Lines Matching refs:Q_1d
51 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()
165 const CeedInt opt_elems = block_size / (Q_1d * Q_1d); in CeedBasisApplyTensorCore_Cuda_shared()
169 …CeedCallBackend(CeedRunKernelDim_Cuda(ceed, data->Weight, grid_size, Q_1d, Q_1d, elems_per_block, … in CeedBasisApplyTensorCore_Cuda_shared()
171 const CeedInt opt_elems = block_size / (Q_1d * Q_1d); in CeedBasisApplyTensorCore_Cuda_shared()
175 …CeedCallBackend(CeedRunKernelDim_Cuda(ceed, data->Weight, grid_size, Q_1d, Q_1d, elems_per_block, … in CeedBasisApplyTensorCore_Cuda_shared()
214 CeedInt Q_1d, dim, num_comp, max_num_points = num_points[0]; in CeedBasisApplyAtPointsCore_Cuda_shared() local
221 CeedCallBackend(CeedBasisGetNumQuadraturePoints1D(basis, &Q_1d)); in CeedBasisApplyAtPointsCore_Cuda_shared()
278 interp_bytes = P_1d * Q_1d * sizeof(CeedScalar); in CeedBasisApplyAtPointsCore_Cuda_shared()
279 CeedCallBackend(CeedCalloc(P_1d * Q_1d, &chebyshev_interp_1d)); in CeedBasisApplyAtPointsCore_Cuda_shared()
292 …e_Cuda(ceed, basis_kernel_source, &data->moduleAtPoints, 8, "BASIS_Q_1D", Q_1d, "BASIS_P_1D", P_1d… in CeedBasisApplyAtPointsCore_Cuda_shared()
293 …CeedIntMax(Q_1d, P_1d), "BASIS_DIM", dim, "BASIS_NUM_COMP", num_comp, "BASIS_NUM_NODES", CeedIntPo… in CeedBasisApplyAtPointsCore_Cuda_shared()
294 … "BASIS_NUM_QPTS", CeedIntPow(Q_1d, dim), "BASIS_NUM_PTS", max_num_points)); in CeedBasisApplyAtPointsCore_Cuda_shared()
316 CeedInt P_1d, Q_1d; in CeedBasisApplyAtPointsCore_Cuda_shared() local
319 CeedCallBackend(CeedBasisGetNumQuadraturePoints1D(basis, &Q_1d)); in CeedBasisApplyAtPointsCore_Cuda_shared()
320 CeedInt thread_1d = CeedIntMax(Q_1d, P_1d); in CeedBasisApplyAtPointsCore_Cuda_shared()
366 CeedInt P_1d, Q_1d; in CeedBasisApplyAtPointsCore_Cuda_shared() local
369 CeedCallBackend(CeedBasisGetNumQuadraturePoints1D(basis, &Q_1d)); in CeedBasisApplyAtPointsCore_Cuda_shared()
370 CeedInt thread_1d = CeedIntMax(Q_1d, P_1d); in CeedBasisApplyAtPointsCore_Cuda_shared()
594 int CeedBasisCreateTensorH1_Cuda_shared(CeedInt dim, CeedInt P_1d, CeedInt Q_1d, const CeedScalar *… in CeedBasisCreateTensorH1_Cuda_shared() argument
598 const CeedInt q_bytes = Q_1d * sizeof(CeedScalar); in CeedBasisCreateTensorH1_Cuda_shared()
617 bool has_collocated_grad = dim == 3 && Q_1d >= P_1d; in CeedBasisCreateTensorH1_Cuda_shared()
622 CeedCallBackend(CeedMalloc(Q_1d * Q_1d, &collo_grad_1d)); in CeedBasisCreateTensorH1_Cuda_shared()
624 CeedCallCuda(ceed, cudaMalloc((void **)&data->d_collo_grad_1d, q_bytes * Q_1d)); in CeedBasisCreateTensorH1_Cuda_shared()
625 …CeedCallCuda(ceed, cudaMemcpy(data->d_collo_grad_1d, collo_grad_1d, q_bytes * Q_1d, cudaMemcpyHost… in CeedBasisCreateTensorH1_Cuda_shared()
634 …CeedCallBackend(CeedCompile_Cuda(ceed, basis_kernel_source, &data->module, 8, "BASIS_Q_1D", Q_1d, … in CeedBasisCreateTensorH1_Cuda_shared()
635 …CeedIntMax(Q_1d, P_1d), "BASIS_DIM", dim, "BASIS_NUM_COMP", num_comp, "BASIS_NUM_NODES", CeedIntPo… in CeedBasisCreateTensorH1_Cuda_shared()
636 … "BASIS_NUM_QPTS", CeedIntPow(Q_1d, dim), "BASIS_HAS_COLLOCATED_GRAD", has_collocated_grad)); in CeedBasisCreateTensorH1_Cuda_shared()