Lines Matching full:ceed
1 // Copyright (c) 2017-2026, Lawrence Livermore National Security, LLC and other CEED contributors.
6 // This file is part of CEED: http://github.com/ceed
8 #include <ceed.h>
9 #include <ceed/backend.h>
10 #include <ceed/jit-tools.h>
14 #include "../hip/ceed-hip-common.h"
15 #include "../hip/ceed-hip-compile.h"
16 #include "ceed-hip-ref.h"
23 Ceed ceed; in CeedBasisApplyCore_Hip() local
31 CeedCallBackend(CeedBasisGetCeed(basis, &ceed)); in CeedBasisApplyCore_Hip()
36 …else CeedCheck(eval_mode == CEED_EVAL_WEIGHT, ceed, CEED_ERROR_BACKEND, "An input vector is requir… in CeedBasisApplyCore_Hip()
54 CeedCallBackend(CeedRunKernel_Hip(ceed, data->Interp, num_elem, block_size, interp_args)); in CeedBasisApplyCore_Hip()
60 CeedCallBackend(CeedRunKernel_Hip(ceed, data->Grad, num_elem, block_size, grad_args)); in CeedBasisApplyCore_Hip()
63 …CeedCheck(data->d_q_weight_1d, ceed, CEED_ERROR_BACKEND, "%s not supported; q_weights_1d not set",… in CeedBasisApplyCore_Hip()
68 …CeedCallBackend(CeedRunKernelDim_Hip(ceed, data->Weight, num_elem, block_size_x, block_size_y, 1, … in CeedBasisApplyCore_Hip()
75 return CeedError(ceed, CEED_ERROR_BACKEND, "%s not supported", CeedEvalModes[eval_mode]); in CeedBasisApplyCore_Hip()
83 CeedCallBackend(CeedDestroy(&ceed)); in CeedBasisApplyCore_Hip()
103 Ceed ceed; in CeedBasisApplyAtPointsCore_Hip() local
121 CeedCallBackend(CeedBasisGetCeed(basis, &ceed)); in CeedBasisApplyAtPointsCore_Hip()
133 CeedCheck(len >= len_required, ceed, CEED_ERROR_BACKEND, in CeedBasisApplyAtPointsCore_Hip()
146 if (data->d_points_per_elem) CeedCallHip(ceed, hipFree(data->d_points_per_elem)); in CeedBasisApplyAtPointsCore_Hip()
147 CeedCallHip(ceed, hipMalloc((void **)&data->d_points_per_elem, num_bytes)); in CeedBasisApplyAtPointsCore_Hip()
153 …CeedCallHip(ceed, hipMemcpy(data->d_points_per_elem, num_points, num_bytes, hipMemcpyHostToDevice)… in CeedBasisApplyAtPointsCore_Hip()
172 CeedCallHip(ceed, hipMalloc((void **)&data->d_chebyshev_interp_1d, interp_bytes)); in CeedBasisApplyAtPointsCore_Hip()
173 …CeedCallHip(ceed, hipMemcpy(data->d_chebyshev_interp_1d, chebyshev_interp_1d, interp_bytes, hipMem… in CeedBasisApplyAtPointsCore_Hip()
178 …const char basis_kernel_source[] = "// AtPoints basis source\n#include <ceed/jit-source/hip/hip-re… in CeedBasisApplyAtPointsCore_Hip()
181 if (data->moduleAtPoints) CeedCallHip(ceed, hipModuleUnload(data->moduleAtPoints)); in CeedBasisApplyAtPointsCore_Hip()
183 …CeedCallBackend(CeedCompile_Hip(ceed, basis_kernel_source, &data->moduleAtPoints, 9, "BASIS_Q_1D",… in CeedBasisApplyAtPointsCore_Hip()
187 …CeedCallBackend(CeedGetKernel_Hip(ceed, data->moduleAtPoints, "InterpAtPoints", &data->InterpAtPoi… in CeedBasisApplyAtPointsCore_Hip()
188 …CeedCallBackend(CeedGetKernel_Hip(ceed, data->moduleAtPoints, "InterpTransposeAtPoints", &data->In… in CeedBasisApplyAtPointsCore_Hip()
189 …CeedCallBackend(CeedGetKernel_Hip(ceed, data->moduleAtPoints, "GradAtPoints", &data->GradAtPoints)… in CeedBasisApplyAtPointsCore_Hip()
190 …CeedCallBackend(CeedGetKernel_Hip(ceed, data->moduleAtPoints, "GradTransposeAtPoints", &data->Grad… in CeedBasisApplyAtPointsCore_Hip()
196 …else CeedCheck(eval_mode == CEED_EVAL_WEIGHT, ceed, CEED_ERROR_BACKEND, "An input vector is requir… in CeedBasisApplyAtPointsCore_Hip()
211 …CeedCallBackend(CeedRunKernel_Hip(ceed, is_transpose ? data->InterpTransposeAtPoints : data->Inter… in CeedBasisApplyAtPointsCore_Hip()
218 …CeedCallBackend(CeedRunKernel_Hip(ceed, is_transpose ? data->GradTransposeAtPoints : data->GradAtP… in CeedBasisApplyAtPointsCore_Hip()
226 return CeedError(ceed, CEED_ERROR_BACKEND, "%s not supported", CeedEvalModes[eval_mode]); in CeedBasisApplyAtPointsCore_Hip()
235 CeedCallBackend(CeedDestroy(&ceed)); in CeedBasisApplyAtPointsCore_Hip()
256 Ceed ceed; in CeedBasisApplyNonTensorCore_Hip() local
265 CeedCallBackend(CeedBasisGetCeed(basis, &ceed)); in CeedBasisApplyNonTensorCore_Hip()
272 …else CeedCheck(eval_mode == CEED_EVAL_WEIGHT, ceed, CEED_ERROR_BACKEND, "An input vector is requir… in CeedBasisApplyNonTensorCore_Hip()
288 …CeedCallBackend(CeedRunKernelDim_Hip(ceed, data->InterpTranspose, grid, block_size_x, 1, elems_per… in CeedBasisApplyNonTensorCore_Hip()
290 …CeedCallBackend(CeedRunKernelDim_Hip(ceed, data->Interp, grid, block_size_x, 1, elems_per_block, i… in CeedBasisApplyNonTensorCore_Hip()
298 …CeedCallBackend(CeedRunKernelDim_Hip(ceed, data->DerivTranspose, grid, block_size_x, 1, elems_per_… in CeedBasisApplyNonTensorCore_Hip()
300 …CeedCallBackend(CeedRunKernelDim_Hip(ceed, data->Deriv, grid, block_size_x, 1, elems_per_block, gr… in CeedBasisApplyNonTensorCore_Hip()
308 …CeedCallBackend(CeedRunKernelDim_Hip(ceed, data->DerivTranspose, grid, block_size_x, 1, elems_per_… in CeedBasisApplyNonTensorCore_Hip()
310 …CeedCallBackend(CeedRunKernelDim_Hip(ceed, data->Deriv, grid, block_size_x, 1, elems_per_block, di… in CeedBasisApplyNonTensorCore_Hip()
318 …CeedCallBackend(CeedRunKernelDim_Hip(ceed, data->DerivTranspose, grid, block_size_x, 1, elems_per_… in CeedBasisApplyNonTensorCore_Hip()
320 …CeedCallBackend(CeedRunKernelDim_Hip(ceed, data->Deriv, grid, block_size_x, 1, elems_per_block, cu… in CeedBasisApplyNonTensorCore_Hip()
324 …CeedCheck(data->d_q_weight, ceed, CEED_ERROR_BACKEND, "%s not supported; q_weights not set", CeedE… in CeedBasisApplyNonTensorCore_Hip()
327 …CeedCallBackend(CeedRunKernelDim_Hip(ceed, data->Weight, grid, num_qpts, 1, elems_per_block, weigh… in CeedBasisApplyNonTensorCore_Hip()
337 CeedCallBackend(CeedDestroy(&ceed)); in CeedBasisApplyNonTensorCore_Hip()
357 Ceed ceed; in CeedBasisDestroy_Hip() local
360 CeedCallBackend(CeedBasisGetCeed(basis, &ceed)); in CeedBasisDestroy_Hip()
362 CeedCallHip(ceed, hipModuleUnload(data->module)); in CeedBasisDestroy_Hip()
363 if (data->moduleAtPoints) CeedCallHip(ceed, hipModuleUnload(data->moduleAtPoints)); in CeedBasisDestroy_Hip()
364 if (data->d_q_weight_1d) CeedCallHip(ceed, hipFree(data->d_q_weight_1d)); in CeedBasisDestroy_Hip()
366 if (data->d_points_per_elem) CeedCallHip(ceed, hipFree(data->d_points_per_elem)); in CeedBasisDestroy_Hip()
367 CeedCallHip(ceed, hipFree(data->d_interp_1d)); in CeedBasisDestroy_Hip()
368 CeedCallHip(ceed, hipFree(data->d_grad_1d)); in CeedBasisDestroy_Hip()
369 CeedCallHip(ceed, hipFree(data->d_chebyshev_interp_1d)); in CeedBasisDestroy_Hip()
371 CeedCallBackend(CeedDestroy(&ceed)); in CeedBasisDestroy_Hip()
379 Ceed ceed; in CeedBasisDestroyNonTensor_Hip() local
382 CeedCallBackend(CeedBasisGetCeed(basis, &ceed)); in CeedBasisDestroyNonTensor_Hip()
384 CeedCallHip(ceed, hipModuleUnload(data->module)); in CeedBasisDestroyNonTensor_Hip()
385 if (data->d_q_weight) CeedCallHip(ceed, hipFree(data->d_q_weight)); in CeedBasisDestroyNonTensor_Hip()
386 CeedCallHip(ceed, hipFree(data->d_interp)); in CeedBasisDestroyNonTensor_Hip()
387 CeedCallHip(ceed, hipFree(data->d_grad)); in CeedBasisDestroyNonTensor_Hip()
388 CeedCallHip(ceed, hipFree(data->d_div)); in CeedBasisDestroyNonTensor_Hip()
389 CeedCallHip(ceed, hipFree(data->d_curl)); in CeedBasisDestroyNonTensor_Hip()
391 CeedCallBackend(CeedDestroy(&ceed)); in CeedBasisDestroyNonTensor_Hip()
400 Ceed ceed; in CeedBasisCreateTensorH1_Hip() local
406 CeedCallBackend(CeedBasisGetCeed(basis, &ceed)); in CeedBasisCreateTensorH1_Hip()
411 CeedCallHip(ceed, hipMalloc((void **)&data->d_q_weight_1d, q_bytes)); in CeedBasisCreateTensorH1_Hip()
412 CeedCallHip(ceed, hipMemcpy(data->d_q_weight_1d, q_weight_1d, q_bytes, hipMemcpyHostToDevice)); in CeedBasisCreateTensorH1_Hip()
414 CeedCallHip(ceed, hipMalloc((void **)&data->d_interp_1d, interp_bytes)); in CeedBasisCreateTensorH1_Hip()
415 CeedCallHip(ceed, hipMemcpy(data->d_interp_1d, interp_1d, interp_bytes, hipMemcpyHostToDevice)); in CeedBasisCreateTensorH1_Hip()
416 CeedCallHip(ceed, hipMalloc((void **)&data->d_grad_1d, interp_bytes)); in CeedBasisCreateTensorH1_Hip()
417 CeedCallHip(ceed, hipMemcpy(data->d_grad_1d, grad_1d, interp_bytes, hipMemcpyHostToDevice)); in CeedBasisCreateTensorH1_Hip()
420 …const char basis_kernel_source[] = "// Tensor basis source\n#include <ceed/jit-source/hip/hip-ref-… in CeedBasisCreateTensorH1_Hip()
423 …CeedCallBackend(CeedCompile_Hip(ceed, basis_kernel_source, &data->module, 7, "BASIS_Q_1D", Q_1d, "… in CeedBasisCreateTensorH1_Hip()
426 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "Interp", &data->Interp)); in CeedBasisCreateTensorH1_Hip()
427 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "Grad", &data->Grad)); in CeedBasisCreateTensorH1_Hip()
428 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "Weight", &data->Weight)); in CeedBasisCreateTensorH1_Hip()
433 CeedCallBackend(CeedSetBackendFunction(ceed, "Basis", basis, "Apply", CeedBasisApply_Hip)); in CeedBasisCreateTensorH1_Hip()
434 CeedCallBackend(CeedSetBackendFunction(ceed, "Basis", basis, "ApplyAdd", CeedBasisApplyAdd_Hip)); in CeedBasisCreateTensorH1_Hip()
435 …CeedCallBackend(CeedSetBackendFunction(ceed, "Basis", basis, "ApplyAtPoints", CeedBasisApplyAtPoin… in CeedBasisCreateTensorH1_Hip()
436 …CeedCallBackend(CeedSetBackendFunction(ceed, "Basis", basis, "ApplyAddAtPoints", CeedBasisApplyAdd… in CeedBasisCreateTensorH1_Hip()
437 CeedCallBackend(CeedSetBackendFunction(ceed, "Basis", basis, "Destroy", CeedBasisDestroy_Hip)); in CeedBasisCreateTensorH1_Hip()
438 CeedCallBackend(CeedDestroy(&ceed)); in CeedBasisCreateTensorH1_Hip()
447 Ceed ceed; in CeedBasisCreateH1_Hip() local
452 CeedCallBackend(CeedBasisGetCeed(basis, &ceed)); in CeedBasisCreateH1_Hip()
459 CeedCallHip(ceed, hipMalloc((void **)&data->d_q_weight, q_bytes)); in CeedBasisCreateH1_Hip()
460 CeedCallHip(ceed, hipMemcpy(data->d_q_weight, q_weight, q_bytes, hipMemcpyHostToDevice)); in CeedBasisCreateH1_Hip()
465 CeedCallHip(ceed, hipMalloc((void **)&data->d_interp, interp_bytes)); in CeedBasisCreateH1_Hip()
466 CeedCallHip(ceed, hipMemcpy(data->d_interp, interp, interp_bytes, hipMemcpyHostToDevice)); in CeedBasisCreateH1_Hip()
471 CeedCallHip(ceed, hipMalloc((void **)&data->d_grad, grad_bytes)); in CeedBasisCreateH1_Hip()
472 CeedCallHip(ceed, hipMemcpy(data->d_grad, grad, grad_bytes, hipMemcpyHostToDevice)); in CeedBasisCreateH1_Hip()
476 …const char basis_kernel_source[] = "// Nontensor basis source\n#include <ceed/jit-source/hip/hip-r… in CeedBasisCreateH1_Hip()
479 …CeedCallBackend(CeedCompile_Hip(ceed, basis_kernel_source, &data->module, 5, "BASIS_Q", num_qpts, … in CeedBasisCreateH1_Hip()
481 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "Interp", &data->Interp)); in CeedBasisCreateH1_Hip()
482 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "InterpTranspose", &data->InterpTranspose)); in CeedBasisCreateH1_Hip()
483 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "Deriv", &data->Deriv)); in CeedBasisCreateH1_Hip()
484 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "DerivTranspose", &data->DerivTranspose)); in CeedBasisCreateH1_Hip()
485 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "Weight", &data->Weight)); in CeedBasisCreateH1_Hip()
490 …CeedCallBackend(CeedSetBackendFunction(ceed, "Basis", basis, "Apply", CeedBasisApplyNonTensor_Hip)… in CeedBasisCreateH1_Hip()
491 …CeedCallBackend(CeedSetBackendFunction(ceed, "Basis", basis, "ApplyAdd", CeedBasisApplyAddNonTenso… in CeedBasisCreateH1_Hip()
492 …CeedCallBackend(CeedSetBackendFunction(ceed, "Basis", basis, "Destroy", CeedBasisDestroyNonTensor_… in CeedBasisCreateH1_Hip()
493 CeedCallBackend(CeedDestroy(&ceed)); in CeedBasisCreateH1_Hip()
502 Ceed ceed; in CeedBasisCreateHdiv_Hip() local
507 CeedCallBackend(CeedBasisGetCeed(basis, &ceed)); in CeedBasisCreateHdiv_Hip()
514 CeedCallHip(ceed, hipMalloc((void **)&data->d_q_weight, q_bytes)); in CeedBasisCreateHdiv_Hip()
515 CeedCallHip(ceed, hipMemcpy(data->d_q_weight, q_weight, q_bytes, hipMemcpyHostToDevice)); in CeedBasisCreateHdiv_Hip()
520 CeedCallHip(ceed, hipMalloc((void **)&data->d_interp, interp_bytes)); in CeedBasisCreateHdiv_Hip()
521 CeedCallHip(ceed, hipMemcpy(data->d_interp, interp, interp_bytes, hipMemcpyHostToDevice)); in CeedBasisCreateHdiv_Hip()
526 CeedCallHip(ceed, hipMalloc((void **)&data->d_div, div_bytes)); in CeedBasisCreateHdiv_Hip()
527 CeedCallHip(ceed, hipMemcpy(data->d_div, div, div_bytes, hipMemcpyHostToDevice)); in CeedBasisCreateHdiv_Hip()
531 …const char basis_kernel_source[] = "// Nontensor basis source\n#include <ceed/jit-source/hip/hip-r… in CeedBasisCreateHdiv_Hip()
534 …CeedCallBackend(CeedCompile_Hip(ceed, basis_kernel_source, &data->module, 5, "BASIS_Q", num_qpts, … in CeedBasisCreateHdiv_Hip()
536 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "Interp", &data->Interp)); in CeedBasisCreateHdiv_Hip()
537 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "InterpTranspose", &data->InterpTranspose)); in CeedBasisCreateHdiv_Hip()
538 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "Deriv", &data->Deriv)); in CeedBasisCreateHdiv_Hip()
539 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "DerivTranspose", &data->DerivTranspose)); in CeedBasisCreateHdiv_Hip()
540 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "Weight", &data->Weight)); in CeedBasisCreateHdiv_Hip()
545 …CeedCallBackend(CeedSetBackendFunction(ceed, "Basis", basis, "Apply", CeedBasisApplyNonTensor_Hip)… in CeedBasisCreateHdiv_Hip()
546 …CeedCallBackend(CeedSetBackendFunction(ceed, "Basis", basis, "ApplyAdd", CeedBasisApplyAddNonTenso… in CeedBasisCreateHdiv_Hip()
547 …CeedCallBackend(CeedSetBackendFunction(ceed, "Basis", basis, "Destroy", CeedBasisDestroyNonTensor_… in CeedBasisCreateHdiv_Hip()
548 CeedCallBackend(CeedDestroy(&ceed)); in CeedBasisCreateHdiv_Hip()
557 Ceed ceed; in CeedBasisCreateHcurl_Hip() local
562 CeedCallBackend(CeedBasisGetCeed(basis, &ceed)); in CeedBasisCreateHcurl_Hip()
569 CeedCallHip(ceed, hipMalloc((void **)&data->d_q_weight, q_bytes)); in CeedBasisCreateHcurl_Hip()
570 CeedCallHip(ceed, hipMemcpy(data->d_q_weight, q_weight, q_bytes, hipMemcpyHostToDevice)); in CeedBasisCreateHcurl_Hip()
575 CeedCallHip(ceed, hipMalloc((void **)&data->d_interp, interp_bytes)); in CeedBasisCreateHcurl_Hip()
576 CeedCallHip(ceed, hipMemcpy(data->d_interp, interp, interp_bytes, hipMemcpyHostToDevice)); in CeedBasisCreateHcurl_Hip()
581 CeedCallHip(ceed, hipMalloc((void **)&data->d_curl, curl_bytes)); in CeedBasisCreateHcurl_Hip()
582 CeedCallHip(ceed, hipMemcpy(data->d_curl, curl, curl_bytes, hipMemcpyHostToDevice)); in CeedBasisCreateHcurl_Hip()
586 …const char basis_kernel_source[] = "// Nontensor basis source\n#include <ceed/jit-source/hip/hip-r… in CeedBasisCreateHcurl_Hip()
589 …CeedCallBackend(CeedCompile_Hip(ceed, basis_kernel_source, &data->module, 5, "BASIS_Q", num_qpts, … in CeedBasisCreateHcurl_Hip()
591 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "Interp", &data->Interp)); in CeedBasisCreateHcurl_Hip()
592 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "InterpTranspose", &data->InterpTranspose)); in CeedBasisCreateHcurl_Hip()
593 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "Deriv", &data->Deriv)); in CeedBasisCreateHcurl_Hip()
594 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "DerivTranspose", &data->DerivTranspose)); in CeedBasisCreateHcurl_Hip()
595 CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, "Weight", &data->Weight)); in CeedBasisCreateHcurl_Hip()
600 …CeedCallBackend(CeedSetBackendFunction(ceed, "Basis", basis, "Apply", CeedBasisApplyNonTensor_Hip)… in CeedBasisCreateHcurl_Hip()
601 …CeedCallBackend(CeedSetBackendFunction(ceed, "Basis", basis, "ApplyAdd", CeedBasisApplyAddNonTenso… in CeedBasisCreateHcurl_Hip()
602 …CeedCallBackend(CeedSetBackendFunction(ceed, "Basis", basis, "Destroy", CeedBasisDestroyNonTensor_… in CeedBasisCreateHcurl_Hip()
603 CeedCallBackend(CeedDestroy(&ceed)); in CeedBasisCreateHcurl_Hip()