1 // Copyright (c) 2017-2022, 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/ceed.h> 9 #include <ceed/backend.h> 10 #include <ceed/jit-tools.h> 11 #include <cuda.h> 12 #include <cuda_runtime.h> 13 #include <nvrtc.h> 14 #include <sstream> 15 #include <stdarg.h> 16 #include <string.h> 17 #include "ceed-cuda-common.h" 18 #include "ceed-cuda-compile.h" 19 20 #define CeedChk_Nvrtc(ceed, x) \ 21 do { \ 22 nvrtcResult result = static_cast<nvrtcResult>(x); \ 23 if (result != NVRTC_SUCCESS) \ 24 return CeedError((ceed), CEED_ERROR_BACKEND, nvrtcGetErrorString(result)); \ 25 } while (0) 26 27 //------------------------------------------------------------------------------ 28 // Compile CUDA kernel 29 //------------------------------------------------------------------------------ 30 int CeedCompileCuda(Ceed ceed, const char *source, CUmodule *module, 31 const CeedInt num_defines, ...) { 32 int ierr; 33 cudaFree(0); // Make sure a Context exists for nvrtc 34 nvrtcProgram prog; 35 36 std::ostringstream code; 37 38 // Get kernel specific options, such as kernel constants 39 if (num_defines > 0) { 40 va_list args; 41 va_start(args, num_defines); 42 char *name; 43 int val; 44 for (int i = 0; i < num_defines; i++) { 45 name = va_arg(args, char *); 46 val = va_arg(args, int); 47 code << "#define " << name << " " << val << "\n"; 48 } 49 va_end(args); 50 } 51 52 // Standard libCEED definitions for CUDA backends 53 char *jit_defs_path, *jit_defs_source; 54 ierr = CeedGetJitAbsolutePath(ceed, 55 "ceed/jit-source/cuda/cuda-jit.h", 56 &jit_defs_path); CeedChkBackend(ierr); 57 ierr = CeedLoadSourceToBuffer(ceed, jit_defs_path, &jit_defs_source); 58 CeedChkBackend(ierr); 59 code << jit_defs_source; 60 code << "\n\n"; 61 ierr = CeedFree(&jit_defs_path); CeedChkBackend(ierr); 62 ierr = CeedFree(&jit_defs_source); CeedChkBackend(ierr); 63 64 // Non-macro options 65 const int num_opts = 3; 66 const char *opts[num_opts]; 67 opts[0] = "-default-device"; 68 struct cudaDeviceProp prop; 69 Ceed_Cuda *ceed_data; 70 ierr = CeedGetData(ceed, &ceed_data); CeedChkBackend(ierr); 71 ierr = cudaGetDeviceProperties(&prop, ceed_data->device_id); 72 CeedChk_Cu(ceed, ierr); 73 std::string arch_arg = "-arch=compute_" + std::to_string(prop.major) + std::to_string(prop.minor); 74 opts[1] = arch_arg.c_str(); 75 opts[2] = "-Dint32_t=int"; 76 77 // Add string source argument provided in call 78 code << source; 79 80 // Create Program 81 CeedChk_Nvrtc(ceed, nvrtcCreateProgram(&prog, code.str().c_str(), NULL, 0, NULL, NULL)); 82 83 // Compile kernel 84 nvrtcResult result = nvrtcCompileProgram(prog, num_opts, opts); 85 if (result != NVRTC_SUCCESS) { 86 size_t log_size; 87 CeedChk_Nvrtc(ceed, nvrtcGetProgramLogSize(prog, &log_size)); 88 char *log; 89 ierr = CeedMalloc(log_size, &log); CeedChkBackend(ierr); 90 CeedChk_Nvrtc(ceed, nvrtcGetProgramLog(prog, log)); 91 return CeedError(ceed, CEED_ERROR_BACKEND, "%s\n%s", 92 nvrtcGetErrorString(result), log); 93 } 94 95 size_t ptx_size; 96 CeedChk_Nvrtc(ceed, nvrtcGetPTXSize(prog, &ptx_size)); 97 char *ptx; 98 ierr = CeedMalloc(ptx_size, &ptx); CeedChkBackend(ierr); 99 CeedChk_Nvrtc(ceed, nvrtcGetPTX(prog, ptx)); 100 CeedChk_Nvrtc(ceed, nvrtcDestroyProgram(&prog)); 101 102 CeedChk_Cu(ceed, cuModuleLoadData(module, ptx)); 103 ierr = CeedFree(&ptx); CeedChkBackend(ierr); 104 return CEED_ERROR_SUCCESS; 105 } 106 107 //------------------------------------------------------------------------------ 108 // Get CUDA kernel 109 //------------------------------------------------------------------------------ 110 int CeedGetKernelCuda(Ceed ceed, CUmodule module, const char *name, 111 CUfunction *kernel) { 112 CeedChk_Cu(ceed, cuModuleGetFunction(kernel, module, name)); 113 return CEED_ERROR_SUCCESS; 114 } 115 116 // Run kernel with block size selected automatically based on the kernel (which 117 // may use enough registers to require a smaller block size than the hardware is 118 // capable). 119 int CeedRunKernelAutoblockCuda(Ceed ceed, CUfunction kernel, size_t points, 120 void **args) { 121 int min_grid_size, max_block_size; 122 CeedChk_Cu(ceed, cuOccupancyMaxPotentialBlockSize(&min_grid_size, 123 &max_block_size, kernel, NULL, 0, 0x10000)); 124 CeedChkBackend(CeedRunKernelCuda(ceed, kernel, CeedDivUpInt(points, 125 max_block_size), max_block_size, args)); 126 return 0; 127 } 128 129 //------------------------------------------------------------------------------ 130 // Run CUDA kernel 131 //------------------------------------------------------------------------------ 132 int CeedRunKernelCuda(Ceed ceed, CUfunction kernel, const int grid_size, 133 const int block_size, void **args) { 134 CeedChkBackend(CeedRunKernelDimSharedCuda(ceed, kernel, grid_size, 135 block_size, 1, 1, 0, args)); 136 return CEED_ERROR_SUCCESS; 137 } 138 139 //------------------------------------------------------------------------------ 140 // Run CUDA kernel for spatial dimension 141 //------------------------------------------------------------------------------ 142 int CeedRunKernelDimCuda(Ceed ceed, CUfunction kernel, const int grid_size, 143 const int block_size_x, const int block_size_y, 144 const int block_size_z, void **args) { 145 CeedChkBackend(CeedRunKernelDimSharedCuda(ceed, kernel, grid_size, 146 block_size_x, block_size_y, block_size_z, 147 0, args)); 148 return CEED_ERROR_SUCCESS; 149 } 150 151 //------------------------------------------------------------------------------ 152 // Run CUDA kernel for spatial dimension with sharde memory 153 //------------------------------------------------------------------------------ 154 int CeedRunKernelDimSharedCuda(Ceed ceed, CUfunction kernel, 155 const int grid_size, const int block_size_x, 156 const int block_size_y, const int block_size_z, 157 const int shared_mem_size, void **args) { 158 CUresult result = cuLaunchKernel(kernel, grid_size, 1, 1, 159 block_size_x, block_size_y, block_size_z, 160 shared_mem_size, NULL, args, NULL); 161 if (result == CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES) { 162 int max_threads_per_block, shared_size_bytes, num_regs; 163 cuFuncGetAttribute(&max_threads_per_block, 164 CU_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK, kernel); 165 cuFuncGetAttribute(&shared_size_bytes, CU_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES, 166 kernel); 167 cuFuncGetAttribute(&num_regs, CU_FUNC_ATTRIBUTE_NUM_REGS, kernel); 168 return CeedError(ceed, CEED_ERROR_BACKEND, 169 "CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES: max_threads_per_block %d on block size (%d,%d,%d), shared_size %d, num_regs %d", 170 max_threads_per_block, block_size_x, block_size_y, block_size_z, 171 shared_size_bytes, num_regs); 172 } else CeedChk_Cu(ceed, result); 173 return CEED_ERROR_SUCCESS; 174 } 175 176 //------------------------------------------------------------------------------ 177