Skip to content

Commit

Permalink
[IMP] move core CUDA RT macros to cuda_rt_essentials.hpp (#1584)
Browse files Browse the repository at this point in the history
The reasoning behind this PR is as follows:
for now, anyone wanting to use `RAFT_CUDA_TRY_NO_THROW` still needs to include `cudart_utils.hpp` which can be costly (compilation) due to the include of `memory_pool.hpp`.
By moving the macros to the essentials, we should not break anything for anyone, but allow anyone to improve compilation times by including the essentials only. At the same time, it should add minimal overhead to the compilation time of the essentials file since the pre-processor is (usually) fast compared to the rest of the compilation pipeline.

Authors:
  - Matt Joux (https://github.com/MatthiasKohl)

Approvers:
  - Allard Hendriksen (https://github.com/ahendriksen)
  - Corey J. Nolet (https://github.com/cjnolet)

URL: #1584
  • Loading branch information
MatthiasKohl authored Jun 10, 2023
1 parent ad0c1c1 commit 5dd7017
Show file tree
Hide file tree
Showing 2 changed files with 37 additions and 35 deletions.
37 changes: 37 additions & 0 deletions cpp/include/raft/util/cuda_rt_essentials.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -23,6 +23,8 @@
#include <cuda_runtime.h>
#include <raft/core/error.hpp>

#include <cstdio>

namespace raft {

/**
Expand Down Expand Up @@ -58,3 +60,38 @@ struct cuda_error : public raft::exception {
throw raft::cuda_error(msg); \
} \
} while (0)

/**
* @brief Debug macro to check for CUDA errors
*
* In a non-release build, this macro will synchronize the specified stream
* before error checking. In both release and non-release builds, this macro
* checks for any pending CUDA errors from previous calls. If an error is
* reported, an exception is thrown detailing the CUDA error that occurred.
*
* The intent of this macro is to provide a mechanism for synchronous and
* deterministic execution for debugging asynchronous CUDA execution. It should
* be used after any asynchronous CUDA call, e.g., cudaMemcpyAsync, or an
* asynchronous kernel launch.
*/
#ifndef NDEBUG
#define RAFT_CHECK_CUDA(stream) RAFT_CUDA_TRY(cudaStreamSynchronize(stream));
#else
#define RAFT_CHECK_CUDA(stream) RAFT_CUDA_TRY(cudaPeekAtLastError());
#endif

// /**
// * @brief check for cuda runtime API errors but log error instead of raising
// * exception.
// */
#define RAFT_CUDA_TRY_NO_THROW(call) \
do { \
cudaError_t const status = call; \
if (cudaSuccess != status) { \
printf("CUDA call='%s' at file=%s line=%d failed with %s\n", \
#call, \
__FILE__, \
__LINE__, \
cudaGetErrorString(status)); \
} \
} while (0)
35 changes: 0 additions & 35 deletions cpp/include/raft/util/cudart_utils.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -34,41 +34,6 @@
#include <mutex>
#include <string>

/**
* @brief Debug macro to check for CUDA errors
*
* In a non-release build, this macro will synchronize the specified stream
* before error checking. In both release and non-release builds, this macro
* checks for any pending CUDA errors from previous calls. If an error is
* reported, an exception is thrown detailing the CUDA error that occurred.
*
* The intent of this macro is to provide a mechanism for synchronous and
* deterministic execution for debugging asynchronous CUDA execution. It should
* be used after any asynchronous CUDA call, e.g., cudaMemcpyAsync, or an
* asynchronous kernel launch.
*/
#ifndef NDEBUG
#define RAFT_CHECK_CUDA(stream) RAFT_CUDA_TRY(cudaStreamSynchronize(stream));
#else
#define RAFT_CHECK_CUDA(stream) RAFT_CUDA_TRY(cudaPeekAtLastError());
#endif

// /**
// * @brief check for cuda runtime API errors but log error instead of raising
// * exception.
// */
#define RAFT_CUDA_TRY_NO_THROW(call) \
do { \
cudaError_t const status = call; \
if (cudaSuccess != status) { \
printf("CUDA call='%s' at file=%s line=%d failed with %s\n", \
#call, \
__FILE__, \
__LINE__, \
cudaGetErrorString(status)); \
} \
} while (0)

namespace raft {

/** Helper method to get to know warp size in device code */
Expand Down

0 comments on commit 5dd7017

Please sign in to comment.