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 <stdbool.h> 12 #include <stddef.h> 13 #include <hip/hip_runtime.h> 14 15 #include "../hip/ceed-hip-common.h" 16 #include "../hip/ceed-hip-compile.h" 17 #include "ceed-hip-shared.h" 18 19 //------------------------------------------------------------------------------ 20 // Compute a block size based on required minimum threads 21 //------------------------------------------------------------------------------ 22 static CeedInt ComputeBlockSizeFromRequirement(const CeedInt required) { 23 CeedInt maxSize = 1024; // Max total threads per block 24 CeedInt currentSize = 64; // Start with one group 25 26 while (currentSize < maxSize) { 27 if (currentSize > required) break; 28 else currentSize = currentSize * 2; 29 } 30 return currentSize; 31 } 32 33 //------------------------------------------------------------------------------ 34 // Compute required thread block sizes for basis kernels given P, Q, dim, and 35 // num_comp (num_comp not currently used, but may be again in other basis 36 // parallelization options) 37 //------------------------------------------------------------------------------ 38 static int ComputeBasisThreadBlockSizes(const CeedInt dim, const CeedInt P_1d, const CeedInt Q_1d, const CeedInt num_comp, CeedInt *block_sizes) { 39 // Note that this will use the same block sizes for all dimensions when compiling, 40 // but as each basis object is defined for a particular dimension, we will never 41 // call any kernels except the ones for the dimension for which we have computed the 42 // block sizes. 43 const CeedInt thread_1d = CeedIntMax(P_1d, Q_1d); 44 45 switch (dim) { 46 case 1: { 47 // Interp kernels: 48 block_sizes[0] = 256; 49 50 // Grad kernels: 51 block_sizes[1] = 256; 52 53 // Weight kernels: 54 block_sizes[2] = 256; 55 } break; 56 case 2: { 57 // Interp kernels: 58 CeedInt required = thread_1d * thread_1d; 59 60 block_sizes[0] = CeedIntMax(256, ComputeBlockSizeFromRequirement(required)); 61 62 // Grad kernels: currently use same required minimum threads 63 block_sizes[1] = CeedIntMax(256, ComputeBlockSizeFromRequirement(required)); 64 65 // Weight kernels: 66 required = CeedIntMax(64, Q_1d * Q_1d); 67 block_sizes[2] = CeedIntMax(256, ComputeBlockSizeFromRequirement(required)); 68 69 } break; 70 case 3: { 71 // Interp kernels: 72 CeedInt required = thread_1d * thread_1d; 73 74 block_sizes[0] = CeedIntMax(256, ComputeBlockSizeFromRequirement(required)); 75 76 // Grad kernels: currently use same required minimum threads 77 block_sizes[1] = CeedIntMax(256, ComputeBlockSizeFromRequirement(required)); 78 79 // Weight kernels: 80 required = Q_1d * Q_1d * Q_1d; 81 block_sizes[2] = CeedIntMax(256, ComputeBlockSizeFromRequirement(required)); 82 } 83 } 84 return CEED_ERROR_SUCCESS; 85 } 86 87 //------------------------------------------------------------------------------ 88 // Apply basis 89 //------------------------------------------------------------------------------ 90 static int CeedBasisApplyTensorCore_Hip_shared(CeedBasis basis, bool apply_add, const CeedInt num_elem, CeedTransposeMode t_mode, 91 CeedEvalMode eval_mode, CeedVector u, CeedVector v) { 92 Ceed ceed; 93 Ceed_Hip *ceed_Hip; 94 CeedInt dim, num_comp; 95 const CeedScalar *d_u; 96 CeedScalar *d_v; 97 CeedBasis_Hip_shared *data; 98 99 CeedCallBackend(CeedBasisGetCeed(basis, &ceed)); 100 CeedCallBackend(CeedGetData(ceed, &ceed_Hip)); 101 CeedCallBackend(CeedBasisGetData(basis, &data)); 102 CeedCallBackend(CeedBasisGetDimension(basis, &dim)); 103 CeedCallBackend(CeedBasisGetNumComponents(basis, &num_comp)); 104 105 // Get read/write access to u, v 106 if (u != CEED_VECTOR_NONE) CeedCallBackend(CeedVectorGetArrayRead(u, CEED_MEM_DEVICE, &d_u)); 107 else CeedCheck(eval_mode == CEED_EVAL_WEIGHT, ceed, CEED_ERROR_BACKEND, "An input vector is required for this CeedEvalMode"); 108 if (apply_add) CeedCallBackend(CeedVectorGetArray(v, CEED_MEM_DEVICE, &d_v)); 109 else CeedCallBackend(CeedVectorGetArrayWrite(v, CEED_MEM_DEVICE, &d_v)); 110 111 // Apply basis operation 112 switch (eval_mode) { 113 case CEED_EVAL_INTERP: { 114 CeedInt P_1d, Q_1d; 115 CeedInt block_size = data->block_sizes[0]; 116 117 CeedCallBackend(CeedBasisGetNumNodes1D(basis, &P_1d)); 118 CeedCallBackend(CeedBasisGetNumQuadraturePoints1D(basis, &Q_1d)); 119 CeedInt thread_1d = CeedIntMax(Q_1d, P_1d); 120 void *interp_args[] = {(void *)&num_elem, &data->d_interp_1d, &d_u, &d_v}; 121 122 if (dim == 1) { 123 CeedInt elems_per_block = 64 * thread_1d > 256 ? 256 / thread_1d : 64; 124 elems_per_block = elems_per_block > 0 ? elems_per_block : 1; 125 CeedInt grid = num_elem / elems_per_block + ((num_elem / elems_per_block * elems_per_block < num_elem) ? 1 : 0); 126 CeedInt shared_mem = elems_per_block * thread_1d * sizeof(CeedScalar); 127 128 if (t_mode == CEED_TRANSPOSE) { 129 CeedCallBackend(CeedRunKernelDimShared_Hip(ceed, apply_add ? data->InterpTransposeAdd : data->InterpTranspose, grid, thread_1d, 1, 130 elems_per_block, shared_mem, interp_args)); 131 } else { 132 CeedCallBackend(CeedRunKernelDimShared_Hip(ceed, data->Interp, grid, thread_1d, 1, elems_per_block, shared_mem, interp_args)); 133 } 134 } else if (dim == 2) { 135 // Check if required threads is small enough to do multiple elems 136 const CeedInt elems_per_block = CeedIntMax(block_size / (thread_1d * thread_1d), 1); 137 CeedInt grid = num_elem / elems_per_block + ((num_elem / elems_per_block * elems_per_block < num_elem) ? 1 : 0); 138 CeedInt shared_mem = elems_per_block * thread_1d * thread_1d * sizeof(CeedScalar); 139 140 if (t_mode == CEED_TRANSPOSE) { 141 CeedCallBackend(CeedRunKernelDimShared_Hip(ceed, apply_add ? data->InterpTransposeAdd : data->InterpTranspose, grid, thread_1d, thread_1d, 142 elems_per_block, shared_mem, interp_args)); 143 } else { 144 CeedCallBackend(CeedRunKernelDimShared_Hip(ceed, data->Interp, grid, thread_1d, thread_1d, elems_per_block, shared_mem, interp_args)); 145 } 146 } else if (dim == 3) { 147 const CeedInt elems_per_block = CeedIntMax(block_size / (thread_1d * thread_1d), 1); 148 CeedInt grid = num_elem / elems_per_block + ((num_elem / elems_per_block * elems_per_block < num_elem) ? 1 : 0); 149 CeedInt shared_mem = elems_per_block * thread_1d * thread_1d * sizeof(CeedScalar); 150 151 if (t_mode == CEED_TRANSPOSE) { 152 CeedCallBackend(CeedRunKernelDimShared_Hip(ceed, apply_add ? data->InterpTransposeAdd : data->InterpTranspose, grid, thread_1d, thread_1d, 153 elems_per_block, shared_mem, interp_args)); 154 } else { 155 CeedCallBackend(CeedRunKernelDimShared_Hip(ceed, data->Interp, grid, thread_1d, thread_1d, elems_per_block, shared_mem, interp_args)); 156 } 157 } 158 } break; 159 case CEED_EVAL_GRAD: { 160 CeedInt P_1d, Q_1d; 161 CeedInt block_size = data->block_sizes[1]; 162 163 CeedCallBackend(CeedBasisGetNumNodes1D(basis, &P_1d)); 164 CeedCallBackend(CeedBasisGetNumQuadraturePoints1D(basis, &Q_1d)); 165 CeedInt thread_1d = CeedIntMax(Q_1d, P_1d); 166 CeedScalar *d_grad_1d = data->d_grad_1d; 167 168 if (data->d_collo_grad_1d) { 169 d_grad_1d = data->d_collo_grad_1d; 170 } 171 void *grad_args[] = {(void *)&num_elem, &data->d_interp_1d, &d_grad_1d, &d_u, &d_v}; 172 if (dim == 1) { 173 CeedInt elems_per_block = 64 * thread_1d > 256 ? 256 / thread_1d : 64; 174 elems_per_block = elems_per_block > 0 ? elems_per_block : 1; 175 CeedInt grid = num_elem / elems_per_block + ((num_elem / elems_per_block * elems_per_block < num_elem) ? 1 : 0); 176 CeedInt shared_mem = elems_per_block * thread_1d * sizeof(CeedScalar); 177 178 if (t_mode == CEED_TRANSPOSE) { 179 CeedCallBackend(CeedRunKernelDimShared_Hip(ceed, apply_add ? data->GradTransposeAdd : data->GradTranspose, grid, thread_1d, 1, 180 elems_per_block, shared_mem, grad_args)); 181 } else { 182 CeedCallBackend(CeedRunKernelDimShared_Hip(ceed, data->Grad, grid, thread_1d, 1, elems_per_block, shared_mem, grad_args)); 183 } 184 } else if (dim == 2) { 185 // Check if required threads is small enough to do multiple elems 186 const CeedInt elems_per_block = CeedIntMax(block_size / (thread_1d * thread_1d), 1); 187 CeedInt grid = num_elem / elems_per_block + ((num_elem / elems_per_block * elems_per_block < num_elem) ? 1 : 0); 188 CeedInt shared_mem = elems_per_block * thread_1d * thread_1d * sizeof(CeedScalar); 189 190 if (t_mode == CEED_TRANSPOSE) { 191 CeedCallBackend(CeedRunKernelDimShared_Hip(ceed, apply_add ? data->GradTransposeAdd : data->GradTranspose, grid, thread_1d, thread_1d, 192 elems_per_block, shared_mem, grad_args)); 193 } else { 194 CeedCallBackend(CeedRunKernelDimShared_Hip(ceed, data->Grad, grid, thread_1d, thread_1d, elems_per_block, shared_mem, grad_args)); 195 } 196 } else if (dim == 3) { 197 const CeedInt elems_per_block = CeedIntMax(block_size / (thread_1d * thread_1d), 1); 198 CeedInt grid = num_elem / elems_per_block + ((num_elem / elems_per_block * elems_per_block < num_elem) ? 1 : 0); 199 CeedInt shared_mem = elems_per_block * thread_1d * thread_1d * sizeof(CeedScalar); 200 201 if (t_mode == CEED_TRANSPOSE) { 202 CeedCallBackend(CeedRunKernelDimShared_Hip(ceed, apply_add ? data->GradTransposeAdd : data->GradTranspose, grid, thread_1d, thread_1d, 203 elems_per_block, shared_mem, grad_args)); 204 } else { 205 CeedCallBackend(CeedRunKernelDimShared_Hip(ceed, data->Grad, grid, thread_1d, thread_1d, elems_per_block, shared_mem, grad_args)); 206 } 207 } 208 } break; 209 case CEED_EVAL_WEIGHT: { 210 CeedInt Q_1d; 211 CeedInt block_size = data->block_sizes[2]; 212 213 CeedCheck(data->d_q_weight_1d, ceed, CEED_ERROR_BACKEND, "%s not supported; q_weights_1d not set", CeedEvalModes[eval_mode]); 214 CeedCallBackend(CeedBasisGetNumQuadraturePoints1D(basis, &Q_1d)); 215 void *weight_args[] = {(void *)&num_elem, (void *)&data->d_q_weight_1d, &d_v}; 216 217 if (dim == 1) { 218 const CeedInt opt_elems = block_size / Q_1d; 219 const CeedInt elems_per_block = opt_elems > 0 ? opt_elems : 1; 220 const CeedInt grid_size = num_elem / elems_per_block + ((num_elem / elems_per_block * elems_per_block < num_elem) ? 1 : 0); 221 222 CeedCallBackend(CeedRunKernelDim_Hip(ceed, data->Weight, grid_size, Q_1d, elems_per_block, 1, weight_args)); 223 } else if (dim == 2) { 224 const CeedInt opt_elems = block_size / (Q_1d * Q_1d); 225 const CeedInt elems_per_block = opt_elems > 0 ? opt_elems : 1; 226 const CeedInt grid_size = num_elem / elems_per_block + ((num_elem / elems_per_block * elems_per_block < num_elem) ? 1 : 0); 227 228 CeedCallBackend(CeedRunKernelDim_Hip(ceed, data->Weight, grid_size, Q_1d, Q_1d, elems_per_block, weight_args)); 229 } else if (dim == 3) { 230 const CeedInt opt_elems = block_size / (Q_1d * Q_1d); 231 const CeedInt elems_per_block = opt_elems > 0 ? opt_elems : 1; 232 const CeedInt grid_size = num_elem / elems_per_block + ((num_elem / elems_per_block * elems_per_block < num_elem) ? 1 : 0); 233 234 CeedCallBackend(CeedRunKernelDim_Hip(ceed, data->Weight, grid_size, Q_1d, Q_1d, elems_per_block, weight_args)); 235 } 236 } break; 237 case CEED_EVAL_NONE: /* handled separately below */ 238 break; 239 // LCOV_EXCL_START 240 case CEED_EVAL_DIV: 241 case CEED_EVAL_CURL: 242 return CeedError(ceed, CEED_ERROR_BACKEND, "%s not supported", CeedEvalModes[eval_mode]); 243 // LCOV_EXCL_STOP 244 } 245 246 // Restore vectors, cover CEED_EVAL_NONE 247 CeedCallBackend(CeedVectorRestoreArray(v, &d_v)); 248 if (eval_mode == CEED_EVAL_NONE) CeedCallBackend(CeedVectorSetArray(v, CEED_MEM_DEVICE, CEED_COPY_VALUES, (CeedScalar *)d_u)); 249 if (eval_mode != CEED_EVAL_WEIGHT) CeedCallBackend(CeedVectorRestoreArrayRead(u, &d_u)); 250 return CEED_ERROR_SUCCESS; 251 } 252 253 int CeedBasisApplyTensor_Hip_shared(CeedBasis basis, const CeedInt num_elem, CeedTransposeMode t_mode, CeedEvalMode eval_mode, CeedVector u, 254 CeedVector v) { 255 CeedCallBackend(CeedBasisApplyTensorCore_Hip_shared(basis, false, num_elem, t_mode, eval_mode, u, v)); 256 return CEED_ERROR_SUCCESS; 257 } 258 259 int CeedBasisApplyAddTensor_Hip_shared(CeedBasis basis, const CeedInt num_elem, CeedTransposeMode t_mode, CeedEvalMode eval_mode, CeedVector u, 260 CeedVector v) { 261 CeedCallBackend(CeedBasisApplyTensorCore_Hip_shared(basis, true, num_elem, t_mode, eval_mode, u, v)); 262 return CEED_ERROR_SUCCESS; 263 } 264 265 //------------------------------------------------------------------------------ 266 // Basis apply - tensor AtPoints 267 //------------------------------------------------------------------------------ 268 static int CeedBasisApplyAtPointsCore_Hip_shared(CeedBasis basis, bool apply_add, const CeedInt num_elem, const CeedInt *num_points, 269 CeedTransposeMode t_mode, CeedEvalMode eval_mode, CeedVector x_ref, CeedVector u, CeedVector v) { 270 Ceed ceed; 271 CeedInt Q_1d, dim, max_num_points = num_points[0]; 272 const CeedInt is_transpose = t_mode == CEED_TRANSPOSE; 273 const int max_block_size = 32; 274 const CeedScalar *d_x, *d_u; 275 CeedScalar *d_v; 276 CeedBasis_Hip_shared *data; 277 278 CeedCallBackend(CeedBasisGetCeed(basis, &ceed)); 279 CeedCallBackend(CeedBasisGetData(basis, &data)); 280 CeedCallBackend(CeedBasisGetNumQuadraturePoints1D(basis, &Q_1d)); 281 CeedCallBackend(CeedBasisGetDimension(basis, &dim)); 282 283 // Check uniform number of points per elem 284 for (CeedInt i = 1; i < num_elem; i++) { 285 CeedCheck(max_num_points == num_points[i], ceed, CEED_ERROR_BACKEND, 286 "BasisApplyAtPoints only supported for the same number of points in each element"); 287 } 288 289 // Weight handled separately 290 if (eval_mode == CEED_EVAL_WEIGHT) { 291 CeedCallBackend(CeedVectorSetValue(v, 1.0)); 292 return CEED_ERROR_SUCCESS; 293 } 294 295 // Build kernels if needed 296 if (data->num_points != max_num_points) { 297 CeedInt P_1d; 298 299 CeedCallBackend(CeedBasisGetNumNodes1D(basis, &P_1d)); 300 data->num_points = max_num_points; 301 302 // -- Create interp matrix to Chebyshev coefficients 303 if (!data->d_chebyshev_interp_1d) { 304 CeedSize interp_bytes; 305 CeedScalar *chebyshev_interp_1d; 306 307 interp_bytes = P_1d * Q_1d * sizeof(CeedScalar); 308 CeedCallBackend(CeedCalloc(P_1d * Q_1d, &chebyshev_interp_1d)); 309 CeedCallBackend(CeedBasisGetChebyshevInterp1D(basis, chebyshev_interp_1d)); 310 CeedCallHip(ceed, hipMalloc((void **)&data->d_chebyshev_interp_1d, interp_bytes)); 311 CeedCallHip(ceed, hipMemcpy(data->d_chebyshev_interp_1d, chebyshev_interp_1d, interp_bytes, hipMemcpyHostToDevice)); 312 CeedCallBackend(CeedFree(&chebyshev_interp_1d)); 313 } 314 315 // -- Compile kernels 316 char *basis_kernel_source; 317 const char *basis_kernel_path; 318 CeedInt num_comp; 319 320 if (data->moduleAtPoints) CeedCallHip(ceed, hipModuleUnload(data->moduleAtPoints)); 321 CeedCallBackend(CeedBasisGetNumComponents(basis, &num_comp)); 322 CeedCallBackend(CeedGetJitAbsolutePath(ceed, "ceed/jit-source/hip/hip-ref-basis-tensor-at-points.h", &basis_kernel_path)); 323 CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "----- Loading Basis Kernel Source -----\n"); 324 CeedCallBackend(CeedLoadSourceToBuffer(ceed, basis_kernel_path, &basis_kernel_source)); 325 CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "----- Loading Basis Kernel Source Complete! -----\n"); 326 CeedCallBackend(CeedCompile_Hip(ceed, basis_kernel_source, &data->moduleAtPoints, 9, "BASIS_Q_1D", Q_1d, "BASIS_P_1D", P_1d, "BASIS_BUF_LEN", 327 Q_1d * CeedIntPow(Q_1d > P_1d ? Q_1d : P_1d, dim - 1), "BASIS_DIM", dim, "BASIS_NUM_COMP", num_comp, 328 "BASIS_NUM_NODES", CeedIntPow(P_1d, dim), "BASIS_NUM_QPTS", CeedIntPow(Q_1d, dim), "BASIS_NUM_PTS", 329 max_num_points, "POINTS_BUFF_LEN", CeedIntPow(Q_1d, dim - 1))); 330 CeedCallBackend(CeedGetKernel_Hip(ceed, data->moduleAtPoints, "InterpAtPoints", &data->InterpAtPoints)); 331 CeedCallBackend(CeedGetKernel_Hip(ceed, data->moduleAtPoints, "GradAtPoints", &data->GradAtPoints)); 332 CeedCallBackend(CeedFree(&basis_kernel_path)); 333 CeedCallBackend(CeedFree(&basis_kernel_source)); 334 } 335 336 // Get read/write access to u, v 337 CeedCallBackend(CeedVectorGetArrayRead(x_ref, CEED_MEM_DEVICE, &d_x)); 338 if (u != CEED_VECTOR_NONE) CeedCallBackend(CeedVectorGetArrayRead(u, CEED_MEM_DEVICE, &d_u)); 339 else CeedCheck(eval_mode == CEED_EVAL_WEIGHT, ceed, CEED_ERROR_BACKEND, "An input vector is required for this CeedEvalMode"); 340 if (apply_add) CeedCallBackend(CeedVectorGetArray(v, CEED_MEM_DEVICE, &d_v)); 341 else CeedCallBackend(CeedVectorGetArrayWrite(v, CEED_MEM_DEVICE, &d_v)); 342 343 // Clear v for transpose operation 344 if (is_transpose && !apply_add) { 345 CeedSize length; 346 347 CeedCallBackend(CeedVectorGetLength(v, &length)); 348 CeedCallHip(ceed, hipMemset(d_v, 0, length * sizeof(CeedScalar))); 349 } 350 351 // Basis action 352 switch (eval_mode) { 353 case CEED_EVAL_INTERP: { 354 void *interp_args[] = {(void *)&num_elem, (void *)&is_transpose, &data->d_chebyshev_interp_1d, &d_x, &d_u, &d_v}; 355 const CeedInt block_size = CeedIntMin(CeedIntPow(Q_1d, dim), max_block_size); 356 357 CeedCallBackend(CeedRunKernel_Hip(ceed, data->InterpAtPoints, num_elem, block_size, interp_args)); 358 } break; 359 case CEED_EVAL_GRAD: { 360 void *grad_args[] = {(void *)&num_elem, (void *)&is_transpose, &data->d_chebyshev_interp_1d, &d_x, &d_u, &d_v}; 361 const CeedInt block_size = CeedIntMin(CeedIntPow(Q_1d, dim), max_block_size); 362 363 CeedCallBackend(CeedRunKernel_Hip(ceed, data->GradAtPoints, num_elem, block_size, grad_args)); 364 } break; 365 case CEED_EVAL_WEIGHT: 366 case CEED_EVAL_NONE: /* handled separately below */ 367 break; 368 // LCOV_EXCL_START 369 case CEED_EVAL_DIV: 370 case CEED_EVAL_CURL: 371 return CeedError(ceed, CEED_ERROR_BACKEND, "%s not supported", CeedEvalModes[eval_mode]); 372 // LCOV_EXCL_STOP 373 } 374 375 // Restore vectors, cover CEED_EVAL_NONE 376 CeedCallBackend(CeedVectorRestoreArrayRead(x_ref, &d_x)); 377 CeedCallBackend(CeedVectorRestoreArray(v, &d_v)); 378 if (eval_mode == CEED_EVAL_NONE) CeedCallBackend(CeedVectorSetArray(v, CEED_MEM_DEVICE, CEED_COPY_VALUES, (CeedScalar *)d_u)); 379 if (eval_mode != CEED_EVAL_WEIGHT) CeedCallBackend(CeedVectorRestoreArrayRead(u, &d_u)); 380 return CEED_ERROR_SUCCESS; 381 } 382 383 static int CeedBasisApplyAtPoints_Hip_shared(CeedBasis basis, const CeedInt num_elem, const CeedInt *num_points, CeedTransposeMode t_mode, 384 CeedEvalMode eval_mode, CeedVector x_ref, CeedVector u, CeedVector v) { 385 CeedCallBackend(CeedBasisApplyAtPointsCore_Hip_shared(basis, false, num_elem, num_points, t_mode, eval_mode, x_ref, u, v)); 386 return CEED_ERROR_SUCCESS; 387 } 388 389 static int CeedBasisApplyAddAtPoints_Hip_shared(CeedBasis basis, const CeedInt num_elem, const CeedInt *num_points, CeedTransposeMode t_mode, 390 CeedEvalMode eval_mode, CeedVector x_ref, CeedVector u, CeedVector v) { 391 CeedCallBackend(CeedBasisApplyAtPointsCore_Hip_shared(basis, true, num_elem, num_points, t_mode, eval_mode, x_ref, u, v)); 392 return CEED_ERROR_SUCCESS; 393 } 394 395 //------------------------------------------------------------------------------ 396 // Destroy basis 397 //------------------------------------------------------------------------------ 398 static int CeedBasisDestroy_Hip_shared(CeedBasis basis) { 399 Ceed ceed; 400 CeedBasis_Hip_shared *data; 401 402 CeedCallBackend(CeedBasisGetCeed(basis, &ceed)); 403 CeedCallBackend(CeedBasisGetData(basis, &data)); 404 CeedCallHip(ceed, hipModuleUnload(data->module)); 405 if (data->moduleAtPoints) CeedCallHip(ceed, hipModuleUnload(data->moduleAtPoints)); 406 if (data->d_q_weight_1d) CeedCallHip(ceed, hipFree(data->d_q_weight_1d)); 407 CeedCallHip(ceed, hipFree(data->d_interp_1d)); 408 CeedCallHip(ceed, hipFree(data->d_grad_1d)); 409 CeedCallHip(ceed, hipFree(data->d_collo_grad_1d)); 410 CeedCallHip(ceed, hipFree(data->d_chebyshev_interp_1d)); 411 CeedCallBackend(CeedFree(&data)); 412 return CEED_ERROR_SUCCESS; 413 } 414 415 //------------------------------------------------------------------------------ 416 // Create tensor basis 417 //------------------------------------------------------------------------------ 418 int CeedBasisCreateTensorH1_Hip_shared(CeedInt dim, CeedInt P_1d, CeedInt Q_1d, const CeedScalar *interp_1d, const CeedScalar *grad_1d, 419 const CeedScalar *q_ref_1d, const CeedScalar *q_weight_1d, CeedBasis basis) { 420 Ceed ceed; 421 char *basis_kernel_source; 422 const char *basis_kernel_path; 423 CeedInt num_comp; 424 const CeedInt q_bytes = Q_1d * sizeof(CeedScalar); 425 const CeedInt interp_bytes = q_bytes * P_1d; 426 CeedBasis_Hip_shared *data; 427 428 CeedCallBackend(CeedBasisGetCeed(basis, &ceed)); 429 CeedCallBackend(CeedCalloc(1, &data)); 430 431 // Copy basis data to GPU 432 if (q_weight_1d) { 433 CeedCallHip(ceed, hipMalloc((void **)&data->d_q_weight_1d, q_bytes)); 434 CeedCallHip(ceed, hipMemcpy(data->d_q_weight_1d, q_weight_1d, q_bytes, hipMemcpyHostToDevice)); 435 } 436 CeedCallHip(ceed, hipMalloc((void **)&data->d_interp_1d, interp_bytes)); 437 CeedCallHip(ceed, hipMemcpy(data->d_interp_1d, interp_1d, interp_bytes, hipMemcpyHostToDevice)); 438 CeedCallHip(ceed, hipMalloc((void **)&data->d_grad_1d, interp_bytes)); 439 CeedCallHip(ceed, hipMemcpy(data->d_grad_1d, grad_1d, interp_bytes, hipMemcpyHostToDevice)); 440 441 // Compute collocated gradient and copy to GPU 442 data->d_collo_grad_1d = NULL; 443 bool has_collocated_grad = dim == 3 && Q_1d >= P_1d; 444 445 if (has_collocated_grad) { 446 CeedScalar *collo_grad_1d; 447 448 CeedCallBackend(CeedMalloc(Q_1d * Q_1d, &collo_grad_1d)); 449 CeedCallBackend(CeedBasisGetCollocatedGrad(basis, collo_grad_1d)); 450 CeedCallHip(ceed, hipMalloc((void **)&data->d_collo_grad_1d, q_bytes * Q_1d)); 451 CeedCallHip(ceed, hipMemcpy(data->d_collo_grad_1d, collo_grad_1d, q_bytes * Q_1d, hipMemcpyHostToDevice)); 452 CeedCallBackend(CeedFree(&collo_grad_1d)); 453 } 454 455 // Set number of threads per block for basis kernels 456 CeedCallBackend(CeedBasisGetNumComponents(basis, &num_comp)); 457 CeedCallBackend(ComputeBasisThreadBlockSizes(dim, P_1d, Q_1d, num_comp, data->block_sizes)); 458 459 // Compile basis kernels 460 CeedCallBackend(CeedGetJitAbsolutePath(ceed, "ceed/jit-source/hip/hip-shared-basis-tensor.h", &basis_kernel_path)); 461 CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "----- Loading Basis Kernel Source -----\n"); 462 CeedCallBackend(CeedLoadSourceToBuffer(ceed, basis_kernel_path, &basis_kernel_source)); 463 CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "----- Loading Basis Kernel Source Complete! -----\n"); 464 CeedCallBackend(CeedCompile_Hip(ceed, basis_kernel_source, &data->module, 11, "BASIS_Q_1D", Q_1d, "BASIS_P_1D", P_1d, "T_1D", 465 CeedIntMax(Q_1d, P_1d), "BASIS_DIM", dim, "BASIS_NUM_COMP", num_comp, "BASIS_NUM_NODES", CeedIntPow(P_1d, dim), 466 "BASIS_NUM_QPTS", CeedIntPow(Q_1d, dim), "BASIS_INTERP_BLOCK_SIZE", data->block_sizes[0], "BASIS_GRAD_BLOCK_SIZE", 467 data->block_sizes[1], "BASIS_WEIGHT_BLOCK_SIZE", data->block_sizes[2], "BASIS_HAS_COLLOCATED_GRAD", 468 has_collocated_grad)); 469 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "Interp", &data->Interp)); 470 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "InterpTranspose", &data->InterpTranspose)); 471 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "InterpTransposeAdd", &data->InterpTransposeAdd)); 472 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "Grad", &data->Grad)); 473 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "GradTranspose", &data->GradTranspose)); 474 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "GradTransposeAdd", &data->GradTransposeAdd)); 475 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "Weight", &data->Weight)); 476 CeedCallBackend(CeedFree(&basis_kernel_path)); 477 CeedCallBackend(CeedFree(&basis_kernel_source)); 478 479 CeedCallBackend(CeedBasisSetData(basis, data)); 480 481 // Register backend functions 482 CeedCallBackend(CeedSetBackendFunction(ceed, "Basis", basis, "Apply", CeedBasisApplyTensor_Hip_shared)); 483 CeedCallBackend(CeedSetBackendFunction(ceed, "Basis", basis, "ApplyAdd", CeedBasisApplyAddTensor_Hip_shared)); 484 CeedCallBackend(CeedSetBackendFunction(ceed, "Basis", basis, "ApplyAtPoints", CeedBasisApplyAtPoints_Hip_shared)); 485 CeedCallBackend(CeedSetBackendFunction(ceed, "Basis", basis, "ApplyAddAtPoints", CeedBasisApplyAddAtPoints_Hip_shared)); 486 CeedCallBackend(CeedSetBackendFunction(ceed, "Basis", basis, "Destroy", CeedBasisDestroy_Hip_shared)); 487 return CEED_ERROR_SUCCESS; 488 } 489 490 //------------------------------------------------------------------------------ 491