Fix unchecked CUDA calls in noexcept functions

include/matx/core/allocator.h

@@ -141,21 +141,21 @@ 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);
       break;
     case MATX_ASYNC_DEVICE_MEMORY:
       if (is_cuda_free()) {
         if constexpr (cuda::std::is_same_v<no_stream_t, StreamType>) {
-          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;

include/matx/core/cache.h

@@ -235,7 +235,7 @@ class matxCache_t {
 
     auto &cval = cache[id];
     if constexpr (is_cuda_executor_v<Executor>) {
-      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<std::recursive_mutex> lock(stream_alloc_mutex);
 

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<int>(e_));                            \
+      } catch (...) {}                                                    \
+      (void)cudaGetLastError();                                           \
+    }                                                                     \
+  } while (0)
+
 // Macro for checking CUDA driver API (CUresult) errors
 #define MATX_CUDA_DRIVER_CHECK(e)                                          \
   do {                                                                     \

include/matx/core/make_sparse_tensor.h

@@ -42,7 +42,7 @@ template <typename T>
 __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<T> makeEmptyStorage() {
 template <typename T>
 __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));
   }

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<Op>) {
         using Format = typename Op::Format;
         fprintf(fp, "format = ");
@@ -630,7 +630,7 @@ namespace matx {
             for (int i = 0; i < R; i++) shape[i] = op.Size(i);
             auto tmpv = make_tensor<typename Op::value_type>(shape);
             (tmpv = op).run(CUDAJITExecutor{});
-            cudaStreamSynchronize(0);
+            MATX_CUDA_CHECK(cudaStreamSynchronize(0));
             detail::InternalPrint(fp, tmpv, dims...);
           };
           switch (r) {
@@ -648,7 +648,7 @@ namespace matx {
         } else {
           auto tmpv = make_tensor<typename Op::value_type>(op.Shape());
           (tmpv = op).run();
-          cudaStreamSynchronize(0);
+          MATX_CUDA_CHECK(cudaStreamSynchronize(0));
           detail::InternalPrint(fp, tmpv, dims...);
         }
       }

include/matx/core/pybind.h

@@ -525,7 +525,7 @@ class MATX_PYBIND_VISIBILITY MatXPybind {
     auto ften = pybind11::array_t<ntype>(resobj);
     constexpr int RANK = TensorType::Rank();
 
-    cudaDeviceSynchronize();
+    MATX_CUDA_CHECK(cudaDeviceSynchronize());
 
     if constexpr (RANK == 0) {
       auto file_val = ften.at();

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 <memory>
+#include <type_traits>
+
+#include <driver_types.h>
+#include <cuda_runtime_api.h>
+
+#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<std::remove_pointer_t<cudaStream_t>, decltype(CudaStreamDeleter)>;
+using CudaEventGuard = std::unique_ptr<std::remove_pointer_t<cudaEvent_t>, decltype(CudaEventDeleter)>;
+
+} // namespace detail
+} // namespace matx

include/matx/core/tensor.h

@@ -745,14 +745,14 @@ class tensor_t : public detail::tensor_impl_t<T,RANK,Desc> {
     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<T,RANK,Desc> {
     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
   }
 

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;
       }
 

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();
   }

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<void*>(tp->get_rhs().Data()), reinterpret_cast<void*>(tp->get_lhs().Data()));
-                cudaMemcpyAsync(reinterpret_cast<void*>(tp->get_lhs().Data()),
+                MATX_CUDA_CHECK(cudaMemcpyAsync(reinterpret_cast<void*>(tp->get_lhs().Data()),
                                 reinterpret_cast<void*>(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<T *>(this)->str(), matx::MATX_NVTX_LOG_API)
 
           run(cudaExecutor{stream, false});
-          cudaEventRecord(ev, stream);
+          MATX_CUDA_CHECK(cudaEventRecord(ev, stream));
         }
 
         /**