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 CeedInt Q_1d, dim; 23 const CeedInt transpose = t_mode == CEED_TRANSPOSE; 24 const int max_block_size = 64; 25 const CeedScalar *d_u; 26 CeedScalar *d_v; 27 CeedBasis_Hip *data; 28 29 CeedCallBackend(CeedBasisGetCeed(basis, &ceed)); 30 CeedCallBackend(CeedBasisGetData(basis, &data)); 31 32 // Read vectors 33 if (u != CEED_VECTOR_NONE) CeedCallBackend(CeedVectorGetArrayRead(u, CEED_MEM_DEVICE, &d_u)); 34 else CeedCheck(eval_mode == CEED_EVAL_WEIGHT, ceed, CEED_ERROR_BACKEND, "An input vector is required for this CeedEvalMode"); 35 CeedCallBackend(CeedVectorGetArrayWrite(v, CEED_MEM_DEVICE, &d_v)); 36 37 // Clear v for transpose operation 38 if (transpose) { 39 CeedSize length; 40 41 CeedCallBackend(CeedVectorGetLength(v, &length)); 42 CeedCallHip(ceed, hipMemset(d_v, 0, length * sizeof(CeedScalar))); 43 } 44 CeedCallBackend(CeedBasisGetNumQuadraturePoints1D(basis, &Q_1d)); 45 CeedCallBackend(CeedBasisGetDimension(basis, &dim)); 46 47 // Basis action 48 switch (eval_mode) { 49 case CEED_EVAL_INTERP: { 50 void *interp_args[] = {(void *)&num_elem, (void *)&transpose, &data->d_interp_1d, &d_u, &d_v}; 51 const CeedInt block_size = CeedIntMin(CeedIntPow(Q_1d, dim), max_block_size); 52 53 CeedCallBackend(CeedRunKernel_Hip(ceed, data->Interp, num_elem, block_size, interp_args)); 54 } break; 55 case CEED_EVAL_GRAD: { 56 void *grad_args[] = {(void *)&num_elem, (void *)&transpose, &data->d_interp_1d, &data->d_grad_1d, &d_u, &d_v}; 57 const CeedInt block_size = max_block_size; 58 59 CeedCallBackend(CeedRunKernel_Hip(ceed, data->Grad, num_elem, block_size, grad_args)); 60 } break; 61 case CEED_EVAL_WEIGHT: { 62 void *weight_args[] = {(void *)&num_elem, (void *)&data->d_q_weight_1d, &d_v}; 63 const int block_size_x = Q_1d; 64 const int block_size_y = dim >= 2 ? Q_1d : 1; 65 66 CeedCallBackend(CeedRunKernelDim_Hip(ceed, data->Weight, num_elem, block_size_x, block_size_y, 1, weight_args)); 67 } break; 68 // LCOV_EXCL_START 69 // Evaluate the divergence to/from the quadrature points 70 case CEED_EVAL_DIV: 71 return CeedError(ceed, CEED_ERROR_BACKEND, "CEED_EVAL_DIV not supported"); 72 // Evaluate the curl to/from the quadrature points 73 case CEED_EVAL_CURL: 74 return CeedError(ceed, CEED_ERROR_BACKEND, "CEED_EVAL_CURL not supported"); 75 // Take no action, BasisApply should not have been called 76 case CEED_EVAL_NONE: 77 return CeedError(ceed, CEED_ERROR_BACKEND, "CEED_EVAL_NONE does not make sense in this context"); 78 // LCOV_EXCL_STOP 79 } 80 81 // Restore vectors 82 if (eval_mode != CEED_EVAL_WEIGHT) { 83 CeedCallBackend(CeedVectorRestoreArrayRead(u, &d_u)); 84 } 85 CeedCallBackend(CeedVectorRestoreArray(v, &d_v)); 86 return CEED_ERROR_SUCCESS; 87 } 88 89 //------------------------------------------------------------------------------ 90 // Basis apply - non-tensor 91 //------------------------------------------------------------------------------ 92 int CeedBasisApplyNonTensor_Hip(CeedBasis basis, const CeedInt num_elem, CeedTransposeMode t_mode, CeedEvalMode eval_mode, CeedVector u, 93 CeedVector v) { 94 Ceed ceed; 95 CeedInt num_nodes, num_qpts; 96 const CeedInt transpose = t_mode == CEED_TRANSPOSE; 97 const int elems_per_block = 1; 98 const int grid = CeedDivUpInt(num_elem, elems_per_block); 99 const CeedScalar *d_u; 100 CeedScalar *d_v; 101 CeedBasisNonTensor_Hip *data; 102 103 CeedCallBackend(CeedBasisGetCeed(basis, &ceed)); 104 CeedCallBackend(CeedBasisGetData(basis, &data)); 105 CeedCallBackend(CeedBasisGetNumQuadraturePoints(basis, &num_qpts)); 106 CeedCallBackend(CeedBasisGetNumNodes(basis, &num_nodes)); 107 108 // Read vectors 109 if (eval_mode != CEED_EVAL_WEIGHT) { 110 CeedCallBackend(CeedVectorGetArrayRead(u, CEED_MEM_DEVICE, &d_u)); 111 } 112 CeedCallBackend(CeedVectorGetArrayWrite(v, CEED_MEM_DEVICE, &d_v)); 113 114 // Clear v for transpose operation 115 if (transpose) { 116 CeedSize length; 117 118 CeedCallBackend(CeedVectorGetLength(v, &length)); 119 CeedCallHip(ceed, hipMemset(d_v, 0, length * sizeof(CeedScalar))); 120 } 121 122 // Apply basis operation 123 switch (eval_mode) { 124 case CEED_EVAL_INTERP: { 125 void *interp_args[] = {(void *)&num_elem, &data->d_interp, &d_u, &d_v}; 126 const int block_size_x = transpose ? num_nodes : num_qpts; 127 128 if (transpose) { 129 CeedCallBackend(CeedRunKernelDim_Hip(ceed, data->InterpTranspose, grid, block_size_x, 1, elems_per_block, interp_args)); 130 } else { 131 CeedCallBackend(CeedRunKernelDim_Hip(ceed, data->Interp, grid, block_size_x, 1, elems_per_block, interp_args)); 132 } 133 } break; 134 case CEED_EVAL_GRAD: { 135 void *grad_args[] = {(void *)&num_elem, &data->d_grad, &d_u, &d_v}; 136 const int block_size_x = transpose ? num_nodes : num_qpts; 137 138 if (transpose) { 139 CeedCallBackend(CeedRunKernelDim_Hip(ceed, data->DerivTranspose, grid, block_size_x, 1, elems_per_block, grad_args)); 140 } else { 141 CeedCallBackend(CeedRunKernelDim_Hip(ceed, data->Deriv, grid, block_size_x, 1, elems_per_block, grad_args)); 142 } 143 } break; 144 case CEED_EVAL_DIV: { 145 void *div_args[] = {(void *)&num_elem, &data->d_div, &d_u, &d_v}; 146 const int block_size_x = transpose ? num_nodes : num_qpts; 147 148 if (transpose) { 149 CeedCallBackend(CeedRunKernelDim_Hip(ceed, data->DerivTranspose, grid, block_size_x, 1, elems_per_block, div_args)); 150 } else { 151 CeedCallBackend(CeedRunKernelDim_Hip(ceed, data->Deriv, grid, block_size_x, 1, elems_per_block, div_args)); 152 } 153 } break; 154 case CEED_EVAL_CURL: { 155 void *curl_args[] = {(void *)&num_elem, &data->d_curl, &d_u, &d_v}; 156 const int block_size_x = transpose ? num_nodes : num_qpts; 157 158 if (transpose) { 159 CeedCallBackend(CeedRunKernelDim_Hip(ceed, data->DerivTranspose, grid, block_size_x, 1, elems_per_block, curl_args)); 160 } else { 161 CeedCallBackend(CeedRunKernelDim_Hip(ceed, data->Deriv, grid, block_size_x, 1, elems_per_block, curl_args)); 162 } 163 } break; 164 case CEED_EVAL_WEIGHT: { 165 void *weight_args[] = {(void *)&num_elem, (void *)&data->d_q_weight, &d_v}; 166 167 CeedCallBackend(CeedRunKernelDim_Hip(ceed, data->Weight, grid, num_qpts, 1, elems_per_block, weight_args)); 168 } break; 169 // LCOV_EXCL_START 170 // Take no action, BasisApply should not have been called 171 case CEED_EVAL_NONE: 172 return CeedError(ceed, CEED_ERROR_BACKEND, "CEED_EVAL_NONE does not make sense in this context"); 173 // LCOV_EXCL_STOP 174 } 175 176 // Restore vectors 177 if (eval_mode != CEED_EVAL_WEIGHT) { 178 CeedCallBackend(CeedVectorRestoreArrayRead(u, &d_u)); 179 } 180 CeedCallBackend(CeedVectorRestoreArray(v, &d_v)); 181 return CEED_ERROR_SUCCESS; 182 } 183 184 //------------------------------------------------------------------------------ 185 // Destroy tensor basis 186 //------------------------------------------------------------------------------ 187 static int CeedBasisDestroy_Hip(CeedBasis basis) { 188 Ceed ceed; 189 CeedBasis_Hip *data; 190 191 CeedCallBackend(CeedBasisGetCeed(basis, &ceed)); 192 CeedCallBackend(CeedBasisGetData(basis, &data)); 193 CeedCallHip(ceed, hipModuleUnload(data->module)); 194 CeedCallHip(ceed, hipFree(data->d_q_weight_1d)); 195 CeedCallHip(ceed, hipFree(data->d_interp_1d)); 196 CeedCallHip(ceed, hipFree(data->d_grad_1d)); 197 CeedCallBackend(CeedFree(&data)); 198 return CEED_ERROR_SUCCESS; 199 } 200 201 //------------------------------------------------------------------------------ 202 // Destroy non-tensor basis 203 //------------------------------------------------------------------------------ 204 static int CeedBasisDestroyNonTensor_Hip(CeedBasis basis) { 205 Ceed ceed; 206 CeedBasisNonTensor_Hip *data; 207 208 CeedCallBackend(CeedBasisGetCeed(basis, &ceed)); 209 CeedCallBackend(CeedBasisGetData(basis, &data)); 210 CeedCallHip(ceed, hipModuleUnload(data->module)); 211 CeedCallHip(ceed, hipFree(data->d_q_weight)); 212 CeedCallHip(ceed, hipFree(data->d_interp)); 213 CeedCallHip(ceed, hipFree(data->d_grad)); 214 CeedCallHip(ceed, hipFree(data->d_div)); 215 CeedCallHip(ceed, hipFree(data->d_curl)); 216 CeedCallBackend(CeedFree(&data)); 217 return CEED_ERROR_SUCCESS; 218 } 219 220 //------------------------------------------------------------------------------ 221 // Create tensor 222 //------------------------------------------------------------------------------ 223 int CeedBasisCreateTensorH1_Hip(CeedInt dim, CeedInt P_1d, CeedInt Q_1d, const CeedScalar *interp_1d, const CeedScalar *grad_1d, 224 const CeedScalar *q_ref_1d, const CeedScalar *q_weight_1d, CeedBasis basis) { 225 Ceed ceed; 226 char *basis_kernel_path, *basis_kernel_source; 227 CeedInt num_comp; 228 const CeedInt q_bytes = Q_1d * sizeof(CeedScalar); 229 const CeedInt interp_bytes = q_bytes * P_1d; 230 CeedBasis_Hip *data; 231 232 CeedCallBackend(CeedBasisGetCeed(basis, &ceed)); 233 CeedCallBackend(CeedCalloc(1, &data)); 234 235 // Copy data to GPU 236 CeedCallHip(ceed, hipMalloc((void **)&data->d_q_weight_1d, q_bytes)); 237 CeedCallHip(ceed, hipMemcpy(data->d_q_weight_1d, q_weight_1d, q_bytes, hipMemcpyHostToDevice)); 238 CeedCallHip(ceed, hipMalloc((void **)&data->d_interp_1d, interp_bytes)); 239 CeedCallHip(ceed, hipMemcpy(data->d_interp_1d, interp_1d, interp_bytes, hipMemcpyHostToDevice)); 240 CeedCallHip(ceed, hipMalloc((void **)&data->d_grad_1d, interp_bytes)); 241 CeedCallHip(ceed, hipMemcpy(data->d_grad_1d, grad_1d, interp_bytes, hipMemcpyHostToDevice)); 242 243 // Compile basis kernels 244 CeedCallBackend(CeedBasisGetNumComponents(basis, &num_comp)); 245 CeedCallBackend(CeedGetJitAbsolutePath(ceed, "ceed/jit-source/hip/hip-ref-basis-tensor.h", &basis_kernel_path)); 246 CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "----- Loading Basis Kernel Source -----\n"); 247 CeedCallBackend(CeedLoadSourceToBuffer(ceed, basis_kernel_path, &basis_kernel_source)); 248 CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "----- Loading Basis Kernel Source Complete! -----\n"); 249 CeedCallBackend(CeedCompile_Hip(ceed, basis_kernel_source, &data->module, 7, "BASIS_Q_1D", Q_1d, "BASIS_P_1D", P_1d, "BASIS_BUF_LEN", 250 num_comp * CeedIntPow(Q_1d > P_1d ? Q_1d : P_1d, dim), "BASIS_DIM", dim, "BASIS_NUM_COMP", num_comp, 251 "BASIS_NUM_NODES", CeedIntPow(P_1d, dim), "BASIS_NUM_QPTS", CeedIntPow(Q_1d, dim))); 252 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "Interp", &data->Interp)); 253 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "Grad", &data->Grad)); 254 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "Weight", &data->Weight)); 255 CeedCallBackend(CeedFree(&basis_kernel_path)); 256 CeedCallBackend(CeedFree(&basis_kernel_source)); 257 258 CeedCallBackend(CeedBasisSetData(basis, data)); 259 260 // Register backend functions 261 CeedCallBackend(CeedSetBackendFunction(ceed, "Basis", basis, "Apply", CeedBasisApply_Hip)); 262 CeedCallBackend(CeedSetBackendFunction(ceed, "Basis", basis, "Destroy", CeedBasisDestroy_Hip)); 263 return CEED_ERROR_SUCCESS; 264 } 265 266 //------------------------------------------------------------------------------ 267 // Create non-tensor H^1 268 //------------------------------------------------------------------------------ 269 int CeedBasisCreateH1_Hip(CeedElemTopology topo, CeedInt dim, CeedInt num_nodes, CeedInt num_qpts, const CeedScalar *interp, const CeedScalar *grad, 270 const CeedScalar *q_ref, const CeedScalar *q_weight, CeedBasis basis) { 271 Ceed ceed; 272 char *basis_kernel_path, *basis_kernel_source; 273 CeedInt num_comp, q_comp_interp, q_comp_grad; 274 const CeedInt q_bytes = num_qpts * sizeof(CeedScalar); 275 CeedBasisNonTensor_Hip *data; 276 277 CeedCallBackend(CeedBasisGetCeed(basis, &ceed)); 278 CeedCallBackend(CeedCalloc(1, &data)); 279 280 // Copy basis data to GPU 281 CeedCallBackend(CeedBasisGetNumQuadratureComponents(basis, CEED_EVAL_INTERP, &q_comp_interp)); 282 CeedCallBackend(CeedBasisGetNumQuadratureComponents(basis, CEED_EVAL_GRAD, &q_comp_grad)); 283 CeedCallHip(ceed, hipMalloc((void **)&data->d_q_weight, q_bytes)); 284 CeedCallHip(ceed, hipMemcpy(data->d_q_weight, q_weight, q_bytes, hipMemcpyHostToDevice)); 285 if (interp) { 286 const CeedInt interp_bytes = q_bytes * num_nodes * q_comp_interp; 287 288 CeedCallHip(ceed, hipMalloc((void **)&data->d_interp, interp_bytes)); 289 CeedCallHip(ceed, hipMemcpy(data->d_interp, interp, interp_bytes, hipMemcpyHostToDevice)); 290 } 291 if (grad) { 292 const CeedInt grad_bytes = q_bytes * num_nodes * q_comp_grad; 293 294 CeedCallHip(ceed, hipMalloc((void **)&data->d_grad, grad_bytes)); 295 CeedCallHip(ceed, hipMemcpy(data->d_grad, grad, grad_bytes, hipMemcpyHostToDevice)); 296 } 297 298 // Compile basis kernels 299 CeedCallBackend(CeedBasisGetNumComponents(basis, &num_comp)); 300 CeedCallBackend(CeedGetJitAbsolutePath(ceed, "ceed/jit-source/hip/hip-ref-basis-nontensor.h", &basis_kernel_path)); 301 CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "----- Loading Basis Kernel Source -----\n"); 302 CeedCallBackend(CeedLoadSourceToBuffer(ceed, basis_kernel_path, &basis_kernel_source)); 303 CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "----- Loading Basis Kernel Source Complete! -----\n"); 304 CeedCallBackend(CeedCompile_Hip(ceed, basis_kernel_source, &data->module, 5, "BASIS_Q", num_qpts, "BASIS_P", num_nodes, "BASIS_Q_COMP_INTERP", 305 q_comp_interp, "BASIS_Q_COMP_DERIV", q_comp_grad, "BASIS_NUM_COMP", num_comp)); 306 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "Interp", &data->Interp)); 307 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "InterpTranspose", &data->InterpTranspose)); 308 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "Deriv", &data->Deriv)); 309 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "DerivTranspose", &data->DerivTranspose)); 310 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "Weight", &data->Weight)); 311 CeedCallBackend(CeedFree(&basis_kernel_path)); 312 CeedCallBackend(CeedFree(&basis_kernel_source)); 313 314 CeedCallBackend(CeedBasisSetData(basis, data)); 315 316 // Register backend functions 317 CeedCallBackend(CeedSetBackendFunction(ceed, "Basis", basis, "Apply", CeedBasisApplyNonTensor_Hip)); 318 CeedCallBackend(CeedSetBackendFunction(ceed, "Basis", basis, "Destroy", CeedBasisDestroyNonTensor_Hip)); 319 return CEED_ERROR_SUCCESS; 320 } 321 322 //------------------------------------------------------------------------------ 323 // Create non-tensor H(div) 324 //------------------------------------------------------------------------------ 325 int CeedBasisCreateHdiv_Hip(CeedElemTopology topo, CeedInt dim, CeedInt num_nodes, CeedInt num_qpts, const CeedScalar *interp, const CeedScalar *div, 326 const CeedScalar *q_ref, const CeedScalar *q_weight, CeedBasis basis) { 327 Ceed ceed; 328 char *basis_kernel_path, *basis_kernel_source; 329 CeedInt num_comp, q_comp_interp, q_comp_div; 330 const CeedInt q_bytes = num_qpts * sizeof(CeedScalar); 331 CeedBasisNonTensor_Hip *data; 332 333 CeedCallBackend(CeedBasisGetCeed(basis, &ceed)); 334 CeedCallBackend(CeedCalloc(1, &data)); 335 336 // Copy basis data to GPU 337 CeedCallBackend(CeedBasisGetNumQuadratureComponents(basis, CEED_EVAL_INTERP, &q_comp_interp)); 338 CeedCallBackend(CeedBasisGetNumQuadratureComponents(basis, CEED_EVAL_DIV, &q_comp_div)); 339 CeedCallHip(ceed, hipMalloc((void **)&data->d_q_weight, q_bytes)); 340 CeedCallHip(ceed, hipMemcpy(data->d_q_weight, q_weight, q_bytes, hipMemcpyHostToDevice)); 341 if (interp) { 342 const CeedInt interp_bytes = q_bytes * num_nodes * q_comp_interp; 343 344 CeedCallHip(ceed, hipMalloc((void **)&data->d_interp, interp_bytes)); 345 CeedCallHip(ceed, hipMemcpy(data->d_interp, interp, interp_bytes, hipMemcpyHostToDevice)); 346 } 347 if (div) { 348 const CeedInt div_bytes = q_bytes * num_nodes * q_comp_div; 349 350 CeedCallHip(ceed, hipMalloc((void **)&data->d_div, div_bytes)); 351 CeedCallHip(ceed, hipMemcpy(data->d_div, div, div_bytes, hipMemcpyHostToDevice)); 352 } 353 354 // Compile basis kernels 355 CeedCallBackend(CeedBasisGetNumComponents(basis, &num_comp)); 356 CeedCallBackend(CeedGetJitAbsolutePath(ceed, "ceed/jit-source/hip/hip-ref-basis-nontensor.h", &basis_kernel_path)); 357 CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "----- Loading Basis Kernel Source -----\n"); 358 CeedCallBackend(CeedLoadSourceToBuffer(ceed, basis_kernel_path, &basis_kernel_source)); 359 CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "----- Loading Basis Kernel Source Complete! -----\n"); 360 CeedCallBackend(CeedCompile_Hip(ceed, basis_kernel_source, &data->module, 5, "BASIS_Q", num_qpts, "BASIS_P", num_nodes, "BASIS_Q_COMP_INTERP", 361 q_comp_interp, "BASIS_Q_COMP_DERIV", q_comp_div, "BASIS_NUM_COMP", num_comp)); 362 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "Interp", &data->Interp)); 363 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "InterpTranspose", &data->InterpTranspose)); 364 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "Deriv", &data->Deriv)); 365 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "DerivTranspose", &data->DerivTranspose)); 366 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "Weight", &data->Weight)); 367 CeedCallBackend(CeedFree(&basis_kernel_path)); 368 CeedCallBackend(CeedFree(&basis_kernel_source)); 369 370 CeedCallBackend(CeedBasisSetData(basis, data)); 371 372 // Register backend functions 373 CeedCallBackend(CeedSetBackendFunction(ceed, "Basis", basis, "Apply", CeedBasisApplyNonTensor_Hip)); 374 CeedCallBackend(CeedSetBackendFunction(ceed, "Basis", basis, "Destroy", CeedBasisDestroyNonTensor_Hip)); 375 return CEED_ERROR_SUCCESS; 376 } 377 378 //------------------------------------------------------------------------------ 379 // Create non-tensor H(curl) 380 //------------------------------------------------------------------------------ 381 int CeedBasisCreateHcurl_Hip(CeedElemTopology topo, CeedInt dim, CeedInt num_nodes, CeedInt num_qpts, const CeedScalar *interp, 382 const CeedScalar *curl, const CeedScalar *q_ref, const CeedScalar *q_weight, CeedBasis basis) { 383 Ceed ceed; 384 char *basis_kernel_path, *basis_kernel_source; 385 CeedInt num_comp, q_comp_interp, q_comp_curl; 386 const CeedInt q_bytes = num_qpts * sizeof(CeedScalar); 387 CeedBasisNonTensor_Hip *data; 388 389 CeedCallBackend(CeedBasisGetCeed(basis, &ceed)); 390 CeedCallBackend(CeedCalloc(1, &data)); 391 392 // Copy basis data to GPU 393 CeedCallBackend(CeedBasisGetNumQuadratureComponents(basis, CEED_EVAL_INTERP, &q_comp_interp)); 394 CeedCallBackend(CeedBasisGetNumQuadratureComponents(basis, CEED_EVAL_CURL, &q_comp_curl)); 395 CeedCallHip(ceed, hipMalloc((void **)&data->d_q_weight, q_bytes)); 396 CeedCallHip(ceed, hipMemcpy(data->d_q_weight, q_weight, q_bytes, hipMemcpyHostToDevice)); 397 if (interp) { 398 const CeedInt interp_bytes = q_bytes * num_nodes * q_comp_interp; 399 400 CeedCallHip(ceed, hipMalloc((void **)&data->d_interp, interp_bytes)); 401 CeedCallHip(ceed, hipMemcpy(data->d_interp, interp, interp_bytes, hipMemcpyHostToDevice)); 402 } 403 if (curl) { 404 const CeedInt curl_bytes = q_bytes * num_nodes * q_comp_curl; 405 406 CeedCallHip(ceed, hipMalloc((void **)&data->d_curl, curl_bytes)); 407 CeedCallHip(ceed, hipMemcpy(data->d_curl, curl, curl_bytes, hipMemcpyHostToDevice)); 408 } 409 410 // Compile basis kernels 411 CeedCallBackend(CeedBasisGetNumComponents(basis, &num_comp)); 412 CeedCallBackend(CeedGetJitAbsolutePath(ceed, "ceed/jit-source/hip/hip-ref-basis-nontensor.h", &basis_kernel_path)); 413 CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "----- Loading Basis Kernel Source -----\n"); 414 CeedCallBackend(CeedLoadSourceToBuffer(ceed, basis_kernel_path, &basis_kernel_source)); 415 CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "----- Loading Basis Kernel Source Complete! -----\n"); 416 CeedCallBackend(CeedCompile_Hip(ceed, basis_kernel_source, &data->module, 5, "BASIS_Q", num_qpts, "BASIS_P", num_nodes, "BASIS_Q_COMP_INTERP", 417 q_comp_interp, "BASIS_Q_COMP_DERIV", q_comp_curl, "BASIS_NUM_COMP", num_comp)); 418 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "Interp", &data->Interp)); 419 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "InterpTranspose", &data->InterpTranspose)); 420 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "Deriv", &data->Deriv)); 421 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "DerivTranspose", &data->DerivTranspose)); 422 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "Weight", &data->Weight)); 423 CeedCallBackend(CeedFree(&basis_kernel_path)); 424 CeedCallBackend(CeedFree(&basis_kernel_source)); 425 426 CeedCallBackend(CeedBasisSetData(basis, data)); 427 428 // Register backend functions 429 CeedCallBackend(CeedSetBackendFunction(ceed, "Basis", basis, "Apply", CeedBasisApplyNonTensor_Hip)); 430 CeedCallBackend(CeedSetBackendFunction(ceed, "Basis", basis, "Destroy", CeedBasisDestroyNonTensor_Hip)); 431 return CEED_ERROR_SUCCESS; 432 } 433 434 //------------------------------------------------------------------------------ 435