1 // Copyright (c) 2017-2018, Lawrence Livermore National Security, LLC. 2 // Produced at the Lawrence Livermore National Laboratory. LLNL-CODE-734707. 3 // All Rights reserved. See files LICENSE and NOTICE for details. 4 // 5 // This file is part of CEED, a collection of benchmarks, miniapps, software 6 // libraries and APIs for efficient high-order finite element and spectral 7 // element discretizations for exascale applications. For more information and 8 // source code availability see http://github.com/ceed. 9 // 10 // The CEED research is supported by the Exascale Computing Project 17-SC-20-SC, 11 // a collaborative effort of two U.S. Department of Energy organizations (Office 12 // of Science and the National Nuclear Security Administration) responsible for 13 // the planning and preparation of a capable exascale ecosystem, including 14 // software, applications, hardware, advanced system engineering and early 15 // testbed platforms, in support of the nation's exascale computing imperative. 16 17 #include <ceed/ceed.h> 18 #include <ceed/backend.h> 19 #include <sstream> 20 #include <stdarg.h> 21 #include <string.h> 22 #include <hip/hiprtc.h> 23 #include "ceed-hip-common.h" 24 #include "ceed-hip-compile.h" 25 26 #define CeedChk_hiprtc(ceed, x) \ 27 do { \ 28 hiprtcResult result = static_cast<hiprtcResult>(x); \ 29 if (result != HIPRTC_SUCCESS) \ 30 return CeedError((ceed), CEED_ERROR_BACKEND, hiprtcGetErrorString(result)); \ 31 } while (0) 32 33 //------------------------------------------------------------------------------ 34 // Compile HIP kernel 35 //------------------------------------------------------------------------------ 36 int CeedCompileHip(Ceed ceed, const char *source, hipModule_t *module, 37 const CeedInt num_opts, ...) { 38 int ierr; 39 hipFree(0); // Make sure a Context exists for hiprtc 40 hiprtcProgram prog; 41 42 std::ostringstream code; 43 // Add hip runtime include statement for generation if runtime < 40400000 44 // (implies ROCm < 4.5) 45 int runtime_version; 46 CeedChk_Hip(ceed, hipRuntimeGetVersion(&runtime_version)); 47 if (runtime_version < 40400000) { 48 code << "\n#include <hip/hip_runtime.h>\n"; 49 } 50 // With ROCm 4.5, need to include these definitions specifically for hiprtc 51 // (but cannot include the runtime header) 52 else { 53 code << "#include <stddef.h>\n"; 54 code << "#define __forceinline__ inline __attribute__((always_inline))\n"; 55 code << "#define HIP_DYNAMIC_SHARED(type, var) extern __shared__ type var[];\n"; 56 } 57 58 // Macro definitions 59 // Get kernel specific options, such as kernel constants 60 const int opts_size = 3; 61 const char *opts[opts_size]; 62 if (num_opts > 0) { 63 va_list args; 64 va_start(args, num_opts); 65 char *name; 66 int val; 67 for (int i = 0; i < num_opts; i++) { 68 name = va_arg(args, char *); 69 val = va_arg(args, int); 70 code << "#define " << name << " " << val << "\n"; 71 } 72 va_end(args); 73 } 74 75 // Standard backend options 76 if (CEED_SCALAR_TYPE == CEED_SCALAR_FP32) { 77 code << "#define CeedScalar float\n"; 78 } 79 else { 80 code << "#define CeedScalar double\n"; 81 } 82 code << "#define CeedInt int\n"; 83 code << "#define CEED_ERROR_SUCCESS 0\n\n"; 84 85 // Non-macro options 86 opts[0] = "-default-device"; 87 struct hipDeviceProp_t prop; 88 Ceed_Hip *ceed_data; 89 ierr = CeedGetData(ceed, (void **)&ceed_data); CeedChkBackend(ierr); 90 CeedChk_Hip(ceed, hipGetDeviceProperties(&prop, ceed_data->device_id)); 91 std::string arch_arg = "--gpu-architecture=" + std::string(prop.gcnArchName); 92 opts[1] = arch_arg.c_str(); 93 opts[2] = "-munsafe-fp-atomics"; 94 95 // Add string source argument provided in call 96 code << source; 97 98 // Create Program 99 CeedChk_hiprtc(ceed, hiprtcCreateProgram(&prog, code.str().c_str(), NULL, 0, NULL, NULL)); 100 101 // Compile kernel 102 hiprtcResult result = hiprtcCompileProgram(prog, opts_size, opts); 103 if (result != HIPRTC_SUCCESS) { 104 size_t log_size; 105 CeedChk_hiprtc(ceed, hiprtcGetProgramLogSize(prog, &log_size)); 106 char *log; 107 ierr = CeedMalloc(log_size, &log); CeedChkBackend(ierr); 108 CeedChk_hiprtc(ceed, hiprtcGetProgramLog(prog, log)); 109 return CeedError(ceed, CEED_ERROR_BACKEND, "%s\n%s", 110 hiprtcGetErrorString(result), log); 111 } 112 113 size_t ptx_size; 114 CeedChk_hiprtc(ceed, hiprtcGetCodeSize(prog, &ptx_size)); 115 char *ptx; 116 ierr = CeedMalloc(ptx_size, &ptx); CeedChkBackend(ierr); 117 CeedChk_hiprtc(ceed, hiprtcGetCode(prog, ptx)); 118 CeedChk_hiprtc(ceed, hiprtcDestroyProgram(&prog)); 119 120 CeedChk_Hip(ceed, hipModuleLoadData(module, ptx)); 121 ierr = CeedFree(&ptx); CeedChkBackend(ierr); 122 123 return CEED_ERROR_SUCCESS; 124 } 125 126 //------------------------------------------------------------------------------ 127 // Get HIP kernel 128 //------------------------------------------------------------------------------ 129 int CeedGetKernelHip(Ceed ceed, hipModule_t module, const char *name, 130 hipFunction_t *kernel) { 131 132 CeedChk_Hip(ceed, hipModuleGetFunction(kernel, module, name)); 133 return CEED_ERROR_SUCCESS; 134 } 135 136 //------------------------------------------------------------------------------ 137 // Run HIP kernel 138 //------------------------------------------------------------------------------ 139 int CeedRunKernelHip(Ceed ceed, hipFunction_t kernel, const int grid_size, 140 const int block_size, void **args) { 141 CeedChk_Hip(ceed, hipModuleLaunchKernel(kernel, grid_size, 1, 1, block_size, 1, 142 1, 0, NULL, args, NULL)); 143 return CEED_ERROR_SUCCESS; 144 } 145 146 //------------------------------------------------------------------------------ 147 // Run HIP kernel for spatial dimension 148 //------------------------------------------------------------------------------ 149 int CeedRunKernelDimHip(Ceed ceed, hipFunction_t kernel, const int grid_size, 150 const int block_size_x, const int block_size_y, 151 const int block_size_z, void **args) { 152 CeedChk_Hip(ceed, hipModuleLaunchKernel(kernel, grid_size, 1, 1, 153 block_size_x, block_size_y, block_size_z, 154 0, NULL, args, NULL)); 155 return CEED_ERROR_SUCCESS; 156 } 157 158 //------------------------------------------------------------------------------ 159 // Run HIP kernel for spatial dimension with shared memory 160 //------------------------------------------------------------------------------ 161 int CeedRunKernelDimSharedHip(Ceed ceed, hipFunction_t kernel, const int grid_size, 162 const int block_size_x, const int block_size_y, 163 const int block_size_z, const int shared_mem_size, 164 void **args) { 165 CeedChk_Hip(ceed, hipModuleLaunchKernel(kernel, grid_size, 1, 1, 166 block_size_x, block_size_y, block_size_z, 167 shared_mem_size, NULL, args, NULL)); 168 return CEED_ERROR_SUCCESS; 169 } 170