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