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