| /libCEED/include/ceed/jit-source/hip/ |
| H A D | hip-ref-basis-tensor.h | 142 __device__ void Weight1d(const CeedInt num_elem, const CeedScalar *q_weight_1d, CeedScalar *w) { in Weight1d() argument 148 if (elem < num_elem) w[elem * BASIS_Q_1D + i] = q_weight_1d[i]; in Weight1d() 155 __device__ void Weight2d(const CeedInt num_elem, const CeedScalar *q_weight_1d, CeedScalar *w) { in Weight2d() argument 165 w[ind] = q_weight_1d[i] * q_weight_1d[j]; in Weight2d() 173 __device__ void Weight3d(const CeedInt num_elem, const CeedScalar *q_weight_1d, CeedScalar *w) { in Weight3d() argument 184 w[ind] = q_weight_1d[i] * q_weight_1d[j] * q_weight_1d[k]; in Weight3d() 193 …l__ void Weight(const CeedInt num_elem, const CeedScalar *__restrict__ q_weight_1d, CeedScalar *__… in Weight() argument 194 if (BASIS_DIM == 1) Weight1d(num_elem, q_weight_1d, v); in Weight() 195 else if (BASIS_DIM == 2) Weight2d(num_elem, q_weight_1d, v); in Weight() 196 else if (BASIS_DIM == 3) Weight3d(num_elem, q_weight_1d, v); in Weight()
|
| H A D | hip-shared-basis-tensor-templates.h | 117 inline __device__ void Weight1d(SharedData_Hip &data, const CeedScalar *__restrict__ q_weight_1d, C… in Weight1d() argument 118 *w = (data.t_id_x < Q_1D) ? q_weight_1d[data.t_id_x] : 0.0; in Weight1d() 309 …id WeightTensor2d(SharedData_Hip &data, const CeedScalar *__restrict__ q_weight_1d, CeedScalar *w)… in WeightTensor2d() argument 310 …*w = (data.t_id_x < Q_1D && data.t_id_y < Q_1D) ? q_weight_1d[data.t_id_x] * q_weight_1d[data.t_id… in WeightTensor2d() 646 …id WeightTensor3d(SharedData_Hip &data, const CeedScalar *__restrict__ q_weight_1d, CeedScalar *w)… in WeightTensor3d() argument 648 const CeedScalar pw = quad ? q_weight_1d[data.t_id_x] * q_weight_1d[data.t_id_y] : 0.0; in WeightTensor3d() 650 w[q] = quad ? pw * q_weight_1d[q] : 0.0; in WeightTensor3d()
|
| H A D | hip-shared-basis-tensor-flattened-templates.h | 277 …Tensor2dFlattened(SharedData_Hip &data, const CeedScalar *__restrict__ q_weight_1d, CeedScalar *w)… in WeightTensor2dFlattened() argument 280 *w = (t_id_x < Q_1D && t_id_y < Q_1D) ? q_weight_1d[t_id_x] * q_weight_1d[t_id_y] : 0.0; in WeightTensor2dFlattened() 673 …Tensor3dFlattened(SharedData_Hip &data, const CeedScalar *__restrict__ q_weight_1d, CeedScalar *w)… in WeightTensor3dFlattened() argument 676 …id_x < Q_1D && t_id_y < Q_1D && t_id_z < Q_1D) ? q_weight_1d[t_id_x] * q_weight_1d[t_id_y] * q_wei… in WeightTensor3dFlattened()
|
| /libCEED/include/ceed/jit-source/cuda/ |
| H A D | cuda-ref-basis-tensor.h | 142 __device__ void Weight1d(const CeedInt num_elem, const CeedScalar *q_weight_1d, CeedScalar *w) { in Weight1d() argument 148 if (elem < num_elem) w[elem * BASIS_Q_1D + i] = q_weight_1d[i]; in Weight1d() 155 __device__ void Weight2d(const CeedInt num_elem, const CeedScalar *q_weight_1d, CeedScalar *w) { in Weight2d() argument 165 w[ind] = q_weight_1d[i] * q_weight_1d[j]; in Weight2d() 173 __device__ void Weight3d(const CeedInt num_elem, const CeedScalar *q_weight_1d, CeedScalar *w) { in Weight3d() argument 184 w[ind] = q_weight_1d[i] * q_weight_1d[j] * q_weight_1d[k]; in Weight3d() 193 …l__ void Weight(const CeedInt num_elem, const CeedScalar *__restrict__ q_weight_1d, CeedScalar *__… in Weight() argument 194 if (BASIS_DIM == 1) Weight1d(num_elem, q_weight_1d, v); in Weight() 195 else if (BASIS_DIM == 2) Weight2d(num_elem, q_weight_1d, v); in Weight() 196 else if (BASIS_DIM == 3) Weight3d(num_elem, q_weight_1d, v); in Weight()
|
| H A D | cuda-shared-basis-tensor-templates.h | 117 inline __device__ void Weight1d(SharedData_Cuda &data, const CeedScalar *__restrict__ q_weight_1d, … in Weight1d() argument 118 *w = (data.t_id_x < Q_1D) ? q_weight_1d[data.t_id_x] : 0.0; in Weight1d() 310 …d WeightTensor2d(SharedData_Cuda &data, const CeedScalar *__restrict__ q_weight_1d, CeedScalar *w)… in WeightTensor2d() argument 311 …*w = (data.t_id_x < Q_1D && data.t_id_y < Q_1D) ? q_weight_1d[data.t_id_x] * q_weight_1d[data.t_id… in WeightTensor2d() 648 …d WeightTensor3d(SharedData_Cuda &data, const CeedScalar *__restrict__ q_weight_1d, CeedScalar *w)… in WeightTensor3d() argument 650 const CeedScalar pw = quad ? q_weight_1d[data.t_id_x] * q_weight_1d[data.t_id_y] : 0.0; in WeightTensor3d() 652 w[q] = quad ? pw * q_weight_1d[q] : 0.0; in WeightTensor3d()
|
| H A D | cuda-shared-basis-tensor-flattened-templates.h | 277 …ensor2dFlattened(SharedData_Cuda &data, const CeedScalar *__restrict__ q_weight_1d, CeedScalar *w)… in WeightTensor2dFlattened() argument 280 *w = (t_id_x < Q_1D && t_id_y < Q_1D) ? q_weight_1d[t_id_x] * q_weight_1d[t_id_y] : 0.0; in WeightTensor2dFlattened() 676 …ensor3dFlattened(SharedData_Cuda &data, const CeedScalar *__restrict__ q_weight_1d, CeedScalar *w)… in WeightTensor3dFlattened() argument 679 …id_x < Q_1D && t_id_y < Q_1D && t_id_z < Q_1D) ? q_weight_1d[t_id_x] * q_weight_1d[t_id_y] * q_wei… in WeightTensor3dFlattened()
|
| /libCEED/tests/ |
| H A D | t330-basis.h | 46 CeedScalar q_ref_1d[q], q_weight_1d[q]; in BuildHdivQuadrilateral() local 49 CeedGaussQuadrature(q, q_ref_1d, q_weight_1d); in BuildHdivQuadrilateral() 53 CeedLobattoQuadrature(q, q_ref_1d, q_weight_1d); in BuildHdivQuadrilateral() 69 q_weights[k1] = q_weight_1d[j] * q_weight_1d[i]; in BuildHdivQuadrilateral()
|
| /libCEED/interface/ |
| H A D | ceed-basis.c | 1537 …nst CeedScalar *grad_1d, const CeedScalar *q_ref_1d, const CeedScalar *q_weight_1d, CeedBasis *bas… in CeedBasisCreateTensorH1() argument 1543 …teTensorH1(delegate, dim, num_comp, P_1d, Q_1d, interp_1d, grad_1d, q_ref_1d, q_weight_1d, basis)); in CeedBasisCreateTensorH1() 1567 CeedCall(CeedCalloc(Q_1d, &(*basis)->q_weight_1d)); in CeedBasisCreateTensorH1() 1569 if (q_weight_1d) memcpy((*basis)->q_weight_1d, q_weight_1d, Q_1d * sizeof(q_weight_1d[0])); in CeedBasisCreateTensorH1() 1574 …CeedCall(ceed->BasisCreateTensorH1(dim, P_1d, Q_1d, interp_1d, grad_1d, q_ref_1d, q_weight_1d, *ba… in CeedBasisCreateTensorH1() 1597 CeedScalar c1, c2, c3, c4, dx, *nodes, *interp_1d, *grad_1d, *q_ref_1d, *q_weight_1d; in CeedBasisCreateTensorH1Lagrange() local 1609 CeedCall(CeedCalloc(Q, &q_weight_1d)); in CeedBasisCreateTensorH1Lagrange() 1613 ierr = CeedGaussQuadrature(Q, q_ref_1d, q_weight_1d); in CeedBasisCreateTensorH1Lagrange() 1616 ierr = CeedLobattoQuadrature(Q, q_ref_1d, q_weight_1d); in CeedBasisCreateTensorH1Lagrange() 1645 …dBasisCreateTensorH1(ceed, dim, num_comp, P, Q, interp_1d, grad_1d, q_ref_1d, q_weight_1d, basis)); in CeedBasisCreateTensorH1Lagrange() [all …]
|
| H A D | ceed-fortran.c | 522 … const CeedScalar *grad_1d, const CeedScalar *q_ref_1d, const CeedScalar *q_weight_1d, int *basis, in fCeedBasisCreateTensorH1() argument 529 …ensorH1(Ceed_dict[*ceed], *dim, *num_comp, *P_1d, *Q_1d, interp_1d, grad_1d, q_ref_1d, q_weight_1d, in fCeedBasisCreateTensorH1() 653 CEED_EXTERN void fCeedGaussQuadrature(int *Q, CeedScalar *q_ref_1d, CeedScalar *q_weight_1d, int *e… in fCeedGaussQuadrature() argument 654 *err = CeedGaussQuadrature(*Q, q_ref_1d, q_weight_1d); in fCeedGaussQuadrature() 658 CEED_EXTERN void fCeedLobattoQuadrature(int *Q, CeedScalar *q_ref_1d, CeedScalar *q_weight_1d, int … in fCeedLobattoQuadrature() argument 659 *err = CeedLobattoQuadrature(*Q, q_ref_1d, q_weight_1d); in fCeedLobattoQuadrature()
|
| /libCEED/include/ceed/jit-source/sycl/ |
| H A D | sycl-shared-basis-tensor.h | 163 kernel void Weight(const CeedInt num_elem, global const CeedScalar *restrict q_weight_1d, global Ce… in Weight() argument 170 Weight1d(BASIS_Q_1D, q_weight_1d, r_W); in Weight() 174 WeightTensor2d(BASIS_Q_1D, q_weight_1d, r_W); in Weight() 178 WeightTensor3d(BASIS_Q_1D, q_weight_1d, r_W); in Weight()
|
| H A D | sycl-shared-basis-tensor-templates.h | 97 inline void Weight1d(const CeedInt Q_1D, const CeedScalar *restrict q_weight_1d, CeedScalar *restri… in Weight1d() argument 99 *w = (item_id_x < Q_1D) ? q_weight_1d[item_id_x] : 0.0; in Weight1d() 266 inline void WeightTensor2d(const CeedInt Q_1D, const CeedScalar *restrict q_weight_1d, CeedScalar *… in WeightTensor2d() argument 270 …*w = (item_id_x < Q_1D && item_id_y < Q_1D) ? q_weight_1d[item_id_x] * q_weight_1d[item_id_y] : 0.… in WeightTensor2d() 587 inline void WeightTensor3d(const CeedInt Q_1D, const CeedScalar *restrict q_weight_1d, CeedScalar *… in WeightTensor3d() argument 592 const CeedScalar w_xy = q_weight_1d[item_id_x] * q_weight_1d[item_id_y]; in WeightTensor3d() 593 for (CeedInt q = 0; q < Q_1D; ++q) w[q] = w_xy * q_weight_1d[q]; in WeightTensor3d()
|
| /libCEED/backends/sycl-shared/ |
| H A D | ceed-sycl-shared.hpp | 34 … const CeedScalar *q_ref_1d, const CeedScalar *q_weight_1d, CeedBasis basis);
|
| H A D | ceed-sycl-shared-basis.sycl.cpp | 169 … const CeedScalar *q_ref_1d, const CeedScalar *q_weight_1d, CeedBasis basis) { in CeedBasisCreateTensorH1_Sycl_shared() argument 204 if (q_weight_1d) { in CeedBasisCreateTensorH1_Sycl_shared() 206 …sycl::event copy_weight = data->sycl_queue.copy<CeedScalar>(q_weight_1d, impl->d_q_weight_1d, Q_1d… in CeedBasisCreateTensorH1_Sycl_shared()
|
| /libCEED/backends/sycl-ref/ |
| H A D | ceed-sycl-ref-basis.sycl.cpp | 248 const CeedScalar *q_weight_1d = impl->d_q_weight_1d; in CeedBasisApplyWeight_Sycl() local 260 if (dim == 1) w[work_item.get_linear_id()] = q_weight_1d[work_item[2]]; in CeedBasisApplyWeight_Sycl() 261 …if (dim == 2) w[work_item.get_linear_id()] = q_weight_1d[work_item[2]] * q_weight_1d[work_item[1]]; in CeedBasisApplyWeight_Sycl() 262 … (dim == 3) w[work_item.get_linear_id()] = q_weight_1d[work_item[2]] * q_weight_1d[work_item[1]] *… in CeedBasisApplyWeight_Sycl() 562 … const CeedScalar *q_ref_1d, const CeedScalar *q_weight_1d, CeedBasis basis) { in CeedBasisCreateTensorH1_Sycl() argument 591 if (q_weight_1d) { in CeedBasisCreateTensorH1_Sycl() 593 …sycl::event copy_weight = data->sycl_queue.copy<CeedScalar>(q_weight_1d, impl->d_q_weight_1d, Q_1d… in CeedBasisCreateTensorH1_Sycl()
|
| H A D | ceed-sycl-ref.hpp | 118 … const CeedScalar *q_ref_1d, const CeedScalar *q_weight_1d, CeedBasis basis);
|
| /libCEED/examples/mfem/ |
| H A D | bp1.hpp | 60 mfem::Vector q_ref_1d(ir.GetNPoints()), q_weight_1d(ir.GetNPoints()); in FESpace2Ceed() local 69 q_weight_1d(i) = ip.weight; in FESpace2Ceed() 78 q_ref_1d.GetData(), q_weight_1d.GetData(), basis); in FESpace2Ceed()
|
| H A D | bp3.hpp | 61 mfem::Vector q_ref_1d(ir.GetNPoints()), q_weight_1d(ir.GetNPoints()); in FESpace2Ceed() local 70 q_weight_1d(i) = ip.weight; in FESpace2Ceed() 79 q_ref_1d.GetData(), q_weight_1d.GetData(), basis); in FESpace2Ceed()
|
| /libCEED/backends/hip-shared/ |
| H A D | ceed-hip-shared.h | 42 … const CeedScalar *q_ref_1d, const CeedScalar *q_weight_1d, CeedBasis basis);
|
| /libCEED/backends/cuda-shared/ |
| H A D | ceed-cuda-shared.h | 41 … const CeedScalar *q_ref_1d, const CeedScalar *q_weight_1d, CeedBasis basis);
|
| /libCEED/backends/ref/ |
| H A D | ceed-ref-basis.c | 171 const CeedScalar *q_weight_1d; in CeedBasisApplyCore_Ref() local 174 CeedCallBackend(CeedBasisGetQWeights(basis, &q_weight_1d)); in CeedBasisApplyCore_Ref() 181 … const CeedScalar w = q_weight_1d[j] * (d == 0 ? 1 : v[((i * Q + j) * post + k) * num_elem]); in CeedBasisApplyCore_Ref() 279 … const CeedScalar *q_ref_1d, const CeedScalar *q_weight_1d, CeedBasis basis) { in CeedBasisCreateTensorH1_Ref() argument
|
| H A D | ceed-ref.h | 71 … const CeedScalar *q_ref_1d, const CeedScalar *q_weight_1d, CeedBasis basis);
|
| /libCEED/backends/magma/ |
| H A D | ceed-magma.h | 79 … const CeedScalar *q_ref_1d, const CeedScalar *q_weight_1d, CeedBasis basis);
|
| H A D | ceed-magma-basis.c | 532 … const CeedScalar *q_ref_1d, const CeedScalar *q_weight_1d, CeedBasis basis) { in CeedBasisCreateTensorH1_Magma() argument 547 if (q_weight_1d) { in CeedBasisCreateTensorH1_Magma() 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()
|
| /libCEED/backends/cuda-ref/ |
| H A D | ceed-cuda-ref.h | 158 … const CeedScalar *q_ref_1d, const CeedScalar *q_weight_1d, CeedBasis basis);
|
| /libCEED/backends/hip-ref/ |
| H A D | ceed-hip-ref.h | 163 … const CeedScalar *q_ref_1d, const CeedScalar *q_weight_1d, CeedBasis basis);
|