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 <sstream> 11 #include <stdarg.h> 12 #include <string.h> 13 #include <hip/hiprtc.h> 14 #include "ceed-hip-common.h" 15 #include "ceed-hip-compile.h" 16 17 #define CeedChk_hiprtc(ceed, x) \ 18 do { \ 19 hiprtcResult result = static_cast<hiprtcResult>(x); \ 20 if (result != HIPRTC_SUCCESS) \ 21 return CeedError((ceed), CEED_ERROR_BACKEND, hiprtcGetErrorString(result)); \ 22 } while (0) 23 24 //------------------------------------------------------------------------------ 25 // Compile HIP kernel 26 //------------------------------------------------------------------------------ 27 int CeedCompileHip(Ceed ceed, const char *source, hipModule_t *module, 28 const CeedInt num_opts, ...) { 29 int ierr; 30 hipFree(0); // Make sure a Context exists for hiprtc 31 hiprtcProgram prog; 32 33 std::ostringstream code; 34 // Add hip runtime include statement for generation if runtime < 40400000 35 // (implies ROCm < 4.5) 36 int runtime_version; 37 CeedChk_Hip(ceed, hipRuntimeGetVersion(&runtime_version)); 38 if (runtime_version < 40400000) { 39 code << "\n#include <hip/hip_runtime.h>\n"; 40 } 41 // With ROCm 4.5, need to include these definitions specifically for hiprtc 42 // (but cannot include the runtime header) 43 else { 44 code << "#include <stddef.h>\n"; 45 code << "#define __forceinline__ inline __attribute__((always_inline))\n"; 46 code << "#define HIP_DYNAMIC_SHARED(type, var) extern __shared__ type var[];\n"; 47 } 48 49 // Macro definitions 50 // Get kernel specific options, such as kernel constants 51 const int opts_size = 3; 52 const char *opts[opts_size]; 53 if (num_opts > 0) { 54 va_list args; 55 va_start(args, num_opts); 56 char *name; 57 int val; 58 for (int i = 0; i < num_opts; 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 backend options 67 if (CEED_SCALAR_TYPE == CEED_SCALAR_FP32) { 68 code << "#define CeedScalar float\n"; 69 } 70 else { 71 code << "#define CeedScalar double\n"; 72 } 73 code << "#define CeedInt int\n"; 74 code << "#define CEED_ERROR_SUCCESS 0\n\n"; 75 76 // Non-macro options 77 opts[0] = "-default-device"; 78 struct hipDeviceProp_t prop; 79 Ceed_Hip *ceed_data; 80 ierr = CeedGetData(ceed, (void **)&ceed_data); CeedChkBackend(ierr); 81 CeedChk_Hip(ceed, hipGetDeviceProperties(&prop, ceed_data->device_id)); 82 std::string arch_arg = "--gpu-architecture=" + std::string(prop.gcnArchName); 83 opts[1] = arch_arg.c_str(); 84 opts[2] = "-munsafe-fp-atomics"; 85 86 // Add string source argument provided in call 87 code << source; 88 89 // Create Program 90 CeedChk_hiprtc(ceed, hiprtcCreateProgram(&prog, code.str().c_str(), NULL, 0, NULL, NULL)); 91 92 // Compile kernel 93 hiprtcResult result = hiprtcCompileProgram(prog, opts_size, opts); 94 if (result != HIPRTC_SUCCESS) { 95 size_t log_size; 96 CeedChk_hiprtc(ceed, hiprtcGetProgramLogSize(prog, &log_size)); 97 char *log; 98 ierr = CeedMalloc(log_size, &log); CeedChkBackend(ierr); 99 CeedChk_hiprtc(ceed, hiprtcGetProgramLog(prog, log)); 100 return CeedError(ceed, CEED_ERROR_BACKEND, "%s\n%s", 101 hiprtcGetErrorString(result), log); 102 } 103 104 size_t ptx_size; 105 CeedChk_hiprtc(ceed, hiprtcGetCodeSize(prog, &ptx_size)); 106 char *ptx; 107 ierr = CeedMalloc(ptx_size, &ptx); CeedChkBackend(ierr); 108 CeedChk_hiprtc(ceed, hiprtcGetCode(prog, ptx)); 109 CeedChk_hiprtc(ceed, hiprtcDestroyProgram(&prog)); 110 111 CeedChk_Hip(ceed, hipModuleLoadData(module, ptx)); 112 ierr = CeedFree(&ptx); CeedChkBackend(ierr); 113 114 return CEED_ERROR_SUCCESS; 115 } 116 117 //------------------------------------------------------------------------------ 118 // Get HIP kernel 119 //------------------------------------------------------------------------------ 120 int CeedGetKernelHip(Ceed ceed, hipModule_t module, const char *name, 121 hipFunction_t *kernel) { 122 123 CeedChk_Hip(ceed, hipModuleGetFunction(kernel, module, name)); 124 return CEED_ERROR_SUCCESS; 125 } 126 127 //------------------------------------------------------------------------------ 128 // Run HIP kernel 129 //------------------------------------------------------------------------------ 130 int CeedRunKernelHip(Ceed ceed, hipFunction_t kernel, const int grid_size, 131 const int block_size, void **args) { 132 CeedChk_Hip(ceed, hipModuleLaunchKernel(kernel, grid_size, 1, 1, block_size, 1, 133 1, 0, NULL, args, NULL)); 134 return CEED_ERROR_SUCCESS; 135 } 136 137 //------------------------------------------------------------------------------ 138 // Run HIP kernel for spatial dimension 139 //------------------------------------------------------------------------------ 140 int CeedRunKernelDimHip(Ceed ceed, hipFunction_t kernel, const int grid_size, 141 const int block_size_x, const int block_size_y, 142 const int block_size_z, void **args) { 143 CeedChk_Hip(ceed, hipModuleLaunchKernel(kernel, grid_size, 1, 1, 144 block_size_x, block_size_y, block_size_z, 145 0, NULL, args, NULL)); 146 return CEED_ERROR_SUCCESS; 147 } 148 149 //------------------------------------------------------------------------------ 150 // Run HIP kernel for spatial dimension with shared memory 151 //------------------------------------------------------------------------------ 152 int CeedRunKernelDimSharedHip(Ceed ceed, hipFunction_t kernel, const int grid_size, 153 const int block_size_x, const int block_size_y, 154 const int block_size_z, const int shared_mem_size, 155 void **args) { 156 CeedChk_Hip(ceed, hipModuleLaunchKernel(kernel, grid_size, 1, 1, 157 block_size_x, block_size_y, block_size_z, 158 shared_mem_size, NULL, args, NULL)); 159 return CEED_ERROR_SUCCESS; 160 } 161