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.h> 9 #include <ceed/backend.h> 10 #include <ceed/jit-tools.h> 11 #include <hip/hip_runtime.h> 12 13 #include "../hip/ceed-hip-common.h" 14 #include "../hip/ceed-hip-compile.h" 15 #include "ceed-hip-ref.h" 16 17 //------------------------------------------------------------------------------ 18 // Basis apply - tensor 19 //------------------------------------------------------------------------------ 20 int CeedBasisApply_Hip(CeedBasis basis, const CeedInt num_elem, CeedTransposeMode t_mode, CeedEvalMode eval_mode, CeedVector u, CeedVector v) { 21 Ceed ceed; 22 CeedCallBackend(CeedBasisGetCeed(basis, &ceed)); 23 Ceed_Hip *ceed_Hip; 24 CeedCallBackend(CeedGetData(ceed, &ceed_Hip)); 25 CeedBasis_Hip *data; 26 CeedCallBackend(CeedBasisGetData(basis, &data)); 27 const CeedInt transpose = t_mode == CEED_TRANSPOSE; 28 const int max_block_size = 64; 29 30 // Read vectors 31 const CeedScalar *d_u; 32 CeedScalar *d_v; 33 if (eval_mode != CEED_EVAL_WEIGHT) { 34 CeedCallBackend(CeedVectorGetArrayRead(u, CEED_MEM_DEVICE, &d_u)); 35 } 36 CeedCallBackend(CeedVectorGetArrayWrite(v, CEED_MEM_DEVICE, &d_v)); 37 38 // Clear v for transpose operation 39 if (t_mode == CEED_TRANSPOSE) { 40 CeedSize length; 41 CeedCallBackend(CeedVectorGetLength(v, &length)); 42 CeedCallHip(ceed, hipMemset(d_v, 0, length * sizeof(CeedScalar))); 43 } 44 45 // Basis action 46 switch (eval_mode) { 47 case CEED_EVAL_INTERP: { 48 void *interp_args[] = {(void *)&num_elem, (void *)&transpose, &data->d_interp_1d, &d_u, &d_v}; 49 CeedInt Q_1d, dim; 50 CeedCallBackend(CeedBasisGetNumQuadraturePoints1D(basis, &Q_1d)); 51 CeedCallBackend(CeedBasisGetDimension(basis, &dim)); 52 CeedInt block_size = CeedIntMin(CeedIntPow(Q_1d, dim), max_block_size); 53 54 CeedCallBackend(CeedRunKernelHip(ceed, data->Interp, num_elem, block_size, interp_args)); 55 } break; 56 case CEED_EVAL_GRAD: { 57 void *grad_args[] = {(void *)&num_elem, (void *)&transpose, &data->d_interp_1d, &data->d_grad_1d, &d_u, &d_v}; 58 CeedInt block_size = max_block_size; 59 60 CeedCallBackend(CeedRunKernelHip(ceed, data->Grad, num_elem, block_size, grad_args)); 61 } break; 62 case CEED_EVAL_WEIGHT: { 63 void *weight_args[] = {(void *)&num_elem, (void *)&data->d_q_weight_1d, &d_v}; 64 const int block_size = 64; 65 int grid_size = num_elem / block_size; 66 if (block_size * grid_size < num_elem) grid_size += 1; 67 68 CeedCallBackend(CeedRunKernelHip(ceed, data->Weight, grid_size, block_size, weight_args)); 69 } break; 70 // LCOV_EXCL_START 71 // Evaluate the divergence to/from the quadrature points 72 case CEED_EVAL_DIV: 73 return CeedError(ceed, CEED_ERROR_BACKEND, "CEED_EVAL_DIV not supported"); 74 // Evaluate the curl to/from the quadrature points 75 case CEED_EVAL_CURL: 76 return CeedError(ceed, CEED_ERROR_BACKEND, "CEED_EVAL_CURL not supported"); 77 // Take no action, BasisApply should not have been called 78 case CEED_EVAL_NONE: 79 return CeedError(ceed, CEED_ERROR_BACKEND, "CEED_EVAL_NONE does not make sense in this context"); 80 // LCOV_EXCL_STOP 81 } 82 83 // Restore vectors 84 if (eval_mode != CEED_EVAL_WEIGHT) { 85 CeedCallBackend(CeedVectorRestoreArrayRead(u, &d_u)); 86 } 87 CeedCallBackend(CeedVectorRestoreArray(v, &d_v)); 88 return CEED_ERROR_SUCCESS; 89 } 90 91 //------------------------------------------------------------------------------ 92 // Basis apply - non-tensor 93 //------------------------------------------------------------------------------ 94 int CeedBasisApplyNonTensor_Hip(CeedBasis basis, const CeedInt num_elem, CeedTransposeMode t_mode, CeedEvalMode eval_mode, CeedVector u, 95 CeedVector v) { 96 Ceed ceed; 97 CeedCallBackend(CeedBasisGetCeed(basis, &ceed)); 98 Ceed_Hip *ceed_Hip; 99 CeedCallBackend(CeedGetData(ceed, &ceed_Hip)); 100 CeedBasisNonTensor_Hip *data; 101 CeedCallBackend(CeedBasisGetData(basis, &data)); 102 CeedInt num_nodes, num_qpts; 103 CeedCallBackend(CeedBasisGetNumQuadraturePoints(basis, &num_qpts)); 104 CeedCallBackend(CeedBasisGetNumNodes(basis, &num_nodes)); 105 const CeedInt transpose = t_mode == CEED_TRANSPOSE; 106 int elemsPerBlock = 1; 107 int grid = num_elem / elemsPerBlock + ((num_elem / elemsPerBlock * elemsPerBlock < num_elem) ? 1 : 0); 108 109 // Read vectors 110 const CeedScalar *d_u; 111 CeedScalar *d_v; 112 if (eval_mode != CEED_EVAL_WEIGHT) { 113 CeedCallBackend(CeedVectorGetArrayRead(u, CEED_MEM_DEVICE, &d_u)); 114 } 115 CeedCallBackend(CeedVectorGetArrayWrite(v, CEED_MEM_DEVICE, &d_v)); 116 117 // Clear v for transpose operation 118 if (t_mode == CEED_TRANSPOSE) { 119 CeedSize length; 120 CeedCallBackend(CeedVectorGetLength(v, &length)); 121 CeedCallHip(ceed, hipMemset(d_v, 0, length * sizeof(CeedScalar))); 122 } 123 124 // Apply basis operation 125 switch (eval_mode) { 126 case CEED_EVAL_INTERP: { 127 void *interp_args[] = {(void *)&num_elem, (void *)&transpose, &data->d_interp, &d_u, &d_v}; 128 const int block_size_x = transpose ? num_nodes : num_qpts; 129 CeedCallBackend(CeedRunKernelDimHip(ceed, data->Interp, grid, block_size_x, 1, elemsPerBlock, interp_args)); 130 } break; 131 case CEED_EVAL_GRAD: { 132 void *grad_args[] = {(void *)&num_elem, (void *)&transpose, &data->d_grad, &d_u, &d_v}; 133 const int block_size_x = transpose ? num_nodes : num_qpts; 134 CeedCallBackend(CeedRunKernelDimHip(ceed, data->Grad, grid, block_size_x, 1, elemsPerBlock, grad_args)); 135 } break; 136 case CEED_EVAL_WEIGHT: { 137 void *weight_args[] = {(void *)&num_elem, (void *)&data->d_q_weight, &d_v}; 138 CeedCallBackend(CeedRunKernelDimHip(ceed, data->Weight, grid, num_qpts, 1, elemsPerBlock, weight_args)); 139 } break; 140 // LCOV_EXCL_START 141 // Evaluate the divergence to/from the quadrature points 142 case CEED_EVAL_DIV: 143 return CeedError(ceed, CEED_ERROR_BACKEND, "CEED_EVAL_DIV not supported"); 144 // Evaluate the curl to/from the quadrature points 145 case CEED_EVAL_CURL: 146 return CeedError(ceed, CEED_ERROR_BACKEND, "CEED_EVAL_CURL not supported"); 147 // Take no action, BasisApply should not have been called 148 case CEED_EVAL_NONE: 149 return CeedError(ceed, CEED_ERROR_BACKEND, "CEED_EVAL_NONE does not make sense in this context"); 150 // LCOV_EXCL_STOP 151 } 152 153 // Restore vectors 154 if (eval_mode != CEED_EVAL_WEIGHT) { 155 CeedCallBackend(CeedVectorRestoreArrayRead(u, &d_u)); 156 } 157 CeedCallBackend(CeedVectorRestoreArray(v, &d_v)); 158 return CEED_ERROR_SUCCESS; 159 } 160 161 //------------------------------------------------------------------------------ 162 // Destroy tensor basis 163 //------------------------------------------------------------------------------ 164 static int CeedBasisDestroy_Hip(CeedBasis basis) { 165 Ceed ceed; 166 CeedCallBackend(CeedBasisGetCeed(basis, &ceed)); 167 168 CeedBasis_Hip *data; 169 CeedCallBackend(CeedBasisGetData(basis, &data)); 170 171 CeedCallHip(ceed, hipModuleUnload(data->module)); 172 173 CeedCallHip(ceed, hipFree(data->d_q_weight_1d)); 174 CeedCallHip(ceed, hipFree(data->d_interp_1d)); 175 CeedCallHip(ceed, hipFree(data->d_grad_1d)); 176 CeedCallBackend(CeedFree(&data)); 177 178 return CEED_ERROR_SUCCESS; 179 } 180 181 //------------------------------------------------------------------------------ 182 // Destroy non-tensor basis 183 //------------------------------------------------------------------------------ 184 static int CeedBasisDestroyNonTensor_Hip(CeedBasis basis) { 185 Ceed ceed; 186 CeedCallBackend(CeedBasisGetCeed(basis, &ceed)); 187 188 CeedBasisNonTensor_Hip *data; 189 CeedCallBackend(CeedBasisGetData(basis, &data)); 190 191 CeedCallHip(ceed, hipModuleUnload(data->module)); 192 193 CeedCallHip(ceed, hipFree(data->d_q_weight)); 194 CeedCallHip(ceed, hipFree(data->d_interp)); 195 CeedCallHip(ceed, hipFree(data->d_grad)); 196 CeedCallBackend(CeedFree(&data)); 197 198 return CEED_ERROR_SUCCESS; 199 } 200 201 //------------------------------------------------------------------------------ 202 // Create tensor 203 //------------------------------------------------------------------------------ 204 int CeedBasisCreateTensorH1_Hip(CeedInt dim, CeedInt P_1d, CeedInt Q_1d, const CeedScalar *interp_1d, const CeedScalar *grad_1d, 205 const CeedScalar *qref1d, const CeedScalar *q_weight_1d, CeedBasis basis) { 206 Ceed ceed; 207 CeedCallBackend(CeedBasisGetCeed(basis, &ceed)); 208 CeedBasis_Hip *data; 209 CeedCallBackend(CeedCalloc(1, &data)); 210 211 // Copy data to GPU 212 const CeedInt q_bytes = Q_1d * sizeof(CeedScalar); 213 CeedCallHip(ceed, hipMalloc((void **)&data->d_q_weight_1d, q_bytes)); 214 CeedCallHip(ceed, hipMemcpy(data->d_q_weight_1d, q_weight_1d, q_bytes, hipMemcpyHostToDevice)); 215 216 const CeedInt interp_bytes = q_bytes * P_1d; 217 CeedCallHip(ceed, hipMalloc((void **)&data->d_interp_1d, interp_bytes)); 218 CeedCallHip(ceed, hipMemcpy(data->d_interp_1d, interp_1d, interp_bytes, hipMemcpyHostToDevice)); 219 220 CeedCallHip(ceed, hipMalloc((void **)&data->d_grad_1d, interp_bytes)); 221 CeedCallHip(ceed, hipMemcpy(data->d_grad_1d, grad_1d, interp_bytes, hipMemcpyHostToDevice)); 222 223 // Compile basis kernels 224 CeedInt ncomp; 225 CeedCallBackend(CeedBasisGetNumComponents(basis, &ncomp)); 226 char *basis_kernel_path, *basis_kernel_source; 227 CeedCallBackend(CeedGetJitAbsolutePath(ceed, "ceed/jit-source/hip/hip-ref-basis-tensor.h", &basis_kernel_path)); 228 CeedDebug256(ceed, 2, "----- Loading Basis Kernel Source -----\n"); 229 CeedCallBackend(CeedLoadSourceToBuffer(ceed, basis_kernel_path, &basis_kernel_source)); 230 CeedDebug256(ceed, 2, "----- Loading Basis Kernel Source Complete! -----\n"); 231 CeedCallBackend(CeedCompileHip(ceed, basis_kernel_source, &data->module, 7, "BASIS_Q_1D", Q_1d, "BASIS_P_1D", P_1d, "BASIS_BUF_LEN", 232 ncomp * CeedIntPow(Q_1d > P_1d ? Q_1d : P_1d, dim), "BASIS_DIM", dim, "BASIS_NUM_COMP", ncomp, "BASIS_NUM_NODES", 233 CeedIntPow(P_1d, dim), "BASIS_NUM_QPTS", CeedIntPow(Q_1d, dim))); 234 CeedCallBackend(CeedGetKernelHip(ceed, data->module, "Interp", &data->Interp)); 235 CeedCallBackend(CeedGetKernelHip(ceed, data->module, "Grad", &data->Grad)); 236 CeedCallBackend(CeedGetKernelHip(ceed, data->module, "Weight", &data->Weight)); 237 CeedCallBackend(CeedFree(&basis_kernel_path)); 238 CeedCallBackend(CeedFree(&basis_kernel_source)); 239 240 CeedCallBackend(CeedBasisSetData(basis, data)); 241 242 CeedCallBackend(CeedSetBackendFunction(ceed, "Basis", basis, "Apply", CeedBasisApply_Hip)); 243 CeedCallBackend(CeedSetBackendFunction(ceed, "Basis", basis, "Destroy", CeedBasisDestroy_Hip)); 244 return CEED_ERROR_SUCCESS; 245 } 246 247 //------------------------------------------------------------------------------ 248 // Create non-tensor 249 //------------------------------------------------------------------------------ 250 int CeedBasisCreateH1_Hip(CeedElemTopology topo, CeedInt dim, CeedInt num_nodes, CeedInt num_qpts, const CeedScalar *interp, const CeedScalar *grad, 251 const CeedScalar *qref, const CeedScalar *q_weight, CeedBasis basis) { 252 Ceed ceed; 253 CeedCallBackend(CeedBasisGetCeed(basis, &ceed)); 254 CeedBasisNonTensor_Hip *data; 255 CeedCallBackend(CeedCalloc(1, &data)); 256 257 // Copy basis data to GPU 258 const CeedInt q_bytes = num_qpts * sizeof(CeedScalar); 259 CeedCallHip(ceed, hipMalloc((void **)&data->d_q_weight, q_bytes)); 260 CeedCallHip(ceed, hipMemcpy(data->d_q_weight, q_weight, q_bytes, hipMemcpyHostToDevice)); 261 262 const CeedInt interp_bytes = q_bytes * num_nodes; 263 CeedCallHip(ceed, hipMalloc((void **)&data->d_interp, interp_bytes)); 264 CeedCallHip(ceed, hipMemcpy(data->d_interp, interp, interp_bytes, hipMemcpyHostToDevice)); 265 266 const CeedInt grad_bytes = q_bytes * num_nodes * dim; 267 CeedCallHip(ceed, hipMalloc((void **)&data->d_grad, grad_bytes)); 268 CeedCallHip(ceed, hipMemcpy(data->d_grad, grad, grad_bytes, hipMemcpyHostToDevice)); 269 270 // Compile basis kernels 271 CeedInt ncomp; 272 CeedCallBackend(CeedBasisGetNumComponents(basis, &ncomp)); 273 char *basis_kernel_path, *basis_kernel_source; 274 CeedCallBackend(CeedGetJitAbsolutePath(ceed, "ceed/jit-source/hip/hip-ref-basis-nontensor.h", &basis_kernel_path)); 275 CeedDebug256(ceed, 2, "----- Loading Basis Kernel Source -----\n"); 276 CeedCallBackend(CeedLoadSourceToBuffer(ceed, basis_kernel_path, &basis_kernel_source)); 277 CeedDebug256(ceed, 2, "----- Loading Basis Kernel Source Complete! -----\n"); 278 CeedCallBackend(CeedCompileHip(ceed, basis_kernel_source, &data->module, 4, "BASIS_Q", num_qpts, "BASIS_P", num_nodes, "BASIS_DIM", dim, 279 "BASIS_NUM_COMP", ncomp)); 280 CeedCallBackend(CeedGetKernelHip(ceed, data->module, "Interp", &data->Interp)); 281 CeedCallBackend(CeedGetKernelHip(ceed, data->module, "Grad", &data->Grad)); 282 CeedCallBackend(CeedGetKernelHip(ceed, data->module, "Weight", &data->Weight)); 283 CeedCallBackend(CeedFree(&basis_kernel_path)); 284 CeedCallBackend(CeedFree(&basis_kernel_source)); 285 CeedCallBackend(CeedBasisSetData(basis, data)); 286 287 // Register backend functions 288 CeedCallBackend(CeedSetBackendFunction(ceed, "Basis", basis, "Apply", CeedBasisApplyNonTensor_Hip)); 289 CeedCallBackend(CeedSetBackendFunction(ceed, "Basis", basis, "Destroy", CeedBasisDestroyNonTensor_Hip)); 290 return CEED_ERROR_SUCCESS; 291 } 292 //------------------------------------------------------------------------------ 293