TEST/DEVICE: Print kernel launch failure instead of exception (#10923)

Throwing an exception triggers cleanup flow that also calls cudaFree(),
so the error is reported from cudaFree() instead of the actual place it
happened.

Author: yosefe

test/gtest/common/cuda.h

@@ -12,7 +12,7 @@
 #include <sstream>
 #include <memory>
 #include <vector>
-
+#include <ucs/debug/log.h>
 
 namespace ucx_cuda {
 
@@ -97,21 +97,21 @@ device_vector<T> make_device_vector(const std::vector<T> &vec)
     return device_vector<T>(vec);
 }
 
-static inline void synchronize()
+static inline ucs_status_t synchronize()
 {
     cudaError_t err = cudaGetLastError();
     if (err != cudaSuccess) {
-        std::stringstream ss;
-        ss << "kernel launch failure: " << cudaGetErrorString(err);
-        throw std::runtime_error(ss.str());
+        ucs_error("kernel launch failure: %s", cudaGetErrorString(err));
+        return UCS_ERR_IO_ERROR;
     }
 
     err = cudaDeviceSynchronize();
     if (err != cudaSuccess) {
-        std::stringstream ss;
-        ss << "cudaDeviceSynchronize(): " << cudaGetErrorString(err);
-        throw std::runtime_error(ss.str());
+        ucs_error("cudaDeviceSynchronize(): %s", cudaGetErrorString(err));
+        return UCS_ERR_IO_ERROR;
     }
+
+    return UCS_OK;
 }
 
 } // namespace ucx_cuda

test/gtest/ucp/cuda/test_kernels.cu

@@ -208,6 +208,10 @@ launch_test_ucp_device_kernel(const test_ucp_device_kernel_params_t &params)
         return UCS_ERR_INVALID_PARAM;
     }
 
-    ucx_cuda::synchronize();
+    ucs_status_t sync_status = ucx_cuda::synchronize();
+    if (sync_status != UCS_OK) {
+        return sync_status;
+    }
+
     return *status;
 }