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 <cuda.h> 12 #include <cuda_runtime.h> 13 14 #include "../cuda/ceed-cuda-common.h" 15 #include "../cuda/ceed-cuda-compile.h" 16 #include "ceed-cuda-ref.h" 17 18 //------------------------------------------------------------------------------ 19 // Basis apply - tensor 20 //------------------------------------------------------------------------------ 21 int CeedBasisApply_Cuda(CeedBasis basis, const CeedInt num_elem, CeedTransposeMode t_mode, CeedEvalMode eval_mode, CeedVector u, CeedVector v) { 22 Ceed ceed; 23 CeedInt Q_1d, dim; 24 const CeedInt is_transpose = t_mode == CEED_TRANSPOSE; 25 const int max_block_size = 32; 26 const CeedScalar *d_u; 27 CeedScalar *d_v; 28 CeedBasis_Cuda *data; 29 30 CeedCallBackend(CeedBasisGetCeed(basis, &ceed)); 31 CeedCallBackend(CeedBasisGetData(basis, &data)); 32 33 // Get read/write access to u, v 34 if (u != CEED_VECTOR_NONE) CeedCallBackend(CeedVectorGetArrayRead(u, CEED_MEM_DEVICE, &d_u)); 35 else CeedCheck(eval_mode == CEED_EVAL_WEIGHT, ceed, CEED_ERROR_BACKEND, "An input vector is required for this CeedEvalMode"); 36 CeedCallBackend(CeedVectorGetArrayWrite(v, CEED_MEM_DEVICE, &d_v)); 37 38 // Clear v for transpose operation 39 if (is_transpose) { 40 CeedSize length; 41 42 CeedCallBackend(CeedVectorGetLength(v, &length)); 43 CeedCallCuda(ceed, cudaMemset(d_v, 0, length * sizeof(CeedScalar))); 44 } 45 CeedCallBackend(CeedBasisGetNumQuadraturePoints1D(basis, &Q_1d)); 46 CeedCallBackend(CeedBasisGetDimension(basis, &dim)); 47 48 // Basis action 49 switch (eval_mode) { 50 case CEED_EVAL_INTERP: { 51 void *interp_args[] = {(void *)&num_elem, (void *)&is_transpose, &data->d_interp_1d, &d_u, &d_v}; 52 const CeedInt block_size = CeedIntMin(CeedIntPow(Q_1d, dim), max_block_size); 53 54 CeedCallBackend(CeedRunKernel_Cuda(ceed, data->Interp, num_elem, block_size, interp_args)); 55 } break; 56 case CEED_EVAL_GRAD: { 57 void *grad_args[] = {(void *)&num_elem, (void *)&is_transpose, &data->d_interp_1d, &data->d_grad_1d, &d_u, &d_v}; 58 const CeedInt block_size = max_block_size; 59 60 CeedCallBackend(CeedRunKernel_Cuda(ceed, data->Grad, num_elem, block_size, grad_args)); 61 } break; 62 case CEED_EVAL_WEIGHT: { 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_Cuda(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_Cuda(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_Cuda *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 CeedCallCuda(ceed, cudaMemset(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_Cuda(ceed, data->InterpTranspose, grid, block_size_x, 1, elems_per_block, interp_args)); 125 } else { 126 CeedCallBackend(CeedRunKernelDim_Cuda(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_Cuda(ceed, data->DerivTranspose, grid, block_size_x, 1, elems_per_block, grad_args)); 135 } else { 136 CeedCallBackend(CeedRunKernelDim_Cuda(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_Cuda(ceed, data->DerivTranspose, grid, block_size_x, 1, elems_per_block, div_args)); 145 } else { 146 CeedCallBackend(CeedRunKernelDim_Cuda(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_Cuda(ceed, data->DerivTranspose, grid, block_size_x, 1, elems_per_block, curl_args)); 155 } else { 156 CeedCallBackend(CeedRunKernelDim_Cuda(ceed, data->Deriv, grid, block_size_x, 1, elems_per_block, curl_args)); 157 } 158 } break; 159 case CEED_EVAL_WEIGHT: { 160 void *weight_args[] = {(void *)&num_elem, (void *)&data->d_q_weight, &d_v}; 161 162 CeedCallBackend(CeedRunKernelDim_Cuda(ceed, data->Weight, grid, num_qpts, 1, elems_per_block, weight_args)); 163 } break; 164 case CEED_EVAL_NONE: /* handled separately below */ 165 break; 166 } 167 168 // Restore vectors, cover CEED_EVAL_NONE 169 CeedCallBackend(CeedVectorRestoreArray(v, &d_v)); 170 if (eval_mode == CEED_EVAL_NONE) CeedCallBackend(CeedVectorSetArray(v, CEED_MEM_DEVICE, CEED_COPY_VALUES, (CeedScalar *)d_u)); 171 if (eval_mode != CEED_EVAL_WEIGHT) CeedCallBackend(CeedVectorRestoreArrayRead(u, &d_u)); 172 return CEED_ERROR_SUCCESS; 173 } 174 175 //------------------------------------------------------------------------------ 176 // Destroy tensor basis 177 //------------------------------------------------------------------------------ 178 static int CeedBasisDestroy_Cuda(CeedBasis basis) { 179 Ceed ceed; 180 CeedBasis_Cuda *data; 181 182 CeedCallBackend(CeedBasisGetCeed(basis, &ceed)); 183 CeedCallBackend(CeedBasisGetData(basis, &data)); 184 CeedCallCuda(ceed, cuModuleUnload(data->module)); 185 CeedCallCuda(ceed, cudaFree(data->d_q_weight_1d)); 186 CeedCallCuda(ceed, cudaFree(data->d_interp_1d)); 187 CeedCallCuda(ceed, cudaFree(data->d_grad_1d)); 188 CeedCallBackend(CeedFree(&data)); 189 return CEED_ERROR_SUCCESS; 190 } 191 192 //------------------------------------------------------------------------------ 193 // Destroy non-tensor basis 194 //------------------------------------------------------------------------------ 195 static int CeedBasisDestroyNonTensor_Cuda(CeedBasis basis) { 196 Ceed ceed; 197 CeedBasisNonTensor_Cuda *data; 198 199 CeedCallBackend(CeedBasisGetCeed(basis, &ceed)); 200 CeedCallBackend(CeedBasisGetData(basis, &data)); 201 CeedCallCuda(ceed, cuModuleUnload(data->module)); 202 CeedCallCuda(ceed, cudaFree(data->d_q_weight)); 203 CeedCallCuda(ceed, cudaFree(data->d_interp)); 204 CeedCallCuda(ceed, cudaFree(data->d_grad)); 205 CeedCallCuda(ceed, cudaFree(data->d_div)); 206 CeedCallCuda(ceed, cudaFree(data->d_curl)); 207 CeedCallBackend(CeedFree(&data)); 208 return CEED_ERROR_SUCCESS; 209 } 210 211 //------------------------------------------------------------------------------ 212 // Create tensor 213 //------------------------------------------------------------------------------ 214 int CeedBasisCreateTensorH1_Cuda(CeedInt dim, CeedInt P_1d, CeedInt Q_1d, const CeedScalar *interp_1d, const CeedScalar *grad_1d, 215 const CeedScalar *q_ref_1d, const CeedScalar *q_weight_1d, CeedBasis basis) { 216 Ceed ceed; 217 char *basis_kernel_source; 218 const char *basis_kernel_path; 219 CeedInt num_comp; 220 const CeedInt q_bytes = Q_1d * sizeof(CeedScalar); 221 const CeedInt interp_bytes = q_bytes * P_1d; 222 CeedBasis_Cuda *data; 223 224 CeedCallBackend(CeedBasisGetCeed(basis, &ceed)); 225 CeedCallBackend(CeedCalloc(1, &data)); 226 227 // Copy data to GPU 228 CeedCallCuda(ceed, cudaMalloc((void **)&data->d_q_weight_1d, q_bytes)); 229 CeedCallCuda(ceed, cudaMemcpy(data->d_q_weight_1d, q_weight_1d, q_bytes, cudaMemcpyHostToDevice)); 230 CeedCallCuda(ceed, cudaMalloc((void **)&data->d_interp_1d, interp_bytes)); 231 CeedCallCuda(ceed, cudaMemcpy(data->d_interp_1d, interp_1d, interp_bytes, cudaMemcpyHostToDevice)); 232 CeedCallCuda(ceed, cudaMalloc((void **)&data->d_grad_1d, interp_bytes)); 233 CeedCallCuda(ceed, cudaMemcpy(data->d_grad_1d, grad_1d, interp_bytes, cudaMemcpyHostToDevice)); 234 235 // Compile basis kernels 236 CeedCallBackend(CeedBasisGetNumComponents(basis, &num_comp)); 237 CeedCallBackend(CeedGetJitAbsolutePath(ceed, "ceed/jit-source/cuda/cuda-ref-basis-tensor.h", &basis_kernel_path)); 238 CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "----- Loading Basis Kernel Source -----\n"); 239 CeedCallBackend(CeedLoadSourceToBuffer(ceed, basis_kernel_path, &basis_kernel_source)); 240 CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "----- Loading Basis Kernel Source Complete! -----\n"); 241 CeedCallBackend(CeedCompile_Cuda(ceed, basis_kernel_source, &data->module, 7, "BASIS_Q_1D", Q_1d, "BASIS_P_1D", P_1d, "BASIS_BUF_LEN", 242 num_comp * CeedIntPow(Q_1d > P_1d ? Q_1d : P_1d, dim), "BASIS_DIM", dim, "BASIS_NUM_COMP", num_comp, 243 "BASIS_NUM_NODES", CeedIntPow(P_1d, dim), "BASIS_NUM_QPTS", CeedIntPow(Q_1d, dim))); 244 CeedCallBackend(CeedGetKernel_Cuda(ceed, data->module, "Interp", &data->Interp)); 245 CeedCallBackend(CeedGetKernel_Cuda(ceed, data->module, "Grad", &data->Grad)); 246 CeedCallBackend(CeedGetKernel_Cuda(ceed, data->module, "Weight", &data->Weight)); 247 CeedCallBackend(CeedFree(&basis_kernel_path)); 248 CeedCallBackend(CeedFree(&basis_kernel_source)); 249 250 CeedCallBackend(CeedBasisSetData(basis, data)); 251 252 // Register backend functions 253 CeedCallBackend(CeedSetBackendFunction(ceed, "Basis", basis, "Apply", CeedBasisApply_Cuda)); 254 CeedCallBackend(CeedSetBackendFunction(ceed, "Basis", basis, "Destroy", CeedBasisDestroy_Cuda)); 255 return CEED_ERROR_SUCCESS; 256 } 257 258 //------------------------------------------------------------------------------ 259 // Create non-tensor H^1 260 //------------------------------------------------------------------------------ 261 int CeedBasisCreateH1_Cuda(CeedElemTopology topo, CeedInt dim, CeedInt num_nodes, CeedInt num_qpts, const CeedScalar *interp, const CeedScalar *grad, 262 const CeedScalar *q_ref, const CeedScalar *q_weight, CeedBasis basis) { 263 Ceed ceed; 264 char *basis_kernel_source; 265 const char *basis_kernel_path; 266 CeedInt num_comp, q_comp_interp, q_comp_grad; 267 const CeedInt q_bytes = num_qpts * sizeof(CeedScalar); 268 CeedBasisNonTensor_Cuda *data; 269 270 CeedCallBackend(CeedBasisGetCeed(basis, &ceed)); 271 CeedCallBackend(CeedCalloc(1, &data)); 272 273 // Copy basis data to GPU 274 CeedCallBackend(CeedBasisGetNumQuadratureComponents(basis, CEED_EVAL_INTERP, &q_comp_interp)); 275 CeedCallBackend(CeedBasisGetNumQuadratureComponents(basis, CEED_EVAL_GRAD, &q_comp_grad)); 276 CeedCallCuda(ceed, cudaMalloc((void **)&data->d_q_weight, q_bytes)); 277 CeedCallCuda(ceed, cudaMemcpy(data->d_q_weight, q_weight, q_bytes, cudaMemcpyHostToDevice)); 278 if (interp) { 279 const CeedInt interp_bytes = q_bytes * num_nodes * q_comp_interp; 280 281 CeedCallCuda(ceed, cudaMalloc((void **)&data->d_interp, interp_bytes)); 282 CeedCallCuda(ceed, cudaMemcpy(data->d_interp, interp, interp_bytes, cudaMemcpyHostToDevice)); 283 } 284 if (grad) { 285 const CeedInt grad_bytes = q_bytes * num_nodes * q_comp_grad; 286 287 CeedCallCuda(ceed, cudaMalloc((void **)&data->d_grad, grad_bytes)); 288 CeedCallCuda(ceed, cudaMemcpy(data->d_grad, grad, grad_bytes, cudaMemcpyHostToDevice)); 289 } 290 291 // Compile basis kernels 292 CeedCallBackend(CeedBasisGetNumComponents(basis, &num_comp)); 293 CeedCallBackend(CeedGetJitAbsolutePath(ceed, "ceed/jit-source/cuda/cuda-ref-basis-nontensor.h", &basis_kernel_path)); 294 CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "----- Loading Basis Kernel Source -----\n"); 295 CeedCallBackend(CeedLoadSourceToBuffer(ceed, basis_kernel_path, &basis_kernel_source)); 296 CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "----- Loading Basis Kernel Source Complete! -----\n"); 297 CeedCallBackend(CeedCompile_Cuda(ceed, basis_kernel_source, &data->module, 5, "BASIS_Q", num_qpts, "BASIS_P", num_nodes, "BASIS_Q_COMP_INTERP", 298 q_comp_interp, "BASIS_Q_COMP_DERIV", q_comp_grad, "BASIS_NUM_COMP", num_comp)); 299 CeedCallBackend(CeedGetKernel_Cuda(ceed, data->module, "Interp", &data->Interp)); 300 CeedCallBackend(CeedGetKernel_Cuda(ceed, data->module, "InterpTranspose", &data->InterpTranspose)); 301 CeedCallBackend(CeedGetKernel_Cuda(ceed, data->module, "Deriv", &data->Deriv)); 302 CeedCallBackend(CeedGetKernel_Cuda(ceed, data->module, "DerivTranspose", &data->DerivTranspose)); 303 CeedCallBackend(CeedGetKernel_Cuda(ceed, data->module, "Weight", &data->Weight)); 304 CeedCallBackend(CeedFree(&basis_kernel_path)); 305 CeedCallBackend(CeedFree(&basis_kernel_source)); 306 307 CeedCallBackend(CeedBasisSetData(basis, data)); 308 309 // Register backend functions 310 CeedCallBackend(CeedSetBackendFunction(ceed, "Basis", basis, "Apply", CeedBasisApplyNonTensor_Cuda)); 311 CeedCallBackend(CeedSetBackendFunction(ceed, "Basis", basis, "Destroy", CeedBasisDestroyNonTensor_Cuda)); 312 return CEED_ERROR_SUCCESS; 313 } 314 315 //------------------------------------------------------------------------------ 316 // Create non-tensor H(div) 317 //------------------------------------------------------------------------------ 318 int CeedBasisCreateHdiv_Cuda(CeedElemTopology topo, CeedInt dim, CeedInt num_nodes, CeedInt num_qpts, const CeedScalar *interp, const CeedScalar *div, 319 const CeedScalar *q_ref, const CeedScalar *q_weight, CeedBasis basis) { 320 Ceed ceed; 321 char *basis_kernel_source; 322 const char *basis_kernel_path; 323 CeedInt num_comp, q_comp_interp, q_comp_div; 324 const CeedInt q_bytes = num_qpts * sizeof(CeedScalar); 325 CeedBasisNonTensor_Cuda *data; 326 327 CeedCallBackend(CeedBasisGetCeed(basis, &ceed)); 328 CeedCallBackend(CeedCalloc(1, &data)); 329 330 // Copy basis data to GPU 331 CeedCallBackend(CeedBasisGetNumQuadratureComponents(basis, CEED_EVAL_INTERP, &q_comp_interp)); 332 CeedCallBackend(CeedBasisGetNumQuadratureComponents(basis, CEED_EVAL_DIV, &q_comp_div)); 333 CeedCallCuda(ceed, cudaMalloc((void **)&data->d_q_weight, q_bytes)); 334 CeedCallCuda(ceed, cudaMemcpy(data->d_q_weight, q_weight, q_bytes, cudaMemcpyHostToDevice)); 335 if (interp) { 336 const CeedInt interp_bytes = q_bytes * num_nodes * q_comp_interp; 337 338 CeedCallCuda(ceed, cudaMalloc((void **)&data->d_interp, interp_bytes)); 339 CeedCallCuda(ceed, cudaMemcpy(data->d_interp, interp, interp_bytes, cudaMemcpyHostToDevice)); 340 } 341 if (div) { 342 const CeedInt div_bytes = q_bytes * num_nodes * q_comp_div; 343 344 CeedCallCuda(ceed, cudaMalloc((void **)&data->d_div, div_bytes)); 345 CeedCallCuda(ceed, cudaMemcpy(data->d_div, div, div_bytes, cudaMemcpyHostToDevice)); 346 } 347 348 // Compile basis kernels 349 CeedCallBackend(CeedBasisGetNumComponents(basis, &num_comp)); 350 CeedCallBackend(CeedGetJitAbsolutePath(ceed, "ceed/jit-source/cuda/cuda-ref-basis-nontensor.h", &basis_kernel_path)); 351 CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "----- Loading Basis Kernel Source -----\n"); 352 CeedCallBackend(CeedLoadSourceToBuffer(ceed, basis_kernel_path, &basis_kernel_source)); 353 CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "----- Loading Basis Kernel Source Complete! -----\n"); 354 CeedCallBackend(CeedCompile_Cuda(ceed, basis_kernel_source, &data->module, 5, "BASIS_Q", num_qpts, "BASIS_P", num_nodes, "BASIS_Q_COMP_INTERP", 355 q_comp_interp, "BASIS_Q_COMP_DERIV", q_comp_div, "BASIS_NUM_COMP", num_comp)); 356 CeedCallBackend(CeedGetKernel_Cuda(ceed, data->module, "Interp", &data->Interp)); 357 CeedCallBackend(CeedGetKernel_Cuda(ceed, data->module, "InterpTranspose", &data->InterpTranspose)); 358 CeedCallBackend(CeedGetKernel_Cuda(ceed, data->module, "Deriv", &data->Deriv)); 359 CeedCallBackend(CeedGetKernel_Cuda(ceed, data->module, "DerivTranspose", &data->DerivTranspose)); 360 CeedCallBackend(CeedGetKernel_Cuda(ceed, data->module, "Weight", &data->Weight)); 361 CeedCallBackend(CeedFree(&basis_kernel_path)); 362 CeedCallBackend(CeedFree(&basis_kernel_source)); 363 364 CeedCallBackend(CeedBasisSetData(basis, data)); 365 366 // Register backend functions 367 CeedCallBackend(CeedSetBackendFunction(ceed, "Basis", basis, "Apply", CeedBasisApplyNonTensor_Cuda)); 368 CeedCallBackend(CeedSetBackendFunction(ceed, "Basis", basis, "Destroy", CeedBasisDestroyNonTensor_Cuda)); 369 return CEED_ERROR_SUCCESS; 370 } 371 372 //------------------------------------------------------------------------------ 373 // Create non-tensor H(curl) 374 //------------------------------------------------------------------------------ 375 int CeedBasisCreateHcurl_Cuda(CeedElemTopology topo, CeedInt dim, CeedInt num_nodes, CeedInt num_qpts, const CeedScalar *interp, 376 const CeedScalar *curl, const CeedScalar *q_ref, const CeedScalar *q_weight, CeedBasis basis) { 377 Ceed ceed; 378 char *basis_kernel_source; 379 const char *basis_kernel_path; 380 CeedInt num_comp, q_comp_interp, q_comp_curl; 381 const CeedInt q_bytes = num_qpts * sizeof(CeedScalar); 382 CeedBasisNonTensor_Cuda *data; 383 384 CeedCallBackend(CeedBasisGetCeed(basis, &ceed)); 385 CeedCallBackend(CeedCalloc(1, &data)); 386 387 // Copy basis data to GPU 388 CeedCallBackend(CeedBasisGetNumQuadratureComponents(basis, CEED_EVAL_INTERP, &q_comp_interp)); 389 CeedCallBackend(CeedBasisGetNumQuadratureComponents(basis, CEED_EVAL_CURL, &q_comp_curl)); 390 CeedCallCuda(ceed, cudaMalloc((void **)&data->d_q_weight, q_bytes)); 391 CeedCallCuda(ceed, cudaMemcpy(data->d_q_weight, q_weight, q_bytes, cudaMemcpyHostToDevice)); 392 if (interp) { 393 const CeedInt interp_bytes = q_bytes * num_nodes * q_comp_interp; 394 395 CeedCallCuda(ceed, cudaMalloc((void **)&data->d_interp, interp_bytes)); 396 CeedCallCuda(ceed, cudaMemcpy(data->d_interp, interp, interp_bytes, cudaMemcpyHostToDevice)); 397 } 398 if (curl) { 399 const CeedInt curl_bytes = q_bytes * num_nodes * q_comp_curl; 400 401 CeedCallCuda(ceed, cudaMalloc((void **)&data->d_curl, curl_bytes)); 402 CeedCallCuda(ceed, cudaMemcpy(data->d_curl, curl, curl_bytes, cudaMemcpyHostToDevice)); 403 } 404 405 // Compile basis kernels 406 CeedCallBackend(CeedBasisGetNumComponents(basis, &num_comp)); 407 CeedCallBackend(CeedGetJitAbsolutePath(ceed, "ceed/jit-source/cuda/cuda-ref-basis-nontensor.h", &basis_kernel_path)); 408 CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "----- Loading Basis Kernel Source -----\n"); 409 CeedCallBackend(CeedLoadSourceToBuffer(ceed, basis_kernel_path, &basis_kernel_source)); 410 CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "----- Loading Basis Kernel Source Complete! -----\n"); 411 CeedCallBackend(CeedCompile_Cuda(ceed, basis_kernel_source, &data->module, 5, "BASIS_Q", num_qpts, "BASIS_P", num_nodes, "BASIS_Q_COMP_INTERP", 412 q_comp_interp, "BASIS_Q_COMP_DERIV", q_comp_curl, "BASIS_NUM_COMP", num_comp)); 413 CeedCallBackend(CeedGetKernel_Cuda(ceed, data->module, "Interp", &data->Interp)); 414 CeedCallBackend(CeedGetKernel_Cuda(ceed, data->module, "InterpTranspose", &data->InterpTranspose)); 415 CeedCallBackend(CeedGetKernel_Cuda(ceed, data->module, "Deriv", &data->Deriv)); 416 CeedCallBackend(CeedGetKernel_Cuda(ceed, data->module, "DerivTranspose", &data->DerivTranspose)); 417 CeedCallBackend(CeedGetKernel_Cuda(ceed, data->module, "Weight", &data->Weight)); 418 CeedCallBackend(CeedFree(&basis_kernel_path)); 419 CeedCallBackend(CeedFree(&basis_kernel_source)); 420 421 CeedCallBackend(CeedBasisSetData(basis, data)); 422 423 // Register backend functions 424 CeedCallBackend(CeedSetBackendFunction(ceed, "Basis", basis, "Apply", CeedBasisApplyNonTensor_Cuda)); 425 CeedCallBackend(CeedSetBackendFunction(ceed, "Basis", basis, "Destroy", CeedBasisDestroyNonTensor_Cuda)); 426 return CEED_ERROR_SUCCESS; 427 } 428 429 //------------------------------------------------------------------------------ 430