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 - tensor AtPoints 87 //------------------------------------------------------------------------------ 88 int CeedBasisApplyAtPoints_Hip(CeedBasis basis, const CeedInt num_elem, const CeedInt *num_points, CeedTransposeMode t_mode, CeedEvalMode eval_mode, 89 CeedVector x_ref, CeedVector u, CeedVector v) { 90 Ceed ceed; 91 CeedInt Q_1d, dim, max_num_points = num_points[0]; 92 const CeedInt is_transpose = t_mode == CEED_TRANSPOSE; 93 const int max_block_size = 32; 94 const CeedScalar *d_x, *d_u; 95 CeedScalar *d_v; 96 CeedBasis_Hip *data; 97 98 CeedCallBackend(CeedBasisGetCeed(basis, &ceed)); 99 CeedCallBackend(CeedBasisGetData(basis, &data)); 100 CeedCallBackend(CeedBasisGetNumQuadraturePoints1D(basis, &Q_1d)); 101 CeedCallBackend(CeedBasisGetDimension(basis, &dim)); 102 103 // Check uniform number of points per elem 104 for (CeedInt i = 1; i < num_elem; i++) { 105 CeedCheck(max_num_points == num_points[i], ceed, CEED_ERROR_BACKEND, 106 "BasisApplyAtPoints only supported for the same number of points in each element"); 107 } 108 109 // Weight handled separately 110 if (eval_mode == CEED_EVAL_WEIGHT) { 111 CeedCall(CeedVectorSetValue(v, 1.0)); 112 return CEED_ERROR_SUCCESS; 113 } 114 115 // Build kernels if needed 116 if (data->num_points != max_num_points) { 117 CeedInt P_1d; 118 119 CeedCallBackend(CeedBasisGetNumNodes1D(basis, &P_1d)); 120 data->num_points = max_num_points; 121 122 // -- Create interp matrix to Chebyshev coefficients 123 if (!data->d_chebyshev_interp_1d) { 124 CeedSize interp_bytes; 125 CeedScalar *chebyshev_interp_1d; 126 127 interp_bytes = P_1d * Q_1d * sizeof(CeedScalar); 128 CeedCallBackend(CeedCalloc(P_1d * Q_1d, &chebyshev_interp_1d)); 129 CeedCall(CeedBasisGetChebyshevInterp1D(basis, chebyshev_interp_1d)); 130 CeedCallHip(ceed, hipMalloc((void **)&data->d_chebyshev_interp_1d, interp_bytes)); 131 CeedCallHip(ceed, hipMemcpy(data->d_chebyshev_interp_1d, chebyshev_interp_1d, interp_bytes, hipMemcpyHostToDevice)); 132 CeedCallBackend(CeedFree(&chebyshev_interp_1d)); 133 } 134 135 // -- Compile kernels 136 char *basis_kernel_source; 137 const char *basis_kernel_path; 138 CeedInt num_comp; 139 140 if (data->moduleAtPoints) CeedCallHip(ceed, hipModuleUnload(data->moduleAtPoints)); 141 CeedCallBackend(CeedBasisGetNumComponents(basis, &num_comp)); 142 CeedCallBackend(CeedGetJitAbsolutePath(ceed, "ceed/jit-source/hip/hip-ref-basis-tensor-at-points.h", &basis_kernel_path)); 143 CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "----- Loading Basis Kernel Source -----\n"); 144 CeedCallBackend(CeedLoadSourceToBuffer(ceed, basis_kernel_path, &basis_kernel_source)); 145 CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "----- Loading Basis Kernel Source Complete! -----\n"); 146 CeedCallBackend(CeedCompile_Hip(ceed, basis_kernel_source, &data->moduleAtPoints, 9, "BASIS_Q_1D", Q_1d, "BASIS_P_1D", P_1d, "BASIS_BUF_LEN", 147 Q_1d * CeedIntPow(Q_1d > P_1d ? Q_1d : P_1d, dim - 1), "BASIS_DIM", dim, "BASIS_NUM_COMP", num_comp, 148 "BASIS_NUM_NODES", CeedIntPow(P_1d, dim), "BASIS_NUM_QPTS", CeedIntPow(Q_1d, dim), "BASIS_NUM_PTS", 149 max_num_points, "POINTS_BUFF_LEN", CeedIntPow(Q_1d, dim - 1))); 150 CeedCallBackend(CeedGetKernel_Hip(ceed, data->moduleAtPoints, "InterpAtPoints", &data->InterpAtPoints)); 151 CeedCallBackend(CeedGetKernel_Hip(ceed, data->moduleAtPoints, "GradAtPoints", &data->GradAtPoints)); 152 CeedCallBackend(CeedFree(&basis_kernel_path)); 153 CeedCallBackend(CeedFree(&basis_kernel_source)); 154 } 155 156 // Get read/write access to u, v 157 CeedCallBackend(CeedVectorGetArrayRead(x_ref, CEED_MEM_DEVICE, &d_x)); 158 if (u != CEED_VECTOR_NONE) CeedCallBackend(CeedVectorGetArrayRead(u, CEED_MEM_DEVICE, &d_u)); 159 else CeedCheck(eval_mode == CEED_EVAL_WEIGHT, ceed, CEED_ERROR_BACKEND, "An input vector is required for this CeedEvalMode"); 160 CeedCallBackend(CeedVectorGetArrayWrite(v, CEED_MEM_DEVICE, &d_v)); 161 162 // Clear v for transpose operation 163 if (is_transpose) { 164 CeedSize length; 165 166 CeedCallBackend(CeedVectorGetLength(v, &length)); 167 CeedCallHip(ceed, hipMemset(d_v, 0, length * sizeof(CeedScalar))); 168 } 169 170 // Basis action 171 switch (eval_mode) { 172 case CEED_EVAL_INTERP: { 173 void *interp_args[] = {(void *)&num_elem, (void *)&is_transpose, &data->d_chebyshev_interp_1d, &d_x, &d_u, &d_v}; 174 const CeedInt block_size = CeedIntMin(CeedIntPow(Q_1d, dim), max_block_size); 175 176 CeedCallBackend(CeedRunKernel_Hip(ceed, data->InterpAtPoints, num_elem, block_size, interp_args)); 177 } break; 178 case CEED_EVAL_GRAD: { 179 void *grad_args[] = {(void *)&num_elem, (void *)&is_transpose, &data->d_chebyshev_interp_1d, &d_x, &d_u, &d_v}; 180 const CeedInt block_size = CeedIntMin(CeedIntPow(Q_1d, dim), max_block_size); 181 182 CeedCallBackend(CeedRunKernel_Hip(ceed, data->GradAtPoints, num_elem, block_size, grad_args)); 183 } break; 184 case CEED_EVAL_WEIGHT: 185 case CEED_EVAL_NONE: /* handled separately below */ 186 break; 187 // LCOV_EXCL_START 188 case CEED_EVAL_DIV: 189 case CEED_EVAL_CURL: 190 return CeedError(ceed, CEED_ERROR_BACKEND, "%s not supported", CeedEvalModes[eval_mode]); 191 // LCOV_EXCL_STOP 192 } 193 194 // Restore vectors, cover CEED_EVAL_NONE 195 CeedCallBackend(CeedVectorRestoreArrayRead(x_ref, &d_x)); 196 CeedCallBackend(CeedVectorRestoreArray(v, &d_v)); 197 if (eval_mode == CEED_EVAL_NONE) CeedCallBackend(CeedVectorSetArray(v, CEED_MEM_DEVICE, CEED_COPY_VALUES, (CeedScalar *)d_u)); 198 if (eval_mode != CEED_EVAL_WEIGHT) CeedCallBackend(CeedVectorRestoreArrayRead(u, &d_u)); 199 return CEED_ERROR_SUCCESS; 200 } 201 202 //------------------------------------------------------------------------------ 203 // Basis apply - non-tensor 204 //------------------------------------------------------------------------------ 205 int CeedBasisApplyNonTensor_Hip(CeedBasis basis, const CeedInt num_elem, CeedTransposeMode t_mode, CeedEvalMode eval_mode, CeedVector u, 206 CeedVector v) { 207 Ceed ceed; 208 CeedInt num_nodes, num_qpts; 209 const CeedInt is_transpose = t_mode == CEED_TRANSPOSE; 210 const int elems_per_block = 1; 211 const int grid = CeedDivUpInt(num_elem, elems_per_block); 212 const CeedScalar *d_u; 213 CeedScalar *d_v; 214 CeedBasisNonTensor_Hip *data; 215 216 CeedCallBackend(CeedBasisGetCeed(basis, &ceed)); 217 CeedCallBackend(CeedBasisGetData(basis, &data)); 218 CeedCallBackend(CeedBasisGetNumQuadraturePoints(basis, &num_qpts)); 219 CeedCallBackend(CeedBasisGetNumNodes(basis, &num_nodes)); 220 221 // Get read/write access to u, v 222 if (u != CEED_VECTOR_NONE) CeedCallBackend(CeedVectorGetArrayRead(u, CEED_MEM_DEVICE, &d_u)); 223 else CeedCheck(eval_mode == CEED_EVAL_WEIGHT, ceed, CEED_ERROR_BACKEND, "An input vector is required for this CeedEvalMode"); 224 CeedCallBackend(CeedVectorGetArrayWrite(v, CEED_MEM_DEVICE, &d_v)); 225 226 // Clear v for transpose operation 227 if (is_transpose) { 228 CeedSize length; 229 230 CeedCallBackend(CeedVectorGetLength(v, &length)); 231 CeedCallHip(ceed, hipMemset(d_v, 0, length * sizeof(CeedScalar))); 232 } 233 234 // Apply basis operation 235 switch (eval_mode) { 236 case CEED_EVAL_INTERP: { 237 void *interp_args[] = {(void *)&num_elem, &data->d_interp, &d_u, &d_v}; 238 const int block_size_x = is_transpose ? num_nodes : num_qpts; 239 240 if (is_transpose) { 241 CeedCallBackend(CeedRunKernelDim_Hip(ceed, data->InterpTranspose, grid, block_size_x, 1, elems_per_block, interp_args)); 242 } else { 243 CeedCallBackend(CeedRunKernelDim_Hip(ceed, data->Interp, grid, block_size_x, 1, elems_per_block, interp_args)); 244 } 245 } break; 246 case CEED_EVAL_GRAD: { 247 void *grad_args[] = {(void *)&num_elem, &data->d_grad, &d_u, &d_v}; 248 const int block_size_x = is_transpose ? num_nodes : num_qpts; 249 250 if (is_transpose) { 251 CeedCallBackend(CeedRunKernelDim_Hip(ceed, data->DerivTranspose, grid, block_size_x, 1, elems_per_block, grad_args)); 252 } else { 253 CeedCallBackend(CeedRunKernelDim_Hip(ceed, data->Deriv, grid, block_size_x, 1, elems_per_block, grad_args)); 254 } 255 } break; 256 case CEED_EVAL_DIV: { 257 void *div_args[] = {(void *)&num_elem, &data->d_div, &d_u, &d_v}; 258 const int block_size_x = is_transpose ? num_nodes : num_qpts; 259 260 if (is_transpose) { 261 CeedCallBackend(CeedRunKernelDim_Hip(ceed, data->DerivTranspose, grid, block_size_x, 1, elems_per_block, div_args)); 262 } else { 263 CeedCallBackend(CeedRunKernelDim_Hip(ceed, data->Deriv, grid, block_size_x, 1, elems_per_block, div_args)); 264 } 265 } break; 266 case CEED_EVAL_CURL: { 267 void *curl_args[] = {(void *)&num_elem, &data->d_curl, &d_u, &d_v}; 268 const int block_size_x = is_transpose ? num_nodes : num_qpts; 269 270 if (is_transpose) { 271 CeedCallBackend(CeedRunKernelDim_Hip(ceed, data->DerivTranspose, grid, block_size_x, 1, elems_per_block, curl_args)); 272 } else { 273 CeedCallBackend(CeedRunKernelDim_Hip(ceed, data->Deriv, grid, block_size_x, 1, elems_per_block, curl_args)); 274 } 275 } break; 276 case CEED_EVAL_WEIGHT: { 277 CeedCheck(data->d_q_weight, ceed, CEED_ERROR_BACKEND, "%s not supported; q_weights not set", CeedEvalModes[eval_mode]); 278 void *weight_args[] = {(void *)&num_elem, (void *)&data->d_q_weight, &d_v}; 279 280 CeedCallBackend(CeedRunKernelDim_Hip(ceed, data->Weight, grid, num_qpts, 1, elems_per_block, weight_args)); 281 } break; 282 case CEED_EVAL_NONE: /* handled separately below */ 283 break; 284 } 285 286 // Restore vectors, cover CEED_EVAL_NONE 287 CeedCallBackend(CeedVectorRestoreArray(v, &d_v)); 288 if (eval_mode == CEED_EVAL_NONE) CeedCallBackend(CeedVectorSetArray(v, CEED_MEM_DEVICE, CEED_COPY_VALUES, (CeedScalar *)d_u)); 289 if (eval_mode != CEED_EVAL_WEIGHT) CeedCallBackend(CeedVectorRestoreArrayRead(u, &d_u)); 290 return CEED_ERROR_SUCCESS; 291 } 292 293 //------------------------------------------------------------------------------ 294 // Destroy tensor basis 295 //------------------------------------------------------------------------------ 296 static int CeedBasisDestroy_Hip(CeedBasis basis) { 297 Ceed ceed; 298 CeedBasis_Hip *data; 299 300 CeedCallBackend(CeedBasisGetCeed(basis, &ceed)); 301 CeedCallBackend(CeedBasisGetData(basis, &data)); 302 CeedCallHip(ceed, hipModuleUnload(data->module)); 303 if (data->moduleAtPoints) CeedCallHip(ceed, hipModuleUnload(data->moduleAtPoints)); 304 if (data->d_q_weight_1d) CeedCallHip(ceed, hipFree(data->d_q_weight_1d)); 305 CeedCallHip(ceed, hipFree(data->d_interp_1d)); 306 CeedCallHip(ceed, hipFree(data->d_grad_1d)); 307 CeedCallHip(ceed, hipFree(data->d_chebyshev_interp_1d)); 308 CeedCallBackend(CeedFree(&data)); 309 return CEED_ERROR_SUCCESS; 310 } 311 312 //------------------------------------------------------------------------------ 313 // Destroy non-tensor basis 314 //------------------------------------------------------------------------------ 315 static int CeedBasisDestroyNonTensor_Hip(CeedBasis basis) { 316 Ceed ceed; 317 CeedBasisNonTensor_Hip *data; 318 319 CeedCallBackend(CeedBasisGetCeed(basis, &ceed)); 320 CeedCallBackend(CeedBasisGetData(basis, &data)); 321 CeedCallHip(ceed, hipModuleUnload(data->module)); 322 if (data->d_q_weight) CeedCallHip(ceed, hipFree(data->d_q_weight)); 323 CeedCallHip(ceed, hipFree(data->d_interp)); 324 CeedCallHip(ceed, hipFree(data->d_grad)); 325 CeedCallHip(ceed, hipFree(data->d_div)); 326 CeedCallHip(ceed, hipFree(data->d_curl)); 327 CeedCallBackend(CeedFree(&data)); 328 return CEED_ERROR_SUCCESS; 329 } 330 331 //------------------------------------------------------------------------------ 332 // Create tensor 333 //------------------------------------------------------------------------------ 334 int CeedBasisCreateTensorH1_Hip(CeedInt dim, CeedInt P_1d, CeedInt Q_1d, const CeedScalar *interp_1d, const CeedScalar *grad_1d, 335 const CeedScalar *q_ref_1d, const CeedScalar *q_weight_1d, CeedBasis basis) { 336 Ceed ceed; 337 char *basis_kernel_source; 338 const char *basis_kernel_path; 339 CeedInt num_comp; 340 const CeedInt q_bytes = Q_1d * sizeof(CeedScalar); 341 const CeedInt interp_bytes = q_bytes * P_1d; 342 CeedBasis_Hip *data; 343 344 CeedCallBackend(CeedBasisGetCeed(basis, &ceed)); 345 CeedCallBackend(CeedCalloc(1, &data)); 346 347 // Copy data to GPU 348 if (q_weight_1d) { 349 CeedCallHip(ceed, hipMalloc((void **)&data->d_q_weight_1d, q_bytes)); 350 CeedCallHip(ceed, hipMemcpy(data->d_q_weight_1d, q_weight_1d, q_bytes, hipMemcpyHostToDevice)); 351 } 352 CeedCallHip(ceed, hipMalloc((void **)&data->d_interp_1d, interp_bytes)); 353 CeedCallHip(ceed, hipMemcpy(data->d_interp_1d, interp_1d, interp_bytes, hipMemcpyHostToDevice)); 354 CeedCallHip(ceed, hipMalloc((void **)&data->d_grad_1d, interp_bytes)); 355 CeedCallHip(ceed, hipMemcpy(data->d_grad_1d, grad_1d, interp_bytes, hipMemcpyHostToDevice)); 356 357 // Compile basis kernels 358 CeedCallBackend(CeedBasisGetNumComponents(basis, &num_comp)); 359 CeedCallBackend(CeedGetJitAbsolutePath(ceed, "ceed/jit-source/hip/hip-ref-basis-tensor.h", &basis_kernel_path)); 360 CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "----- Loading Basis Kernel Source -----\n"); 361 CeedCallBackend(CeedLoadSourceToBuffer(ceed, basis_kernel_path, &basis_kernel_source)); 362 CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "----- Loading Basis Kernel Source Complete! -----\n"); 363 CeedCallBackend(CeedCompile_Hip(ceed, basis_kernel_source, &data->module, 7, "BASIS_Q_1D", Q_1d, "BASIS_P_1D", P_1d, "BASIS_BUF_LEN", 364 Q_1d * CeedIntPow(Q_1d > P_1d ? Q_1d : P_1d, dim - 1), "BASIS_DIM", dim, "BASIS_NUM_COMP", num_comp, 365 "BASIS_NUM_NODES", CeedIntPow(P_1d, dim), "BASIS_NUM_QPTS", CeedIntPow(Q_1d, dim))); 366 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "Interp", &data->Interp)); 367 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "Grad", &data->Grad)); 368 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "Weight", &data->Weight)); 369 CeedCallBackend(CeedFree(&basis_kernel_path)); 370 CeedCallBackend(CeedFree(&basis_kernel_source)); 371 372 CeedCallBackend(CeedBasisSetData(basis, data)); 373 374 // Register backend functions 375 CeedCallBackend(CeedSetBackendFunction(ceed, "Basis", basis, "Apply", CeedBasisApply_Hip)); 376 CeedCallBackend(CeedSetBackendFunction(ceed, "Basis", basis, "ApplyAtPoints", CeedBasisApplyAtPoints_Hip)); 377 CeedCallBackend(CeedSetBackendFunction(ceed, "Basis", basis, "Destroy", CeedBasisDestroy_Hip)); 378 return CEED_ERROR_SUCCESS; 379 } 380 381 //------------------------------------------------------------------------------ 382 // Create non-tensor H^1 383 //------------------------------------------------------------------------------ 384 int CeedBasisCreateH1_Hip(CeedElemTopology topo, CeedInt dim, CeedInt num_nodes, CeedInt num_qpts, const CeedScalar *interp, const CeedScalar *grad, 385 const CeedScalar *q_ref, const CeedScalar *q_weight, CeedBasis basis) { 386 Ceed ceed; 387 char *basis_kernel_source; 388 const char *basis_kernel_path; 389 CeedInt num_comp, q_comp_interp, q_comp_grad; 390 const CeedInt q_bytes = num_qpts * sizeof(CeedScalar); 391 CeedBasisNonTensor_Hip *data; 392 393 CeedCallBackend(CeedBasisGetCeed(basis, &ceed)); 394 CeedCallBackend(CeedCalloc(1, &data)); 395 396 // Copy basis data to GPU 397 CeedCallBackend(CeedBasisGetNumQuadratureComponents(basis, CEED_EVAL_INTERP, &q_comp_interp)); 398 CeedCallBackend(CeedBasisGetNumQuadratureComponents(basis, CEED_EVAL_GRAD, &q_comp_grad)); 399 if (q_weight) { 400 CeedCallHip(ceed, hipMalloc((void **)&data->d_q_weight, q_bytes)); 401 CeedCallHip(ceed, hipMemcpy(data->d_q_weight, q_weight, q_bytes, hipMemcpyHostToDevice)); 402 } 403 if (interp) { 404 const CeedInt interp_bytes = q_bytes * num_nodes * q_comp_interp; 405 406 CeedCallHip(ceed, hipMalloc((void **)&data->d_interp, interp_bytes)); 407 CeedCallHip(ceed, hipMemcpy(data->d_interp, interp, interp_bytes, hipMemcpyHostToDevice)); 408 } 409 if (grad) { 410 const CeedInt grad_bytes = q_bytes * num_nodes * q_comp_grad; 411 412 CeedCallHip(ceed, hipMalloc((void **)&data->d_grad, grad_bytes)); 413 CeedCallHip(ceed, hipMemcpy(data->d_grad, grad, grad_bytes, hipMemcpyHostToDevice)); 414 } 415 416 // Compile basis kernels 417 CeedCallBackend(CeedBasisGetNumComponents(basis, &num_comp)); 418 CeedCallBackend(CeedGetJitAbsolutePath(ceed, "ceed/jit-source/hip/hip-ref-basis-nontensor.h", &basis_kernel_path)); 419 CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "----- Loading Basis Kernel Source -----\n"); 420 CeedCallBackend(CeedLoadSourceToBuffer(ceed, basis_kernel_path, &basis_kernel_source)); 421 CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "----- Loading Basis Kernel Source Complete! -----\n"); 422 CeedCallBackend(CeedCompile_Hip(ceed, basis_kernel_source, &data->module, 5, "BASIS_Q", num_qpts, "BASIS_P", num_nodes, "BASIS_Q_COMP_INTERP", 423 q_comp_interp, "BASIS_Q_COMP_DERIV", q_comp_grad, "BASIS_NUM_COMP", num_comp)); 424 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "Interp", &data->Interp)); 425 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "InterpTranspose", &data->InterpTranspose)); 426 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "Deriv", &data->Deriv)); 427 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "DerivTranspose", &data->DerivTranspose)); 428 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "Weight", &data->Weight)); 429 CeedCallBackend(CeedFree(&basis_kernel_path)); 430 CeedCallBackend(CeedFree(&basis_kernel_source)); 431 432 CeedCallBackend(CeedBasisSetData(basis, data)); 433 434 // Register backend functions 435 CeedCallBackend(CeedSetBackendFunction(ceed, "Basis", basis, "Apply", CeedBasisApplyNonTensor_Hip)); 436 CeedCallBackend(CeedSetBackendFunction(ceed, "Basis", basis, "Destroy", CeedBasisDestroyNonTensor_Hip)); 437 return CEED_ERROR_SUCCESS; 438 } 439 440 //------------------------------------------------------------------------------ 441 // Create non-tensor H(div) 442 //------------------------------------------------------------------------------ 443 int CeedBasisCreateHdiv_Hip(CeedElemTopology topo, CeedInt dim, CeedInt num_nodes, CeedInt num_qpts, const CeedScalar *interp, const CeedScalar *div, 444 const CeedScalar *q_ref, const CeedScalar *q_weight, CeedBasis basis) { 445 Ceed ceed; 446 char *basis_kernel_source; 447 const char *basis_kernel_path; 448 CeedInt num_comp, q_comp_interp, q_comp_div; 449 const CeedInt q_bytes = num_qpts * sizeof(CeedScalar); 450 CeedBasisNonTensor_Hip *data; 451 452 CeedCallBackend(CeedBasisGetCeed(basis, &ceed)); 453 CeedCallBackend(CeedCalloc(1, &data)); 454 455 // Copy basis data to GPU 456 CeedCallBackend(CeedBasisGetNumQuadratureComponents(basis, CEED_EVAL_INTERP, &q_comp_interp)); 457 CeedCallBackend(CeedBasisGetNumQuadratureComponents(basis, CEED_EVAL_DIV, &q_comp_div)); 458 if (q_weight) { 459 CeedCallHip(ceed, hipMalloc((void **)&data->d_q_weight, q_bytes)); 460 CeedCallHip(ceed, hipMemcpy(data->d_q_weight, q_weight, q_bytes, hipMemcpyHostToDevice)); 461 } 462 if (interp) { 463 const CeedInt interp_bytes = q_bytes * num_nodes * q_comp_interp; 464 465 CeedCallHip(ceed, hipMalloc((void **)&data->d_interp, interp_bytes)); 466 CeedCallHip(ceed, hipMemcpy(data->d_interp, interp, interp_bytes, hipMemcpyHostToDevice)); 467 } 468 if (div) { 469 const CeedInt div_bytes = q_bytes * num_nodes * q_comp_div; 470 471 CeedCallHip(ceed, hipMalloc((void **)&data->d_div, div_bytes)); 472 CeedCallHip(ceed, hipMemcpy(data->d_div, div, div_bytes, hipMemcpyHostToDevice)); 473 } 474 475 // Compile basis kernels 476 CeedCallBackend(CeedBasisGetNumComponents(basis, &num_comp)); 477 CeedCallBackend(CeedGetJitAbsolutePath(ceed, "ceed/jit-source/hip/hip-ref-basis-nontensor.h", &basis_kernel_path)); 478 CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "----- Loading Basis Kernel Source -----\n"); 479 CeedCallBackend(CeedLoadSourceToBuffer(ceed, basis_kernel_path, &basis_kernel_source)); 480 CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "----- Loading Basis Kernel Source Complete! -----\n"); 481 CeedCallBackend(CeedCompile_Hip(ceed, basis_kernel_source, &data->module, 5, "BASIS_Q", num_qpts, "BASIS_P", num_nodes, "BASIS_Q_COMP_INTERP", 482 q_comp_interp, "BASIS_Q_COMP_DERIV", q_comp_div, "BASIS_NUM_COMP", num_comp)); 483 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "Interp", &data->Interp)); 484 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "InterpTranspose", &data->InterpTranspose)); 485 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "Deriv", &data->Deriv)); 486 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "DerivTranspose", &data->DerivTranspose)); 487 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "Weight", &data->Weight)); 488 CeedCallBackend(CeedFree(&basis_kernel_path)); 489 CeedCallBackend(CeedFree(&basis_kernel_source)); 490 491 CeedCallBackend(CeedBasisSetData(basis, data)); 492 493 // Register backend functions 494 CeedCallBackend(CeedSetBackendFunction(ceed, "Basis", basis, "Apply", CeedBasisApplyNonTensor_Hip)); 495 CeedCallBackend(CeedSetBackendFunction(ceed, "Basis", basis, "Destroy", CeedBasisDestroyNonTensor_Hip)); 496 return CEED_ERROR_SUCCESS; 497 } 498 499 //------------------------------------------------------------------------------ 500 // Create non-tensor H(curl) 501 //------------------------------------------------------------------------------ 502 int CeedBasisCreateHcurl_Hip(CeedElemTopology topo, CeedInt dim, CeedInt num_nodes, CeedInt num_qpts, const CeedScalar *interp, 503 const CeedScalar *curl, const CeedScalar *q_ref, const CeedScalar *q_weight, CeedBasis basis) { 504 Ceed ceed; 505 char *basis_kernel_source; 506 const char *basis_kernel_path; 507 CeedInt num_comp, q_comp_interp, q_comp_curl; 508 const CeedInt q_bytes = num_qpts * sizeof(CeedScalar); 509 CeedBasisNonTensor_Hip *data; 510 511 CeedCallBackend(CeedBasisGetCeed(basis, &ceed)); 512 CeedCallBackend(CeedCalloc(1, &data)); 513 514 // Copy basis data to GPU 515 CeedCallBackend(CeedBasisGetNumQuadratureComponents(basis, CEED_EVAL_INTERP, &q_comp_interp)); 516 CeedCallBackend(CeedBasisGetNumQuadratureComponents(basis, CEED_EVAL_CURL, &q_comp_curl)); 517 if (q_weight) { 518 CeedCallHip(ceed, hipMalloc((void **)&data->d_q_weight, q_bytes)); 519 CeedCallHip(ceed, hipMemcpy(data->d_q_weight, q_weight, q_bytes, hipMemcpyHostToDevice)); 520 } 521 if (interp) { 522 const CeedInt interp_bytes = q_bytes * num_nodes * q_comp_interp; 523 524 CeedCallHip(ceed, hipMalloc((void **)&data->d_interp, interp_bytes)); 525 CeedCallHip(ceed, hipMemcpy(data->d_interp, interp, interp_bytes, hipMemcpyHostToDevice)); 526 } 527 if (curl) { 528 const CeedInt curl_bytes = q_bytes * num_nodes * q_comp_curl; 529 530 CeedCallHip(ceed, hipMalloc((void **)&data->d_curl, curl_bytes)); 531 CeedCallHip(ceed, hipMemcpy(data->d_curl, curl, curl_bytes, hipMemcpyHostToDevice)); 532 } 533 534 // Compile basis kernels 535 CeedCallBackend(CeedBasisGetNumComponents(basis, &num_comp)); 536 CeedCallBackend(CeedGetJitAbsolutePath(ceed, "ceed/jit-source/hip/hip-ref-basis-nontensor.h", &basis_kernel_path)); 537 CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "----- Loading Basis Kernel Source -----\n"); 538 CeedCallBackend(CeedLoadSourceToBuffer(ceed, basis_kernel_path, &basis_kernel_source)); 539 CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "----- Loading Basis Kernel Source Complete! -----\n"); 540 CeedCallBackend(CeedCompile_Hip(ceed, basis_kernel_source, &data->module, 5, "BASIS_Q", num_qpts, "BASIS_P", num_nodes, "BASIS_Q_COMP_INTERP", 541 q_comp_interp, "BASIS_Q_COMP_DERIV", q_comp_curl, "BASIS_NUM_COMP", num_comp)); 542 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "Interp", &data->Interp)); 543 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "InterpTranspose", &data->InterpTranspose)); 544 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "Deriv", &data->Deriv)); 545 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "DerivTranspose", &data->DerivTranspose)); 546 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "Weight", &data->Weight)); 547 CeedCallBackend(CeedFree(&basis_kernel_path)); 548 CeedCallBackend(CeedFree(&basis_kernel_source)); 549 550 CeedCallBackend(CeedBasisSetData(basis, data)); 551 552 // Register backend functions 553 CeedCallBackend(CeedSetBackendFunction(ceed, "Basis", basis, "Apply", CeedBasisApplyNonTensor_Hip)); 554 CeedCallBackend(CeedSetBackendFunction(ceed, "Basis", basis, "Destroy", CeedBasisDestroyNonTensor_Hip)); 555 return CEED_ERROR_SUCCESS; 556 } 557 558 //------------------------------------------------------------------------------ 559