Refactor driver APIs

Author: davebayer

cudax/include/cuda/experimental/__container/async_buffer.cuh

@@ -177,8 +177,8 @@ private:
     static_assert(::cuda::std::contiguous_iterator<_Iter>, "Non contiguous iterators are not supported");
     // TODO use batched memcpy for non-contiguous iterators, it allows to
     // specify stream ordered access
-    ::cuda::__driver::__memcpyAsync(
-      __dest, ::cuda::std::to_address(__first), sizeof(_Tp) * __count, __buf_.stream().get());
+    _CCCL_TRY_DRIVER_API(
+      __memcpyAsync(__dest, ::cuda::std::to_address(__first), sizeof(_Tp) * __count, __buf_.stream().get()));
   }
 
 public:
@@ -635,11 +635,11 @@ template <typename _BufferTo, typename _BufferFrom>
 void __copy_cross_buffers(stream_ref __stream, _BufferTo& __to, const _BufferFrom& __from)
 {
   __stream.wait(__from.stream());
-  ::cuda::__driver::__memcpyAsync(
+  _CCCL_TRY_DRIVER_API(__memcpyAsync(
     __to.__unwrapped_begin(),
     __from.__unwrapped_begin(),
     sizeof(typename _BufferTo::value_type) * __from.size(),
-    __stream.get());
+    __stream.get()));
 }
 
 _CCCL_BEGIN_NAMESPACE_ARCH_DEPENDENT
