1 // Copyright (c) 2017-2024, Lawrence Livermore National Security, LLC and other CEED contributors. 2 // All Rights Reserved. See the top-level LICENSE and NOTICE files for details. 3 // 4 // SPDX-License-Identifier: BSD-2-Clause 5 // 6 // This file is part of CEED: http://github.com/ceed 7 8 #include "ceed-cuda-compile.h" 9 10 #include <ceed.h> 11 #include <ceed/backend.h> 12 #include <ceed/jit-tools.h> 13 #include <cuda_runtime.h> 14 #include <nvrtc.h> 15 #include <stdarg.h> 16 #include <string.h> 17 18 #include <sstream> 19 20 #include "ceed-cuda-common.h" 21 22 #define CeedChk_Nvrtc(ceed, x) \ 23 do { \ 24 nvrtcResult result = static_cast<nvrtcResult>(x); \ 25 if (result != NVRTC_SUCCESS) return CeedError((ceed), CEED_ERROR_BACKEND, nvrtcGetErrorString(result)); \ 26 } while (0) 27 28 #define CeedCallNvrtc(ceed, ...) \ 29 do { \ 30 int ierr_q_ = __VA_ARGS__; \ 31 CeedChk_Nvrtc(ceed, ierr_q_); \ 32 } while (0) 33 34 //------------------------------------------------------------------------------ 35 // Compile CUDA kernel 36 //------------------------------------------------------------------------------ 37 int CeedCompile_Cuda(Ceed ceed, const char *source, CUmodule *module, const CeedInt num_defines, ...) { 38 size_t ptx_size; 39 char *ptx; 40 const int num_opts = 4; 41 CeedInt num_jit_source_dirs = 0, num_jit_defines = 0; 42 const char **opts; 43 nvrtcProgram prog; 44 struct cudaDeviceProp prop; 45 Ceed_Cuda *ceed_data; 46 47 cudaFree(0); // Make sure a Context exists for nvrtc 48 49 std::ostringstream code; 50 51 // Get kernel specific options, such as kernel constants 52 if (num_defines > 0) { 53 va_list args; 54 va_start(args, num_defines); 55 char *name; 56 int val; 57 58 for (int i = 0; i < num_defines; i++) { 59 name = va_arg(args, char *); 60 val = va_arg(args, int); 61 code << "#define " << name << " " << val << "\n"; 62 } 63 va_end(args); 64 } 65 66 // Standard libCEED definitions for CUDA backends 67 code << "#include <ceed/jit-source/cuda/cuda-jit.h>\n\n"; 68 69 // Non-macro options 70 CeedCallBackend(CeedCalloc(num_opts, &opts)); 71 opts[0] = "-default-device"; 72 CeedCallBackend(CeedGetData(ceed, &ceed_data)); 73 CeedCallCuda(ceed, cudaGetDeviceProperties(&prop, ceed_data->device_id)); 74 std::string arch_arg = 75 #if CUDA_VERSION >= 11010 76 // NVRTC used to support only virtual architectures through the option 77 // -arch, since it was only emitting PTX. It will now support actual 78 // architectures as well to emit SASS. 79 // https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html#dynamic-code-generation 80 "-arch=sm_" 81 #else 82 "-arch=compute_" 83 #endif 84 + std::to_string(prop.major) + std::to_string(prop.minor); 85 opts[1] = arch_arg.c_str(); 86 opts[2] = "-Dint32_t=int"; 87 opts[3] = "-DCEED_RUNNING_JIT_PASS=1"; 88 // Additional include dirs 89 { 90 const char **jit_source_dirs; 91 92 CeedCallBackend(CeedGetJitSourceRoots(ceed, &num_jit_source_dirs, &jit_source_dirs)); 93 CeedCallBackend(CeedRealloc(num_opts + num_jit_source_dirs, &opts)); 94 for (CeedInt i = 0; i < num_jit_source_dirs; i++) { 95 std::ostringstream include_dir_arg; 96 97 include_dir_arg << "-I" << jit_source_dirs[i]; 98 CeedCallBackend(CeedStringAllocCopy(include_dir_arg.str().c_str(), (char **)&opts[num_opts + i])); 99 } 100 CeedCallBackend(CeedRestoreJitSourceRoots(ceed, &jit_source_dirs)); 101 } 102 // User defines 103 { 104 const char **jit_defines; 105 106 CeedCallBackend(CeedGetJitDefines(ceed, &num_jit_defines, &jit_defines)); 107 CeedCallBackend(CeedRealloc(num_opts + num_jit_source_dirs + num_jit_defines, &opts)); 108 for (CeedInt i = 0; i < num_jit_defines; i++) { 109 std::ostringstream define_arg; 110 111 define_arg << "-D" << jit_defines[i]; 112 CeedCallBackend(CeedStringAllocCopy(define_arg.str().c_str(), (char **)&opts[num_opts + num_jit_source_dirs + i])); 113 } 114 CeedCallBackend(CeedRestoreJitDefines(ceed, &jit_defines)); 115 } 116 117 // Add string source argument provided in call 118 code << source; 119 120 // Create Program 121 CeedCallNvrtc(ceed, nvrtcCreateProgram(&prog, code.str().c_str(), NULL, 0, NULL, NULL)); 122 123 // Compile kernel 124 nvrtcResult result = nvrtcCompileProgram(prog, num_opts + num_jit_source_dirs + num_jit_defines, opts); 125 126 for (CeedInt i = 0; i < num_jit_source_dirs; i++) { 127 CeedCallBackend(CeedFree(&opts[num_opts + i])); 128 } 129 for (CeedInt i = 0; i < num_jit_defines; i++) { 130 CeedCallBackend(CeedFree(&opts[num_opts + num_jit_source_dirs + i])); 131 } 132 CeedCallBackend(CeedFree(&opts)); 133 if (result != NVRTC_SUCCESS) { 134 char *log; 135 size_t log_size; 136 137 CeedDebug256(ceed, CEED_DEBUG_COLOR_ERROR, "---------- CEED JIT SOURCE FAILED TO COMPILE ----------\n"); 138 CeedDebug(ceed, "Source:\n%s\n", code.str().c_str()); 139 CeedDebug256(ceed, CEED_DEBUG_COLOR_ERROR, "---------- CEED JIT SOURCE FAILED TO COMPILE ----------\n"); 140 CeedCallNvrtc(ceed, nvrtcGetProgramLogSize(prog, &log_size)); 141 CeedCallBackend(CeedMalloc(log_size, &log)); 142 CeedCallNvrtc(ceed, nvrtcGetProgramLog(prog, log)); 143 return CeedError(ceed, CEED_ERROR_BACKEND, "%s\n%s", nvrtcGetErrorString(result), log); 144 } 145 146 #if CUDA_VERSION >= 11010 147 CeedCallNvrtc(ceed, nvrtcGetCUBINSize(prog, &ptx_size)); 148 CeedCallBackend(CeedMalloc(ptx_size, &ptx)); 149 CeedCallNvrtc(ceed, nvrtcGetCUBIN(prog, ptx)); 150 #else 151 CeedCallNvrtc(ceed, nvrtcGetPTXSize(prog, &ptx_size)); 152 CeedCallBackend(CeedMalloc(ptx_size, &ptx)); 153 CeedCallNvrtc(ceed, nvrtcGetPTX(prog, ptx)); 154 #endif 155 CeedCallNvrtc(ceed, nvrtcDestroyProgram(&prog)); 156 157 CeedCallCuda(ceed, cuModuleLoadData(module, ptx)); 158 CeedCallBackend(CeedFree(&ptx)); 159 return CEED_ERROR_SUCCESS; 160 } 161 162 //------------------------------------------------------------------------------ 163 // Get CUDA kernel 164 //------------------------------------------------------------------------------ 165 int CeedGetKernel_Cuda(Ceed ceed, CUmodule module, const char *name, CUfunction *kernel) { 166 CeedCallCuda(ceed, cuModuleGetFunction(kernel, module, name)); 167 return CEED_ERROR_SUCCESS; 168 } 169 170 //------------------------------------------------------------------------------ 171 // Run CUDA kernel with block size selected automatically based on the kernel 172 // (which may use enough registers to require a smaller block size than the 173 // hardware is capable) 174 //------------------------------------------------------------------------------ 175 int CeedRunKernelAutoblockCuda(Ceed ceed, CUfunction kernel, size_t points, void **args) { 176 int min_grid_size, max_block_size; 177 178 CeedCallCuda(ceed, cuOccupancyMaxPotentialBlockSize(&min_grid_size, &max_block_size, kernel, NULL, 0, 0x10000)); 179 CeedCallBackend(CeedRunKernel_Cuda(ceed, kernel, CeedDivUpInt(points, max_block_size), max_block_size, args)); 180 return CEED_ERROR_SUCCESS; 181 } 182 183 //------------------------------------------------------------------------------ 184 // Run CUDA kernel 185 //------------------------------------------------------------------------------ 186 int CeedRunKernel_Cuda(Ceed ceed, CUfunction kernel, const int grid_size, const int block_size, void **args) { 187 CeedCallBackend(CeedRunKernelDimShared_Cuda(ceed, kernel, grid_size, block_size, 1, 1, 0, args)); 188 return CEED_ERROR_SUCCESS; 189 } 190 191 //------------------------------------------------------------------------------ 192 // Run CUDA kernel for spatial dimension 193 //------------------------------------------------------------------------------ 194 int CeedRunKernelDim_Cuda(Ceed ceed, CUfunction kernel, const int grid_size, const int block_size_x, const int block_size_y, const int block_size_z, 195 void **args) { 196 CeedCallBackend(CeedRunKernelDimShared_Cuda(ceed, kernel, grid_size, block_size_x, block_size_y, block_size_z, 0, args)); 197 return CEED_ERROR_SUCCESS; 198 } 199 200 //------------------------------------------------------------------------------ 201 // Run CUDA kernel for spatial dimension with shared memory 202 //------------------------------------------------------------------------------ 203 int CeedRunKernelDimShared_Cuda(Ceed ceed, CUfunction kernel, const int grid_size, const int block_size_x, const int block_size_y, 204 const int block_size_z, const int shared_mem_size, void **args) { 205 #if CUDA_VERSION >= 9000 206 cuFuncSetAttribute(kernel, CU_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES, shared_mem_size); 207 #endif 208 CUresult result = cuLaunchKernel(kernel, grid_size, 1, 1, block_size_x, block_size_y, block_size_z, shared_mem_size, NULL, args, NULL); 209 210 if (result == CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES) { 211 int max_threads_per_block, shared_size_bytes, num_regs; 212 213 cuFuncGetAttribute(&max_threads_per_block, CU_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK, kernel); 214 cuFuncGetAttribute(&shared_size_bytes, CU_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES, kernel); 215 cuFuncGetAttribute(&num_regs, CU_FUNC_ATTRIBUTE_NUM_REGS, kernel); 216 return CeedError(ceed, CEED_ERROR_BACKEND, 217 "CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES: max_threads_per_block %d on block size (%d,%d,%d), shared_size %d, num_regs %d", 218 max_threads_per_block, block_size_x, block_size_y, block_size_z, shared_size_bytes, num_regs); 219 } else CeedChk_Cu(ceed, result); 220 return CEED_ERROR_SUCCESS; 221 } 222 223 //------------------------------------------------------------------------------ 224