From 56254839bef10809f1a9276995cf642e78aac530 Mon Sep 17 00:00:00 2001 From: Artem Chirkin <9253178+achirkin@users.noreply.github.com> Date: Wed, 19 Aug 2026 07:34:13 -0700 Subject: [PATCH 1/4] Better error reporting with std::source_location --- cpp/include/raft/common/nccl_macros.hpp | 26 ++- cpp/include/raft/core/cublas_macros.hpp | 27 ++- cpp/include/raft/core/cusolver_macros.hpp | 27 ++- cpp/include/raft/core/cusparse_macros.hpp | 27 ++- cpp/include/raft/core/detail/macros.hpp | 16 +- cpp/include/raft/core/device_resources.hpp | 17 +- cpp/include/raft/core/error.hpp | 161 +++++++++++++++--- cpp/include/raft/core/interruptible.hpp | 45 +++-- .../raft/core/resource/cuda_stream.hpp | 20 ++- .../raft/core/resource/stream_view.hpp | 20 ++- cpp/include/raft/core/stream_view.hpp | 13 +- cpp/include/raft/util/cuda_rt_essentials.hpp | 47 +++-- cpp/include/raft/util/kernel_launch.hpp | 47 +---- cpp/tests/CMakeLists.txt | 3 +- cpp/tests/core/error.cpp | 154 +++++++++++++++++ cpp/tests/core/interruptible.cu | 119 +++++++++++++ cpp/tests/util/cudart_utils.cpp | 31 +++- 17 files changed, 628 insertions(+), 172 deletions(-) create mode 100644 cpp/tests/core/error.cpp diff --git a/cpp/include/raft/common/nccl_macros.hpp b/cpp/include/raft/common/nccl_macros.hpp index 0f1defbac3..c53846b745 100644 --- a/cpp/include/raft/common/nccl_macros.hpp +++ b/cpp/include/raft/common/nccl_macros.hpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2025, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -9,19 +9,17 @@ * Invokes a NCCL runtime API function call, if the call does not return ncclSuccess, throws an * exception detailing the NCCL error that occurred */ -#define RAFT_NCCL_TRY(call) \ - do { \ - ncclResult_t const status = (call); \ - if (ncclSuccess != status) { \ - std::string msg{}; \ - SET_ERROR_MSG(msg, \ - "NCCL error encountered at: ", \ - "call='%s', Reason=%d:%s", \ - #call, \ - status, \ - ncclGetErrorString(status)); \ - throw raft::logic_error(msg); \ - } \ +#define RAFT_NCCL_TRY(call) \ + do { \ + ncclResult_t const status = (call); \ + if (ncclSuccess != status) { \ + throw raft::logic_error(raft::format_error_message(std::source_location::current(), \ + "NCCL error encountered at: ", \ + "call='%s', Reason=%d:%s", \ + #call, \ + status, \ + ncclGetErrorString(status))); \ + } \ } while (0); // FIXME: Remove after consumer rename diff --git a/cpp/include/raft/core/cublas_macros.hpp b/cpp/include/raft/core/cublas_macros.hpp index 4d9cf7043d..e94703514b 100644 --- a/cpp/include/raft/core/cublas_macros.hpp +++ b/cpp/include/raft/core/cublas_macros.hpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -77,19 +77,18 @@ inline const char* cublas_error_to_string(cublasStatus_t err) * Invokes a cuBLAS runtime API function call, if the call does not return * CUBLAS_STATUS_SUCCESS, throws an exception detailing the cuBLAS error that occurred */ -#define RAFT_CUBLAS_TRY(call) \ - do { \ - cublasStatus_t const status = (call); \ - if (CUBLAS_STATUS_SUCCESS != status) { \ - std::string msg{}; \ - SET_ERROR_MSG(msg, \ - "cuBLAS error encountered at: ", \ - "call='%s', Reason=%d:%s", \ - #call, \ - status, \ - raft::linalg::detail::cublas_error_to_string(status)); \ - throw raft::cublas_error(msg); \ - } \ +#define RAFT_CUBLAS_TRY(call) \ + do { \ + cublasStatus_t const status = (call); \ + if (CUBLAS_STATUS_SUCCESS != status) { \ + throw raft::cublas_error( \ + raft::format_error_message(std::source_location::current(), \ + "cuBLAS error encountered at: ", \ + "call='%s', Reason=%d:%s", \ + #call, \ + status, \ + raft::linalg::detail::cublas_error_to_string(status))); \ + } \ } while (0) // FIXME: Remove after consumers rename diff --git a/cpp/include/raft/core/cusolver_macros.hpp b/cpp/include/raft/core/cusolver_macros.hpp index 7496feb058..b957f171eb 100644 --- a/cpp/include/raft/core/cusolver_macros.hpp +++ b/cpp/include/raft/core/cusolver_macros.hpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -78,19 +78,18 @@ inline const char* cusolver_error_to_string(cusolverStatus_t err) * Invokes a cuSOLVER runtime API function call, if the call does not return * CUSolver_STATUS_SUCCESS, throws an exception detailing the cuSOLVER error that occurred */ -#define RAFT_CUSOLVER_TRY(call) \ - do { \ - cusolverStatus_t const status = (call); \ - if (CUSOLVER_STATUS_SUCCESS != status) { \ - std::string msg{}; \ - SET_ERROR_MSG(msg, \ - "cuSOLVER error encountered at: ", \ - "call='%s', Reason=%d:%s", \ - #call, \ - status, \ - raft::linalg::detail::cusolver_error_to_string(status)); \ - throw raft::cusolver_error(msg); \ - } \ +#define RAFT_CUSOLVER_TRY(call) \ + do { \ + cusolverStatus_t const status = (call); \ + if (CUSOLVER_STATUS_SUCCESS != status) { \ + throw raft::cusolver_error( \ + raft::format_error_message(std::source_location::current(), \ + "cuSOLVER error encountered at: ", \ + "call='%s', Reason=%d:%s", \ + #call, \ + status, \ + raft::linalg::detail::cusolver_error_to_string(status))); \ + } \ } while (0) // FIXME: remove after consumer rename diff --git a/cpp/include/raft/core/cusparse_macros.hpp b/cpp/include/raft/core/cusparse_macros.hpp index d1a8495d12..1eeb26472c 100644 --- a/cpp/include/raft/core/cusparse_macros.hpp +++ b/cpp/include/raft/core/cusparse_macros.hpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -70,19 +70,18 @@ inline const char* cusparse_error_to_string(cusparseStatus_t err) * Invokes a cuSparse runtime API function call, if the call does not return * CUSPARSE_STATUS_SUCCESS, throws an exception detailing the cuSparse error that occurred */ -#define RAFT_CUSPARSE_TRY(call) \ - do { \ - cusparseStatus_t const status = (call); \ - if (CUSPARSE_STATUS_SUCCESS != status) { \ - std::string msg{}; \ - SET_ERROR_MSG(msg, \ - "cuSparse error encountered at: ", \ - "call='%s', Reason=%d:%s", \ - #call, \ - status, \ - raft::sparse::detail::cusparse_error_to_string(status)); \ - throw raft::cusparse_error(msg); \ - } \ +#define RAFT_CUSPARSE_TRY(call) \ + do { \ + cusparseStatus_t const status = (call); \ + if (CUSPARSE_STATUS_SUCCESS != status) { \ + throw raft::cusparse_error( \ + raft::format_error_message(std::source_location::current(), \ + "cuSparse error encountered at: ", \ + "call='%s', Reason=%d:%s", \ + #call, \ + status, \ + raft::sparse::detail::cusparse_error_to_string(status))); \ + } \ } while (0) /** diff --git a/cpp/include/raft/core/detail/macros.hpp b/cpp/include/raft/core/detail/macros.hpp index 776d1dd971..aa69e38438 100644 --- a/cpp/include/raft/core/detail/macros.hpp +++ b/cpp/include/raft/core/detail/macros.hpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -131,3 +131,17 @@ #define RAFT_STRINGIFY_DETAIL(...) #__VA_ARGS__ #define RAFT_STRINGIFY(...) RAFT_STRINGIFY_DETAIL(__VA_ARGS__) #endif + +// Let the compiler check the arguments of a printf-style function against its format string, the +// same way it checks calls to printf itself. +// +// `fmt_index` is the one-based index of the format string parameter and `first_arg_index` the index +// of the first variadic argument (for a non-static member function, `this` is parameter one). +#ifndef RAFT_FORMAT_PRINTF +#if defined(__GNUC__) +#define RAFT_FORMAT_PRINTF(fmt_index, first_arg_index) \ + __attribute__((format(printf, fmt_index, first_arg_index))) +#else +#define RAFT_FORMAT_PRINTF(fmt_index, first_arg_index) +#endif +#endif diff --git a/cpp/include/raft/core/device_resources.hpp b/cpp/include/raft/core/device_resources.hpp index 2cfad4f700..9471f738d9 100644 --- a/cpp/include/raft/core/device_resources.hpp +++ b/cpp/include/raft/core/device_resources.hpp @@ -38,6 +38,7 @@ #include #include +#include #include #include #include @@ -113,13 +114,25 @@ class device_resources : public resources { /** * @brief synchronize a stream on the current container + * + * @param[in] stream stream to synchronize + * @param[in] location the call site to blame for the errors; leave at its default */ - void sync_stream(rmm::cuda_stream_view stream) const { resource::sync_stream(*this, stream); } + void sync_stream(rmm::cuda_stream_view stream, + std::source_location location = std::source_location::current()) const + { + resource::sync_stream(*this, stream, location); + } /** * @brief synchronize main stream on the current container + * + * @param[in] location the call site to blame for the errors; leave at its default */ - void sync_stream() const { resource::sync_stream(*this); } + void sync_stream(std::source_location location = std::source_location::current()) const + { + resource::sync_stream(*this, location); + } /** * @brief returns main stream on the current container diff --git a/cpp/include/raft/core/error.hpp b/cpp/include/raft/core/error.hpp index 9cfbe699d0..633dcfaa1e 100644 --- a/cpp/include/raft/core/error.hpp +++ b/cpp/include/raft/core/error.hpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -13,9 +13,12 @@ #define ENABLE_COLLECT_CALLSTACK #endif +#include +#include #include #include #include +#include #include #include #include @@ -29,6 +32,51 @@ namespace RAFT_EXPORT raft { +namespace detail { + +/** + * @brief Write into @p out the message described by @p location, @p location_prefix and @p fmt. + * + * Does not throw, so that the caller can `va_end` its arguments before reporting a failure. + * + * @param[out] out receives the formatted message; untouched on failure + * @param[in] location the call site to blame + * @param[in] location_prefix text placed in front of the location, e.g. "RAFT failure at " + * @param[in] fmt printf-style format string of the reason + * @param[in] args arguments of @p fmt, started and ended by the caller + * @return whether the message could be formatted + */ +[[nodiscard]] inline auto try_vformat_error_message(std::string& out, + std::source_location location, + char const* location_prefix, + char const* fmt, + std::va_list args) -> bool +{ + char const* location_fmt = "file=%s line=%d function=%s: "; + char const* file = location.file_name(); + auto line = static_cast(location.line()); + char const* function = location.function_name(); + + std::va_list args_measure; + va_copy(args_measure, args); + int size1 = std::snprintf(nullptr, 0, "%s", location_prefix); + int size2 = std::snprintf(nullptr, 0, location_fmt, file, line, function); + int size3 = std::vsnprintf(nullptr, 0, fmt, args_measure); + va_end(args_measure); + if (size1 < 0 || size2 < 0 || size3 < 0) { return false; } + + auto size = static_cast(size1 + size2 + size3 + 1); /* +1 for final '\0' */ + std::vector buf(size); + std::snprintf(buf.data(), static_cast(size1) + 1, "%s", location_prefix); + std::snprintf( + buf.data() + size1, static_cast(size2) + 1, location_fmt, file, line, function); + std::vsnprintf(buf.data() + size1 + size2, static_cast(size3) + 1, fmt, args); + out.assign(buf.data(), buf.data() + size - 1); /* -1 to remove final '\0' */ + return true; +} + +} // namespace detail + /** * @defgroup error_handling Exceptions & Error Handling * @{ @@ -149,10 +197,78 @@ struct non_cuda_build_error : public raft::exception { explicit non_cuda_build_error(std::string const& message) : raft::exception(message) {} }; +/** + * @brief Format an error message that blames the call site given by @p location. + * + * A function that reports errors on behalf of its caller takes a `std::source_location` defaulted + * to `std::source_location::current()` and forwards it here, so that the message points at the + * caller rather than at the implementation: + * @code + * void sync(cudaStream_t stream, std::source_location loc = std::source_location::current()) + * { + * auto status = cudaStreamSynchronize(stream); + * if (status != cudaSuccess) { + * throw raft::cuda_error(raft::format_error_message( + * loc, "CUDA error encountered at: ", "Reason=%s", cudaGetErrorString(status))); + * } + * } + * @endcode + * + * The enclosing function is reported next to the file and the line, since it names the template + * instantiation that the file and the line alone cannot. + * + * @param[in] location the call site to blame + * @param[in] location_prefix text placed in front of the location, e.g. "RAFT failure at " + * @param[in] fmt printf-style format string describing the reason of the error + * @param[in] ... arguments of @p fmt; only types that may be passed through `...` are allowed + * @return the message, of the form "file= line= function=: " + */ +[[nodiscard]] RAFT_FORMAT_PRINTF(3, 4) inline auto format_error_message( + std::source_location location, char const* location_prefix, char const* fmt, ...) -> std::string +{ + std::string msg{}; + std::va_list args; + va_start(args, fmt); + bool const formatted = + detail::try_vformat_error_message(msg, location, location_prefix, fmt, args); + va_end(args); + if (!formatted) { throw raft::exception("Error in snprintf, cannot handle raft exception."); } + return msg; +} + /** * @} */ +namespace detail { + +/** + * @brief The implementation of the deprecated SET_ERROR_MSG macro. + * + * Same as raft::format_error_message; exists as a separate entry point only to carry the + * deprecation warning of the macro to its call sites. + */ +#ifndef RAFT_HIDE_DEPRECATION_WARNINGS +[[deprecated( + "SET_ERROR_MSG is deprecated, use raft::format_error_message(" + "std::source_location::current(), location_prefix, fmt, ...) instead")]] +#endif +RAFT_FORMAT_PRINTF(3, 4) inline auto format_error_message_deprecated(std::source_location location, + char const* location_prefix, + char const* fmt, + ...) -> std::string +{ + std::string msg{}; + std::va_list args; + va_start(args, fmt); + bool const formatted = try_vformat_error_message(msg, location, location_prefix, fmt, args); + va_end(args); + if (!formatted) { throw raft::exception("Error in snprintf, cannot handle raft exception."); } + return msg; +} + +} // namespace detail + } // namespace RAFT_EXPORT raft // FIXME: Need to be replaced with RAFT_FAIL @@ -186,21 +302,14 @@ struct non_cuda_build_error : public raft::exception { /** * Macro to append error message to first argument. * This should only be called in contexts where it is OK to throw exceptions! + * + * @deprecated use raft::format_error_message instead, which does not need a macro to know the call + * site and can be forwarded a location captured elsewhere. */ -#define SET_ERROR_MSG(msg, location_prefix, fmt, ...) \ - do { \ - int size1 = std::snprintf(nullptr, 0, "%s", location_prefix); \ - int size2 = std::snprintf(nullptr, 0, "file=%s line=%d: ", __FILE__, __LINE__); \ - int size3 = std::snprintf(nullptr, 0, fmt, ##__VA_ARGS__); \ - if (size1 < 0 || size2 < 0 || size3 < 0) \ - throw raft::exception("Error in snprintf, cannot handle raft exception."); \ - auto size = size1 + size2 + size3 + 1; /* +1 for final '\0' */ \ - std::vector buf(size); \ - std::snprintf(buf.data(), size1 + 1 /* +1 for '\0' */, "%s", location_prefix); \ - std::snprintf( \ - buf.data() + size1, size2 + 1 /* +1 for '\0' */, "file=%s line=%d: ", __FILE__, __LINE__); \ - std::snprintf(buf.data() + size1 + size2, size3 + 1 /* +1 for '\0' */, fmt, ##__VA_ARGS__); \ - msg += std::string(buf.data(), buf.data() + size - 1); /* -1 to remove final '\0' */ \ +#define SET_ERROR_MSG(msg, location_prefix, fmt, ...) \ + do { \ + msg += raft::detail::format_error_message_deprecated( \ + std::source_location::current(), location_prefix, fmt, ##__VA_ARGS__); \ } while (0) /** @@ -216,13 +325,12 @@ struct non_cuda_build_error : public raft::exception { * optional format tagas * @throw raft::logic_error if the condition evaluates to false. */ -#define RAFT_EXPECTS(cond, fmt, ...) \ - do { \ - if (!(cond)) { \ - std::string msg{}; \ - SET_ERROR_MSG(msg, "RAFT failure at ", fmt, ##__VA_ARGS__); \ - throw raft::logic_error(msg); \ - } \ +#define RAFT_EXPECTS(cond, fmt, ...) \ + do { \ + if (!(cond)) { \ + throw raft::logic_error(raft::format_error_message( \ + std::source_location::current(), "RAFT failure at ", fmt, ##__VA_ARGS__)); \ + } \ } while (0) /** @@ -232,11 +340,10 @@ struct non_cuda_build_error : public raft::exception { * optional format tagas * @throw always throws raft::logic_error */ -#define RAFT_FAIL(fmt, ...) \ - do { \ - std::string msg{}; \ - SET_ERROR_MSG(msg, "RAFT failure at ", fmt, ##__VA_ARGS__); \ - throw raft::logic_error(msg); \ +#define RAFT_FAIL(fmt, ...) \ + do { \ + throw raft::logic_error(raft::format_error_message( \ + std::source_location::current(), "RAFT failure at ", fmt, ##__VA_ARGS__)); \ } while (0) /** diff --git a/cpp/include/raft/core/interruptible.hpp b/cpp/include/raft/core/interruptible.hpp index 9b04d4a64d..0ce89b8238 100644 --- a/cpp/include/raft/core/interruptible.hpp +++ b/cpp/include/raft/core/interruptible.hpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -18,6 +18,7 @@ #include #include #include +#include #include #include @@ -68,14 +69,17 @@ class interruptible { * called on this CPU thread. * * @param [in] stream a CUDA stream. + * @param [in] location the call site to blame for the errors; leave at its default unless + * synchronizing on behalf of a caller, in which case forward the caller's location. * * @throw raft::interrupted_exception if interruptible::cancel() was called on the current CPU * thread before the currently captured work has been finished. * @throw raft::cuda_error if another CUDA error happens. */ - static inline void synchronize(rmm::cuda_stream_view stream) + static inline void synchronize(rmm::cuda_stream_view stream, + std::source_location location = std::source_location::current()) { - get_token()->synchronize_impl(cudaStreamQuery, stream); + get_token()->synchronize_impl(cudaStreamQuery, stream, "cudaStreamQuery", location); } /** @@ -83,14 +87,17 @@ class interruptible { * called on this CPU thread. * * @param [in] event a CUDA event. + * @param [in] location the call site to blame for the errors; leave at its default unless + * synchronizing on behalf of a caller, in which case forward the caller's location. * * @throw raft::interrupted_exception if interruptible::cancel() was called on the current CPU * thread before the currently captured work has been finished. * @throw raft::cuda_error if another CUDA error happens. */ - static inline void synchronize(cudaEvent_t event) + static inline void synchronize(cudaEvent_t event, + std::source_location location = std::source_location::current()) { - get_token()->synchronize_impl(cudaEventQuery, event); + get_token()->synchronize_impl(cudaEventQuery, event, "cudaEventQuery", location); } /** @@ -104,10 +111,16 @@ class interruptible { * * Both `yield` and `yield_no_throw` reset the state to non-cancelled after execution. * + * @param [in] location the call site to blame for the interruption; leave at its default unless + * yielding on behalf of a caller, in which case forward the caller's location. + * * @throw raft::interrupted_exception if interruptible::cancel() was called on the current CPU * thread. */ - static inline void yield() { get_token()->yield_impl(); } + static inline void yield(std::source_location location = std::source_location::current()) + { + get_token()->yield_impl(location); + } /** * @brief Check the thread state, whether the thread can continue execution or is interrupted by @@ -269,10 +282,11 @@ class interruptible { interruptible() noexcept { yield_no_throw_impl(); } - void yield_impl() + void yield_impl(std::source_location location) { if (!yield_no_throw_impl()) { - throw interrupted_exception("The work in this thread was cancelled."); + throw interrupted_exception(raft::format_error_message( + location, "RAFT failure at ", "The work in this thread was cancelled.")); } } @@ -281,17 +295,26 @@ class interruptible { return continue_.test_and_set(std::memory_order_relaxed); } + /** + * @param [in] query the CUDA API function polling the state of @p object + * @param [in] object the stream or the event to wait for + * @param [in] call the name of @p query, as it should appear in an error message + * @param [in] location the call site to blame for the errors + */ template - inline void synchronize_impl(Query query, Object object) + inline void synchronize_impl(Query query, + Object object, + char const* call, + std::source_location location) { cudaError_t query_result; while (true) { - yield_impl(); + yield_impl(location); query_result = query(object); if (query_result != cudaErrorNotReady) { break; } std::this_thread::yield(); } - RAFT_CUDA_TRY(query_result); + raft::check_cuda_error(query_result, call, location); } }; diff --git a/cpp/include/raft/core/resource/cuda_stream.hpp b/cpp/include/raft/core/resource/cuda_stream.hpp index a2c6036abb..8f0e4aec32 100644 --- a/cpp/include/raft/core/resource/cuda_stream.hpp +++ b/cpp/include/raft/core/resource/cuda_stream.hpp @@ -14,6 +14,8 @@ #include +#include + namespace RAFT_EXPORT raft { namespace resource { class cuda_stream_resource : public resource { @@ -81,16 +83,28 @@ inline void set_cuda_stream(resources& res, rmm::cuda_stream_view stream_view) * * @param[in] res the raft resources object * @param[in] stream stream to synchronize + * @param[in] location the call site to blame for the errors; leave at its default unless + * synchronizing on behalf of a caller, in which case forward the caller's location. */ -inline void sync_stream(const resources& res, rmm::cuda_stream_view stream) +inline void sync_stream(const resources& res, + rmm::cuda_stream_view stream, + std::source_location location = std::source_location::current()) { - interruptible::synchronize(stream); + interruptible::synchronize(stream, location); } /** * @brief synchronize main stream on the resources instance + * + * @param[in] res the raft resources object + * @param[in] location the call site to blame for the errors; leave at its default unless + * synchronizing on behalf of a caller, in which case forward the caller's location. */ -inline void sync_stream(const resources& res) { sync_stream(res, get_cuda_stream(res)); } +inline void sync_stream(const resources& res, + std::source_location location = std::source_location::current()) +{ + sync_stream(res, get_cuda_stream(res), location); +} /** * @} diff --git a/cpp/include/raft/core/resource/stream_view.hpp b/cpp/include/raft/core/resource/stream_view.hpp index 86cdbf6e33..c5eb25f82a 100644 --- a/cpp/include/raft/core/resource/stream_view.hpp +++ b/cpp/include/raft/core/resource/stream_view.hpp @@ -11,6 +11,8 @@ #include #endif +#include + namespace RAFT_EXPORT raft { namespace resource { struct stream_view_resource : public resource { @@ -74,16 +76,28 @@ inline void set_stream_view(resources& res, raft::stream_view view) * * @param[in] res the raft resources object * @param[in] stream stream to synchronize + * @param[in] location the call site to blame for the errors; leave at its default unless + * synchronizing on behalf of a caller, in which case forward the caller's location. */ -inline void sync_stream_view(const resources& res, raft::stream_view stream) +inline void sync_stream_view(const resources& res, + raft::stream_view stream, + std::source_location location = std::source_location::current()) { - stream.interruptible_synchronize(); + stream.interruptible_synchronize(location); } /** * @brief synchronize main stream on the resources instance + * + * @param[in] res the raft resources object + * @param[in] location the call site to blame for the errors; leave at its default unless + * synchronizing on behalf of a caller, in which case forward the caller's location. */ -inline void sync_stream_view(const resources& res) { sync_stream_view(res, get_stream_view(res)); } +inline void sync_stream_view(const resources& res, + std::source_location location = std::source_location::current()) +{ + sync_stream_view(res, get_stream_view(res), location); +} /** * @} diff --git a/cpp/include/raft/core/stream_view.hpp b/cpp/include/raft/core/stream_view.hpp index 9a9180da50..5d4a8b6936 100644 --- a/cpp/include/raft/core/stream_view.hpp +++ b/cpp/include/raft/core/stream_view.hpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once @@ -13,6 +13,8 @@ #include #endif +#include + namespace RAFT_EXPORT raft { } // namespace RAFT_EXPORT raft @@ -72,10 +74,15 @@ struct stream_view { [[nodiscard]] auto is_default() const { return base_view_.is_default(); } void synchronize() const { base_view_.synchronize(); } void synchronize_no_throw() const { base_view_.synchronize_no_throw(); } - void interruptible_synchronize() const + /** + * @param[in] location the call site to blame for the errors; leave at its default unless + * synchronizing on behalf of a caller, in which case forward the caller's location. + */ + void interruptible_synchronize( + [[maybe_unused]] std::source_location location = std::source_location::current()) const { #ifndef RAFT_DISABLE_CUDA - interruptible::synchronize(base_view_); + interruptible::synchronize(base_view_, location); #else synchronize(); #endif diff --git a/cpp/include/raft/util/cuda_rt_essentials.hpp b/cpp/include/raft/util/cuda_rt_essentials.hpp index 3ce7636ab4..96c3863c30 100644 --- a/cpp/include/raft/util/cuda_rt_essentials.hpp +++ b/cpp/include/raft/util/cuda_rt_essentials.hpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2023-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2023-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -15,6 +15,7 @@ #include #include +#include namespace raft { @@ -26,6 +27,33 @@ struct cuda_error : public raft::exception { explicit cuda_error(std::string const& message) : raft::exception(message) {} }; +/** + * @brief Throw raft::cuda_error blaming @p location, unless @p status is cudaSuccess. + * + * The function form of RAFT_CUDA_TRY: a utility that checks a CUDA call on behalf of its caller + * forwards the location it received, so that the reported location is the caller's rather than the + * utility's. + * + * @param[in] status the status returned by a CUDA runtime API call + * @param[in] call the name of the CUDA runtime API call, as it should appear in the message + * @param[in] location the call site to blame; leave at its default unless forwarding one + * + * @throw raft::cuda_error if @p status is not cudaSuccess + */ +inline void check_cuda_error(cudaError_t status, + char const* call, + std::source_location location = std::source_location::current()) +{ + if (status == cudaSuccess) { return; } + cudaGetLastError(); // clear the error, so that it does not affect the subsequent calls + throw cuda_error(format_error_message(location, + "CUDA error encountered at: ", + "call='%s', Reason=%s:%s", + call, + cudaGetErrorName(status), + cudaGetErrorString(status))); +} + } // namespace raft /** @@ -36,20 +64,9 @@ struct cuda_error : public raft::exception { * exception detailing the CUDA error that occurred * */ -#define RAFT_CUDA_TRY(call) \ - do { \ - cudaError_t const status = call; \ - if (status != cudaSuccess) { \ - cudaGetLastError(); \ - std::string msg{}; \ - SET_ERROR_MSG(msg, \ - "CUDA error encountered at: ", \ - "call='%s', Reason=%s:%s", \ - #call, \ - cudaGetErrorName(status), \ - cudaGetErrorString(status)); \ - throw raft::cuda_error(msg); \ - } \ +#define RAFT_CUDA_TRY(call) \ + do { \ + raft::check_cuda_error((call), #call); \ } while (0) /** diff --git a/cpp/include/raft/util/kernel_launch.hpp b/cpp/include/raft/util/kernel_launch.hpp index 7fb1a064a2..8172392f4d 100644 --- a/cpp/include/raft/util/kernel_launch.hpp +++ b/cpp/include/raft/util/kernel_launch.hpp @@ -15,56 +15,15 @@ #include #include -#include #include #include -#include #include #include -#include namespace raft { namespace detail { -/** - * @brief Format a cuda_error message with an explicit call-site location. - * - * Mirrors SET_ERROR_MSG / RAFT_CUDA_TRY formatting but does not use those macros, so the reported - * location is the caller's rather than this header. The enclosing function is reported too, since - * it names the template instantiation that the file and line alone cannot. - */ -inline std::string format_cuda_launch_error(cudaError_t status, std::source_location location) -{ - char const* location_prefix = "CUDA error encountered at: "; - char const* location_fmt = "file=%s line=%d function=%s: "; - char const* fmt = "call='%s', Reason=%s:%s"; - char const* call = "cudaLaunchKernelExC"; - char const* file = location.file_name(); - auto line = static_cast(location.line()); - char const* function = location.function_name(); - - int size1 = std::snprintf(nullptr, 0, "%s", location_prefix); - int size2 = std::snprintf(nullptr, 0, location_fmt, file, line, function); - int size3 = - std::snprintf(nullptr, 0, fmt, call, cudaGetErrorName(status), cudaGetErrorString(status)); - if (size1 < 0 || size2 < 0 || size3 < 0) { - throw raft::exception("Error in snprintf, cannot handle raft exception."); - } - auto size = static_cast(size1 + size2 + size3 + 1); - std::vector buf(size); - std::snprintf(buf.data(), static_cast(size1) + 1, "%s", location_prefix); - std::snprintf( - buf.data() + size1, static_cast(size2) + 1, location_fmt, file, line, function); - std::snprintf(buf.data() + size1 + size2, - static_cast(size3) + 1, - fmt, - call, - cudaGetErrorName(status), - cudaGetErrorString(status)); - return std::string(buf.data(), buf.data() + size - 1); -} - /** * @brief Launch a kernel, copying the launch arguments into parameters first. * @@ -79,10 +38,8 @@ void dispatch(cudaLaunchConfig_t const& config, { std::array arg_ptrs{ {const_cast(static_cast(std::addressof(params)))...}}; - cudaError_t status = cudaLaunchKernelExC(&config, kernel, arg_ptrs.data()); - if (status == cudaSuccess) { return; } - cudaGetLastError(); // clear sticky error - throw raft::cuda_error(format_cuda_launch_error(status, location)); + raft::check_cuda_error( + cudaLaunchKernelExC(&config, kernel, arg_ptrs.data()), "cudaLaunchKernelExC", location); } } // namespace detail diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index 9207a64cb3..54fd6f13d8 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -95,6 +95,7 @@ if(BUILD_TESTS) core/bitset.cu core/device_resources_manager.cpp core/device_setter.cpp + core/error.cpp core/math_device.cu core/math_host.cpp core/operators_device.cu @@ -121,7 +122,7 @@ if(BUILD_TESTS) ConfigureTest( NAME CORE_TEST PATH core/stream_view.cpp core/mdspan_copy.cpp core/host_memory_resource.cpp - core/allocation_tracking.cpp LIB EXPLICIT_INSTANTIATE_ONLY NOCUDA + core/allocation_tracking.cpp core/error.cpp LIB EXPLICIT_INSTANTIATE_ONLY NOCUDA ) list(APPEND EXT_HEADER_TEST_SOURCES ext_headers/raft_core_logger.cpp) diff --git a/cpp/tests/core/error.cpp b/cpp/tests/core/error.cpp new file mode 100644 index 0000000000..cea729537b --- /dev/null +++ b/cpp/tests/core/error.cpp @@ -0,0 +1,154 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +#include + +#include + +#include +#include +#include + +namespace raft { + +namespace { + +/** The first line of the message, i.e. the message without the appended call stack. */ +auto first_line(std::string const& msg) -> std::string { return msg.substr(0, msg.find('\n')); } + +auto reg_escape(std::string const& s) -> std::string +{ + static std::regex const special_chars{R"([-[\]{}()*+?.,\^$|#\s])"}; + return std::regex_replace(s, special_chars, R"(\$&)"); +} + +/** + * Expected message of an error reported in this file, at a line matching @p line_pattern, in a + * function whose name contains @p function. + */ +auto expected_message(char const* prefix, + std::string const& line_pattern, + char const* function, + std::string const& reason) -> std::regex +{ + std::string re{"^"}; + re += reg_escape(prefix); + re += R"(file=.*error\.cpp line=)"; + re += line_pattern; + re += " function=.*"; + re += function; + re += ".*: "; + re += reg_escape(reason); + re += "$"; + return std::regex{re}; +} + +auto line_of(std::source_location location) -> std::string +{ + return std::to_string(location.line()); +} + +/** Any line of this file; used where the exact line is not the point of the test. */ +constexpr char const* k_any_line = R"(\d+)"; + +/** A utility reporting an error on behalf of its caller. */ +auto report_on_behalf_of_caller(std::source_location location = std::source_location::current()) + -> std::string +{ + return format_error_message(location, "RAFT failure at ", "%s", "reported by a utility"); +} + +} // namespace + +TEST(Error, FormatBlamesTheGivenLocation) +{ + auto const loc = std::source_location::current(); + auto msg = format_error_message(loc, "RAFT failure at ", "value=%d, name='%s'", 42, "answer"); + + EXPECT_TRUE(std::regex_match( + msg, + expected_message( + "RAFT failure at ", line_of(loc), "FormatBlamesTheGivenLocation", "value=42, name='answer'"))) + << "message:'" << msg << "'"; +} + +TEST(Error, FormatBlamesTheCallerOfAUtility) +{ + auto msg = report_on_behalf_of_caller(); + + // This test function is blamed, not report_on_behalf_of_caller and not error.hpp. + EXPECT_TRUE(std::regex_match( + msg, + expected_message( + "RAFT failure at ", k_any_line, "FormatBlamesTheCallerOfAUtility", "reported by a utility"))) + << "message:'" << msg << "'"; + EXPECT_EQ(msg.find("error.hpp"), std::string::npos) << "message:'" << msg << "'"; + EXPECT_EQ(msg.find("report_on_behalf_of_caller"), std::string::npos) << "message:'" << msg << "'"; +} + +TEST(Error, FormatHandlesLongReasons) +{ + std::string reason{"This is a test string repeated many times. "}; + for (size_t i = 0; i < 6; ++i) { + reason += reason; + } + EXPECT_TRUE(reason.size() > 2048) << "size of the test string is: " << reason.size(); + + auto const loc = std::source_location::current(); + auto msg = format_error_message(loc, "RAFT failure at ", (reason + "%d").c_str(), 121); + + EXPECT_TRUE(std::regex_match( + msg, + expected_message("RAFT failure at ", line_of(loc), "FormatHandlesLongReasons", reason + "121"))) + << "message:'" << msg << "'"; +} + +TEST(Error, ExpectsBlamesTheCallSite) +{ + int const x = -1; + try { + RAFT_EXPECTS(x > 0, "x must be positive, got %d", x); + FAIL() << "Expected logic_error from a violated expectation"; + } catch (raft::logic_error const& e) { + auto msg = first_line(e.what()); + EXPECT_TRUE(std::regex_match( + msg, + expected_message( + "RAFT failure at ", k_any_line, "ExpectsBlamesTheCallSite", "x must be positive, got -1"))) + << "message:'" << msg << "'"; + } +} + +TEST(Error, FailBlamesTheCallSite) +{ + try { + RAFT_FAIL("cannot do %s", "that"); + FAIL() << "Expected logic_error from RAFT_FAIL"; + } catch (raft::logic_error const& e) { + auto msg = first_line(e.what()); + EXPECT_TRUE(std::regex_match( + msg, + expected_message("RAFT failure at ", k_any_line, "FailBlamesTheCallSite", "cannot do that"))) + << "message:'" << msg << "'"; + } +} + +// The deprecated macro must keep working for the projects that have not switched yet. +#pragma GCC diagnostic push +#pragma GCC diagnostic ignored "-Wdeprecated-declarations" +TEST(Error, DeprecatedSetErrorMsgAppends) +{ + std::string msg{"prefix:"}; + ASSERT_NO_THROW(SET_ERROR_MSG(msg, "location prefix:", "value=%d", 123)); + + EXPECT_TRUE(std::regex_match( + msg, + expected_message( + "prefix:location prefix:", k_any_line, "DeprecatedSetErrorMsgAppends", "value=123"))) + << "message:'" << msg << "'"; +} +#pragma GCC diagnostic pop + +} // namespace raft diff --git a/cpp/tests/core/interruptible.cu b/cpp/tests/core/interruptible.cu index cdcdd67d96..be039b02d2 100644 --- a/cpp/tests/core/interruptible.cu +++ b/cpp/tests/core/interruptible.cu @@ -6,9 +6,12 @@ #include #include #include +#include +#include #include #include +#include #include #include @@ -16,6 +19,9 @@ #include #include #include +#include +#include +#include #include #include @@ -87,6 +93,119 @@ TEST(Raft, InterruptibleDelayedInit) }).join(); } +namespace { + +/** + * A stream that cannot be synchronized: querying a capturing stream fails with a non-sticky error, + * which is a safe way to exercise the error reporting; the stream is fully usable again afterwards. + */ +class capturing_stream { + public: + capturing_stream() + { + RAFT_CUDA_TRY(cudaStreamBeginCapture(stream_.value(), cudaStreamCaptureModeRelaxed)); + } + + ~capturing_stream() + { + // Querying the stream has invalidated the capture, so ending it is expected to fail. + cudaGraph_t graph{}; + cudaStreamEndCapture(stream_.value(), &graph); + cudaGetLastError(); + } + + capturing_stream(capturing_stream const&) = delete; + capturing_stream(capturing_stream&&) = delete; + auto operator=(capturing_stream const&) -> capturing_stream& = delete; + auto operator=(capturing_stream&&) -> capturing_stream& = delete; + + [[nodiscard]] auto view() const -> rmm::cuda_stream_view { return stream_.view(); } + + private: + rmm::cuda_stream stream_{}; +}; + +/** A utility synchronizing on behalf of its caller. */ +void sync_on_behalf_of_caller(raft::resources const& res, + std::source_location location = std::source_location::current()) +{ + resource::sync_stream(res, location); +} + +} // namespace + +TEST(Raft, InterruptibleSynchronizeBlamesTheCallSite) +{ + capturing_stream stream{}; + std::string caught{}; + auto sync_line = __LINE__ + 2; + try { + interruptible::synchronize(stream.view()); + FAIL() << "Expected cuda_error from synchronizing a capturing stream"; + } catch (raft::cuda_error const& e) { + caught = e.what(); + } + + // Must blame this test translation unit, not the interruptible implementation. + EXPECT_EQ(caught.find("interruptible.hpp"), std::string::npos) << caught; + std::string re_exp{R"(CUDA error encountered at: file=.*interruptible\.cu line=)"}; + re_exp += std::to_string(sync_line); + re_exp += R"( function=.*InterruptibleSynchronizeBlamesTheCallSite.*: )"; + re_exp += R"(call='cudaStreamQuery', Reason=.*)"; + EXPECT_TRUE(std::regex_search(caught, std::regex(re_exp))) + << "message:'" << caught << "'\nexpected regex:'" << re_exp << "'"; +} + +TEST(Raft, SyncStreamBlamesTheCallSite) +{ + capturing_stream stream{}; + raft::resources res; + resource::set_cuda_stream(res, stream.view()); + + std::string caught{}; + try { + sync_on_behalf_of_caller(res); + FAIL() << "Expected cuda_error from synchronizing a capturing stream"; + } catch (raft::cuda_error const& e) { + caught = e.what(); + } + + // The location travels from here through sync_on_behalf_of_caller, resource::sync_stream and + // interruptible::synchronize, so that none of them is blamed. + for (auto const* implementation : + {"interruptible.hpp", "cuda_stream.hpp", "sync_on_behalf_of_caller"}) { + EXPECT_EQ(caught.find(implementation), std::string::npos) << caught; + } + std::string re_exp{R"(CUDA error encountered at: file=.*interruptible\.cu line=\d+)"}; + re_exp += R"( function=.*SyncStreamBlamesTheCallSite.*: )"; + re_exp += R"(call='cudaStreamQuery', Reason=.*)"; + EXPECT_TRUE(std::regex_search(caught, std::regex(re_exp))) + << "message:'" << caught << "'\nexpected regex:'" << re_exp << "'"; +} + +TEST(Raft, InterruptedExceptionBlamesTheCallSite) +{ + interruptible::get_token()->cancel(); + std::string caught{}; + auto yield_line = __LINE__ + 2; + try { + interruptible::yield(); + FAIL() << "Expected interrupted_exception after cancelling this thread"; + } catch (interrupted_exception const& e) { + caught = e.what(); + } + + std::string re_exp{R"(RAFT failure at file=.*interruptible\.cu line=)"}; + re_exp += std::to_string(yield_line); + re_exp += R"( function=.*InterruptedExceptionBlamesTheCallSite.*: )"; + re_exp += R"(The work in this thread was cancelled\.)"; + EXPECT_TRUE(std::regex_search(caught, std::regex(re_exp))) + << "message:'" << caught << "'\nexpected regex:'" << re_exp << "'"; + + // clear the cancellation state to not disrupt other tests + interruptible::yield_no_throw(); +} + TEST(Raft, InterruptibleOpenMP) { // number of threads must be smaller than max number of resident grids for GPU diff --git a/cpp/tests/util/cudart_utils.cpp b/cpp/tests/util/cudart_utils.cpp index 4ecc602201..2d7f56aa46 100644 --- a/cpp/tests/util/cudart_utils.cpp +++ b/cpp/tests/util/cudart_utils.cpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2024, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -13,6 +13,8 @@ #include #include +#include +#include #include namespace raft { @@ -64,18 +66,37 @@ TEST(Raft, Utils) << "expected regex:'" << re_exp << "'"; } - // Now we test SET_ERROR_MSG instead of THROW + // Now we test raft::format_error_message instead of THROW + auto const loc = std::source_location::current(); std::string msg{"prefix:"}; - ASSERT_NO_THROW(SET_ERROR_MSG(msg, "location prefix:", test_format_c, 123)); + msg += raft::format_error_message(loc, "location prefix:", test_format_c, 123); std::string re_exp{"^prefix:location prefix:file="}; re_exp += reg_file; - // test code must be at line >10 (copyright), assume line is never >9999 - re_exp += " line=\\d{2,4}: "; + re_exp += " line="; + re_exp += std::to_string(loc.line()); + re_exp += " function=.*Raft_Utils_Test.*: "; re_exp += reg_escape(test); re_exp += "123$"; EXPECT_TRUE(std::regex_match(msg, std::regex(re_exp))) << "message:'" << msg << "'" << std::endl << "expected regex:'" << re_exp << "'"; + + // The deprecated macro keeps appending the same message, at its own call site + std::string legacy_msg{"prefix:"}; +#pragma GCC diagnostic push +#pragma GCC diagnostic ignored "-Wdeprecated-declarations" + ASSERT_NO_THROW(SET_ERROR_MSG(legacy_msg, "location prefix:", test_format_c, 123)); +#pragma GCC diagnostic pop + + std::string re_exp_legacy{"^prefix:location prefix:file="}; + re_exp_legacy += reg_file; + // test code must be at line >10 (copyright), assume line is never >9999 + re_exp_legacy += " line=\\d{2,4} function=.*Raft_Utils_Test.*: "; + re_exp_legacy += reg_escape(test); + re_exp_legacy += "123$"; + EXPECT_TRUE(std::regex_match(legacy_msg, std::regex(re_exp_legacy))) + << "message:'" << legacy_msg << "'" << std::endl + << "expected regex:'" << re_exp_legacy << "'"; } TEST(Raft, GetDeviceForAddress) From 935a46087bc25a5363f59a3322b8da2491ccbfcc Mon Sep 17 00:00:00 2001 From: Artem Chirkin <9253178+achirkin@users.noreply.github.com> Date: Wed, 19 Aug 2026 07:34:45 -0700 Subject: [PATCH 2/4] Add a note to docs --- docs/source/developer_guide.md | 2 ++ 1 file changed, 2 insertions(+) diff --git a/docs/source/developer_guide.md b/docs/source/developer_guide.md index ae646ce901..1339013520 100644 --- a/docs/source/developer_guide.md +++ b/docs/source/developer_guide.md @@ -227,6 +227,8 @@ Keep in mind that this only applies to files tracked by git that have been modif ## Error handling Call CUDA APIs via the provided helper macros `RAFT_CUDA_TRY`, `RAFT_CUBLAS_TRY` and `RAFT_CUSOLVER_TRY`. These macros take care of checking the return values of the used API calls and generate an exception when the command is not successful. If you need to avoid an exception, e.g. inside a destructor, use `RAFT_CUDA_TRY_NO_THROW`, `RAFT_CUBLAS_TRY_NO_THROW ` and `RAFT_CUSOLVER_TRY_NO_THROW`. These macros log the error but do not throw an exception. +A function that reports an error on behalf of its caller cannot use these macros, because they would blame its own line rather than the caller's. Such a function should take a `std::source_location` parameter defaulted to `std::source_location::current()` and forward it to `raft::format_error_message` (or to `raft::check_cuda_error`, the function form of `RAFT_CUDA_TRY`), so that the reported location is the caller's; `raft::resource::sync_stream` is an example. The `SET_ERROR_MSG` macro is deprecated in favour of `raft::format_error_message`. + ## Logging ### Introduction From 010ff14860baa433d9c7603268ac5f9d73f2cd61 Mon Sep 17 00:00:00 2001 From: Artem Chirkin <9253178+achirkin@users.noreply.github.com> Date: Thu, 20 Aug 2026 05:00:40 -0700 Subject: [PATCH 3/4] Simplify RAFT_CUDA_TRY --- cpp/include/raft/util/cuda_rt_essentials.hpp | 5 +---- 1 file changed, 1 insertion(+), 4 deletions(-) diff --git a/cpp/include/raft/util/cuda_rt_essentials.hpp b/cpp/include/raft/util/cuda_rt_essentials.hpp index 96c3863c30..f5b724d21c 100644 --- a/cpp/include/raft/util/cuda_rt_essentials.hpp +++ b/cpp/include/raft/util/cuda_rt_essentials.hpp @@ -64,10 +64,7 @@ inline void check_cuda_error(cudaError_t status, * exception detailing the CUDA error that occurred * */ -#define RAFT_CUDA_TRY(call) \ - do { \ - raft::check_cuda_error((call), #call); \ - } while (0) + #define RAFT_CUDA_TRY(call) raft::check_cuda_error(call, #call) /** * @brief Debug macro to check for CUDA errors From 30a44d008cc4ac8fcd83bddba85b45cbc91f5fef Mon Sep 17 00:00:00 2001 From: Artem Chirkin <9253178+achirkin@users.noreply.github.com> Date: Thu, 20 Aug 2026 07:08:52 -0700 Subject: [PATCH 4/4] Fix style --- cpp/include/raft/util/cuda_rt_essentials.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/include/raft/util/cuda_rt_essentials.hpp b/cpp/include/raft/util/cuda_rt_essentials.hpp index f5b724d21c..f72ef2da2d 100644 --- a/cpp/include/raft/util/cuda_rt_essentials.hpp +++ b/cpp/include/raft/util/cuda_rt_essentials.hpp @@ -64,7 +64,7 @@ inline void check_cuda_error(cudaError_t status, * exception detailing the CUDA error that occurred * */ - #define RAFT_CUDA_TRY(call) raft::check_cuda_error(call, #call) +#define RAFT_CUDA_TRY(call) raft::check_cuda_error(call, #call) /** * @brief Debug macro to check for CUDA errors