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