diff --git a/include/matx/core/error.h b/include/matx/core/error.h index 26b975a58..432676032 100644 --- a/include/matx/core/error.h +++ b/include/matx/core/error.h @@ -39,6 +39,7 @@ #include #ifdef __CUDACC__ #include +#include #endif #include "matx/core/stacktrace.h" @@ -265,14 +266,64 @@ namespace matx MATX_CUDA_CHECK(e); \ } +namespace detail { +#ifdef __CUDACC__ +// The CUDA driver library (libcuda.so.1 on Linux, nvcuda.dll on Windows) ships with the display +// driver, not with the toolkit. Calling its entry points directly makes every consumer record a +// load-time dependency on it, so a binary using MatX cannot even start on a machine that has no +// NVIDIA driver installed: the loader rejects it before main() runs, leaving the application no +// way to report the problem. The CUDA runtime hands out the same entry points and resolves the +// driver lazily, which keeps that dependency out of the link. +// +// Returns nullptr when the entry point is unavailable. +template +inline Fn DrvEntryPoint(const char *symbol) +{ + // Ask for the API version the installed driver actually implements rather than the toolkit's + // CUDA_VERSION. A driver older than the toolkit is a supported configuration (CUDA minor + // version compatibility), and asking such a driver for a newer version than it provides makes + // the query fail with cudaDriverEntryPointSymbolNotFound. + int driver_version = 0; + if (cudaDriverGetVersion(&driver_version) != cudaSuccess || driver_version <= 0) { + return nullptr; + } + const unsigned int api_version = static_cast(driver_version) < CUDA_VERSION + ? static_cast(driver_version) + : static_cast(CUDA_VERSION); + + void *fn = nullptr; + cudaDriverEntryPointQueryResult qres = cudaDriverEntryPointSymbolNotFound; + const cudaError_t rc = + cudaGetDriverEntryPointByVersion(symbol, &fn, api_version, cudaEnableDefault, &qres); + // An unavailable symbol is reported as a success with a null pointer, so both must be checked + if (rc != cudaSuccess || qres != cudaDriverEntryPointSuccess) { + return nullptr; + } + return reinterpret_cast(fn); +} + +// Deliberately does not throw: this runs on an error path, where failing would replace the +// original error with a less useful one. +inline const char *DrvGetErrorString(CUresult error) +{ + using fn_t = CUresult(CUDAAPI *)(CUresult, const char **); + static const fn_t fn = DrvEntryPoint("cuGetErrorString"); + const char *str = nullptr; + if (fn == nullptr || fn(error, &str) != CUDA_SUCCESS) { + return nullptr; + } + return str; +} +#endif +} + // Macro for checking CUDA driver API (CUresult) errors #define MATX_CUDA_DRIVER_CHECK(e) \ do { \ const CUresult e_ = (e); \ if (e_ != CUDA_SUCCESS) \ { \ - const char *err_str = nullptr; \ - cuGetErrorString(e_, &err_str); \ + const char *err_str = matx::detail::DrvGetErrorString(e_); \ MATX_LOG_ERROR("{}:{} CUDA Driver Error: {} ({})", __FILE__, __LINE__, err_str != nullptr ? err_str : "unknown", static_cast(e_)); \ MATX_THROW(matx::matxCudaError, err_str != nullptr ? err_str : "unknown"); \ } \ diff --git a/include/matx/core/nvrtc_helper.h b/include/matx/core/nvrtc_helper.h index b73563810..a9877b521 100644 --- a/include/matx/core/nvrtc_helper.h +++ b/include/matx/core/nvrtc_helper.h @@ -330,13 +330,50 @@ std::vector __MATX_HOST__ __MATX_INLINE__ get_preprocessor_options( } \ } while (0) +// These have no CUDA runtime API equivalent, so they are resolved through the runtime instead of +// linked against libcuda. See matx/core/error.h for why that matters. +inline CUresult DrvModuleLoadDataEx(CUmodule *module, const void *image, unsigned int numOptions, + CUjit_option *options, void **optionValues) { + using fn_t = CUresult(CUDAAPI *)(CUmodule *, const void *, unsigned int, CUjit_option *, void **); + static const fn_t fn = DrvEntryPoint("cuModuleLoadDataEx"); + MATX_ASSERT_STR(fn != nullptr, matxCudaError, "CUDA driver entry point cuModuleLoadDataEx is unavailable"); + return fn(module, image, numOptions, options, optionValues); +} + +inline CUresult DrvModuleGetFunction(CUfunction *hfunc, CUmodule hmod, const char *name) { + using fn_t = CUresult(CUDAAPI *)(CUfunction *, CUmodule, const char *); + static const fn_t fn = DrvEntryPoint("cuModuleGetFunction"); + MATX_ASSERT_STR(fn != nullptr, matxCudaError, "CUDA driver entry point cuModuleGetFunction is unavailable"); + return fn(hfunc, hmod, name); +} + +inline CUresult DrvFuncSetAttribute(CUfunction hfunc, CUfunction_attribute attrib, int value) { + using fn_t = CUresult(CUDAAPI *)(CUfunction, CUfunction_attribute, int); + static const fn_t fn = DrvEntryPoint("cuFuncSetAttribute"); + MATX_ASSERT_STR(fn != nullptr, matxCudaError, "CUDA driver entry point cuFuncSetAttribute is unavailable"); + return fn(hfunc, attrib, value); +} + +inline CUresult DrvLaunchKernel(CUfunction f, unsigned int gridDimX, unsigned int gridDimY, + unsigned int gridDimZ, unsigned int blockDimX, + unsigned int blockDimY, unsigned int blockDimZ, + unsigned int sharedMemBytes, CUstream hStream, void **kernelParams, + void **extra) { + using fn_t = CUresult(CUDAAPI *)(CUfunction, unsigned int, unsigned int, unsigned int, + unsigned int, unsigned int, unsigned int, unsigned int, CUstream, + void **, void **); + static const fn_t fn = DrvEntryPoint("cuLaunchKernel"); + MATX_ASSERT_STR(fn != nullptr, matxCudaError, "CUDA driver entry point cuLaunchKernel is unavailable"); + return fn(f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, sharedMemBytes, + hStream, kernelParams, extra); +} + // Helper function to check CUDA Driver API errors #define CUDA_CHECK(call) \ do { \ CUresult result = call; \ if (result != CUDA_SUCCESS) { \ - const char* errStr = nullptr; \ - cuGetErrorString(result, &errStr); \ + const char* errStr = DrvGetErrorString(result); \ std::string error_msg = std::string("CUDA error: ") + \ (errStr != nullptr ? errStr : "unknown"); \ MATX_LOG_ERROR("{}", error_msg); \ @@ -881,10 +918,10 @@ auto nvrtc_compile_and_run([[maybe_unused]] const std::string &name, // Load the cached cubin into a CUDA module CUmodule module; - CUDA_CHECK(cuModuleLoadDataEx(&module, cached_cubin_ptr->data, 0, nullptr, nullptr)); + CUDA_CHECK(DrvModuleLoadDataEx(&module, cached_cubin_ptr->data, 0, nullptr, nullptr)); // Get kernel function using the cached lowered name - CUDA_CHECK(cuModuleGetFunction(&kernel_func, module, lowered_name.c_str())); + CUDA_CHECK(DrvModuleGetFunction(&kernel_func, module, lowered_name.c_str())); // Cache both module and function to prevent resource leak // Module must stay loaded for function to remain valid @@ -1063,10 +1100,10 @@ auto nvrtc_compile_and_run([[maybe_unused]] const std::string &name, // Load LTO-IR into CUDA module CUmodule module; - CUDA_CHECK(cuModuleLoadDataEx(&module, cubin.data(), 0, nullptr, nullptr)); + CUDA_CHECK(DrvModuleLoadDataEx(&module, cubin.data(), 0, nullptr, nullptr)); // Get kernel function using the lowered name - CUDA_CHECK(cuModuleGetFunction(&kernel_func, module, lowered_name.c_str())); + CUDA_CHECK(DrvModuleGetFunction(&kernel_func, module, lowered_name.c_str())); // Cache both module and function to prevent resource leak // Module must stay loaded for function to remain valid @@ -1092,7 +1129,7 @@ auto nvrtc_compile_and_run([[maybe_unused]] const std::string &name, if (dynamic_shmem_size > max_shared_memory_per_block) { MATX_LOG_DEBUG("Requested dynamic shared memory ({} bytes) exceeds default per-block shared memory limit ({})", dynamic_shmem_size, max_shared_memory_per_block); - CUDA_CHECK(cuFuncSetAttribute(kernel_func, CU_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES, dynamic_shmem_size)); + CUDA_CHECK(DrvFuncSetAttribute(kernel_func, CU_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES, dynamic_shmem_size)); } auto storage = op.ToJITStorage(); @@ -1111,13 +1148,13 @@ auto nvrtc_compile_and_run([[maybe_unused]] const std::string &name, MATX_LOG_DEBUG("Launching kernel with grid=({}, {}, {}), block=({}, {}, {}), dynamic_shmem_size={} bytes", blocks.x, blocks.y, blocks.z, threads.x, threads.y, threads.z, dynamic_shmem_size); // Launch kernel - CUDA_CHECK(cuLaunchKernel(kernel_func, - blocks.x, blocks.y, blocks.z, - threads.x, threads.y, threads.z, - dynamic_shmem_size, - stream, - args, - nullptr)); + CUDA_CHECK(DrvLaunchKernel(kernel_func, + blocks.x, blocks.y, blocks.z, + threads.x, threads.y, threads.z, + dynamic_shmem_size, + stream, + args, + nullptr)); } else { // Rank 0-4 kernels: Pass individual size parameters @@ -1139,13 +1176,13 @@ auto nvrtc_compile_and_run([[maybe_unused]] const std::string &name, MATX_LOG_DEBUG("Launching kernel with grid=({}, {}, {}), block=({}, {}, {}), dynamic_shmem_size={} bytes", blocks.x, blocks.y, blocks.z, threads.x, threads.y, threads.z, dynamic_shmem_size); // Launch kernel - CUDA_CHECK(cuLaunchKernel(kernel_func, - blocks.x, blocks.y, blocks.z, - threads.x, threads.y, threads.z, - dynamic_shmem_size, - stream, - args, - nullptr)); + CUDA_CHECK(DrvLaunchKernel(kernel_func, + blocks.x, blocks.y, blocks.z, + threads.x, threads.y, threads.z, + dynamic_shmem_size, + stream, + args, + nullptr)); } } diff --git a/include/matx/core/print.h b/include/matx/core/print.h index 0ca94b878..333da890f 100644 --- a/include/matx/core/print.h +++ b/include/matx/core/print.h @@ -597,17 +597,19 @@ namespace matx { // Try to get pointer from cuda if (kind == MATX_INVALID_MEMORY) { - CUmemorytype mtype; - void *data[] = {&mtype}; - CUpointer_attribute attrs[] = {CU_POINTER_ATTRIBUTE_MEMORY_TYPE}; - MATX_CUDA_DRIVER_CHECK(cuPointerGetAttributes(1, - &attrs[0], - data, - reinterpret_cast(op.Data()))); - MATX_ASSERT_STR(mtype == CU_MEMORYTYPE_HOST || mtype == 0 || mtype == CU_MEMORYTYPE_DEVICE, + cudaPointerAttributes ptr_attr{}; + MATX_CUDA_CHECK(cudaPointerGetAttributes(&ptr_attr, op.Data())); + // cudaMemoryTypeUnregistered replaces the zero the driver API left behind for a + // pointer CUDA has no record of + MATX_ASSERT_STR(ptr_attr.type == cudaMemoryTypeHost || + ptr_attr.type == cudaMemoryTypeUnregistered || + ptr_attr.type == cudaMemoryTypeDevice || + ptr_attr.type == cudaMemoryTypeManaged, matxNotSupported, "Invalid memory type for printing"); - if (mtype == CU_MEMORYTYPE_DEVICE) { + // The driver API folded managed memory into CU_MEMORYTYPE_DEVICE, so keeping managed + // pointers on the device path preserves the existing behaviour + if (ptr_attr.type == cudaMemoryTypeDevice || ptr_attr.type == cudaMemoryTypeManaged) { detail::DevicePrint(fp, op, dims...); } else { diff --git a/include/matx/core/tensor.h b/include/matx/core/tensor.h index 4b050632f..7083d0b68 100644 --- a/include/matx/core/tensor.h +++ b/include/matx/core/tensor.h @@ -1143,11 +1143,15 @@ MATX_LOOP_UNROLL } __MATX_INLINE__ __MATX_HOST__ bool IsHostAccessiblePointer() { - void* hostPtr = nullptr; - [[maybe_unused]] const CUresult retval = - cuPointerGetAttribute(&hostPtr, CU_POINTER_ATTRIBUTE_HOST_POINTER, (CUdeviceptr)this->Data()); - MATX_ASSERT_STR_EXP(retval, CUDA_SUCCESS, matxNotSupported, "Pointer is not host-accessible"); - return hostPtr != nullptr; + // cudaPointerAttributes::hostPointer is the runtime equivalent of the driver's + // CU_POINTER_ATTRIBUTE_HOST_POINTER: the address that may be dereferenced on the host, or + // null if there is none. Unlike the driver call it also succeeds for a pointer CUDA has no + // record of, reporting it as cudaMemoryTypeUnregistered with a null hostPointer, so an + // ordinary malloc'd pointer makes this predicate return false rather than throw. + cudaPointerAttributes ptr_attr{}; + [[maybe_unused]] const cudaError_t retval = cudaPointerGetAttributes(&ptr_attr, this->Data()); + MATX_ASSERT_STR_EXP(retval, cudaSuccess, matxNotSupported, "Pointer is not host-accessible"); + return ptr_attr.hostPointer != nullptr; } /** @@ -1518,7 +1522,6 @@ MATX_LOOP_UNROLL // Pass in the base pointer, not a potentially offset pointer void *data_ptr = const_cast(static_cast(this->GetStorage().data())); auto kind = GetPointerKind(data_ptr); - auto cu_ptr = reinterpret_cast(data_ptr); if (kind == MATX_INVALID_MEMORY) { // GetStorage().data() is only guaranteed to be the true allocation base @@ -1527,23 +1530,22 @@ MATX_LOOP_UNROLL // reinterpreted address (e.g. RealView()/ImagView()) lands here instead, // so classify it from the driver's own record of that address. // Managed memory is reported separately from plain device memory since - // DLPack defines a distinct kDLCUDAManaged device type for it. - CUpointer_attribute attr[] = {CU_POINTER_ATTRIBUTE_MEMORY_TYPE, CU_POINTER_ATTRIBUTE_DEVICE_ORDINAL, CU_POINTER_ATTRIBUTE_IS_MANAGED}; - CUmemorytype mem_type; - int dev_ord; - int is_managed; - void *data[3] = {&mem_type, &dev_ord, &is_managed}; - MATX_CUDA_DRIVER_CHECK(cuPointerGetAttributes(sizeof(attr)/sizeof(attr[0]), attr, data, cu_ptr)); - - if (is_managed) { + // DLPack defines a distinct kDLCUDAManaged device type for it. The runtime + // API reports managed memory as its own memory type, so it answers in one + // call what the driver API needed a separate IS_MANAGED query for. + cudaPointerAttributes ptr_attr{}; + MATX_CUDA_CHECK(cudaPointerGetAttributes(&ptr_attr, data_ptr)); + const int dev_ord = ptr_attr.device; + + if (ptr_attr.type == cudaMemoryTypeManaged) { t->device.device_type = kDLCUDAManaged; t->device.device_id = dev_ord; } - else if (mem_type == CU_MEMORYTYPE_DEVICE) { + else if (ptr_attr.type == cudaMemoryTypeDevice) { t->device.device_type = kDLCUDA; t->device.device_id = dev_ord; } - else if (mem_type == CU_MEMORYTYPE_HOST) { + else if (ptr_attr.type == cudaMemoryTypeHost) { t->device.device_type = kDLCUDAHost; t->device.device_id = dev_ord; } @@ -1558,10 +1560,9 @@ MATX_LOOP_UNROLL else { // We have a record of this pointer's memory space; only the device // ordinal still needs to come from the driver - CUpointer_attribute attr = CU_POINTER_ATTRIBUTE_DEVICE_ORDINAL; - int dev_ord; - void *data = &dev_ord; - MATX_CUDA_DRIVER_CHECK(cuPointerGetAttributes(1, &attr, &data, cu_ptr)); + cudaPointerAttributes ptr_attr{}; + MATX_CUDA_CHECK(cudaPointerGetAttributes(&ptr_attr, data_ptr)); + const int dev_ord = ptr_attr.device; switch (kind) { case MATX_MANAGED_MEMORY: