Home
last modified time | relevance | path

Searched refs:Q_1d (Results 1 – 25 of 34) sorted by relevance

12

/libCEED/backends/hip-shared/
H A Dceed-hip-shared-basis.c39 static int ComputeBasisThreadBlockSizes(const CeedInt dim, const CeedInt P_1d, const CeedInt Q_1d, … in ComputeBasisThreadBlockSizes() argument
44 const CeedInt thread_1d = CeedIntMax(P_1d, Q_1d); in ComputeBasisThreadBlockSizes()
67 required = CeedIntMax(64, Q_1d * Q_1d); in ComputeBasisThreadBlockSizes()
81 required = Q_1d * Q_1d * Q_1d; in ComputeBasisThreadBlockSizes()
118 CeedInt P_1d, Q_1d; in CeedBasisApplyTensorCore_Hip_shared() local
123 CeedCallBackend(CeedBasisGetNumQuadraturePoints1D(basis, &Q_1d)); in CeedBasisApplyTensorCore_Hip_shared()
124 CeedInt thread_1d = CeedIntMax(Q_1d, P_1d); in CeedBasisApplyTensorCore_Hip_shared()
165 CeedInt P_1d, Q_1d; in CeedBasisApplyTensorCore_Hip_shared() local
170 CeedCallBackend(CeedBasisGetNumQuadraturePoints1D(basis, &Q_1d)); in CeedBasisApplyTensorCore_Hip_shared()
171 CeedInt thread_1d = CeedIntMax(Q_1d, P_1d); in CeedBasisApplyTensorCore_Hip_shared()
[all …]
H A Dceed-hip-shared.h41 CEED_INTERN int CeedBasisCreateTensorH1_Hip_shared(CeedInt dim, CeedInt P_1d, CeedInt Q_1d, const C…
/libCEED/interface/
H A Dceed-basis.c400 CeedInt dim, num_comp, num_q_comp, num_nodes, P_1d = 1, Q_1d = 1, total_num_points = 0; in CeedBasisApplyAtPointsCheckDims() local
405 CeedCall(CeedBasisGetNumQuadraturePoints1D(basis, &Q_1d)); in CeedBasisApplyAtPointsCheckDims()
478 CeedInt dim, num_comp, P_1d = 1, Q_1d = 1, total_num_points = num_points[0]; in CeedBasisApplyAtPoints_Core() local
484 CeedCall(CeedBasisGetNumQuadraturePoints1D(basis, &Q_1d)); in CeedBasisApplyAtPoints_Core()
507 CeedCall(CeedCalloc(P_1d * Q_1d, &chebyshev_interp_1d)); in CeedBasisApplyAtPoints_Core()
508 CeedCall(CeedCalloc(P_1d * Q_1d, &chebyshev_grad_1d)); in CeedBasisApplyAtPoints_Core()
509 CeedCall(CeedCalloc(Q_1d, &chebyshev_q_weight_1d)); in CeedBasisApplyAtPoints_Core()
514 CeedCall(CeedVectorCreate(ceed, num_comp * CeedIntPow(Q_1d, dim), &basis->vec_chebyshev)); in CeedBasisApplyAtPoints_Core()
515 …CeedCall(CeedBasisCreateTensorH1(ceed, dim, num_comp, P_1d, Q_1d, chebyshev_interp_1d, chebyshev_g… in CeedBasisApplyAtPoints_Core()
532 …CeedCall(CeedBasisCreateTensorH1Lagrange(ceed_ref, dim, num_comp, P_1d, Q_1d, CEED_GAUSS, &basis_r… in CeedBasisApplyAtPoints_Core()
[all …]
/libCEED/backends/cuda-shared/
H A Dceed-cuda-shared-basis.c51 CeedInt P_1d, Q_1d; in CeedBasisApplyTensorCore_Cuda_shared() local
55 CeedCallBackend(CeedBasisGetNumQuadraturePoints1D(basis, &Q_1d)); in CeedBasisApplyTensorCore_Cuda_shared()
56 CeedInt thread_1d = CeedIntMax(Q_1d, P_1d); in CeedBasisApplyTensorCore_Cuda_shared()
101 CeedInt P_1d, Q_1d; in CeedBasisApplyTensorCore_Cuda_shared() local
105 CeedCallBackend(CeedBasisGetNumQuadraturePoints1D(basis, &Q_1d)); in CeedBasisApplyTensorCore_Cuda_shared()
106 CeedInt thread_1d = CeedIntMax(Q_1d, P_1d); in CeedBasisApplyTensorCore_Cuda_shared()
153 CeedInt Q_1d; in CeedBasisApplyTensorCore_Cuda_shared() local
157 CeedCallBackend(CeedBasisGetNumQuadraturePoints1D(basis, &Q_1d)); in CeedBasisApplyTensorCore_Cuda_shared()
160 const CeedInt elems_per_block = block_size / Q_1d; in CeedBasisApplyTensorCore_Cuda_shared()
163 …CeedCallBackend(CeedRunKernelDim_Cuda(ceed, data->Weight, grid_size, Q_1d, elems_per_block, 1, wei… in CeedBasisApplyTensorCore_Cuda_shared()
[all …]
H A Dceed-cuda-shared.h40 CEED_INTERN int CeedBasisCreateTensorH1_Cuda_shared(CeedInt dim, CeedInt P_1d, CeedInt Q_1d, const …
/libCEED/backends/sycl-gen/
H A Dceed-sycl-gen-operator-build.sycl.cpp30 …idCalculate_Sycl_gen(const CeedInt dim, const CeedInt P_1d, const CeedInt Q_1d, CeedInt *block_siz… in BlockGridCalculate_Sycl_gen() argument
31 const CeedInt thread1d = CeedIntMax(Q_1d, P_1d); in BlockGridCalculate_Sycl_gen()
66 …CeedInt Q, P_1d = 0, Q_1d = 0, elem_size, num_input_fields, num_output_fields, n… in CeedOperatorBuildKernel_Sycl_gen() local
90 Q_1d = Q; in CeedOperatorBuildKernel_Sycl_gen()
149 CeedCallBackend(CeedBasisGetNumQuadraturePoints1D(basis, &Q_1d)); in CeedOperatorBuildKernel_Sycl_gen()
174 CeedCallBackend(CeedBasisGetNumQuadraturePoints1D(basis, &Q_1d)); in CeedOperatorBuildKernel_Sycl_gen()
184 impl->Q_1d = Q_1d; in CeedOperatorBuildKernel_Sycl_gen()
216 CeedCallBackend(BlockGridCalculate_Sycl_gen(dim, P_1d, Q_1d, block_sizes)); in CeedOperatorBuildKernel_Sycl_gen()
223 code << "#define CEED_Q_VLA " << Q_1d << "\n\n"; in CeedOperatorBuildKernel_Sycl_gen()
264 code << " const CeedInt Q_1D = " << Q_1d << ";\n"; in CeedOperatorBuildKernel_Sycl_gen()
[all …]
H A Dceed-sycl-gen-operator-build.hpp9 …idCalculate_Sycl_gen(const CeedInt dim, const CeedInt P_1d, const CeedInt Q_1d, CeedInt *block_siz…
H A Dceed-sycl-gen.hpp18 CeedInt Q_1d; member
H A Dceed-sycl-gen-operator.sycl.cpp124 const CeedInt Q_1d = impl->Q_1d; in CeedOperatorApplyAdd_Sycl_gen() local
128 CeedCallBackend(BlockGridCalculate_Sycl_gen(dim, P_1d, Q_1d, block_sizes)); in CeedOperatorApplyAdd_Sycl_gen()
/libCEED/backends/sycl-shared/
H A Dceed-sycl-shared-basis.sycl.cpp168 int CeedBasisCreateTensorH1_Sycl_shared(CeedInt dim, CeedInt P_1d, CeedInt Q_1d, const CeedScalar *… in CeedBasisCreateTensorH1_Sycl_shared() argument
182 const CeedInt thread_1d = CeedIntMax(Q_1d, P_1d); in CeedBasisCreateTensorH1_Sycl_shared()
184 const CeedInt num_qpts = CeedIntPow(Q_1d, dim); in CeedBasisCreateTensorH1_Sycl_shared()
196 CeedCallBackend(ComputeLocalRange(ceed, dim, Q_1d, impl->weight_local_range)); in CeedBasisCreateTensorH1_Sycl_shared()
205 …CeedCallSycl(ceed, impl->d_q_weight_1d = sycl::malloc_device<CeedScalar>(Q_1d, data->sycl_device, … in CeedBasisCreateTensorH1_Sycl_shared()
206 …::event copy_weight = data->sycl_queue.copy<CeedScalar>(q_weight_1d, impl->d_q_weight_1d, Q_1d, e); in CeedBasisCreateTensorH1_Sycl_shared()
210 const CeedInt interp_length = Q_1d * P_1d; in CeedBasisCreateTensorH1_Sycl_shared()
223 const bool has_collocated_grad = (dim == 3) && (Q_1d >= P_1d); in CeedBasisCreateTensorH1_Sycl_shared()
227 const CeedInt cgrad_length = Q_1d * Q_1d; in CeedBasisCreateTensorH1_Sycl_shared()
229 CeedCallBackend(CeedMalloc(Q_1d * Q_1d, &collo_grad_1d)); in CeedBasisCreateTensorH1_Sycl_shared()
[all …]
H A Dceed-sycl-shared.hpp33 CEED_INTERN int CeedBasisCreateTensorH1_Sycl_shared(CeedInt dim, CeedInt P_1d, CeedInt Q_1d, const …
/libCEED/backends/ref/
H A Dceed-ref-basis.c51 CeedInt P_1d, Q_1d; in CeedBasisApplyCore_Ref() local
54 CeedCallBackend(CeedBasisGetNumQuadraturePoints1D(basis, &Q_1d)); in CeedBasisApplyCore_Ref()
61 CeedInt P = P_1d, Q = Q_1d; in CeedBasisApplyCore_Ref()
64 P = Q_1d; in CeedBasisApplyCore_Ref()
86 CeedInt P = P_1d, Q = Q_1d; in CeedBasisApplyCore_Ref()
89 P = Q_1d; in CeedBasisApplyCore_Ref()
90 Q = Q_1d; in CeedBasisApplyCore_Ref()
112 P = Q_1d, Q = Q_1d; in CeedBasisApplyCore_Ref()
114 P = Q_1d; in CeedBasisApplyCore_Ref()
148 P = Q_1d; in CeedBasisApplyCore_Ref()
[all …]
H A Dceed-ref.h70 CEED_INTERN int CeedBasisCreateTensorH1_Ref(CeedInt dim, CeedInt P_1d, CeedInt Q_1d, const CeedScal…
/libCEED/backends/hip-ref/
H A Dceed-hip-ref-basis.c24 CeedInt Q_1d, dim; in CeedBasisApplyCore_Hip() local
45 CeedCallBackend(CeedBasisGetNumQuadraturePoints1D(basis, &Q_1d)); in CeedBasisApplyCore_Hip()
52 const CeedInt block_size = CeedIntMin(CeedIntPow(Q_1d, dim), max_block_size); in CeedBasisApplyCore_Hip()
65 const int block_size_x = Q_1d; in CeedBasisApplyCore_Hip()
66 const int block_size_y = dim >= 2 ? Q_1d : 1; in CeedBasisApplyCore_Hip()
104 CeedInt Q_1d, dim, max_num_points = num_points[0]; in CeedBasisApplyAtPointsCore_Hip() local
112 CeedCallBackend(CeedBasisGetNumQuadraturePoints1D(basis, &Q_1d)); in CeedBasisApplyAtPointsCore_Hip()
169 interp_bytes = P_1d * Q_1d * sizeof(CeedScalar); in CeedBasisApplyAtPointsCore_Hip()
170 CeedCallBackend(CeedCalloc(P_1d * Q_1d, &chebyshev_interp_1d)); in CeedBasisApplyAtPointsCore_Hip()
183 …le_Hip(ceed, basis_kernel_source, &data->moduleAtPoints, 9, "BASIS_Q_1D", Q_1d, "BASIS_P_1D", P_1d… in CeedBasisApplyAtPointsCore_Hip()
[all …]
/libCEED/backends/cuda-ref/
H A Dceed-cuda-ref-basis.c25 CeedInt Q_1d, dim; in CeedBasisApplyCore_Cuda() local
45 CeedCallBackend(CeedBasisGetNumQuadraturePoints1D(basis, &Q_1d)); in CeedBasisApplyCore_Cuda()
52 const CeedInt block_size = CeedIntMin(CeedIntPow(Q_1d, dim), max_block_size); in CeedBasisApplyCore_Cuda()
65 const int block_size_x = Q_1d; in CeedBasisApplyCore_Cuda()
66 const int block_size_y = dim >= 2 ? Q_1d : 1; in CeedBasisApplyCore_Cuda()
105 CeedInt Q_1d, dim, max_num_points = num_points[0]; in CeedBasisApplyAtPointsCore_Cuda() local
113 CeedCallBackend(CeedBasisGetNumQuadraturePoints1D(basis, &Q_1d)); in CeedBasisApplyAtPointsCore_Cuda()
170 interp_bytes = P_1d * Q_1d * sizeof(CeedScalar); in CeedBasisApplyAtPointsCore_Cuda()
171 CeedCallBackend(CeedCalloc(P_1d * Q_1d, &chebyshev_interp_1d)); in CeedBasisApplyAtPointsCore_Cuda()
184 …e_Cuda(ceed, basis_kernel_source, &data->moduleAtPoints, 9, "BASIS_Q_1D", Q_1d, "BASIS_P_1D", P_1d… in CeedBasisApplyAtPointsCore_Cuda()
[all …]
H A Dceed-cuda-ref.h157 CEED_INTERN int CeedBasisCreateTensorH1_Cuda(CeedInt dim, CeedInt P_1d, CeedInt Q_1d, const CeedSca…
/libCEED/backends/hip-gen/
H A Dceed-hip-gen-operator-build.cpp34 …st CeedInt dim, const CeedInt num_elem, const CeedInt P_1d, const CeedInt Q_1d, CeedInt *block_siz… in BlockGridCalculate_Hip_gen() argument
35 const CeedInt thread_1d = CeedIntMax(Q_1d, P_1d); in BlockGridCalculate_Hip_gen()
64 …CeedQFunctionField *qf_output_fields, CeedInt *max_P, CeedInt *max_P_1d, CeedInt *Q, CeedInt *Q_1d, in CeedOperatorBuildKernelData_Hip_gen() argument
99 *Q_1d = 0; in CeedOperatorBuildKernelData_Hip_gen()
126 …CeedCheck(*Q_1d == 0 || field_Q_1d == *Q_1d, ceed, CEED_ERROR_BACKEND, "Quadrature spaces must be … in CeedOperatorBuildKernelData_Hip_gen()
127 *Q_1d = field_Q_1d; in CeedOperatorBuildKernelData_Hip_gen()
158 …CeedCheck(*Q_1d == 0 || field_Q_1d == *Q_1d, ceed, CEED_ERROR_BACKEND, "Quadrature spaces must be … in CeedOperatorBuildKernelData_Hip_gen()
159 *Q_1d = field_Q_1d; in CeedOperatorBuildKernelData_Hip_gen()
209 … CeedInt max_dim, CeedInt Q, CeedInt Q_1d, bool is_input, bool is_all_tensor, bool is_at_points, in CeedOperatorBuildKernelFieldData_Hip_gen() argument
262 …code << tab << "const CeedInt " << P_name << " = " << (basis == CEED_BASIS_NONE ? Q_1d : P_1d) << … in CeedOperatorBuildKernelFieldData_Hip_gen()
[all …]
H A Dceed-hip-gen-operator-build.h9 …ridCalculate_Hip_gen(CeedInt dim, CeedInt num_elem, CeedInt P_1d, CeedInt Q_1d, CeedInt *block_siz…
H A Dceed-hip-gen.h17 CeedInt Q, Q_1d; member
/libCEED/backends/cuda-gen/
H A Dceed-cuda-gen-operator-build.cpp37 …CeedQFunctionField *qf_output_fields, CeedInt *max_P, CeedInt *max_P_1d, CeedInt *Q, CeedInt *Q_1d, in CeedOperatorBuildKernelData_Cuda_gen() argument
72 *Q_1d = 0; in CeedOperatorBuildKernelData_Cuda_gen()
99 …CeedCheck(*Q_1d == 0 || field_Q_1d == *Q_1d, ceed, CEED_ERROR_BACKEND, "Quadrature spaces must be … in CeedOperatorBuildKernelData_Cuda_gen()
100 *Q_1d = field_Q_1d; in CeedOperatorBuildKernelData_Cuda_gen()
131 …CeedCheck(*Q_1d == 0 || field_Q_1d == *Q_1d, ceed, CEED_ERROR_BACKEND, "Quadrature spaces must be … in CeedOperatorBuildKernelData_Cuda_gen()
132 *Q_1d = field_Q_1d; in CeedOperatorBuildKernelData_Cuda_gen()
182 … CeedInt max_dim, CeedInt Q, CeedInt Q_1d, bool is_input, bool is_all_tensor, bool is_at_points, in CeedOperatorBuildKernelFieldData_Cuda_gen() argument
235 …code << tab << "const CeedInt " << P_name << " = " << (basis == CEED_BASIS_NONE ? Q_1d : P_1d) << … in CeedOperatorBuildKernelFieldData_Cuda_gen()
252 interp_bytes = P_1d * Q_1d * sizeof(CeedScalar); in CeedOperatorBuildKernelFieldData_Cuda_gen()
253 CeedCallBackend(CeedCalloc(P_1d * Q_1d, &chebyshev_interp_1d)); in CeedOperatorBuildKernelFieldData_Cuda_gen()
[all …]
H A Dceed-cuda-gen.h17 CeedInt Q, Q_1d; member
/libCEED/backends/sycl-ref/
H A Dceed-sycl-ref-basis.sycl.cpp68 const CeedInt Q_1d = kh.get_specialization_constant<BASIS_Q_1D_ID>(); in CeedBasisApplyInterp_Sycl() local
71 const CeedInt num_qpts = CeedIntPow(Q_1d, dim); 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()
171 const CeedInt Q_1d = kh.get_specialization_constant<BASIS_Q_1D_ID>(); in CeedBasisApplyGrad_Sycl() local
174 const CeedInt num_qpts = CeedIntPow(Q_1d, dim); in CeedBasisApplyGrad_Sycl()
175 const CeedInt P = is_transpose ? Q_1d : P_1d; in CeedBasisApplyGrad_Sycl()
176 const CeedInt Q = is_transpose ? P_1d : Q_1d; in CeedBasisApplyGrad_Sycl()
247 const CeedInt Q_1d = impl->Q_1d; in CeedBasisApplyWeight_Sycl() local
250 const CeedInt num_quad_x = Q_1d; in CeedBasisApplyWeight_Sycl()
[all …]
H A Dceed-sycl-ref.hpp49 CeedInt Q_1d; member
117 CEED_INTERN int CeedBasisCreateTensorH1_Sycl(CeedInt dim, CeedInt P_1d, CeedInt Q_1d, const CeedSca…
/libCEED/backends/magma/
H A Dceed-magma-basis.c33 CeedInt dim, num_comp, num_nodes, P_1d, Q_1d, P, Q; in CeedBasisApplyCore_Magma() local
45 CeedCallBackend(CeedBasisGetNumQuadraturePoints1D(basis, &Q_1d)); in CeedBasisApplyCore_Magma()
47 Q = Q_1d; in CeedBasisApplyCore_Magma()
49 P = Q_1d; in CeedBasisApplyCore_Magma()
63 CeedInt elem_qpts_size = CeedIntPow(Q_1d, dim); in CeedBasisApplyCore_Magma()
128 CeedInt elem_qpts_size = CeedIntPow(Q_1d, dim); in CeedBasisApplyCore_Magma()
531 int CeedBasisCreateTensorH1_Magma(CeedInt dim, CeedInt P_1d, CeedInt Q_1d, const CeedScalar *interp… in CeedBasisCreateTensorH1_Magma() argument
548 CeedCallBackend(magma_malloc((void **)&impl->d_q_weight_1d, Q_1d * sizeof(q_weight_1d[0]))); in CeedBasisCreateTensorH1_Magma()
549 …magma_setvector(Q_1d, sizeof(q_weight_1d[0]), q_weight_1d, 1, impl->d_q_weight_1d, 1, data->queue); in CeedBasisCreateTensorH1_Magma()
551 CeedCallBackend(magma_malloc((void **)&impl->d_interp_1d, Q_1d * P_1d * sizeof(interp_1d[0]))); in CeedBasisCreateTensorH1_Magma()
[all …]
H A Dceed-magma.h78 CEED_INTERN int CeedBasisCreateTensorH1_Magma(CeedInt dim, CeedInt P_1d, CeedInt Q_1d, const CeedSc…

12