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