39 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()
217       CeedInt Q_1d;  in CeedBasisApplyTensorCore_Hip_shared()  local
221       CeedCallBackend(CeedBasisGetNumQuadraturePoints1D(basis, &Q_1d));  in CeedBasisApplyTensorCore_Hip_shared()
225         const CeedInt opt_elems       = block_size / Q_1d;  in CeedBasisApplyTensorCore_Hip_shared()
229 …CeedCallBackend(CeedRunKernelDim_Hip(ceed, data->Weight, grid_size, Q_1d, elems_per_block, 1, weig…  in CeedBasisApplyTensorCore_Hip_shared()
231         const CeedInt opt_elems       = block_size / (Q_1d * Q_1d);  in CeedBasisApplyTensorCore_Hip_shared()
235 …CeedCallBackend(CeedRunKernelDim_Hip(ceed, data->Weight, grid_size, Q_1d, Q_1d, elems_per_block, w…  in CeedBasisApplyTensorCore_Hip_shared()
237         const CeedInt opt_elems       = block_size / (Q_1d * Q_1d);  in CeedBasisApplyTensorCore_Hip_shared()
241 …CeedCallBackend(CeedRunKernelDim_Hip(ceed, data->Weight, grid_size, Q_1d, Q_1d, elems_per_block, w…  in CeedBasisApplyTensorCore_Hip_shared()
279   CeedInt               Q_1d, dim, max_num_points = num_points[0];  in CeedBasisApplyAtPointsCore_Hip_shared()  local
286   CeedCallBackend(CeedBasisGetNumQuadraturePoints1D(basis, &Q_1d));  in CeedBasisApplyAtPointsCore_Hip_shared()
343       interp_bytes = P_1d * Q_1d * sizeof(CeedScalar);  in CeedBasisApplyAtPointsCore_Hip_shared()
344       CeedCallBackend(CeedCalloc(P_1d * Q_1d, &chebyshev_interp_1d));  in CeedBasisApplyAtPointsCore_Hip_shared()
357 …le_Hip(ceed, basis_kernel_source, &data->moduleAtPoints, 9, "BASIS_Q_1D", Q_1d, "BASIS_P_1D", P_1d…  in CeedBasisApplyAtPointsCore_Hip_shared()
358 …CeedIntMax(Q_1d, P_1d), "BASIS_DIM", dim, "BASIS_NUM_COMP", num_comp, "BASIS_NUM_NODES", CeedIntPo…  in CeedBasisApplyAtPointsCore_Hip_shared()
359 …"BASIS_NUM_QPTS", CeedIntPow(Q_1d, dim), "BASIS_NUM_PTS", max_num_points, "BASIS_INTERP_BLOCK_SIZE…  in CeedBasisApplyAtPointsCore_Hip_shared()
382       CeedInt P_1d, Q_1d;  in CeedBasisApplyAtPointsCore_Hip_shared()  local
386       CeedCallBackend(CeedBasisGetNumQuadraturePoints1D(basis, &Q_1d));  in CeedBasisApplyAtPointsCore_Hip_shared()
387       CeedInt thread_1d     = CeedIntMax(Q_1d, P_1d);  in CeedBasisApplyAtPointsCore_Hip_shared()
430       CeedInt P_1d, Q_1d;  in CeedBasisApplyAtPointsCore_Hip_shared()  local
434       CeedCallBackend(CeedBasisGetNumQuadraturePoints1D(basis, &Q_1d));  in CeedBasisApplyAtPointsCore_Hip_shared()
435       CeedInt thread_1d   = CeedIntMax(Q_1d, P_1d);  in CeedBasisApplyAtPointsCore_Hip_shared()
655 int CeedBasisCreateTensorH1_Hip_shared(CeedInt dim, CeedInt P_1d, CeedInt Q_1d, const CeedScalar *i…  in CeedBasisCreateTensorH1_Hip_shared()  argument
659   const CeedInt         q_bytes      = Q_1d * sizeof(CeedScalar);  in CeedBasisCreateTensorH1_Hip_shared()
678   bool has_collocated_grad = dim == 3 && Q_1d >= P_1d;  in CeedBasisCreateTensorH1_Hip_shared()
683     CeedCallBackend(CeedMalloc(Q_1d * Q_1d, &collo_grad_1d));  in CeedBasisCreateTensorH1_Hip_shared()
685     CeedCallHip(ceed, hipMalloc((void **)&data->d_collo_grad_1d, q_bytes * Q_1d));  in CeedBasisCreateTensorH1_Hip_shared()
686 …CeedCallHip(ceed, hipMemcpy(data->d_collo_grad_1d, collo_grad_1d, q_bytes * Q_1d, hipMemcpyHostToD…  in CeedBasisCreateTensorH1_Hip_shared()
692   CeedCallBackend(ComputeBasisThreadBlockSizes(dim, P_1d, Q_1d, num_comp, data->block_sizes));  in CeedBasisCreateTensorH1_Hip_shared()
698 …CeedCallBackend(CeedCompile_Hip(ceed, basis_kernel_source, &data->module, 11, "BASIS_Q_1D", Q_1d, …  in CeedBasisCreateTensorH1_Hip_shared()
699 …CeedIntMax(Q_1d, P_1d), "BASIS_DIM", dim, "BASIS_NUM_COMP", num_comp, "BASIS_NUM_NODES", CeedIntPo…  in CeedBasisCreateTensorH1_Hip_shared()
700 …"BASIS_NUM_QPTS", CeedIntPow(Q_1d, dim), "BASIS_INTERP_BLOCK_SIZE", data->block_sizes[0], "BASIS_G…  in CeedBasisCreateTensorH1_Hip_shared()