@@ -660,16 +660,9 @@ __fill_n(cuda::stream_ref __stream, _Tp* __first, ::cuda::std::size_t __count, c
   // we need to check the attributes
   if constexpr (_Accessability == __memory_accessability::__device_and_host)
   {
-    __driver::__pointer_attribute_value_type_t<CU_POINTER_ATTRIBUTE_MEMORY_TYPE> __type;
-    bool __is_managed{};
-    auto __status1 = ::cuda::__driver::__pointerGetAttributeNoThrow<CU_POINTER_ATTRIBUTE_MEMORY_TYPE>(__type, __first);
-    auto __status2 =
-      ::cuda::__driver::__pointerGetAttributeNoThrow<CU_POINTER_ATTRIBUTE_IS_MANAGED>(__is_managed, __first);
-    if (__status1 != ::cudaSuccess || __status2 != ::cudaSuccess)
-    {
-      __throw_cuda_error(__status1, "Failed to get buffer memory attributes");
-    }
-    if (__type == ::CU_MEMORYTYPE_HOST && !__is_managed)
+    const auto __mem_type   = _CCCL_TRY_DRIVER_API(__pointerGetAttribute<::CU_POINTER_ATTRIBUTE_MEMORY_TYPE>(__first));
+    const auto __is_managed = _CCCL_TRY_DRIVER_API(__pointerGetAttribute<::CU_POINTER_ATTRIBUTE_IS_MANAGED>(__first));
+    if (__mem_type == ::CU_MEMORYTYPE_HOST && !__is_managed)
     {
       __fill_n<_Tp, __memory_accessability::__host>(__stream, __first, __count, __value);
     }
@@ -687,7 +680,7 @@ __fill_n(cuda::stream_ref __stream, _Tp* __first, ::cuda::std::size_t __count, c
   {
     if constexpr (sizeof(_Tp) <= 4)
     {
-      ::cuda::__driver::__memsetAsync(__first, __value, __count, __stream.get());
+      _CCCL_TRY_DRIVER_API(__memsetAsync(__first, __value, __count, __stream.get()));
     }
     else
     {

cudax/include/cuda/experimental/__green_context/green_ctx.cuh

@@ -47,9 +47,9 @@ struct green_context
       : __dev_id(__device.get())
   {
     // TODO get CUdevice from device
-    auto __dev_handle = ::cuda::__driver::__deviceGet(__dev_id);
-    __green_ctx       = ::cuda::__driver::__greenCtxCreate(__dev_handle);
-    __transformed     = ::cuda::__driver::__ctxFromGreenCtx(__green_ctx);
+    auto __dev_handle = _CCCL_TRY_DRIVER_API(__deviceGet(__dev_id));
+    __green_ctx       = _CCCL_TRY_DRIVER_API(__greenCtxCreate(__dev_handle));
+    __transformed     = _CCCL_TRY_DRIVER_API(__ctxFromGreenCtx(__green_ctx));
   }
 
   green_context(const green_context&)            = delete;
@@ -58,17 +58,17 @@ struct green_context
   // TODO this probably should be the runtime equivalent once available
   [[nodiscard]] static green_context from_native_handle(CUgreenCtx __gctx)
   {
-    CUcontext __transformed = ::cuda::__driver::__ctxFromGreenCtx(__gctx);
-    ::cuda::__driver::__ctxPush(__transformed);
-    CUdevice __device = ::cuda::__driver::__ctxGetDevice();
-    ::cuda::__driver::__ctxPop();
+    CUcontext __transformed = _CCCL_TRY_DRIVER_API(__ctxFromGreenCtx(__gctx));
+    (void) _CCCL_TRY_DRIVER_API(__ctxPush(__transformed));
+    CUdevice __device = _CCCL_TRY_DRIVER_API(__ctxGetDevice());
+    (void) _CCCL_TRY_DRIVER_API(__ctxPop());
     return green_context(::cuda::__driver::__cudevice_to_ordinal(__device), __gctx, __transformed);
   }
 
 #  if _CCCL_CTK_AT_LEAST(13, 0)
   [[nodiscard]] _CCCL_HOST_API green_context_id id() const
   {
-    return green_context_id{::cuda::__driver::__greenCtxGetId(__green_ctx)};
+    return green_context_id{_CCCL_TRY_DRIVER_API(__greenCtxGetId(__green_ctx))};
   }
 #  endif // _CCCL_CTK_AT_LEAST(13, 0)
 
@@ -83,7 +83,7 @@ struct green_context
   {
     if (__green_ctx)
     {
-      [[maybe_unused]] cudaError_t __status = ::cuda::__driver::__greenCtxDestroyNoThrow(__green_ctx);
+      _CCCL_ASSERT_DRIVER_API(__greenCtxDestroy(__green_ctx));
     }
   }
 

cudax/include/cuda/experimental/__kernel/attributes.cuh

@@ -47,8 +47,8 @@ struct __kernel_attr_impl
   template <class _Signature>
   [[nodiscard]] type operator()(kernel_ref<_Signature> __kernel, device_ref __dev) const
   {
-    return static_cast<type>(
-      ::cuda::__driver::__kernelGetAttribute(_Attr, __kernel.get(), ::cuda::__driver::__deviceGet(__dev.get())));
+    return static_cast<type>(_CCCL_TRY_DRIVER_API(
+      __kernelGetAttribute(_Attr, __kernel.get(), _CCCL_TRY_DRIVER_API(__deviceGet(__dev.get())))));
   }
 };
 

cudax/include/cuda/experimental/__kernel/kernel_ref.cuh

@@ -95,7 +95,7 @@ public:
   //! @throws cuda_error if the kernel name cannot be obtained
   [[nodiscard]] ::cuda::std::string_view name() const
   {
-    return ::cuda::__driver::__kernelGetName(__kernel_);
+    return _CCCL_TRY_DRIVER_API(__kernelGetName(__kernel_));
   }
 #endif // _CCCL_CTK_AT_LEAST(12, 3)
 

cudax/include/cuda/experimental/__launch/configuration.cuh

@@ -319,32 +319,34 @@ template <class _Tp>
 [[nodiscard]] ::cudaError_t __apply_launch_option(
   const dynamic_shared_memory<_Tp>& __opt, ::CUlaunchConfig& __config, ::CUfunction __kernel) noexcept
 {
-  ::cudaError_t __status = ::cudaSuccess;
-
   // Since CUDA 12.4, querying CU_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES requires the function to be loaded.
-  if (::cuda::__driver::__version_at_least(12, 4))
+  const auto __needs_load = ::cuda::__driver::__version_at_least(12, 4);
+  if (__needs_load.__error_ != ::CUDA_SUCCESS)
+  {
+    return static_cast<::cudaError_t>(__needs_load.__error_);
+  }
+
+  if (__needs_load.__value_)
   {
-    __status = ::cuda::__driver::__functionLoadNoThrow(__kernel);
-    if (__status != ::cudaSuccess)
+    const auto __status = ::cuda::__driver::__functionLoad(__kernel).__error_;
+    if (__status != ::CUDA_SUCCESS)
     {
-      return __status;
+      return static_cast<::cudaError_t>(__status);
     }
   }
 
-  int __static_smem_size{};
-  __status = ::cuda::__driver::__functionGetAttributeNoThrow(
-    __static_smem_size, ::CU_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES, __kernel);
-  if (__status != ::cudaSuccess)
+  const auto __static_smem_size =
+    ::cuda::__driver::__functionGetAttribute(::CU_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES, __kernel);
+  if (__static_smem_size.__error_ != ::CUDA_SUCCESS)
   {
-    return __status;
+    return static_cast<::cudaError_t>(__static_smem_size.__error_);
   }
 
-  int __max_dyn_smem_size{};
-  __status = ::cuda::__driver::__functionGetAttributeNoThrow(
-    __max_dyn_smem_size, ::CU_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES, __kernel);
-  if (__status != ::cudaSuccess)
+  const auto __max_dyn_smem_size =
+    ::cuda::__driver::__functionGetAttribute(::CU_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES, __kernel);
+  if (__max_dyn_smem_size.__error_ != ::CUDA_SUCCESS)
   {
-    return __status;
+    return static_cast<::cudaError_t>(__max_dyn_smem_size.__error_);
   }
 
   const auto __dyn_smem_size = ::cuda::overflow_cast<int>(__opt.size_bytes());
@@ -353,19 +355,21 @@ template <class _Tp>
     return ::cudaErrorInvalidValue;
   }
 
-  const int __smem_size = __static_smem_size + __dyn_smem_size.value;
+  const int __smem_size = __static_smem_size.__value_ + __dyn_smem_size.value;
   if (static_cast<::cuda::std::size_t>(__smem_size) > __max_portable_dyn_smem_size && !__opt.__non_portable_)
   {
     return ::cudaErrorInvalidValue;
   }
 
-  if (__max_dyn_smem_size < __dyn_smem_size.value)
+  if (__max_dyn_smem_size.__value_ < __dyn_smem_size.value)
   {
-    __status = ::cuda::__driver::__functionSetAttributeNoThrow(
-      __kernel, ::CU_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES, __dyn_smem_size.value);
-    if (__status != ::cudaSuccess)
+    const auto __status =
+      ::cuda::__driver::__functionSetAttribute(
+        __kernel, ::CU_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES, __dyn_smem_size.value)
+        .__error_;
+    if (__status != ::CUDA_SUCCESS)
     {
-      return __status;
+      return static_cast<::cudaError_t>(__status);
     }
   }
 

cudax/include/cuda/experimental/__launch/launch.cuh

@@ -60,7 +60,7 @@ __global__ static void __kernel_launcher_no_config(_Kernel __kernel_fn, _Args...
 template <class... _Args>
 [[nodiscard]] _CCCL_HOST_API CUfunction __get_cufunction_of(kernel_ref<void(_Args...)> __kernel)
 {
-  return ::cuda::__driver::__kernelGetFunction(__kernel.get());
+  return _CCCL_TRY_DRIVER_API(__kernelGetFunction(__kernel.get()));
 }
 
 template <class... _Args>
@@ -90,12 +90,12 @@ __do_launch(_GraphInserter&& __inserter, ::CUlaunchConfig& __config, ::CUfunctio
 
   auto __dependencies = __inserter.get_dependencies();
 
-  const auto __node = ::cuda::__driver::__graphAddKernelNode(
-    __inserter.get_graph().get(), __dependencies.data(), __dependencies.size(), __node_params);
+  const auto __node = _CCCL_TRY_DRIVER_API(
+    __graphAddKernelNode(__inserter.get_graph().get(), __dependencies.data(), __dependencies.size(), __node_params));
 
   for (unsigned int __i = 0; __i < __config.numAttrs; ++__i)
   {
-    ::cuda::__driver::__graphKernelNodeSetAttribute(__node, __config.attrs[__i].id, __config.attrs[__i].value);
+    _CCCL_TRY_DRIVER_API(__graphKernelNodeSetAttribute(__node, __config.attrs[__i].id, __config.attrs[__i].value));
   }
 
   // TODO skip the update if called on rvalue?
@@ -111,7 +111,7 @@ _CCCL_HOST_API void inline __do_launch(
 #if defined(_CUDAX_LAUNCH_CONFIG_TEST)
   test_launch_kernel_replacement(__config, __kernel, __args_ptrs);
 #else // ^^^ _CUDAX_LAUNCH_CONFIG_TEST ^^^ / vvv !_CUDAX_LAUNCH_CONFIG_TEST vvv
-  ::cuda::__driver::__launchKernel(__config, __kernel, __args_ptrs);
+  _CCCL_TRY_DRIVER_API(__launchKernel(__config, __kernel, __args_ptrs));
 #endif // ^^^ !_CUDAX_LAUNCH_CONFIG_TEST ^^^
 }
 

cudax/include/cuda/experimental/__library/library.cuh

@@ -79,7 +79,7 @@ struct library : public library_ref
   {
     if (__library_ != value_type{})
     {
-      [[maybe_unused]] const auto __status = ::cuda::__driver::__libraryUnloadNoThrow(__library_);
+      _CCCL_ASSERT_DRIVER_API(__libraryUnload(__library_));
     }
   }
 

cudax/include/cuda/experimental/__library/library_ref.cuh

@@ -72,15 +72,15 @@ public:
   //! @throws cuda_error if the library could not be queried for the kernel
   [[nodiscard]] bool has_kernel(const char* __name) const
   {
-    ::CUkernel __kernel{};
-    switch (const auto __res = ::cuda::__driver::__libraryGetKernelNoThrow(__kernel, __library_, __name))
+    const auto __status = ::cuda::__driver::__libraryGetKernel(__library_, __name).__error_;
+    switch (__status)
     {
-      case ::cudaSuccess:
+      case ::CUDA_SUCCESS:
         return true;
-      case ::cudaErrorSymbolNotFound:
+      case ::CUDA_ERROR_NOT_FOUND:
         return false;
       default:
-        ::cuda::__throw_cuda_error(__res, "Failed to get the kernel from library");
+        ::cuda::__throw_cuda_error(static_cast<::cudaError_t>(__status), "Failed to get the kernel from library");
     }
   }
 
@@ -96,13 +96,7 @@ public:
   template <class _Signature>
   [[nodiscard]] kernel_ref<_Signature> kernel(const char* __name) const
   {
-    ::CUkernel __kernel{};
-    if (const auto __res = ::cuda::__driver::__libraryGetKernelNoThrow(__kernel, __library_, __name);
-        __res != ::cudaSuccess)
-    {
-      ::cuda::__throw_cuda_error(__res, "Failed to get the kernel from the library");
-    }
-    return kernel_ref<_Signature>{__kernel};
+    return kernel_ref<_Signature>{_CCCL_TRY_DRIVER_API(__libraryGetKernel(__library_, __name))};
   }
 
   //! @brief Checks if the library contains a global symbol with the given name on a device
@@ -117,16 +111,15 @@ public:
   {
     ::cuda::__ensure_current_context __ctx_guard(__device);
 
-    ::CUdeviceptr __dptr{};
-    ::cuda::std::size_t __size{};
-    switch (const auto __res = ::cuda::__driver::__libraryGetGlobalNoThrow(__dptr, __size, __library_, __name))
+    const auto __status = ::cuda::__driver::__libraryGetGlobal(__library_, __name).__error_;
+    switch (__status)
     {
-      case ::cudaSuccess:
+      case ::CUDA_SUCCESS:
         return true;
-      case ::cudaErrorSymbolNotFound:
+      case ::CUDA_ERROR_NOT_FOUND:
         return false;
       default:
-        ::cuda::__throw_cuda_error(__res, "Failed to get the global symbol from library");
+        ::cuda::__throw_cuda_error(static_cast<::cudaError_t>(__status), "Failed to get the global symbol from library");
     }
   }
 
@@ -141,15 +134,8 @@ public:
   [[nodiscard]] library_symbol_info global(const char* __name, ::cuda::device_ref __device) const
   {
     ::cuda::__ensure_current_context __ctx_guard(__device);
-
-    ::CUdeviceptr __dptr{};
-    ::cuda::std::size_t __size{};
-    if (const auto __res = ::cuda::__driver::__libraryGetGlobalNoThrow(__dptr, __size, __library_, __name);
-        __res != ::cudaSuccess)
-    {
-      ::cuda::__throw_cuda_error(__res, "Failed to get the global symbol from the library");
-    }
-    return library_symbol_info{reinterpret_cast<void*>(__dptr), __size};
+    auto [__ptr, __size] = _CCCL_TRY_DRIVER_API(__libraryGetGlobal(__library_, __name));
+    return library_symbol_info{__ptr, __size};
   }
 
   //! @brief Checks if the library contains a managed symbol with the given name
@@ -163,16 +149,16 @@ public:
   //! @note Managed memory is shared across devices
   [[nodiscard]] bool has_managed(const char* __name) const
   {
-    ::CUdeviceptr __dptr{};
-    ::cuda::std::size_t __size{};
-    switch (const auto __res = ::cuda::__driver::__libraryGetManagedNoThrow(__dptr, __size, __library_, __name))
+    const auto __status = ::cuda::__driver::__libraryGetManaged(__library_, __name).__error_;
+    switch (__status)
     {
-      case ::cudaSuccess:
+      case ::CUDA_SUCCESS:
         return true;
-      case ::cudaErrorSymbolNotFound:
+      case ::CUDA_ERROR_NOT_FOUND:
         return false;
       default:
-        ::cuda::__throw_cuda_error(__res, "Failed to get the managed symbol from library");
+        ::cuda::__throw_cuda_error(
+          static_cast<::cudaError_t>(__status), "Failed to get the managed symbol from library");
     }
   }
 
@@ -187,14 +173,8 @@ public:
   //! @note Managed memory is shared across devices
   [[nodiscard]] library_symbol_info managed(const char* __name) const
   {
-    ::CUdeviceptr __dptr{};
-    ::cuda::std::size_t __size{};
-    if (const auto __res = ::cuda::__driver::__libraryGetManagedNoThrow(__dptr, __size, __library_, __name);
-        __res != ::cudaSuccess)
-    {
-      ::cuda::__throw_cuda_error(__res, "Failed to get the managed symbol from the library");
-    }
-    return library_symbol_info{reinterpret_cast<void*>(__dptr), __size};
+    auto [__ptr, __size] = _CCCL_TRY_DRIVER_API(__libraryGetManaged(__library_, __name));
+    return library_symbol_info{__ptr, __size};
   }
 
   //! @brief Gets the CUlibrary handle

cudax/include/cuda/experimental/__memory_resource/synchronous_resource_adapter.cuh

@@ -85,7 +85,7 @@ struct synchronous_resource_adapter : ::cuda::mr::__copy_default_queries<_Resour
     }
     else
     {
-      ::cuda::__driver::__streamSynchronizeNoThrow(__stream.get());
+      _CCCL_ASSERT_DRIVER_API(__streamSynchronize(__stream.get()));
       __resource.deallocate_sync(__ptr, __bytes, __alignment);
     }
   }

cudax/include/cuda/experimental/__stream/stream.cuh

@@ -95,7 +95,7 @@ struct stream : stream_ref
     {
       // Needs to call driver API in case current device is not set, runtime version would set dev 0 current
       // Alternative would be to store the device and push/pop here
-      [[maybe_unused]] auto status = ::cuda::__driver::__streamDestroyNoThrow(__stream);
+      _CCCL_ASSERT_DRIVER_API(__streamDestroy(__stream));
     }
   }