Skip to content

Remove cuda driver calls from core #1247

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Draft
andrea-cimatoribus-pix4d wants to merge 1 commit into
base: main
Choose a base branch
from
Draft

Remove cuda driver calls from core #1247

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
55 changes: 53 additions & 2 deletions include/matx/core/error.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -39,6 +39,7 @@
#include <sstream>
#ifdef __CUDACC__
#include <cuda.h>
#include <cuda_runtime_api.h>
#endif

#include "matx/core/stacktrace.h"
Expand Down Expand Up @@ -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 <typename Fn>
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<unsigned int>(driver_version) < CUDA_VERSION
? static_cast<unsigned int>(driver_version)
: static_cast<unsigned int>(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>(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<fn_t>("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<int>(e_)); \
MATX_THROW(matx::matxCudaError, err_str != nullptr ? err_str : "unknown"); \
} \
Expand Down
79 changes: 58 additions & 21 deletions include/matx/core/nvrtc_helper.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -330,13 +330,50 @@ std::vector<std::string> __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<fn_t>("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<fn_t>("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<fn_t>("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<fn_t>("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); \
Expand Down Expand Up @@ -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
Expand Down Expand Up @@ -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
Expand All @@ -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();
Expand All @@ -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
Expand All @@ -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));
}
}

Expand Down
20 changes: 11 additions & 9 deletions include/matx/core/print.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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<CUdeviceptr>(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 {
Expand Down
43 changes: 22 additions & 21 deletions include/matx/core/tensor.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;
}

/**
Expand Down Expand Up @@ -1518,7 +1522,6 @@ MATX_LOOP_UNROLL
// Pass in the base pointer, not a potentially offset pointer
void *data_ptr = const_cast<void *>(static_cast<const void *>(this->GetStorage().data()));
auto kind = GetPointerKind(data_ptr);
auto cu_ptr = reinterpret_cast<CUdeviceptr>(data_ptr);

if (kind == MATX_INVALID_MEMORY) {
// GetStorage().data() is only guaranteed to be the true allocation base
Expand All @@ -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;
}
Expand All @@ -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:
Expand Down