From d748b655e7cbadecb332061b0a22b70cb852691a Mon Sep 17 00:00:00 2001 From: lennartvoelz Date: Fri, 28 Aug 2026 10:24:07 +0200 Subject: [PATCH 1/8] Add MATX_CUDA_CHECK_NOEXCEPT for error checking, logging and draining in noexcept contexts Adds MATX_CUDA_CHECK_NOEXCEPT: - checks the error state against cudaSuccess - on failure, logs the error - logging is guarded by try/catch, to allow safe usage in noexcept contexts (logging with std::format could throw) --- include/matx/core/error.h | 21 +++++++++++++++++++++ 1 file changed, 21 insertions(+) diff --git a/include/matx/core/error.h b/include/matx/core/error.h index 90e872aff..991b079f8 100644 --- a/include/matx/core/error.h +++ b/include/matx/core/error.h @@ -265,6 +265,27 @@ namespace matx MATX_CUDA_CHECK(e); \ } +// Macro for checking cuda errors in noexcept functions and destructors, +// where MATX_CUDA_CHECK's throw would call std::terminate(). Logs the +// error and drains the sticky CUDA error state via cudaGetLastError() +// instead of throwing, so the failure is diagnosed here rather than +// leaking into an unrelated later call. The logging call itself is +// wrapped in try/catch so this macro can be used in noexcept contexts +// (std::format can theoretically throw). +#define MATX_CUDA_CHECK_NOEXCEPT(e) \ + do { \ + const auto e_ = (e); \ + if (e_ != cudaSuccess) \ + { \ + try { \ + MATX_LOG_ERROR("{}:{} CUDA Error (noexcept context): {} ({})", \ + __FILE__, __LINE__, cudaGetErrorString(e_), \ + static_cast(e_)); \ + } catch (...) {} \ + (void)cudaGetLastError(); \ + } \ + } while (0) + // This macro asserts compatible dimensions of current class to an operator. #define MATX_ASSERT_COMPATIBLE_OP_SIZES(op) \ if constexpr (Rank() > 0) { \ From 9ef52fd9bc18086c453d33408275a5d16e5987a6 Mon Sep 17 00:00:00 2001 From: lennartvoelz Date: Fri, 28 Aug 2026 10:56:58 +0200 Subject: [PATCH 2/8] Add MATX_CUDA_CHECK_NOEXCEPT macro around bare CUDA calls in the memory tracker --- include/matx/core/allocator.h | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/include/matx/core/allocator.h b/include/matx/core/allocator.h index c61546720..57c65cfe5 100644 --- a/include/matx/core/allocator.h +++ b/include/matx/core/allocator.h @@ -141,10 +141,10 @@ struct MemTracker { case MATX_MANAGED_MEMORY: [[fallthrough]]; case MATX_DEVICE_MEMORY: - if (is_cuda_free()) cudaFree(ptr); + if (is_cuda_free()) MATX_CUDA_CHECK_NOEXCEPT(cudaFree(ptr)); break; case MATX_HOST_MEMORY: - if (is_cuda_free()) cudaFreeHost(ptr); + if (is_cuda_free()) MATX_CUDA_CHECK_NOEXCEPT(cudaFreeHost(ptr)); break; case MATX_HOST_MALLOC_MEMORY: free(ptr); @@ -152,10 +152,10 @@ struct MemTracker { case MATX_ASYNC_DEVICE_MEMORY: if (is_cuda_free()) { if constexpr (std::is_same_v) { - cudaFreeAsync(ptr, iter->second.stream); + MATX_CUDA_CHECK_NOEXCEPT(cudaFreeAsync(ptr, iter->second.stream)); } else { - cudaFreeAsync(ptr, st.stream); + MATX_CUDA_CHECK_NOEXCEPT(cudaFreeAsync(ptr, st.stream)); } } break; From cfc86b7746ffd2061552b1f12444fe1cff33002a Mon Sep 17 00:00:00 2001 From: lennartvoelz Date: Fri, 28 Aug 2026 11:00:39 +0200 Subject: [PATCH 3/8] Add MATX_CUDA_CHECK around bare CUDA calls in cache.h --- include/matx/core/cache.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/include/matx/core/cache.h b/include/matx/core/cache.h index 29506944c..23ad9c974 100644 --- a/include/matx/core/cache.h +++ b/include/matx/core/cache.h @@ -235,7 +235,7 @@ class matxCache_t { auto &cval = cache[id]; if constexpr (is_cuda_executor_v) { - cudaGetDevice(&key.device_id); + MATX_CUDA_CHECK(cudaGetDevice(&key.device_id)); } else { key.device_id = 0; @@ -262,7 +262,7 @@ class matxCache_t { void *ptr = nullptr; CacheCommonParamsKey key; key.thread_id = std::this_thread::get_id(); - cudaGetDevice(&key.device_id); + MATX_CUDA_CHECK(cudaGetDevice(&key.device_id)); [[maybe_unused]] std::lock_guard lock(stream_alloc_mutex); From 61bf6d73550ea9d76c41527f508ea101d2208f71 Mon Sep 17 00:00:00 2001 From: lennartvoelz Date: Fri, 28 Aug 2026 11:19:35 +0200 Subject: [PATCH 4/8] Add MATX_CUDA_CHECK around bare cuda calls in make_sparse_tensor.h --- include/matx/core/make_sparse_tensor.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/include/matx/core/make_sparse_tensor.h b/include/matx/core/make_sparse_tensor.h index 442eb7e36..a0de745d1 100644 --- a/include/matx/core/make_sparse_tensor.h +++ b/include/matx/core/make_sparse_tensor.h @@ -42,7 +42,7 @@ template __MATX_INLINE__ static void setZero(T *ptr, index_t sz, matxMemorySpace_t space) { if (space == MATX_DEVICE_MEMORY || space == MATX_ASYNC_DEVICE_MEMORY) { - cudaMemset(ptr, 0, sz * sizeof(T)); + MATX_CUDA_CHECK(cudaMemset(ptr, 0, sz * sizeof(T))); } else { memset(ptr, 0, sz * sizeof(T)); } @@ -71,7 +71,7 @@ __MATX_INLINE__ Storage makeEmptyStorage() { template __MATX_INLINE__ static void setVal(T *ptr, T val, matxMemorySpace_t space) { if (space == MATX_DEVICE_MEMORY || space == MATX_ASYNC_DEVICE_MEMORY) { - cudaMemcpy(ptr, &val, sizeof(T), cudaMemcpyHostToDevice); + MATX_CUDA_CHECK(cudaMemcpy(ptr, &val, sizeof(T), cudaMemcpyHostToDevice)); } else { memcpy(ptr, &val, sizeof(T)); } From f631a62c5fbc3ef84cb51e9919d71f9a5f96c867 Mon Sep 17 00:00:00 2001 From: lennartvoelz Date: Fri, 28 Aug 2026 11:55:36 +0200 Subject: [PATCH 5/8] Add MATX_CUDA_CHECK around bare CUDA calls under core/ Covers print.h, pybind.h, and tensor.h's prefetch helpers. --- include/matx/core/print.h | 6 +++--- include/matx/core/pybind.h | 2 +- include/matx/core/tensor.h | 12 ++++++------ 3 files changed, 10 insertions(+), 10 deletions(-) diff --git a/include/matx/core/print.h b/include/matx/core/print.h index 17a03a8d3..939b279c9 100644 --- a/include/matx/core/print.h +++ b/include/matx/core/print.h @@ -554,7 +554,7 @@ namespace matx { MATX_NVTX_START("", matx::MATX_NVTX_LOG_API) #ifdef __CUDACC__ - cudaDeviceSynchronize(); + MATX_CUDA_CHECK(cudaDeviceSynchronize()); if constexpr (is_sparse_tensor_v) { using Format = typename Op::Format; fprintf(fp, "format = "); @@ -631,7 +631,7 @@ namespace matx { for (int i = 0; i < R; i++) shape[i] = op.Size(i); auto tmpv = make_tensor(shape); (tmpv = op).run(CUDAJITExecutor{}); - cudaStreamSynchronize(0); + MATX_CUDA_CHECK(cudaStreamSynchronize(0)); detail::InternalPrint(fp, tmpv, dims...); }; switch (r) { @@ -649,7 +649,7 @@ namespace matx { } else { auto tmpv = make_tensor(op.Shape()); (tmpv = op).run(); - cudaStreamSynchronize(0); + MATX_CUDA_CHECK(cudaStreamSynchronize(0)); detail::InternalPrint(fp, tmpv, dims...); } } diff --git a/include/matx/core/pybind.h b/include/matx/core/pybind.h index 49aca6951..7db0f78ff 100644 --- a/include/matx/core/pybind.h +++ b/include/matx/core/pybind.h @@ -525,7 +525,7 @@ class MATX_PYBIND_VISIBILITY MatXPybind { auto ften = pybind11::array_t(resobj); constexpr int RANK = TensorType::Rank(); - cudaDeviceSynchronize(); + MATX_CUDA_CHECK(cudaDeviceSynchronize()); if constexpr (RANK == 0) { auto file_val = ften.at(); diff --git a/include/matx/core/tensor.h b/include/matx/core/tensor.h index 8c53daac3..f529c6763 100644 --- a/include/matx/core/tensor.h +++ b/include/matx/core/tensor.h @@ -745,14 +745,14 @@ class tensor_t : public detail::tensor_impl_t { MATX_NVTX_START("", matx::MATX_NVTX_LOG_API) int dev; - cudaGetDevice(&dev); + MATX_CUDA_CHECK_NOEXCEPT(cudaGetDevice(&dev)); #if CUDART_VERSION <= 12000 - cudaMemPrefetchAsync(this->Data(), this->desc_.TotalSize() * sizeof(T), dev, stream); + MATX_CUDA_CHECK_NOEXCEPT(cudaMemPrefetchAsync(this->Data(), this->desc_.TotalSize() * sizeof(T), dev, stream)); #else cudaMemLocation loc; loc.id = dev; loc.type = cudaMemLocationTypeDevice; - cudaMemPrefetchAsync(this->Data(), this->desc_.TotalSize() * sizeof(T), loc, 0, stream); + MATX_CUDA_CHECK_NOEXCEPT(cudaMemPrefetchAsync(this->Data(), this->desc_.TotalSize() * sizeof(T), loc, 0, stream)); #endif } @@ -771,13 +771,13 @@ class tensor_t : public detail::tensor_impl_t { MATX_NVTX_START("", matx::MATX_NVTX_LOG_API) #if CUDART_VERSION <= 12000 - cudaMemPrefetchAsync(this->Data(), this->desc_.TotalSize() * sizeof(T), cudaCpuDeviceId, - stream); + MATX_CUDA_CHECK_NOEXCEPT(cudaMemPrefetchAsync(this->Data(), this->desc_.TotalSize() * sizeof(T), cudaCpuDeviceId, + stream)); #else cudaMemLocation loc; loc.id = cudaCpuDeviceId; loc.type = cudaMemLocationTypeHost; - cudaMemPrefetchAsync(this->Data(), this->desc_.TotalSize() * sizeof(T), loc, 0, stream); + MATX_CUDA_CHECK_NOEXCEPT(cudaMemPrefetchAsync(this->Data(), this->desc_.TotalSize() * sizeof(T), loc, 0, stream)); #endif } From f1bab6edf7bf96f907b52c3098b9292352194f9f Mon Sep 17 00:00:00 2001 From: lennartvoelz Date: Fri, 28 Aug 2026 11:58:32 +0200 Subject: [PATCH 6/8] Add MATX_CUDA_CHECK around bare CUDA calls under executors/ Covers cuda_executor_common.h's timing/profiling calls and distributed.h's device-guard and stream-cleanup destructors. --- include/matx/executors/cuda_executor_common.h | 14 +++++++------- include/matx/executors/distributed.h | 11 ++++++----- 2 files changed, 13 insertions(+), 12 deletions(-) diff --git a/include/matx/executors/cuda_executor_common.h b/include/matx/executors/cuda_executor_common.h index 958f68e04..bb0a2decf 100644 --- a/include/matx/executors/cuda_executor_common.h +++ b/include/matx/executors/cuda_executor_common.h @@ -120,8 +120,8 @@ namespace detail ~CudaExecutorBase() { if (profiling_) { - cudaEventDestroy(start_); - cudaEventDestroy(stop_); + MATX_CUDA_CHECK_NOEXCEPT(cudaEventDestroy(start_)); + MATX_CUDA_CHECK_NOEXCEPT(cudaEventDestroy(stop_)); } } @@ -134,14 +134,14 @@ namespace detail * @brief Synchronize the cuda executor's stream * */ - void sync() { cudaStreamSynchronize(stream_); } + void sync() { MATX_CUDA_CHECK(cudaStreamSynchronize(stream_)); } /** * @brief Start a timer for profiling workload */ void start_timer() { if (profiling_) { - cudaEventRecord(start_, stream_); + MATX_CUDA_CHECK(cudaEventRecord(start_, stream_)); } } @@ -150,7 +150,7 @@ namespace detail */ void stop_timer() { if (profiling_) { - cudaEventRecord(stop_, stream_); + MATX_CUDA_CHECK(cudaEventRecord(stop_, stream_)); } } @@ -163,8 +163,8 @@ namespace detail MATX_THROW(matxInvalidParameter, "Profiling not enabled when using get_time_ms()"); } float time; - cudaEventSynchronize(stop_); - cudaEventElapsedTime(&time, start_, stop_); + MATX_CUDA_CHECK(cudaEventSynchronize(stop_)); + MATX_CUDA_CHECK(cudaEventElapsedTime(&time, start_, stop_)); return time; } diff --git a/include/matx/executors/distributed.h b/include/matx/executors/distributed.h index ce00457f6..9f79910f8 100644 --- a/include/matx/executors/distributed.h +++ b/include/matx/executors/distributed.h @@ -150,9 +150,10 @@ class distributed_device_guard { ~distributed_device_guard() { if (changed_) { - // Destructors must not throw. A later CUDA call will report a failure to - // restore the original device if the context has become unusable. - (void)cudaSetDevice(previous_); + // Destructors must not throw. Log and drain the sticky error here + // instead of letting a failure to restore the original device leak + // into an unrelated later call. + MATX_CUDA_CHECK_NOEXCEPT(cudaSetDevice(previous_)); } } @@ -260,11 +261,11 @@ class distributedCUDAExecutor { const bool restore_device = cudaGetDevice(&previous_device) == cudaSuccess; for (const auto &entry : streams_) { if (cudaSetDevice(entry.device_id) == cudaSuccess) { - (void)cudaStreamDestroy(entry.stream); + MATX_CUDA_CHECK_NOEXCEPT(cudaStreamDestroy(entry.stream)); } } if (restore_device) { - (void)cudaSetDevice(previous_device); + MATX_CUDA_CHECK_NOEXCEPT(cudaSetDevice(previous_device)); } streams_.clear(); } From 381af9897af531f095c6a7e4bc38e329ebb4d223 Mon Sep 17 00:00:00 2001 From: lennartvoelz Date: Fri, 28 Aug 2026 12:08:44 +0200 Subject: [PATCH 7/8] Add MATX_CUDA_CHECK around bare CUDA calls under operators/ and transforms/ Covers base_operator.h's copy/event-record paths and the solver transforms (cgsolve, chol, eig, fft, inverse, lu, qr, solve, svd). --- include/matx/operators/base_operator.h | 6 ++--- include/matx/transforms/cgsolve.h | 26 +++++++++++----------- include/matx/transforms/chol/chol_cuda.h | 4 ++-- include/matx/transforms/eig/eig_cuda.h | 6 ++--- include/matx/transforms/fft/fft_cuda.h | 4 ++-- include/matx/transforms/inverse.h | 12 +++++----- include/matx/transforms/lu/lu_cuda.h | 4 ++-- include/matx/transforms/qr/qr_cuda.h | 12 +++++----- include/matx/transforms/solve/solve_cuda.h | 26 +++++++++++----------- include/matx/transforms/svd/svd_cuda.h | 20 ++++++++--------- 10 files changed, 60 insertions(+), 60 deletions(-) diff --git a/include/matx/operators/base_operator.h b/include/matx/operators/base_operator.h index 8a29656f5..7c9df89d6 100644 --- a/include/matx/operators/base_operator.h +++ b/include/matx/operators/base_operator.h @@ -255,11 +255,11 @@ namespace matx MATX_ASSERT_STR(tp->get_lhs().Bytes() >= tp->get_rhs().Bytes(), matxInvalidSize, "LHS tensor is smaller than RHS tensor in assignment"); MATX_LOG_TRACE("Copying {} bytes from {} to {} using cudaMemcpyAsync", tp->get_lhs().Bytes(), reinterpret_cast(tp->get_rhs().Data()), reinterpret_cast(tp->get_lhs().Data())); - cudaMemcpyAsync(reinterpret_cast(tp->get_lhs().Data()), + MATX_CUDA_CHECK(cudaMemcpyAsync(reinterpret_cast(tp->get_lhs().Data()), reinterpret_cast(tp->get_rhs().Data()), tp->get_rhs().Bytes(), cudaMemcpyDefault, - ex.getStream()); + ex.getStream())); } else { MATX_LOG_TRACE("Copying {} bytes from {} to {} using kernel", @@ -321,7 +321,7 @@ namespace matx MATX_NVTX_START(static_cast(this)->str(), matx::MATX_NVTX_LOG_API) run(cudaExecutor{stream, false}); - cudaEventRecord(ev, stream); + MATX_CUDA_CHECK(cudaEventRecord(ev, stream)); } /** diff --git a/include/matx/transforms/cgsolve.h b/include/matx/transforms/cgsolve.h index da4d0192d..ce6039d56 100644 --- a/include/matx/transforms/cgsolve.h +++ b/include/matx/transforms/cgsolve.h @@ -75,8 +75,8 @@ namespace matx if(tol>0.0f) { - cudaStreamCreateWithFlags(&d2h,cudaStreamNonBlocking); - cudaEventCreate(&event); + MATX_CUDA_CHECK(cudaStreamCreateWithFlags(&d2h,cudaStreamNonBlocking)); + MATX_CUDA_CHECK(cudaEventCreate(&event)); } int converged_host = false; @@ -128,9 +128,9 @@ namespace matx if(tol>0.0f) { (converged = matx::all(as_int(sqrt(r0r0) < tol))).run(stream); - - cudaEventRecord(event, stream); - cudaStreamWaitEvent(d2h, event); + + MATX_CUDA_CHECK(cudaEventRecord(event, stream)); + MATX_CUDA_CHECK(cudaStreamWaitEvent(d2h, event)); } int i; @@ -165,17 +165,17 @@ MATX_IGNORE_WARNING_POP_MSVC // copy convergence criteria to host. // This is in unpinned memory and cannot on most systems run asynchronously. // We do this here to hide the copy/sync behind prior launch latency/execution. - cudaMemcpyAsync(&converged_host, converged.Data(), sizeof(int), cudaMemcpyDeviceToHost, d2h); - cudaStreamSynchronize(d2h); + MATX_CUDA_CHECK(cudaMemcpyAsync(&converged_host, converged.Data(), sizeof(int), cudaMemcpyDeviceToHost, d2h)); + MATX_CUDA_CHECK(cudaStreamSynchronize(d2h)); if(converged_host != 0) { // != 0 instead of == true: converged_host is int (CUDA device copy); mixing int and bool triggers MSVC C4805 as an error under /WX break; } (converged = matx::all(as_int(sqrt(r1r1) < tol))).run(stream); - - cudaEventRecord(event, stream); - cudaStreamWaitEvent(d2h, event); + + MATX_CUDA_CHECK(cudaEventRecord(event, stream)); + MATX_CUDA_CHECK(cudaStreamWaitEvent(d2h, event)); } // p = r1 + b * p @@ -188,9 +188,9 @@ MATX_IGNORE_WARNING_POP_MSVC } if(tol>0.0f) { - cudaEventDestroy(event); - cudaStreamDestroy(d2h); + MATX_CUDA_CHECK(cudaEventDestroy(event)); + MATX_CUDA_CHECK(cudaStreamDestroy(d2h)); } } - + } // end namespace matx diff --git a/include/matx/transforms/chol/chol_cuda.h b/include/matx/transforms/chol/chol_cuda.h index 1e2cd999d..1beca033a 100644 --- a/include/matx/transforms/chol/chol_cuda.h +++ b/include/matx/transforms/chol/chol_cuda.h @@ -169,10 +169,10 @@ class matxDnCholCUDAPlan_t : matxDnCUDASolver_t { } std::vector h_info(this->batch_a_ptrs.size()); - cudaMemcpyAsync(h_info.data(), this->d_info, sizeof(int) * this->batch_a_ptrs.size(), cudaMemcpyDeviceToHost, stream); + MATX_CUDA_CHECK(cudaMemcpyAsync(h_info.data(), this->d_info, sizeof(int) * this->batch_a_ptrs.size(), cudaMemcpyDeviceToHost, stream)); // This will block. Figure this out later - cudaStreamSynchronize(stream); + MATX_CUDA_CHECK(cudaStreamSynchronize(stream)); for (const auto& info : h_info) { if (info < 0) { diff --git a/include/matx/transforms/eig/eig_cuda.h b/include/matx/transforms/eig/eig_cuda.h index 3202ea0ba..452e7c6bc 100644 --- a/include/matx/transforms/eig/eig_cuda.h +++ b/include/matx/transforms/eig/eig_cuda.h @@ -204,7 +204,7 @@ class matxDnEigCUDAPlan_t : matxDnCUDASolver_t { ("cusolverDnXsyevBatched failed with error " + std::to_string(ret)).c_str()); std::vector h_info(params.batch_size); - cudaMemcpyAsync(h_info.data(), this->d_info, sizeof(int) * params.batch_size, cudaMemcpyDeviceToHost, stream); + MATX_CUDA_CHECK(cudaMemcpyAsync(h_info.data(), this->d_info, sizeof(int) * params.batch_size, cudaMemcpyDeviceToHost, stream)); #else SetBatchPointers(out, this->batch_a_ptrs); SetBatchPointers(w, this->batch_w_ptrs); @@ -224,11 +224,11 @@ class matxDnEigCUDAPlan_t : matxDnCUDASolver_t { } std::vector h_info(this->batch_a_ptrs.size()); - cudaMemcpyAsync(h_info.data(), this->d_info, sizeof(int) * this->batch_a_ptrs.size(), cudaMemcpyDeviceToHost, stream); + MATX_CUDA_CHECK(cudaMemcpyAsync(h_info.data(), this->d_info, sizeof(int) * this->batch_a_ptrs.size(), cudaMemcpyDeviceToHost, stream)); #endif // This will block. Figure this out later - cudaStreamSynchronize(stream); + MATX_CUDA_CHECK(cudaStreamSynchronize(stream)); for (const auto& info : h_info) { if (info < 0) { diff --git a/include/matx/transforms/fft/fft_cuda.h b/include/matx/transforms/fft/fft_cuda.h index 05dffec2c..7c16e8e9f 100644 --- a/include/matx/transforms/fft/fft_cuda.h +++ b/include/matx/transforms/fft/fft_cuda.h @@ -376,7 +376,7 @@ matxCUDAFFTPlan1D_t(OutTensorType &o, const InTensorType &i) MATX_NVTX_START("", matx::MATX_NVTX_LOG_INTERNAL) int dev; - cudaGetDevice(&dev); + MATX_CUDA_CHECK(cudaGetDevice(&dev)); this->workspace_ = nullptr; this->params_ = this->GetFFTParams(o, i, 1); @@ -509,7 +509,7 @@ class matxCUDAFFTPlan2D_t : public matxCUDAFFTPlan_tworkspace_ = nullptr; this->params_ = this->GetFFTParams(o, i, 2); diff --git a/include/matx/transforms/inverse.h b/include/matx/transforms/inverse.h index 80d5b488a..738d2c76d 100644 --- a/include/matx/transforms/inverse.h +++ b/include/matx/transforms/inverse.h @@ -406,8 +406,8 @@ class matxInversePlan_t { MATX_ASSERT(ret == CUBLAS_STATUS_SUCCESS, matxInverseError); - cudaMemcpyAsync(h_info, d_info, sizeof(int) * params.batch_size, cudaMemcpyDeviceToHost, stream); - cudaStreamSynchronize(stream); + MATX_CUDA_CHECK(cudaMemcpyAsync(h_info, d_info, sizeof(int) * params.batch_size, cudaMemcpyDeviceToHost, stream)); + MATX_CUDA_CHECK(cudaStreamSynchronize(stream)); for (size_t i = 0; i < params.batch_size; i++) { if (h_info[i] != 0) { MATX_THROW(matxLUError, "inverse failed"); @@ -447,8 +447,8 @@ class matxInversePlan_t { } MATX_ASSERT(ret == CUBLAS_STATUS_SUCCESS, matxInverseError); - cudaMemcpyAsync(h_info, d_info, sizeof(int) * params.batch_size, cudaMemcpyDeviceToHost, stream); - cudaStreamSynchronize(stream); + MATX_CUDA_CHECK(cudaMemcpyAsync(h_info, d_info, sizeof(int) * params.batch_size, cudaMemcpyDeviceToHost, stream)); + MATX_CUDA_CHECK(cudaStreamSynchronize(stream)); for (size_t i = 0; i < params.batch_size; i++) { if (h_info[i] != 0) { MATX_THROW(matxLUError, "inverse failed"); @@ -479,8 +479,8 @@ class matxInversePlan_t { d_info); MATX_ASSERT_STR_EXP(solver_ret, CUSOLVER_STATUS_SUCCESS, matxSolverError, "Error in cusolverDnXgetrf"); - cudaMemcpyAsync(h_info, d_info, sizeof(int) * params.batch_size, cudaMemcpyDeviceToHost, stream); - cudaStreamSynchronize(stream); + MATX_CUDA_CHECK(cudaMemcpyAsync(h_info, d_info, sizeof(int) * params.batch_size, cudaMemcpyDeviceToHost, stream)); + MATX_CUDA_CHECK(cudaStreamSynchronize(stream)); for (size_t i = 0; i < params.batch_size; i++) { if (h_info[i] != 0) { MATX_THROW(matxLUError, "inverse failed"); diff --git a/include/matx/transforms/lu/lu_cuda.h b/include/matx/transforms/lu/lu_cuda.h index fb419ef0a..449d39b85 100644 --- a/include/matx/transforms/lu/lu_cuda.h +++ b/include/matx/transforms/lu/lu_cuda.h @@ -177,10 +177,10 @@ class matxDnLUCUDAPlan_t : matxDnCUDASolver_t { } std::vector h_info(this->batch_a_ptrs.size()); - cudaMemcpyAsync(h_info.data(), this->d_info, sizeof(int) * this->batch_a_ptrs.size(), cudaMemcpyDeviceToHost, stream); + MATX_CUDA_CHECK(cudaMemcpyAsync(h_info.data(), this->d_info, sizeof(int) * this->batch_a_ptrs.size(), cudaMemcpyDeviceToHost, stream)); // This will block. Figure this out later - cudaStreamSynchronize(stream); + MATX_CUDA_CHECK(cudaStreamSynchronize(stream)); for (const auto& info : h_info) { if (info < 0) { diff --git a/include/matx/transforms/qr/qr_cuda.h b/include/matx/transforms/qr/qr_cuda.h index 5021a65d2..30c24e6a3 100644 --- a/include/matx/transforms/qr/qr_cuda.h +++ b/include/matx/transforms/qr/qr_cuda.h @@ -367,10 +367,10 @@ class matxDnQRCUDAPlan_t : matxDnCUDASolver_t { } std::vector h_info(this->batch_a_ptrs.size()); - cudaMemcpyAsync(h_info.data(), this->d_info, sizeof(int) * this->batch_a_ptrs.size(), cudaMemcpyDeviceToHost, stream); + MATX_CUDA_CHECK(cudaMemcpyAsync(h_info.data(), this->d_info, sizeof(int) * this->batch_a_ptrs.size(), cudaMemcpyDeviceToHost, stream)); // This will block. Figure this out later - cudaStreamSynchronize(stream); + MATX_CUDA_CHECK(cudaStreamSynchronize(stream)); for ([[maybe_unused]] const auto& info : h_info) { MATX_ASSERT_STR_EXP(info, 0, matxSolverError, @@ -636,10 +636,10 @@ class matxDnEconQRCUDAPlan_t : matxDnCUDASolver_t { } std::vector h_info(this->batch_a_ptrs.size()); - cudaMemcpyAsync(h_info.data(), this->d_info, sizeof(int) * this->batch_a_ptrs.size(), cudaMemcpyDeviceToHost, stream); + MATX_CUDA_CHECK(cudaMemcpyAsync(h_info.data(), this->d_info, sizeof(int) * this->batch_a_ptrs.size(), cudaMemcpyDeviceToHost, stream)); // This will block. Figure this out later - cudaStreamSynchronize(stream); + MATX_CUDA_CHECK(cudaStreamSynchronize(stream)); for ([[maybe_unused]] const auto& info : h_info) { MATX_ASSERT_STR_EXP(info, 0, matxSolverError, @@ -671,10 +671,10 @@ class matxDnEconQRCUDAPlan_t : matxDnCUDASolver_t { MATX_ASSERT(ret == CUSOLVER_STATUS_SUCCESS, matxSolverError); } - cudaMemcpyAsync(h_info.data(), this->d_info, sizeof(int) * this->batch_a_ptrs.size(), cudaMemcpyDeviceToHost, stream); + MATX_CUDA_CHECK(cudaMemcpyAsync(h_info.data(), this->d_info, sizeof(int) * this->batch_a_ptrs.size(), cudaMemcpyDeviceToHost, stream)); // This will block. Figure this out later - cudaStreamSynchronize(stream); + MATX_CUDA_CHECK(cudaStreamSynchronize(stream)); for ([[maybe_unused]] const auto& info : h_info) { MATX_ASSERT_STR_EXP(info, 0, matxSolverError, diff --git a/include/matx/transforms/solve/solve_cuda.h b/include/matx/transforms/solve/solve_cuda.h index 9d967c923..ea6b0382a 100644 --- a/include/matx/transforms/solve/solve_cuda.h +++ b/include/matx/transforms/solve/solve_cuda.h @@ -302,12 +302,12 @@ class DenseSolveCublasBatchedPlan_t { SetBatchPointers(b_col, h_b_array); } - cudaMemcpyAsync(d_a_array_.get(), h_a_array.data(), + MATX_CUDA_CHECK(cudaMemcpyAsync(d_a_array_.get(), h_a_array.data(), h_a_array.size() * sizeof(T *), cudaMemcpyHostToDevice, - stream); - cudaMemcpyAsync(d_b_array_.get(), h_b_array.data(), + stream)); + MATX_CUDA_CHECK(cudaMemcpyAsync(d_b_array_.get(), h_b_array.data(), h_b_array.size() * sizeof(T *), cudaMemcpyHostToDevice, - stream); + stream)); [[maybe_unused]] auto ret = cublasSetStream(handle_.get(), stream); MATX_ASSERT(ret == CUBLAS_STATUS_SUCCESS, matxSolverError); @@ -317,9 +317,9 @@ class DenseSolveCublasBatchedPlan_t { MATX_ASSERT(ret == CUBLAS_STATUS_SUCCESS, matxSolverError); std::vector h_info(batch_size); - cudaMemcpyAsync(h_info.data(), d_info_.get(), sizeof(int) * batch_size, - cudaMemcpyDeviceToHost, stream); - cudaStreamSynchronize(stream); + MATX_CUDA_CHECK(cudaMemcpyAsync(h_info.data(), d_info_.get(), sizeof(int) * batch_size, + cudaMemcpyDeviceToHost, stream)); + MATX_CUDA_CHECK(cudaStreamSynchronize(stream)); CheckDenseSolveInfos(h_info, "cuBLAS", "getrfBatched"); int h_getrs_info = 0; @@ -409,10 +409,10 @@ class DenseSolveCusolverLoopPlan_t { } std::vector h_info(params_.batch_size); - cudaMemcpyAsync(h_info.data(), d_info_.get(), + MATX_CUDA_CHECK(cudaMemcpyAsync(h_info.data(), d_info_.get(), sizeof(int) * params_.batch_size, cudaMemcpyDeviceToHost, - stream); - cudaStreamSynchronize(stream); + stream)); + MATX_CUDA_CHECK(cudaStreamSynchronize(stream)); CheckDenseSolveInfos(h_info, "cuSolver", "Xgetrf"); for (uint32_t i = 0; i < params_.batch_size; i++) { @@ -425,10 +425,10 @@ class DenseSolveCusolverLoopPlan_t { MATX_ASSERT(ret == CUSOLVER_STATUS_SUCCESS, matxSolverError); } - cudaMemcpyAsync(h_info.data(), d_info_.get(), + MATX_CUDA_CHECK(cudaMemcpyAsync(h_info.data(), d_info_.get(), sizeof(int) * params_.batch_size, cudaMemcpyDeviceToHost, - stream); - cudaStreamSynchronize(stream); + stream)); + MATX_CUDA_CHECK(cudaStreamSynchronize(stream)); CheckDenseSolveInfos(h_info, "cuSolver", "Xgetrs"); } diff --git a/include/matx/transforms/svd/svd_cuda.h b/include/matx/transforms/svd/svd_cuda.h index bf0e56e00..6093cefaa 100644 --- a/include/matx/transforms/svd/svd_cuda.h +++ b/include/matx/transforms/svd/svd_cuda.h @@ -425,8 +425,8 @@ inline void svdbpi_impl(UType &U, SType &S, VTType &VT, const AType &A, int max_ cudaStream_t d2h = nullptr; cudaEvent_t event = nullptr; if(tol>0.0f) { - cudaStreamCreateWithFlags(&d2h,cudaStreamNonBlocking); - cudaEventCreate(&event); + MATX_CUDA_CHECK(cudaStreamCreateWithFlags(&d2h,cudaStreamNonBlocking)); + MATX_CUDA_CHECK(cudaEventCreate(&event)); } auto [AT, Q, Qold, R, Z, l2Norm, converged] = detail::svdbpi_impl_workspace(A, stream); @@ -461,7 +461,7 @@ inline void svdbpi_impl(UType &U, SType &S, VTType &VT, const AType &A, int max_ if(tol!=0.0f) { - cudaStreamSynchronize(d2h); // wait for d2h transfer to finish + MATX_CUDA_CHECK(cudaStreamSynchronize(d2h)); // wait for d2h transfer to finish if(converged_host != 0) { // if converged exit loop break; @@ -479,14 +479,14 @@ inline void svdbpi_impl(UType &U, SType &S, VTType &VT, const AType &A, int max_ } // event to record when converged is ready in stream - cudaEventRecord(event, stream); + MATX_CUDA_CHECK(cudaEventRecord(event, stream)); // wait for d2h transfer until converged is ready - cudaStreamWaitEvent(d2h, event); + MATX_CUDA_CHECK(cudaStreamWaitEvent(d2h, event)); // copy convergence criteria to host. // This is in unpinned memory and cannot on most systems run asynchronously. // We do this here to hide the copy/sync behind prior launch latency/execution of next iteration. - cudaMemcpyAsync(&converged_host, converged.Data(), sizeof(int), cudaMemcpyDeviceToHost, d2h); + MATX_CUDA_CHECK(cudaMemcpyAsync(&converged_host, converged.Data(), sizeof(int), cudaMemcpyDeviceToHost, d2h)); } } @@ -520,8 +520,8 @@ inline void svdbpi_impl(UType &U, SType &S, VTType &VT, const AType &A, int max_ } if(tol>0.0f) { - cudaEventDestroy(event); - cudaStreamDestroy(d2h); + MATX_CUDA_CHECK(cudaEventDestroy(event)); + MATX_CUDA_CHECK(cudaStreamDestroy(d2h)); } } @@ -852,10 +852,10 @@ class matxDnSVDCUDAPlan_t : matxDnCUDASolver_t { } std::vector h_info(this->batch_a_ptrs.size()); - cudaMemcpyAsync(h_info.data(), this->d_info, sizeof(int) * this->batch_a_ptrs.size(), cudaMemcpyDeviceToHost, stream); + MATX_CUDA_CHECK(cudaMemcpyAsync(h_info.data(), this->d_info, sizeof(int) * this->batch_a_ptrs.size(), cudaMemcpyDeviceToHost, stream)); // This will block. Figure this out later - cudaStreamSynchronize(stream); + MATX_CUDA_CHECK(cudaStreamSynchronize(stream)); for (const auto& info : h_info) { if (info < 0) { From 31c0b3c35b2cb55d96c547b44e805053f20af7e5 Mon Sep 17 00:00:00 2001 From: lennartvoelz Date: Fri, 28 Aug 2026 13:40:56 +0200 Subject: [PATCH 8/8] Fix stream/event leak on exception in cgsolve and svdbpi The newly-added MATX_CUDA_CHECK calls in cgsolve_impl and svdbpi_impl can now throw between creating their temporary d2h stream/event and the end-of-function cleanup that used to destroy them, leaking both on any such throw. Replace the manual create/destroy with detail::CudaStreamGuard/CudaEventGuard (core/resource_guard.h), a shared unique_ptr-based RAII guard for opaque CUDA handles, so cleanup runs on any exception path. --- include/matx/core/resource_guard.h | 74 ++++++++++++++++++++++++++ include/matx/transforms/cgsolve.h | 36 +++++++------ include/matx/transforms/svd/svd_cuda.h | 31 ++++++----- 3 files changed, 111 insertions(+), 30 deletions(-) create mode 100644 include/matx/core/resource_guard.h diff --git a/include/matx/core/resource_guard.h b/include/matx/core/resource_guard.h new file mode 100644 index 000000000..c82d66dfb --- /dev/null +++ b/include/matx/core/resource_guard.h @@ -0,0 +1,74 @@ +//////////////////////////////////////////////////////////////////////////////// +// BSD 3-Clause License +// +// Copyright (c) 2026, NVIDIA Corporation +// All rights reserved. +// +// Redistribution and use in source and binary forms, with or without +// modification, are permitted provided that the following conditions are met: +// +// 1. Redistributions of source code must retain the above copyright notice, this +// list of conditions and the following disclaimer. +// +// 2. Redistributions in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// 3. Neither the name of the copyright holder nor the names of its +// contributors may be used to endorse or promote products derived from +// this software without specific prior written permission. +// +// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" +// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE +// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE +// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE +// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL +// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR +// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER +// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, +// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +///////////////////////////////////////////////////////////////////////////////// + +#pragma once + +#include +#include + +#include +#include + +#include "matx/core/error.h" + +namespace matx { +namespace detail { + +/** + * @brief RAII guards for opaque CUDA C-API handles, built on + * std::unique_ptr with a stateless deleter. reset(handle) takes ownership; + * the handle is destroyed on the next reset(), on guard destruction, or + * if an exception unwinds through the guard. + * + * Creation and validation stay at the call site, since cudaStreamCreate*, + * cudaEventCreate, etc. each have their own success code/error type: + * + * detail::CudaStreamGuard d2h_guard; + * cudaStream_t handle; + * MATX_CUDA_CHECK(cudaStreamCreateWithFlags(&handle, cudaStreamNonBlocking)); + * d2h_guard.reset(handle); + * ... + * MATX_CUDA_CHECK(cudaStreamSynchronize(d2h_guard.get())); + */ +inline constexpr auto CudaStreamDeleter = [](cudaStream_t stream) noexcept { + MATX_CUDA_CHECK_NOEXCEPT(cudaStreamDestroy(stream)); +}; + +inline constexpr auto CudaEventDeleter = [](cudaEvent_t event) noexcept { + MATX_CUDA_CHECK_NOEXCEPT(cudaEventDestroy(event)); +}; + +using CudaStreamGuard = std::unique_ptr, decltype(CudaStreamDeleter)>; +using CudaEventGuard = std::unique_ptr, decltype(CudaEventDeleter)>; + +} // namespace detail +} // namespace matx diff --git a/include/matx/transforms/cgsolve.h b/include/matx/transforms/cgsolve.h index ce6039d56..4daa7bee1 100644 --- a/include/matx/transforms/cgsolve.h +++ b/include/matx/transforms/cgsolve.h @@ -34,6 +34,7 @@ #include "matx/transforms/reduce.h" #include "matx/core/nvtx.h" +#include "matx/core/resource_guard.h" #include "matx/core/type_utils.h" #include "matx/operators/all.h" #include "matx/operators/if.h" @@ -70,13 +71,19 @@ namespace matx MATX_ASSERT_STR(A.Rank() -1 == X.Rank(), matxInvalidDim, "cgsolve: A rank must be one larger than X rank"); MATX_ASSERT_STR(X.Rank() == B.Rank(), matxInvalidDim, "cgsole: X rank and B rank must match"); - cudaStream_t d2h = nullptr; - cudaEvent_t event = nullptr; - + // Declared in this order so, if tol>0.0f, event_guard (destructed + // first) is torn down before the stream it may still be queued on. + detail::CudaStreamGuard d2h_guard; + detail::CudaEventGuard event_guard; if(tol>0.0f) { - MATX_CUDA_CHECK(cudaStreamCreateWithFlags(&d2h,cudaStreamNonBlocking)); - MATX_CUDA_CHECK(cudaEventCreate(&event)); + cudaStream_t d2h_handle; + MATX_CUDA_CHECK(cudaStreamCreateWithFlags(&d2h_handle,cudaStreamNonBlocking)); + d2h_guard.reset(d2h_handle); + + cudaEvent_t event_handle; + MATX_CUDA_CHECK(cudaEventCreate(&event_handle)); + event_guard.reset(event_handle); } int converged_host = false; @@ -129,8 +136,8 @@ namespace matx if(tol>0.0f) { (converged = matx::all(as_int(sqrt(r0r0) < tol))).run(stream); - MATX_CUDA_CHECK(cudaEventRecord(event, stream)); - MATX_CUDA_CHECK(cudaStreamWaitEvent(d2h, event)); + MATX_CUDA_CHECK(cudaEventRecord(event_guard.get(), stream)); + MATX_CUDA_CHECK(cudaStreamWaitEvent(d2h_guard.get(), event_guard.get())); } int i; @@ -165,8 +172,8 @@ MATX_IGNORE_WARNING_POP_MSVC // copy convergence criteria to host. // This is in unpinned memory and cannot on most systems run asynchronously. // We do this here to hide the copy/sync behind prior launch latency/execution. - MATX_CUDA_CHECK(cudaMemcpyAsync(&converged_host, converged.Data(), sizeof(int), cudaMemcpyDeviceToHost, d2h)); - MATX_CUDA_CHECK(cudaStreamSynchronize(d2h)); + MATX_CUDA_CHECK(cudaMemcpyAsync(&converged_host, converged.Data(), sizeof(int), cudaMemcpyDeviceToHost, d2h_guard.get())); + MATX_CUDA_CHECK(cudaStreamSynchronize(d2h_guard.get())); if(converged_host != 0) { // != 0 instead of == true: converged_host is int (CUDA device copy); mixing int and bool triggers MSVC C4805 as an error under /WX break; @@ -174,8 +181,8 @@ MATX_IGNORE_WARNING_POP_MSVC (converged = matx::all(as_int(sqrt(r1r1) < tol))).run(stream); - MATX_CUDA_CHECK(cudaEventRecord(event, stream)); - MATX_CUDA_CHECK(cudaStreamWaitEvent(d2h, event)); + MATX_CUDA_CHECK(cudaEventRecord(event_guard.get(), stream)); + MATX_CUDA_CHECK(cudaStreamWaitEvent(d2h_guard.get(), event_guard.get())); } // p = r1 + b * p @@ -183,14 +190,9 @@ MATX_IGNORE_WARNING_POP_MSVC (IF( pApc != value_type(0), updateP)).run(stream); // Advance residual - swap(r0r0, r1r1); + swap(r0r0, r1r1); swap(r0r0c, r1r1c); } - - if(tol>0.0f) { - MATX_CUDA_CHECK(cudaEventDestroy(event)); - MATX_CUDA_CHECK(cudaStreamDestroy(d2h)); - } } } // end namespace matx diff --git a/include/matx/transforms/svd/svd_cuda.h b/include/matx/transforms/svd/svd_cuda.h index 6093cefaa..cb671d214 100644 --- a/include/matx/transforms/svd/svd_cuda.h +++ b/include/matx/transforms/svd/svd_cuda.h @@ -37,6 +37,7 @@ #include "matx/core/error.h" #include "matx/core/nvtx.h" +#include "matx/core/resource_guard.h" #include "matx/core/tensor.h" #include "matx/core/cache.h" #include "matx/operators/slice.h" @@ -422,11 +423,20 @@ inline void svdbpi_impl(UType &U, SType &S, VTType &VT, const AType &A, int max_ MATX_ASSERT_STR(S.Size(RANK-2) == d, matxInvalidDim, "svdbpi: S must have Size(RANK-2) == d"); int converged_host = false; - cudaStream_t d2h = nullptr; - cudaEvent_t event = nullptr; + + // Declared in this order so, if tol>0.0f, event_guard (destructed first) + // is torn down before the stream it may still be queued on. + detail::CudaStreamGuard d2h_guard; + detail::CudaEventGuard event_guard; + if(tol>0.0f) { - MATX_CUDA_CHECK(cudaStreamCreateWithFlags(&d2h,cudaStreamNonBlocking)); - MATX_CUDA_CHECK(cudaEventCreate(&event)); + cudaStream_t d2h_handle; + MATX_CUDA_CHECK(cudaStreamCreateWithFlags(&d2h_handle,cudaStreamNonBlocking)); + d2h_guard.reset(d2h_handle); + + cudaEvent_t event_handle; + MATX_CUDA_CHECK(cudaEventCreate(&event_handle)); + event_guard.reset(event_handle); } auto [AT, Q, Qold, R, Z, l2Norm, converged] = detail::svdbpi_impl_workspace(A, stream); @@ -461,7 +471,7 @@ inline void svdbpi_impl(UType &U, SType &S, VTType &VT, const AType &A, int max_ if(tol!=0.0f) { - MATX_CUDA_CHECK(cudaStreamSynchronize(d2h)); // wait for d2h transfer to finish + MATX_CUDA_CHECK(cudaStreamSynchronize(d2h_guard.get())); // wait for d2h transfer to finish if(converged_host != 0) { // if converged exit loop break; @@ -479,14 +489,14 @@ inline void svdbpi_impl(UType &U, SType &S, VTType &VT, const AType &A, int max_ } // event to record when converged is ready in stream - MATX_CUDA_CHECK(cudaEventRecord(event, stream)); + MATX_CUDA_CHECK(cudaEventRecord(event_guard.get(), stream)); // wait for d2h transfer until converged is ready - MATX_CUDA_CHECK(cudaStreamWaitEvent(d2h, event)); + MATX_CUDA_CHECK(cudaStreamWaitEvent(d2h_guard.get(), event_guard.get())); // copy convergence criteria to host. // This is in unpinned memory and cannot on most systems run asynchronously. // We do this here to hide the copy/sync behind prior launch latency/execution of next iteration. - MATX_CUDA_CHECK(cudaMemcpyAsync(&converged_host, converged.Data(), sizeof(int), cudaMemcpyDeviceToHost, d2h)); + MATX_CUDA_CHECK(cudaMemcpyAsync(&converged_host, converged.Data(), sizeof(int), cudaMemcpyDeviceToHost, d2h_guard.get())); } } @@ -518,11 +528,6 @@ inline void svdbpi_impl(UType &U, SType &S, VTType &VT, const AType &A, int max_ // IF required to avoid nans when singular value is 0 (IF(D != STypeS(0), VT = VT / D)).run(stream); } - - if(tol>0.0f) { - MATX_CUDA_CHECK(cudaEventDestroy(event)); - MATX_CUDA_CHECK(cudaStreamDestroy(d2h)); - } }