| /libCEED/include/ceed/jit-source/hip/ |
| H A D | hip-ref-basis-nontensor.h | 21 …erp(const CeedInt num_elem, const CeedScalar *__restrict__ d_B, const CeedScalar *__restrict__ d_U, in Interp() argument 25 … BASIS_NUM_COMP * BASIS_Q * num_elem, d_B, d_U, d_V); in Interp() 29 …ose(const CeedInt num_elem, const CeedScalar *__restrict__ d_B, const CeedScalar *__restrict__ d_U, in InterpTranspose() argument 33 … BASIS_NUM_COMP * BASIS_Q * num_elem, d_B, d_U, d_V); in InterpTranspose() 40 …riv(const CeedInt num_elem, const CeedScalar *__restrict__ d_B, const CeedScalar *__restrict__ d_U, in Deriv() argument 44 … BASIS_NUM_COMP * BASIS_Q * num_elem, d_B, d_U, d_V); in Deriv() 48 …ose(const CeedInt num_elem, const CeedScalar *__restrict__ d_B, const CeedScalar *__restrict__ d_U, in DerivTranspose() argument 52 … BASIS_NUM_COMP * BASIS_Q * num_elem, d_B, d_U, d_V); in DerivTranspose()
|
| H A D | hip-shared-basis-tensor.h | 19 …void Interp(const CeedInt num_elem, const CeedScalar *c_B, const CeedScalar *__restrict__ d_U, Cee… in __launch_bounds__() 40 …tStrided1d<BASIS_NUM_COMP, BASIS_P_1D>(data, elem, 1, BASIS_P_1D * num_elem, BASIS_P_1D, d_U, r_U); in __launch_bounds__() 44 … BASIS_P_1D>(data, elem, 1, BASIS_P_1D * BASIS_P_1D * num_elem, BASIS_P_1D * BASIS_P_1D, d_U, r_U); in __launch_bounds__() 49 … BASIS_P_1D * BASIS_P_1D * BASIS_P_1D, d_U, r_U); in __launch_bounds__() 58 …st CeedInt num_elem, const CeedScalar *c_B, const CeedScalar *__restrict__ d_U, CeedScalar *__rest… in __launch_bounds__() 73 …tStrided1d<BASIS_NUM_COMP, BASIS_P_1D>(data, elem, 1, BASIS_P_1D * num_elem, BASIS_P_1D, d_U, r_U); in __launch_bounds__() 76 … BASIS_P_1D>(data, elem, 1, BASIS_P_1D * BASIS_P_1D * num_elem, BASIS_P_1D * BASIS_P_1D, d_U, r_U); in __launch_bounds__() 80 … BASIS_P_1D * BASIS_P_1D * BASIS_P_1D, d_U, r_U); in __launch_bounds__() 88 …st CeedInt num_elem, const CeedScalar *c_B, const CeedScalar *__restrict__ d_U, CeedScalar *__rest… in __launch_bounds__() 109 …tStrided1d<BASIS_NUM_COMP, BASIS_Q_1D>(data, elem, 1, BASIS_Q_1D * num_elem, BASIS_Q_1D, d_U, r_U); in __launch_bounds__() [all …]
|
| H A D | hip-shared-basis-nontensor.h | 19 …void Interp(const CeedInt num_elem, const CeedScalar *c_B, const CeedScalar *__restrict__ d_U, Cee… in __launch_bounds__() 39 …ReadElementStrided1d<BASIS_NUM_COMP, BASIS_P>(data, elem, 1, BASIS_P * num_elem, BASIS_P, d_U, r_U… in __launch_bounds__() 46 …st CeedInt num_elem, const CeedScalar *c_B, const CeedScalar *__restrict__ d_U, CeedScalar *__rest… in __launch_bounds__() 66 …ReadElementStrided1d<BASIS_NUM_COMP, BASIS_Q>(data, elem, 1, BASIS_Q * num_elem, BASIS_Q, d_U, r_U… in __launch_bounds__() 73 …st CeedInt num_elem, const CeedScalar *c_B, const CeedScalar *__restrict__ d_U, CeedScalar *__rest… in __launch_bounds__() 93 …ReadElementStrided1d<BASIS_NUM_COMP, BASIS_Q>(data, elem, 1, BASIS_Q * num_elem, BASIS_Q, d_U, r_U… in __launch_bounds__() 103 …void Grad(const CeedInt num_elem, const CeedScalar *c_G, const CeedScalar *__restrict__ d_U, CeedS… in __launch_bounds__() 123 …ReadElementStrided1d<BASIS_NUM_COMP, BASIS_P>(data, elem, 1, BASIS_P * num_elem, BASIS_P, d_U, r_U… in __launch_bounds__() 130 …st CeedInt num_elem, const CeedScalar *c_G, const CeedScalar *__restrict__ d_U, CeedScalar *__rest… in __launch_bounds__() 150 …rided1d<BASIS_NUM_COMP * BASIS_DIM, BASIS_Q>(data, elem, 1, BASIS_Q * num_elem, BASIS_Q, d_U, r_U); in __launch_bounds__() [all …]
|
| H A D | hip-ref-basis-nontensor-templates.h | 18 const CeedScalar *__restrict__ d_U, CeedScalar *__restrict__ d_V) { in Contract() argument 26 U = &d_U[elem * strides_elem_U + comp * strides_comp_U]; in Contract() 45 … const CeedScalar *__restrict__ d_U, CeedScalar *__restrict__ d_V) { in ContractTranspose() argument 55 U = &d_U[elem * strides_elem_U + comp * strides_comp_U + d * strides_q_comp_U]; in ContractTranspose()
|
| H A D | hip-shared-basis-tensor-at-points.h | 25 const CeedScalar *__restrict__ d_U, CeedScalar *__restrict__ d_V) { in __launch_bounds__() 49 …tStrided1d<BASIS_NUM_COMP, BASIS_P_1D>(data, elem, 1, BASIS_P_1D * num_elem, BASIS_P_1D, d_U, r_U); in __launch_bounds__() 52 … BASIS_P_1D>(data, elem, 1, BASIS_P_1D * BASIS_P_1D * num_elem, BASIS_P_1D * BASIS_P_1D, d_U, r_U); in __launch_bounds__() 56 … BASIS_P_1D * BASIS_P_1D * BASIS_P_1D, d_U, r_U); in __launch_bounds__() 81 const CeedScalar *__restrict__ d_U, CeedScalar *__restrict__ d_V) { in __launch_bounds__() 124 …M_PTS>(data, elem, i, points_per_elem[elem], 1, num_elem * BASIS_NUM_PTS, BASIS_NUM_PTS, d_U, r_U); in __launch_bounds__() 151 … const CeedScalar *__restrict__ d_U, CeedScalar *__restrict__ d_V) { in __launch_bounds__() 183 …M_PTS>(data, elem, i, points_per_elem[elem], 1, num_elem * BASIS_NUM_PTS, BASIS_NUM_PTS, d_U, r_U); in __launch_bounds__() 213 const CeedScalar *__restrict__ d_U, CeedScalar *__restrict__ d_V) { in __launch_bounds__() 237 …tStrided1d<BASIS_NUM_COMP, BASIS_P_1D>(data, elem, 1, BASIS_P_1D * num_elem, BASIS_P_1D, d_U, r_U); in __launch_bounds__() [all …]
|
| /libCEED/include/ceed/jit-source/cuda/ |
| H A D | cuda-ref-basis-nontensor.h | 21 …erp(const CeedInt num_elem, const CeedScalar *__restrict__ d_B, const CeedScalar *__restrict__ d_U, in Interp() argument 25 … BASIS_NUM_COMP * BASIS_Q * num_elem, d_B, d_U, d_V); in Interp() 29 …ose(const CeedInt num_elem, const CeedScalar *__restrict__ d_B, const CeedScalar *__restrict__ d_U, in InterpTranspose() argument 33 … BASIS_NUM_COMP * BASIS_Q * num_elem, d_B, d_U, d_V); in InterpTranspose() 40 …riv(const CeedInt num_elem, const CeedScalar *__restrict__ d_B, const CeedScalar *__restrict__ d_U, in Deriv() argument 44 … BASIS_NUM_COMP * BASIS_Q * num_elem, d_B, d_U, d_V); in Deriv() 48 …ose(const CeedInt num_elem, const CeedScalar *__restrict__ d_B, const CeedScalar *__restrict__ d_U, in DerivTranspose() argument 52 … BASIS_NUM_COMP * BASIS_Q * num_elem, d_B, d_U, d_V); in DerivTranspose()
|
| H A D | cuda-shared-basis-tensor.h | 18 …st CeedInt num_elem, const CeedScalar *c_B, const CeedScalar *__restrict__ d_U, CeedScalar *__rest… in Interp() argument 39 …tStrided1d<BASIS_NUM_COMP, BASIS_P_1D>(data, elem, 1, BASIS_P_1D * num_elem, BASIS_P_1D, d_U, r_U); in Interp() 43 … BASIS_P_1D>(data, elem, 1, BASIS_P_1D * BASIS_P_1D * num_elem, BASIS_P_1D * BASIS_P_1D, d_U, r_U); in Interp() 48 … BASIS_P_1D * BASIS_P_1D * BASIS_P_1D, d_U, r_U); in Interp() 56 …InterpCollocated(const CeedInt num_elem, const CeedScalar *c_B, const CeedScalar *__restrict__ d_U, in InterpCollocated() argument 72 …tStrided1d<BASIS_NUM_COMP, BASIS_P_1D>(data, elem, 1, BASIS_P_1D * num_elem, BASIS_P_1D, d_U, r_U); in InterpCollocated() 75 … BASIS_P_1D>(data, elem, 1, BASIS_P_1D * BASIS_P_1D * num_elem, BASIS_P_1D * BASIS_P_1D, d_U, r_U); in InterpCollocated() 79 … BASIS_P_1D * BASIS_P_1D * BASIS_P_1D, d_U, r_U); in InterpCollocated() 86 … InterpTranspose(const CeedInt num_elem, const CeedScalar *c_B, const CeedScalar *__restrict__ d_U, in InterpTranspose() argument 108 …tStrided1d<BASIS_NUM_COMP, BASIS_Q_1D>(data, elem, 1, BASIS_Q_1D * num_elem, BASIS_Q_1D, d_U, r_U); in InterpTranspose() [all …]
|
| H A D | cuda-shared-basis-nontensor.h | 18 …st CeedInt num_elem, const CeedScalar *c_B, const CeedScalar *__restrict__ d_U, CeedScalar *__rest… in Interp() argument 38 …ReadElementStrided1d<BASIS_NUM_COMP, BASIS_P>(data, elem, 1, BASIS_P * num_elem, BASIS_P, d_U, r_U… in Interp() 44 … InterpTranspose(const CeedInt num_elem, const CeedScalar *c_B, const CeedScalar *__restrict__ d_U, in InterpTranspose() argument 65 …ReadElementStrided1d<BASIS_NUM_COMP, BASIS_Q>(data, elem, 1, BASIS_Q * num_elem, BASIS_Q, d_U, r_U… in InterpTranspose() 71 …terpTransposeAdd(const CeedInt num_elem, const CeedScalar *c_B, const CeedScalar *__restrict__ d_U, in InterpTransposeAdd() argument 92 …ReadElementStrided1d<BASIS_NUM_COMP, BASIS_Q>(data, elem, 1, BASIS_Q * num_elem, BASIS_Q, d_U, r_U… in InterpTransposeAdd() 101 …st CeedInt num_elem, const CeedScalar *c_G, const CeedScalar *__restrict__ d_U, CeedScalar *__rest… in Grad() argument 121 …ReadElementStrided1d<BASIS_NUM_COMP, BASIS_P>(data, elem, 1, BASIS_P * num_elem, BASIS_P, d_U, r_U… in Grad() 127 …id GradTranspose(const CeedInt num_elem, const CeedScalar *c_G, const CeedScalar *__restrict__ d_U, in GradTranspose() argument 148 …rided1d<BASIS_NUM_COMP * BASIS_DIM, BASIS_Q>(data, elem, 1, BASIS_Q * num_elem, BASIS_Q, d_U, r_U); in GradTranspose() [all …]
|
| H A D | cuda-ref-basis-nontensor-templates.h | 18 const CeedScalar *__restrict__ d_U, CeedScalar *__restrict__ d_V) { in Contract() argument 26 U = d_U + elem * strides_elem_U + comp * strides_comp_U; in Contract() 45 … const CeedScalar *__restrict__ d_U, CeedScalar *__restrict__ d_V) { in ContractTranspose() argument 55 U = &d_U[elem * strides_elem_U + comp * strides_comp_U + d * strides_q_comp_U]; in ContractTranspose()
|
| H A D | cuda-shared-basis-tensor-at-points.h | 24 …const CeedScalar *__restrict__ d_X, const CeedScalar *__restrict__ d_U, CeedScalar *__restrict__ d… in InterpAtPoints() argument 48 …tStrided1d<BASIS_NUM_COMP, BASIS_P_1D>(data, elem, 1, BASIS_P_1D * num_elem, BASIS_P_1D, d_U, r_U); in InterpAtPoints() 51 … BASIS_P_1D>(data, elem, 1, BASIS_P_1D * BASIS_P_1D * num_elem, BASIS_P_1D * BASIS_P_1D, d_U, r_U); in InterpAtPoints() 55 … BASIS_P_1D * BASIS_P_1D * BASIS_P_1D, d_U, r_U); in InterpAtPoints() 80 … const CeedScalar *__restrict__ d_U, CeedScalar *__restrict__ d_V) { in InterpTransposeAtPoints() argument 123 …M_PTS>(data, elem, i, points_per_elem[elem], 1, num_elem * BASIS_NUM_PTS, BASIS_NUM_PTS, d_U, r_U); in InterpTransposeAtPoints() 150 … const CeedScalar *__restrict__ d_U, CeedScalar *__restrict__ d_V) { in InterpTransposeAddAtPoints() argument 182 …M_PTS>(data, elem, i, points_per_elem[elem], 1, num_elem * BASIS_NUM_PTS, BASIS_NUM_PTS, d_U, r_U); in InterpTransposeAddAtPoints() 211 …const CeedScalar *__restrict__ d_X, const CeedScalar *__restrict__ d_U, CeedScalar *__restrict__ d… in GradAtPoints() argument 235 …tStrided1d<BASIS_NUM_COMP, BASIS_P_1D>(data, elem, 1, BASIS_P_1D * num_elem, BASIS_P_1D, d_U, r_U); in GradAtPoints() [all …]
|
| /libCEED/include/ceed/jit-source/sycl/ |
| H A D | sycl-shared-basis-tensor.h | 22 …Int num_elem, global const CeedScalar *restrict d_interp_1d, global const CeedScalar *restrict d_U, in Interp() argument 37 …1d(BASIS_NUM_COMP, BASIS_P_1D, num_elem, 1, BASIS_NUM_NODES * num_elem, BASIS_NUM_NODES, d_U, r_U); in Interp() 42 …2d(BASIS_NUM_COMP, BASIS_P_1D, num_elem, 1, BASIS_NUM_NODES * num_elem, BASIS_NUM_NODES, d_U, r_U); in Interp() 47 …3d(BASIS_NUM_COMP, BASIS_P_1D, num_elem, 1, BASIS_NUM_NODES * num_elem, BASIS_NUM_NODES, d_U, r_U); in Interp() 53 …Int num_elem, global const CeedScalar *restrict d_interp_1d, global const CeedScalar *restrict d_U, in InterpTranspose() argument 71 …ed1d(BASIS_NUM_COMP, BASIS_Q_1D, num_elem, 1, BASIS_NUM_QPTS * num_elem, BASIS_NUM_QPTS, d_U, r_U); in InterpTranspose() 76 …ed2d(BASIS_NUM_COMP, BASIS_Q_1D, num_elem, 1, BASIS_NUM_QPTS * num_elem, BASIS_NUM_QPTS, d_U, r_U); in InterpTranspose() 81 …ed3d(BASIS_NUM_COMP, BASIS_Q_1D, num_elem, 1, BASIS_NUM_QPTS * num_elem, BASIS_NUM_QPTS, d_U, r_U); in InterpTranspose() 91 global const CeedScalar *restrict d_U, global CeedScalar *restrict d_V) { in Grad() argument 108 …1d(BASIS_NUM_COMP, BASIS_P_1D, num_elem, 1, BASIS_NUM_NODES * num_elem, BASIS_NUM_NODES, d_U, r_U); in Grad() [all …]
|
| /libCEED/backends/sycl-ref/ |
| H A D | ceed-sycl-ref-basis.sycl.cpp | 339 const CeedScalar *d_U, CeedScalar *d_V) { in CeedBasisApplyNonTensorInterp_Sycl() argument 364 const CeedScalar *U = d_U + elem * u_stride + comp * u_comp_stride; in CeedBasisApplyNonTensorInterp_Sycl() 380 const CeedScalar *d_U, CeedScalar *d_V) { in CeedBasisApplyNonTensorGrad_Sycl() argument 413 const CeedScalar *U = d_U + elem * u_stride + comp * u_comp_stride + d1 * u_dim_stride; in CeedBasisApplyNonTensorGrad_Sycl()
|