From fd513b48467cd09ecef4d158faecd8644c701026 Mon Sep 17 00:00:00 2001 From: Nikhil Talpallikar Date: Fri, 1 Aug 2025 10:27:09 -0700 Subject: [PATCH 1/9] Fix null pointer refrence issue with cuda driver API function pointers in case cuInit fails --- .../matrixMulDynlinkJIT/cuda_drvapi_dynlink.c | 22 +++++++++---------- 1 file changed, 11 insertions(+), 11 deletions(-) diff --git a/Samples/0_Introduction/matrixMulDynlinkJIT/cuda_drvapi_dynlink.c b/Samples/0_Introduction/matrixMulDynlinkJIT/cuda_drvapi_dynlink.c index 8adc4d87a..a38afc8d8 100644 --- a/Samples/0_Introduction/matrixMulDynlinkJIT/cuda_drvapi_dynlink.c +++ b/Samples/0_Introduction/matrixMulDynlinkJIT/cuda_drvapi_dynlink.c @@ -345,17 +345,6 @@ CUresult CUDAAPI cuInit(unsigned int Flags, int cudaVersion) CHECKED_CALL(LOAD_LIBRARY(&CudaDrvLib)); - // cuInit is required; alias it to _cuInit - GET_PROC_EX(cuInit, _cuInit, 1); - CHECKED_CALL(_cuInit(Flags)); - - // available since 2.2. if not present, version 1.0 is assumed - GET_PROC_OPTIONAL(cuDriverGetVersion); - - if (cuDriverGetVersion) { - CHECKED_CALL(cuDriverGetVersion(&driverVer)); - } - // fetch all function pointers GET_PROC(cuDeviceGet); GET_PROC(cuDeviceGetCount); @@ -620,5 +609,16 @@ CUresult CUDAAPI cuInit(unsigned int Flags, int cudaVersion) #endif } + // cuInit is required; alias it to _cuInit + GET_PROC_EX(cuInit, _cuInit, 1); + CHECKED_CALL(_cuInit(Flags)); + + // available since 2.2. if not present, version 1.0 is assumed + GET_PROC_OPTIONAL(cuDriverGetVersion); + + if (cuDriverGetVersion) { + CHECKED_CALL(cuDriverGetVersion(&driverVer)); + } + return CUDA_SUCCESS; } From f8aab0053f3ea388472e3e574139ec6fa99322f4 Mon Sep 17 00:00:00 2001 From: Nikhil Talpallikar Date: Tue, 5 Aug 2025 13:46:44 -0700 Subject: [PATCH 2/9] Clean implementation for failure path when cuInit fails --- .../matrixMulDynlinkJIT/cuda_drvapi_dynlink.c | 28 +++++++++---------- 1 file changed, 14 insertions(+), 14 deletions(-) diff --git a/Samples/0_Introduction/matrixMulDynlinkJIT/cuda_drvapi_dynlink.c b/Samples/0_Introduction/matrixMulDynlinkJIT/cuda_drvapi_dynlink.c index a38afc8d8..6f1c97383 100644 --- a/Samples/0_Introduction/matrixMulDynlinkJIT/cuda_drvapi_dynlink.c +++ b/Samples/0_Introduction/matrixMulDynlinkJIT/cuda_drvapi_dynlink.c @@ -342,9 +342,21 @@ CUresult CUDAAPI cuInit(unsigned int Flags, int cudaVersion) { CUDADRIVER CudaDrvLib; int driverVer = 1000; + CUresult result = CUDA_SUCCESS; - CHECKED_CALL(LOAD_LIBRARY(&CudaDrvLib)); + result = LOAD_LIBRARY(&CudaDrvLib); + // cuInit is required; alias it to _cuInit + GET_PROC_EX(cuInit, _cuInit, 1); + result = _cuInit(Flags); + + // available since 2.2. if not present, version 1.0 is assumed + GET_PROC_OPTIONAL(cuDriverGetVersion); + + if (cuDriverGetVersion) { + result = cuDriverGetVersion(&driverVer); + } + // fetch all function pointers GET_PROC(cuDeviceGet); GET_PROC(cuDeviceGetCount); @@ -608,17 +620,5 @@ CUresult CUDAAPI cuInit(unsigned int Flags, int cudaVersion) GET_PROC(cuGraphicsD3D9RegisterResource); #endif } - - // cuInit is required; alias it to _cuInit - GET_PROC_EX(cuInit, _cuInit, 1); - CHECKED_CALL(_cuInit(Flags)); - - // available since 2.2. if not present, version 1.0 is assumed - GET_PROC_OPTIONAL(cuDriverGetVersion); - - if (cuDriverGetVersion) { - CHECKED_CALL(cuDriverGetVersion(&driverVer)); - } - - return CUDA_SUCCESS; + return result; } From 527b29dbd0dac5ff100e8946997aaf3d54c33958 Mon Sep 17 00:00:00 2001 From: Nikhil Talpallikar Date: Tue, 5 Aug 2025 13:49:28 -0700 Subject: [PATCH 3/9] Clean implementation for failure path when cuInit fails. Removed CHECKED_CALL macro which returned prematurely --- .../matrixMulDynlinkJIT/cuda_drvapi_dynlink.c | 8 -------- 1 file changed, 8 deletions(-) diff --git a/Samples/0_Introduction/matrixMulDynlinkJIT/cuda_drvapi_dynlink.c b/Samples/0_Introduction/matrixMulDynlinkJIT/cuda_drvapi_dynlink.c index 6f1c97383..05048e15c 100644 --- a/Samples/0_Introduction/matrixMulDynlinkJIT/cuda_drvapi_dynlink.c +++ b/Samples/0_Introduction/matrixMulDynlinkJIT/cuda_drvapi_dynlink.c @@ -324,14 +324,6 @@ static CUresult LOAD_LIBRARY(CUDADRIVER *pInstance) #error unsupported platform #endif -#define CHECKED_CALL(call) \ - do { \ - CUresult result = (call); \ - if (CUDA_SUCCESS != result) { \ - return result; \ - } \ - } while (0) - #define GET_PROC_REQUIRED(name) GET_PROC_EX(name, name, 1) #define GET_PROC_OPTIONAL(name) GET_PROC_EX(name, name, 0) #define GET_PROC(name) GET_PROC_REQUIRED(name) From d2c52db3e0f08095fc2a01d2bfc1712a6ecb8e2d Mon Sep 17 00:00:00 2001 From: Nikhil Talpallikar Date: Wed, 6 Aug 2025 00:29:22 -0700 Subject: [PATCH 4/9] Fixed the error path to initialize error path function pointers. Exit with error in case of LOADLIBRARY failureas initialize of function pointers in case of LOADLIBRARY failure will fail --- .../matrixMulDynlinkJIT/cuda_drvapi_dynlink.c | 68 ++++++++++++++++--- .../matrixMulDynlinkJIT/helper_cuda_drvapi.h | 10 +++ 2 files changed, 70 insertions(+), 8 deletions(-) diff --git a/Samples/0_Introduction/matrixMulDynlinkJIT/cuda_drvapi_dynlink.c b/Samples/0_Introduction/matrixMulDynlinkJIT/cuda_drvapi_dynlink.c index 05048e15c..b4e8b95e7 100644 --- a/Samples/0_Introduction/matrixMulDynlinkJIT/cuda_drvapi_dynlink.c +++ b/Samples/0_Introduction/matrixMulDynlinkJIT/cuda_drvapi_dynlink.c @@ -242,12 +242,23 @@ static CUresult LOAD_LIBRARY(CUDADRIVER *pInstance) if (*pInstance == NULL) { printf("LoadLibrary \"%s\" failed!\n", __CudaLibName); - return CUDA_ERROR_UNKNOWN; + exit(EXIT_FAILURE); } return CUDA_SUCCESS; } +CUresult GET_DRIVER_HANDLE(CUDADRIVER* pInstance) +{ + *pInstance = GetModuleHandle(__CudaLibName); + if (*pInstance) { + return CUDA_SUCCESS; + } + else { + return CUDA_ERROR_UNKNOWN; + } +} + #define GET_PROC_EX(name, alias, required) \ alias = (t##name *)GetProcAddress(CudaDrvLib, #name); \ if (alias == NULL && required) { \ @@ -269,6 +280,13 @@ static CUresult LOAD_LIBRARY(CUDADRIVER *pInstance) return CUDA_ERROR_UNKNOWN; \ } +#define GET_PROC_ERROR_FUNCTIONS(name, alias, required) \ + alias = (t##name *)GetProcAddress(CudaDrvLib, #name); \ + if (alias == NULL && required) { \ + printf("Failed to find error function \"%s\" in %s\n", #name, __CudaLibName); \ + exit(EXIT_FAILURE); \ + } \ + #elif defined(__unix__) || defined(__QNX__) || defined(__APPLE__) || defined(__MACOSX) #include @@ -293,12 +311,23 @@ static CUresult LOAD_LIBRARY(CUDADRIVER *pInstance) if (*pInstance == NULL) { printf("dlopen \"%s\" failed!\n", __CudaLibName); - return CUDA_ERROR_UNKNOWN; + exit(EXIT_FAILURE); } return CUDA_SUCCESS; } +CUresult GET_DRIVER_HANDLE(CUDADRIVER* pInstance) +{ + *pInstance = dlopen(__CudaLibName, RTLD_NOLOAD); + if (*pInstance) { + return CUDA_SUCCESS; + } + else { + return CUDA_ERROR_UNKNOWN; + } +} + #define GET_PROC_EX(name, alias, required) \ alias = (t##name *)dlsym(CudaDrvLib, #name); \ if (alias == NULL && required) { \ @@ -320,33 +349,56 @@ static CUresult LOAD_LIBRARY(CUDADRIVER *pInstance) return CUDA_ERROR_UNKNOWN; \ } +#define GET_PROC_ERROR_FUNCTIONS(name, alias, required) \ + alias = (t##name *)dlsym(CudaDrvLib, #name); \ + if (alias == NULL && required) { \ + printf("Failed to find error function \"%s\" in %s\n", #name, __CudaLibName); \ + exit(EXIT_FAILURE); \ + } + #else #error unsupported platform #endif +#define CHECKED_CALL(call) \ + do { \ + CUresult result = (call); \ + if (CUDA_SUCCESS != result) { \ + return result; \ + } \ + } while (0) + #define GET_PROC_REQUIRED(name) GET_PROC_EX(name, name, 1) #define GET_PROC_OPTIONAL(name) GET_PROC_EX(name, name, 0) #define GET_PROC(name) GET_PROC_REQUIRED(name) #define GET_PROC_V2(name) GET_PROC_EX_V2(name, name, 1) #define GET_PROC_V3(name) GET_PROC_EX_V3(name, name, 1) +CUresult INIT_ERROR_FUNCTIONS(void) +{ + CUDADRIVER CudaDrvLib; + CUresult result = CUDA_SUCCESS; + result = GET_DRIVER_HANDLE(&CudaDrvLib); + GET_PROC_ERROR_FUNCTIONS(cuGetErrorString, cuGetErrorString, 1); + return result; +} + CUresult CUDAAPI cuInit(unsigned int Flags, int cudaVersion) { CUDADRIVER CudaDrvLib; int driverVer = 1000; - CUresult result = CUDA_SUCCESS; - - result = LOAD_LIBRARY(&CudaDrvLib); + + CHECKED_CALL(LOAD_LIBRARY(&CudaDrvLib)); // cuInit is required; alias it to _cuInit GET_PROC_EX(cuInit, _cuInit, 1); - result = _cuInit(Flags); + CHECKED_CALL(_cuInit(Flags)); // available since 2.2. if not present, version 1.0 is assumed GET_PROC_OPTIONAL(cuDriverGetVersion); if (cuDriverGetVersion) { - result = cuDriverGetVersion(&driverVer); + CHECKED_CALL(cuDriverGetVersion(&driverVer)); } // fetch all function pointers @@ -612,5 +664,5 @@ CUresult CUDAAPI cuInit(unsigned int Flags, int cudaVersion) GET_PROC(cuGraphicsD3D9RegisterResource); #endif } - return result; + return CUDA_SUCCESS; } diff --git a/Samples/0_Introduction/matrixMulDynlinkJIT/helper_cuda_drvapi.h b/Samples/0_Introduction/matrixMulDynlinkJIT/helper_cuda_drvapi.h index 7c61ff2a4..7f0e6d78d 100644 --- a/Samples/0_Introduction/matrixMulDynlinkJIT/helper_cuda_drvapi.h +++ b/Samples/0_Introduction/matrixMulDynlinkJIT/helper_cuda_drvapi.h @@ -42,11 +42,21 @@ inline int ftoi(float value) { return (value >= 0 ? static_cast(value + 0.5 #ifndef checkCudaErrors #define checkCudaErrors(err) __checkCudaErrors(err, __FILE__, __LINE__) +extern "C" CUresult INIT_ERROR_FUNCTIONS(void); + // These are the inline versions for all of the SDK helper functions inline void __checkCudaErrors(CUresult err, const char *file, const int line) { if (CUDA_SUCCESS != err) { const char *errorStr = NULL; + + if (!cuGetErrorString) { + CUresult result = INIT_ERROR_FUNCTIONS(); + if (result != CUDA_SUCCESS) { + printf("CUDA driver API failed"); + exit(EXIT_FAILURE); + } + } cuGetErrorString(err, &errorStr); fprintf(stderr, "checkCudaErrors() Driver API error = %04d \"%s\" from file <%s>, " From 6df7127c2371f1193257e0c31551eef4a6aaea94 Mon Sep 17 00:00:00 2001 From: Nikhil Talpallikar Date: Wed, 6 Aug 2025 00:36:35 -0700 Subject: [PATCH 5/9] Fixed dlopen on linux with lazy load flag --- .../0_Introduction/matrixMulDynlinkJIT/cuda_drvapi_dynlink.c | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/Samples/0_Introduction/matrixMulDynlinkJIT/cuda_drvapi_dynlink.c b/Samples/0_Introduction/matrixMulDynlinkJIT/cuda_drvapi_dynlink.c index b4e8b95e7..0983226f5 100644 --- a/Samples/0_Introduction/matrixMulDynlinkJIT/cuda_drvapi_dynlink.c +++ b/Samples/0_Introduction/matrixMulDynlinkJIT/cuda_drvapi_dynlink.c @@ -319,7 +319,7 @@ static CUresult LOAD_LIBRARY(CUDADRIVER *pInstance) CUresult GET_DRIVER_HANDLE(CUDADRIVER* pInstance) { - *pInstance = dlopen(__CudaLibName, RTLD_NOLOAD); + *pInstance = dlopen(__CudaLibName, RTLD_LAZY); if (*pInstance) { return CUDA_SUCCESS; } From 0861db73ad82d635d790ef512215112f8ac4c2f3 Mon Sep 17 00:00:00 2001 From: ntalpallikar Date: Wed, 6 Aug 2025 10:56:14 -0700 Subject: [PATCH 6/9] fixing indentation --- .../matrixMulDynlinkJIT/cuda_drvapi_dynlink.c | 44 +++++++++---------- 1 file changed, 21 insertions(+), 23 deletions(-) diff --git a/Samples/0_Introduction/matrixMulDynlinkJIT/cuda_drvapi_dynlink.c b/Samples/0_Introduction/matrixMulDynlinkJIT/cuda_drvapi_dynlink.c index 0983226f5..cda089efc 100644 --- a/Samples/0_Introduction/matrixMulDynlinkJIT/cuda_drvapi_dynlink.c +++ b/Samples/0_Introduction/matrixMulDynlinkJIT/cuda_drvapi_dynlink.c @@ -280,12 +280,12 @@ CUresult GET_DRIVER_HANDLE(CUDADRIVER* pInstance) return CUDA_ERROR_UNKNOWN; \ } -#define GET_PROC_ERROR_FUNCTIONS(name, alias, required) \ - alias = (t##name *)GetProcAddress(CudaDrvLib, #name); \ - if (alias == NULL && required) { \ - printf("Failed to find error function \"%s\" in %s\n", #name, __CudaLibName); \ - exit(EXIT_FAILURE); \ - } \ +#define GET_PROC_ERROR_FUNCTIONS(name, alias, required) \ + alias = (t##name *)GetProcAddress(CudaDrvLib, #name); \ + if (alias == NULL && required) { \ + printf("Failed to find error function \"%s\" in %s\n", #name, __CudaLibName); \ + exit(EXIT_FAILURE); \ + } \ #elif defined(__unix__) || defined(__QNX__) || defined(__APPLE__) || defined(__MACOSX) @@ -349,23 +349,23 @@ CUresult GET_DRIVER_HANDLE(CUDADRIVER* pInstance) return CUDA_ERROR_UNKNOWN; \ } -#define GET_PROC_ERROR_FUNCTIONS(name, alias, required) \ - alias = (t##name *)dlsym(CudaDrvLib, #name); \ - if (alias == NULL && required) { \ - printf("Failed to find error function \"%s\" in %s\n", #name, __CudaLibName); \ - exit(EXIT_FAILURE); \ +#define GET_PROC_ERROR_FUNCTIONS(name, alias, required) \ + alias = (t##name *)dlsym(CudaDrvLib, #name); \ + if (alias == NULL && required) { \ + printf("Failed to find error function \"%s\" in %s\n", #name, __CudaLibName); \ + exit(EXIT_FAILURE); \ } #else #error unsupported platform #endif -#define CHECKED_CALL(call) \ - do { \ - CUresult result = (call); \ - if (CUDA_SUCCESS != result) { \ - return result; \ - } \ +#define CHECKED_CALL(call) \ + do { \ + CUresult result = (call); \ + if (CUDA_SUCCESS != result) { \ + return result; \ + } \ } while (0) #define GET_PROC_REQUIRED(name) GET_PROC_EX(name, name, 1) @@ -377,8 +377,8 @@ CUresult GET_DRIVER_HANDLE(CUDADRIVER* pInstance) CUresult INIT_ERROR_FUNCTIONS(void) { CUDADRIVER CudaDrvLib; - CUresult result = CUDA_SUCCESS; - result = GET_DRIVER_HANDLE(&CudaDrvLib); + CUresult result = CUDA_SUCCESS; + result = GET_DRIVER_HANDLE(&CudaDrvLib); GET_PROC_ERROR_FUNCTIONS(cuGetErrorString, cuGetErrorString, 1); return result; } @@ -386,8 +386,7 @@ CUresult INIT_ERROR_FUNCTIONS(void) CUresult CUDAAPI cuInit(unsigned int Flags, int cudaVersion) { CUDADRIVER CudaDrvLib; - int driverVer = 1000; - + int driverVer = 1000; CHECKED_CALL(LOAD_LIBRARY(&CudaDrvLib)); // cuInit is required; alias it to _cuInit @@ -399,8 +398,7 @@ CUresult CUDAAPI cuInit(unsigned int Flags, int cudaVersion) if (cuDriverGetVersion) { CHECKED_CALL(cuDriverGetVersion(&driverVer)); - } - + } // fetch all function pointers GET_PROC(cuDeviceGet); GET_PROC(cuDeviceGetCount); From b4aaab387ec48a07994e3c542be52bdd110467e1 Mon Sep 17 00:00:00 2001 From: Nikhil Talpallikar Date: Fri, 29 Aug 2025 12:37:12 -0700 Subject: [PATCH 7/9] fixed formatting --- .../matrixMulDynlinkJIT/cuda_drvapi_dynlink.c | 50 +++++++++---------- .../matrixMulDynlinkJIT/helper_cuda_drvapi.h | 2 +- 2 files changed, 26 insertions(+), 26 deletions(-) diff --git a/Samples/0_Introduction/matrixMulDynlinkJIT/cuda_drvapi_dynlink.c b/Samples/0_Introduction/matrixMulDynlinkJIT/cuda_drvapi_dynlink.c index 0983226f5..6a50a7498 100644 --- a/Samples/0_Introduction/matrixMulDynlinkJIT/cuda_drvapi_dynlink.c +++ b/Samples/0_Introduction/matrixMulDynlinkJIT/cuda_drvapi_dynlink.c @@ -248,8 +248,8 @@ static CUresult LOAD_LIBRARY(CUDADRIVER *pInstance) return CUDA_SUCCESS; } -CUresult GET_DRIVER_HANDLE(CUDADRIVER* pInstance) -{ +CUresult GET_DRIVER_HANDLE(CUDADRIVER *pInstance) +{ *pInstance = GetModuleHandle(__CudaLibName); if (*pInstance) { return CUDA_SUCCESS; @@ -280,12 +280,12 @@ CUresult GET_DRIVER_HANDLE(CUDADRIVER* pInstance) return CUDA_ERROR_UNKNOWN; \ } -#define GET_PROC_ERROR_FUNCTIONS(name, alias, required) \ - alias = (t##name *)GetProcAddress(CudaDrvLib, #name); \ - if (alias == NULL && required) { \ - printf("Failed to find error function \"%s\" in %s\n", #name, __CudaLibName); \ - exit(EXIT_FAILURE); \ - } \ +#define GET_PROC_ERROR_FUNCTIONS(name, alias, required) \ + alias = (t##name *)GetProcAddress(CudaDrvLib, #name); \ + if (alias == NULL && required) { \ + printf("Failed to find error function \"%s\" in %s\n", #name, __CudaLibName); \ + exit(EXIT_FAILURE); \ + } #elif defined(__unix__) || defined(__QNX__) || defined(__APPLE__) || defined(__MACOSX) @@ -317,8 +317,8 @@ static CUresult LOAD_LIBRARY(CUDADRIVER *pInstance) return CUDA_SUCCESS; } -CUresult GET_DRIVER_HANDLE(CUDADRIVER* pInstance) -{ +CUresult GET_DRIVER_HANDLE(CUDADRIVER *pInstance) +{ *pInstance = dlopen(__CudaLibName, RTLD_LAZY); if (*pInstance) { return CUDA_SUCCESS; @@ -349,23 +349,23 @@ CUresult GET_DRIVER_HANDLE(CUDADRIVER* pInstance) return CUDA_ERROR_UNKNOWN; \ } -#define GET_PROC_ERROR_FUNCTIONS(name, alias, required) \ - alias = (t##name *)dlsym(CudaDrvLib, #name); \ - if (alias == NULL && required) { \ - printf("Failed to find error function \"%s\" in %s\n", #name, __CudaLibName); \ - exit(EXIT_FAILURE); \ +#define GET_PROC_ERROR_FUNCTIONS(name, alias, required) \ + alias = (t##name *)dlsym(CudaDrvLib, #name); \ + if (alias == NULL && required) { \ + printf("Failed to find error function \"%s\" in %s\n", #name, __CudaLibName); \ + exit(EXIT_FAILURE); \ } #else #error unsupported platform #endif -#define CHECKED_CALL(call) \ - do { \ - CUresult result = (call); \ - if (CUDA_SUCCESS != result) { \ - return result; \ - } \ +#define CHECKED_CALL(call) \ + do { \ + CUresult result = (call); \ + if (CUDA_SUCCESS != result) { \ + return result; \ + } \ } while (0) #define GET_PROC_REQUIRED(name) GET_PROC_EX(name, name, 1) @@ -377,8 +377,8 @@ CUresult GET_DRIVER_HANDLE(CUDADRIVER* pInstance) CUresult INIT_ERROR_FUNCTIONS(void) { CUDADRIVER CudaDrvLib; - CUresult result = CUDA_SUCCESS; - result = GET_DRIVER_HANDLE(&CudaDrvLib); + CUresult result = CUDA_SUCCESS; + result = GET_DRIVER_HANDLE(&CudaDrvLib); GET_PROC_ERROR_FUNCTIONS(cuGetErrorString, cuGetErrorString, 1); return result; } @@ -387,7 +387,7 @@ CUresult CUDAAPI cuInit(unsigned int Flags, int cudaVersion) { CUDADRIVER CudaDrvLib; int driverVer = 1000; - + CHECKED_CALL(LOAD_LIBRARY(&CudaDrvLib)); // cuInit is required; alias it to _cuInit @@ -400,7 +400,7 @@ CUresult CUDAAPI cuInit(unsigned int Flags, int cudaVersion) if (cuDriverGetVersion) { CHECKED_CALL(cuDriverGetVersion(&driverVer)); } - + // fetch all function pointers GET_PROC(cuDeviceGet); GET_PROC(cuDeviceGetCount); diff --git a/Samples/0_Introduction/matrixMulDynlinkJIT/helper_cuda_drvapi.h b/Samples/0_Introduction/matrixMulDynlinkJIT/helper_cuda_drvapi.h index 7f0e6d78d..befd53bb2 100644 --- a/Samples/0_Introduction/matrixMulDynlinkJIT/helper_cuda_drvapi.h +++ b/Samples/0_Introduction/matrixMulDynlinkJIT/helper_cuda_drvapi.h @@ -49,7 +49,7 @@ inline void __checkCudaErrors(CUresult err, const char *file, const int line) { if (CUDA_SUCCESS != err) { const char *errorStr = NULL; - + if (!cuGetErrorString) { CUresult result = INIT_ERROR_FUNCTIONS(); if (result != CUDA_SUCCESS) { From 3f1efe238f304286052aaa5c0e33c97c063f7f41 Mon Sep 17 00:00:00 2001 From: Nicholas Angle Date: Fri, 5 Sep 2025 14:57:48 -0700 Subject: [PATCH 8/9] Add trim notification sample --- Samples/3_CUDA_Features/CMakeLists.txt | 1 + .../trimNotification/CMakeLists.txt | 36 +++++ .../trimNotification/README.md | 33 +++++ .../trimNotification/trimNotification.cpp | 125 ++++++++++++++++++ 4 files changed, 195 insertions(+) create mode 100644 Samples/3_CUDA_Features/trimNotification/CMakeLists.txt create mode 100644 Samples/3_CUDA_Features/trimNotification/README.md create mode 100644 Samples/3_CUDA_Features/trimNotification/trimNotification.cpp diff --git a/Samples/3_CUDA_Features/CMakeLists.txt b/Samples/3_CUDA_Features/CMakeLists.txt index d7416b242..43df5b289 100644 --- a/Samples/3_CUDA_Features/CMakeLists.txt +++ b/Samples/3_CUDA_Features/CMakeLists.txt @@ -21,4 +21,5 @@ add_subdirectory(newdelete) add_subdirectory(ptxjit) add_subdirectory(simpleCudaGraphs) add_subdirectory(tf32TensorCoreGemm) +add_subdirectory(trimNotification) add_subdirectory(warpAggregatedAtomicsCG) diff --git a/Samples/3_CUDA_Features/trimNotification/CMakeLists.txt b/Samples/3_CUDA_Features/trimNotification/CMakeLists.txt new file mode 100644 index 000000000..391f2e1a2 --- /dev/null +++ b/Samples/3_CUDA_Features/trimNotification/CMakeLists.txt @@ -0,0 +1,36 @@ +cmake_minimum_required(VERSION 3.20) + +list(APPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_SOURCE_DIR}/../../../cmake/Modules") + +project(trimNotification LANGUAGES C CXX CUDA) + +find_package(CUDAToolkit REQUIRED) + +set(CMAKE_POSITION_INDEPENDENT_CODE ON) + +set(CMAKE_CUDA_ARCHITECTURES 75 80 86 87 89 90 100 110 120) +set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -Wno-deprecated-gpu-targets") +if(ENABLE_CUDA_DEBUG) + set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -G") # enable cuda-gdb (may significantly affect performance on some targets) +else() + set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -lineinfo") # add line information to all builds for debug tools (exclusive to -G option) +endif() + +# Include directories and libraries +include_directories(../../../Common) + +# Source file +add_executable(trimNotification trimNotification.cpp) + +target_include_directories(trimNotification PRIVATE ${CUDAToolkit_INCLUDE_DIRS}) + +target_link_libraries(trimNotification PUBLIC CUDA::cudart) + +target_compile_options(trimNotification PRIVATE $<$:--extended-lambda>) + +target_compile_features(trimNotification PRIVATE cxx_std_20 cuda_std_20) + +set_target_properties(trimNotification PROPERTIES CUDA_SEPARABLE_COMPILATION ON) + + + diff --git a/Samples/3_CUDA_Features/trimNotification/README.md b/Samples/3_CUDA_Features/trimNotification/README.md new file mode 100644 index 000000000..edf4168dd --- /dev/null +++ b/Samples/3_CUDA_Features/trimNotification/README.md @@ -0,0 +1,33 @@ +# trimNotification - Trim Notification + +## Description + +This example demonstrates use of cudaDeviceRegisterAsyncNotification to manage memory in CUDA. A GPU running in WDDM mode is required for this sample. + +## Key Concepts + +Memory Management + +## Supported SM Architectures + +## Supported OSes + +Windows + +## Supported CPU Architecture + +x86_64 + +## CUDA APIs involved + +### [CUDA Runtime API](http://docs.nvidia.com/cuda/cuda-runtime-api/index.html) +cudaSetDevice, cudaDeviceRegisterAsyncNotification, cudaDeviceUnregisterAsyncNotification, cudaMalloc, cudaFree + +## Dependencies needed to build/run + +## Prerequisites + +Download and install the [CUDA Toolkit](https://developer.nvidia.com/cuda-downloads) for your corresponding platform. +Make sure the dependencies mentioned in [Dependencies]() section above are installed. + +## References (for more details) diff --git a/Samples/3_CUDA_Features/trimNotification/trimNotification.cpp b/Samples/3_CUDA_Features/trimNotification/trimNotification.cpp new file mode 100644 index 000000000..e538689ec --- /dev/null +++ b/Samples/3_CUDA_Features/trimNotification/trimNotification.cpp @@ -0,0 +1,125 @@ +/* Copyright (c) 2025, 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: + * * Redistributions of source code must retain the above copyright + * notice, this list of conditions and the following disclaimer. + * * 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. + * * Neither the name of NVIDIA CORPORATION 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 ``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 OWNER 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. + */ + +#include + +#include +#include +#include +#include +#include + +constexpr uint32_t memblockSize = 40 * 1024 * 1024; +constexpr uint32_t stepTime = 10; // The time to wait between allocation steps in ms +constexpr uint64_t maxSteps = 500; + +// This sample uses the cudaDeviceRegisterAsyncNotification API to allocate as much video memory as it can. +// A notification to free some video memory is sent whenever the application goes over its budget. If multiple +// instances of this program are launched, all instances will likely end up with similar budgets and in +// steady-state will have similar amounts of video memory allocated. + +int main() +{ + // This structure will be our 'userData' to pass into the callback. This will give the + // callback state that it can affect without needing to use globals. + struct MemoryInfo { + std::vector memblocks; + std::mutex mutex; + bool allowedToGrow = true; + } memoryInfo; + + // The function that we will register as a callback. This could also be a standalone function + // instead of a lambda, but, since it is only referenced here, we can use a lambda to avoid + // using the global namespace. + cudaAsyncCallback basicUserCallback = [](cudaAsyncNotificationInfo_t* notificationInfo, void* userData, cudaAsyncCallbackHandle_t callback) + { + MemoryInfo* memoryInfo = (MemoryInfo*)userData; + + // Must check the type before accessing the info member of cudaAsyncNotificationInfo_t. + // Otherwise, we could misinterpret notificationInfo->info if a different type of + // notification is sent. + if (notificationInfo->type == cudaAsyncNotificationTypeOverBudget) { + printf("Asked to free %lld bytes of video memory\n", notificationInfo->info.overBudget.bytesOverBudget); + uint64_t numBlocksToFree = (notificationInfo->info.overBudget.bytesOverBudget / memblockSize) + 1; + + // The async notification will be on a separate thread, so shared data should be synchronized. + std::scoped_lock lock(memoryInfo->mutex); + memoryInfo->allowedToGrow = false; + + // Free the required number of memblocks + std::span blocksToFree(memoryInfo->memblocks.end() - numBlocksToFree, numBlocksToFree); + + for (auto& block : blocksToFree) { + cudaFree(block); + block = nullptr; + } + + // Shrink the vector to remove the freed blocks + std::erase(memoryInfo->memblocks, nullptr); + } + }; + + // Initialize CUDA + const int cudaDevice = 0; + cudaSetDevice(cudaDevice); + + // This callback handle is an opaque object that can be used to unregister the notification via + // cudaDeviceUnregisterAsyncNotification and to identify which callback registration a given + // notification corresponds to. + cudaAsyncCallbackHandle_t callback; + cudaDeviceRegisterAsyncNotification(cudaDevice, basicUserCallback, (void*)&memoryInfo, &callback); + + // Attempt to allocate a block of memory, then sleep before repeating. + uint64_t stepCounter = 0; + while (stepCounter < maxSteps) + { + { + // The async notification will be on a separate thread, so shared data should be synchronized. + std::scoped_lock lock(memoryInfo.mutex); + + if (memoryInfo.allowedToGrow) + { + int* newMemblock; + cudaError_t cudaStatus = cudaMalloc((void**)&newMemblock, memblockSize); + if (cudaStatus != cudaSuccess) { + fprintf(stderr, "cudaMalloc failed!"); + } + memoryInfo.memblocks.push_back(newMemblock); + } + } + + std::this_thread::sleep_for(std::chrono::milliseconds(stepTime)); + stepCounter++; + } + + printf("Current memory allocated: %lld MB\n", memoryInfo.memblocks.size() * memblockSize / 1024 / 1024); + + // Unregister callback handle in application cleanup + cudaDeviceUnregisterAsyncNotification(cudaDevice, callback); + + return 0; +} From 6c97408018273cda5ea1bef663502b3f9334f47a Mon Sep 17 00:00:00 2001 From: Nicholas Angle Date: Thu, 11 Dec 2025 16:21:08 -0800 Subject: [PATCH 9/9] Fix formatting with pre-commit --- .../trimNotification/CMakeLists.txt | 3 - .../trimNotification/trimNotification.cpp | 78 +++++++++---------- 2 files changed, 38 insertions(+), 43 deletions(-) diff --git a/Samples/3_CUDA_Features/trimNotification/CMakeLists.txt b/Samples/3_CUDA_Features/trimNotification/CMakeLists.txt index 391f2e1a2..d27bc2ccf 100644 --- a/Samples/3_CUDA_Features/trimNotification/CMakeLists.txt +++ b/Samples/3_CUDA_Features/trimNotification/CMakeLists.txt @@ -31,6 +31,3 @@ target_compile_options(trimNotification PRIVATE $<$:--ext target_compile_features(trimNotification PRIVATE cxx_std_20 cuda_std_20) set_target_properties(trimNotification PROPERTIES CUDA_SEPARABLE_COMPILATION ON) - - - diff --git a/Samples/3_CUDA_Features/trimNotification/trimNotification.cpp b/Samples/3_CUDA_Features/trimNotification/trimNotification.cpp index e538689ec..3033fee0f 100644 --- a/Samples/3_CUDA_Features/trimNotification/trimNotification.cpp +++ b/Samples/3_CUDA_Features/trimNotification/trimNotification.cpp @@ -25,9 +25,8 @@ * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. */ -#include - #include +#include #include #include #include @@ -39,49 +38,50 @@ constexpr uint64_t maxSteps = 500; // This sample uses the cudaDeviceRegisterAsyncNotification API to allocate as much video memory as it can. // A notification to free some video memory is sent whenever the application goes over its budget. If multiple -// instances of this program are launched, all instances will likely end up with similar budgets and in +// instances of this program are launched, all instances will likely end up with similar budgets and in // steady-state will have similar amounts of video memory allocated. int main() { // This structure will be our 'userData' to pass into the callback. This will give the // callback state that it can affect without needing to use globals. - struct MemoryInfo { - std::vector memblocks; - std::mutex mutex; - bool allowedToGrow = true; + struct MemoryInfo + { + std::vector memblocks; + std::mutex mutex; + bool allowedToGrow = true; } memoryInfo; // The function that we will register as a callback. This could also be a standalone function // instead of a lambda, but, since it is only referenced here, we can use a lambda to avoid // using the global namespace. - cudaAsyncCallback basicUserCallback = [](cudaAsyncNotificationInfo_t* notificationInfo, void* userData, cudaAsyncCallbackHandle_t callback) - { - MemoryInfo* memoryInfo = (MemoryInfo*)userData; - - // Must check the type before accessing the info member of cudaAsyncNotificationInfo_t. - // Otherwise, we could misinterpret notificationInfo->info if a different type of - // notification is sent. - if (notificationInfo->type == cudaAsyncNotificationTypeOverBudget) { - printf("Asked to free %lld bytes of video memory\n", notificationInfo->info.overBudget.bytesOverBudget); - uint64_t numBlocksToFree = (notificationInfo->info.overBudget.bytesOverBudget / memblockSize) + 1; + cudaAsyncCallback basicUserCallback = + [](cudaAsyncNotificationInfo_t *notificationInfo, void *userData, cudaAsyncCallbackHandle_t callback) { + MemoryInfo *memoryInfo = (MemoryInfo *)userData; + + // Must check the type before accessing the info member of cudaAsyncNotificationInfo_t. + // Otherwise, we could misinterpret notificationInfo->info if a different type of + // notification is sent. + if (notificationInfo->type == cudaAsyncNotificationTypeOverBudget) { + printf("Asked to free %lld bytes of video memory\n", notificationInfo->info.overBudget.bytesOverBudget); + uint64_t numBlocksToFree = (notificationInfo->info.overBudget.bytesOverBudget / memblockSize) + 1; + + // The async notification will be on a separate thread, so shared data should be synchronized. + std::scoped_lock lock(memoryInfo->mutex); + memoryInfo->allowedToGrow = false; + + // Free the required number of memblocks + std::span blocksToFree(memoryInfo->memblocks.end() - numBlocksToFree, numBlocksToFree); + + for (auto &block : blocksToFree) { + cudaFree(block); + block = nullptr; + } - // The async notification will be on a separate thread, so shared data should be synchronized. - std::scoped_lock lock(memoryInfo->mutex); - memoryInfo->allowedToGrow = false; - - // Free the required number of memblocks - std::span blocksToFree(memoryInfo->memblocks.end() - numBlocksToFree, numBlocksToFree); - - for (auto& block : blocksToFree) { - cudaFree(block); - block = nullptr; + // Shrink the vector to remove the freed blocks + std::erase(memoryInfo->memblocks, nullptr); } - - // Shrink the vector to remove the freed blocks - std::erase(memoryInfo->memblocks, nullptr); - } - }; + }; // Initialize CUDA const int cudaDevice = 0; @@ -91,27 +91,25 @@ int main() // cudaDeviceUnregisterAsyncNotification and to identify which callback registration a given // notification corresponds to. cudaAsyncCallbackHandle_t callback; - cudaDeviceRegisterAsyncNotification(cudaDevice, basicUserCallback, (void*)&memoryInfo, &callback); + cudaDeviceRegisterAsyncNotification(cudaDevice, basicUserCallback, (void *)&memoryInfo, &callback); // Attempt to allocate a block of memory, then sleep before repeating. uint64_t stepCounter = 0; - while (stepCounter < maxSteps) - { + while (stepCounter < maxSteps) { { // The async notification will be on a separate thread, so shared data should be synchronized. std::scoped_lock lock(memoryInfo.mutex); - if (memoryInfo.allowedToGrow) - { - int* newMemblock; - cudaError_t cudaStatus = cudaMalloc((void**)&newMemblock, memblockSize); + if (memoryInfo.allowedToGrow) { + int *newMemblock; + cudaError_t cudaStatus = cudaMalloc((void **)&newMemblock, memblockSize); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaMalloc failed!"); } memoryInfo.memblocks.push_back(newMemblock); } } - + std::this_thread::sleep_for(std::chrono::milliseconds(stepTime)); stepCounter++; }