Home
last modified time | relevance | path

Searched refs:is_transpose (Results 1 – 7 of 7) sorted by relevance

/libCEED/include/ceed/jit-source/hip/
H A Dhip-ref-basis-tensor.h19 extern "C" __global__ void Interp(const CeedInt num_elem, const CeedInt is_transpose, const CeedSca… in Interp() argument
31 const CeedInt P = is_transpose ? BASIS_Q_1D : BASIS_P_1D; in Interp()
32 const CeedInt Q = is_transpose ? BASIS_P_1D : BASIS_Q_1D; in Interp()
33 const CeedInt stride_0 = is_transpose ? 1 : BASIS_P_1D; in Interp()
34 const CeedInt stride_1 = is_transpose ? BASIS_P_1D : 1; in Interp()
35 const CeedInt u_stride = is_transpose ? BASIS_NUM_QPTS : BASIS_NUM_NODES; in Interp()
36 const CeedInt v_stride = is_transpose ? BASIS_NUM_NODES : BASIS_NUM_QPTS; in Interp()
37 const CeedInt u_comp_stride = num_elem * (is_transpose ? BASIS_NUM_QPTS : BASIS_NUM_NODES); in Interp()
38 const CeedInt v_comp_stride = num_elem * (is_transpose ? BASIS_NUM_NODES : BASIS_NUM_QPTS); in Interp()
39 const CeedInt u_size = is_transpose ? BASIS_NUM_QPTS : BASIS_NUM_NODES; in Interp()
[all …]
/libCEED/include/ceed/jit-source/cuda/
H A Dcuda-ref-basis-tensor.h19 extern "C" __global__ void Interp(const CeedInt num_elem, const CeedInt is_transpose, const CeedSca… in Interp() argument
31 const CeedInt P = is_transpose ? BASIS_Q_1D : BASIS_P_1D; in Interp()
32 const CeedInt Q = is_transpose ? BASIS_P_1D : BASIS_Q_1D; in Interp()
33 const CeedInt stride_0 = is_transpose ? 1 : BASIS_P_1D; in Interp()
34 const CeedInt stride_1 = is_transpose ? BASIS_P_1D : 1; in Interp()
35 const CeedInt u_stride = is_transpose ? BASIS_NUM_QPTS : BASIS_NUM_NODES; in Interp()
36 const CeedInt v_stride = is_transpose ? BASIS_NUM_NODES : BASIS_NUM_QPTS; in Interp()
37 const CeedInt u_comp_stride = num_elem * (is_transpose ? BASIS_NUM_QPTS : BASIS_NUM_NODES); in Interp()
38 const CeedInt v_comp_stride = num_elem * (is_transpose ? BASIS_NUM_NODES : BASIS_NUM_QPTS); in Interp()
39 const CeedInt u_size = is_transpose ? BASIS_NUM_QPTS : BASIS_NUM_NODES; in Interp()
[all …]
/libCEED/backends/sycl-ref/
H A Dceed-sycl-ref-basis.sycl.cpp38 template <int is_transpose>
62 …cgh.parallel_for<CeedBasisSyclInterp<is_transpose>>(kernel_range, [=](sycl::nd_item<1> work_item, … in CeedBasisApplyInterp_Sycl()
72 const CeedInt P = is_transpose ? Q_1d : P_1d; in CeedBasisApplyInterp_Sycl()
73 const CeedInt Q = is_transpose ? P_1d : Q_1d; in CeedBasisApplyInterp_Sycl()
74 const CeedInt stride_0 = is_transpose ? 1 : P_1d; in CeedBasisApplyInterp_Sycl()
75 const CeedInt stride_1 = is_transpose ? P_1d : 1; in CeedBasisApplyInterp_Sycl()
76 const CeedInt u_stride = is_transpose ? num_qpts : num_nodes; in CeedBasisApplyInterp_Sycl()
77 const CeedInt v_stride = is_transpose ? num_nodes : num_qpts; in CeedBasisApplyInterp_Sycl()
141 template <int is_transpose>
165 …cgh.parallel_for<CeedBasisSyclGrad<is_transpose>>(kernel_range, [=](sycl::nd_item<1> work_item, sy… in CeedBasisApplyGrad_Sycl()
[all …]
/libCEED/backends/hip-ref/
H A Dceed-hip-ref-basis.c25 const CeedInt is_transpose = t_mode == CEED_TRANSPOSE; in CeedBasisApplyCore_Hip() local
41 if (is_transpose) CeedCallBackend(CeedVectorSetValue(v, 0.0)); in CeedBasisApplyCore_Hip()
51 …void *interp_args[] = {(void *)&num_elem, (void *)&is_transpose, &data->d_interp_1d, &d_u,… in CeedBasisApplyCore_Hip()
57 …void *grad_args[] = {(void *)&num_elem, (void *)&is_transpose, &data->d_interp_1d, &data->… in CeedBasisApplyCore_Hip()
105 const CeedInt is_transpose = t_mode == CEED_TRANSPOSE; in CeedBasisApplyAtPointsCore_Hip() local
131 CeedCallBackend(CeedVectorGetLength(is_transpose ? u : v, &len)); in CeedBasisApplyAtPointsCore_Hip()
140 if (is_transpose) { in CeedBasisApplyAtPointsCore_Hip()
201 if (is_transpose) CeedCallBackend(CeedVectorSetValue(v, 0.0)); 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()
[all …]
/libCEED/backends/cuda-ref/
H A Dceed-cuda-ref-basis.c26 const CeedInt is_transpose = t_mode == CEED_TRANSPOSE; in CeedBasisApplyCore_Cuda() local
42 if (is_transpose) CeedCallBackend(CeedVectorSetValue(v, 0.0)); in CeedBasisApplyCore_Cuda()
51 …void *interp_args[] = {(void *)&num_elem, (void *)&is_transpose, &data->d_interp_1d, &d_u,… in CeedBasisApplyCore_Cuda()
57 …void *grad_args[] = {(void *)&num_elem, (void *)&is_transpose, &data->d_interp_1d, &data->… in CeedBasisApplyCore_Cuda()
106 const CeedInt is_transpose = t_mode == CEED_TRANSPOSE; in CeedBasisApplyAtPointsCore_Cuda() local
132 CeedCallBackend(CeedVectorGetLength(is_transpose ? u : v, &len)); in CeedBasisApplyAtPointsCore_Cuda()
141 if (is_transpose) { in CeedBasisApplyAtPointsCore_Cuda()
202 if (is_transpose) CeedCallBackend(CeedVectorSetValue(v, 0.0)); in CeedBasisApplyAtPointsCore_Cuda()
212 …CeedCallBackend(CeedRunKernel_Cuda(ceed, is_transpose ? data->InterpTransposeAtPoints : data->Inte… in CeedBasisApplyAtPointsCore_Cuda()
219 …CeedCallBackend(CeedRunKernel_Cuda(ceed, is_transpose ? data->GradTransposeAtPoints : data->GradAt… in CeedBasisApplyAtPointsCore_Cuda()
[all …]
/libCEED/backends/hip-shared/
H A Dceed-hip-shared-basis.c280 const CeedInt is_transpose = t_mode == CEED_TRANSPOSE; in CeedBasisApplyAtPointsCore_Hip_shared() local
305 CeedCallBackend(CeedVectorGetLength(is_transpose ? u : v, &len)); in CeedBasisApplyAtPointsCore_Hip_shared()
314 if (is_transpose) { in CeedBasisApplyAtPointsCore_Hip_shared()
396 if (is_transpose) { in CeedBasisApplyAtPointsCore_Hip_shared()
408 if (is_transpose) { in CeedBasisApplyAtPointsCore_Hip_shared()
420 if (is_transpose) { in CeedBasisApplyAtPointsCore_Hip_shared()
444 if (is_transpose) { in CeedBasisApplyAtPointsCore_Hip_shared()
456 if (is_transpose) { in CeedBasisApplyAtPointsCore_Hip_shared()
468 if (is_transpose) { in CeedBasisApplyAtPointsCore_Hip_shared()
/libCEED/backends/cuda-shared/
H A Dceed-cuda-shared-basis.c215 const CeedInt is_transpose = t_mode == CEED_TRANSPOSE; in CeedBasisApplyAtPointsCore_Cuda_shared() local
240 CeedCallBackend(CeedVectorGetLength(is_transpose ? u : v, &len)); in CeedBasisApplyAtPointsCore_Cuda_shared()
249 if (is_transpose) { in CeedBasisApplyAtPointsCore_Cuda_shared()
330 if (is_transpose) { in CeedBasisApplyAtPointsCore_Cuda_shared()
344 if (is_transpose) { in CeedBasisApplyAtPointsCore_Cuda_shared()
356 if (is_transpose) { in CeedBasisApplyAtPointsCore_Cuda_shared()
380 if (is_transpose) { in CeedBasisApplyAtPointsCore_Cuda_shared()
393 if (is_transpose) { in CeedBasisApplyAtPointsCore_Cuda_shared()
405 if (is_transpose) { in CeedBasisApplyAtPointsCore_Cuda_shared()