| /libCEED/include/ceed/jit-source/cuda/ |
| H A D | cuda-shared-basis-tensor-templates.h | 19 template <int NUM_COMP, int P_1D, int Q_1D> 35 template <int NUM_COMP, int P_1D, int Q_1D> 51 template <int NUM_COMP, int P_1D, int Q_1D, int T_1D> 53 for (CeedInt comp = 0; comp < NUM_COMP; comp++) { in Interp1d() 54 ContractX1d<NUM_COMP, P_1D, Q_1D>(data, &r_U[comp], c_B, &r_V[comp]); in Interp1d() 61 template <int NUM_COMP, int P_1D, int Q_1D, int T_1D> 64 for (CeedInt comp = 0; comp < NUM_COMP; comp++) { in InterpTranspose1d() 65 ContractTransposeX1d<NUM_COMP, P_1D, Q_1D>(data, &r_U[comp], c_B, &r_V[comp]); in InterpTranspose1d() 72 template <int NUM_COMP, int P_1D, int Q_1D, int T_1D> 75 for (CeedInt comp = 0; comp < NUM_COMP; comp++) { in InterpCollocatedNodes1d() [all …]
|
| H A D | cuda-shared-basis-tensor-flattened-templates.h | 19 template <int NUM_COMP, int P_1D, int Q_1D, int T_1D> 36 template <int NUM_COMP, int P_1D, int Q_1D, int T_1D> 53 template <int NUM_COMP, int P_1D, int Q_1D, int T_1D> 70 template <int NUM_COMP, int P_1D, int Q_1D, int T_1D> 87 template <int NUM_COMP, int P_1D, int Q_1D, int T_1D> 103 template <int NUM_COMP, int Q_1D, int T_1D> 107 for (CeedInt comp = 0; comp < NUM_COMP; comp++) { in QPack2d() 115 template <int NUM_COMP, int Q_1D, int T_1D> 119 for (CeedInt comp = 0; comp < NUM_COMP; comp++) { in QUnpack2d() 130 template <int NUM_COMP, int P_1D, int Q_1D, int T_1D> [all …]
|
| H A D | cuda-shared-basis-nontensor-templates.h | 15 template <int NUM_COMP, int P_1D, int Q_1D> 31 template <int NUM_COMP, int P_1D, int Q_1D> 46 template <int NUM_COMP, int P, int Q, int T_1D> 49 for (CeedInt comp = 0; comp < NUM_COMP; comp++) { in InterpNonTensor() 50 Contract1d<NUM_COMP, P, Q>(data, &r_U[comp], c_B, &r_V[comp]); in InterpNonTensor() 57 template <int NUM_COMP, int P, int Q, int T_1D> 60 for (CeedInt comp = 0; comp < NUM_COMP; comp++) { in InterpTransposeNonTensor() 62 ContractTranspose1d<NUM_COMP, P, Q>(data, &r_U[comp], c_B, &r_V[comp]); in InterpTransposeNonTensor() 69 template <int NUM_COMP, int DIM, int P, int Q, int T_1D> 72 for (CeedInt comp = 0; comp < NUM_COMP; comp++) { in GradNonTensor() [all …]
|
| H A D | cuda-gen-templates.h | 27 template <int NUM_COMP, int COMP_STRIDE, int NUM_PTS> 32 for (CeedInt comp = 0; comp < NUM_COMP; comp++) { in ReadPoint() 40 template <int NUM_COMP, int COMP_STRIDE, int NUM_PTS> 46 for (CeedInt comp = 0; comp < NUM_COMP; comp++) { in WritePoint() 59 template <int NUM_COMP, int P_1D> 72 template <int NUM_COMP, int COMP_STRIDE, int P_1D> 79 for (CeedInt comp = 0; comp < NUM_COMP; comp++) r_u[comp] = d_u[ind + COMP_STRIDE * comp]; in ReadLVecStandard1d() 86 template <int NUM_COMP, int P_1D, int STRIDES_NODE, int STRIDES_COMP, int STRIDES_ELEM> 93 for (CeedInt comp = 0; comp < NUM_COMP; comp++) r_u[comp] = d_u[ind + comp * STRIDES_COMP]; in ReadLVecStrided1d() 100 template <int NUM_COMP, int COMP_STRIDE, int P_1D> [all …]
|
| H A D | cuda-shared-basis-read-write-templates.h | 27 template <int NUM_COMP, int P_1D> 34 for (CeedInt comp = 0; comp < NUM_COMP; comp++) { in ReadElementStrided1d() 43 template <int NUM_COMP, int P_1D> 50 for (CeedInt comp = 0; comp < NUM_COMP; comp++) { in WriteElementStrided1d() 56 template <int NUM_COMP, int P_1D> 63 for (CeedInt comp = 0; comp < NUM_COMP; comp++) { in SumElementStrided1d() 76 template <int NUM_COMP, int P_1D> 83 for (CeedInt comp = 0; comp < NUM_COMP; comp++) { in ReadElementStrided2d() 92 template <int NUM_COMP, int P_1D> 99 for (CeedInt comp = 0; comp < NUM_COMP; comp++) { in WriteElementStrided2d() [all …]
|
| H A D | cuda-shared-basis-tensor-at-points-templates.h | 43 template <int NUM_COMP, int NUM_POINTS, int P_1D, int Q_1D> 48 for (CeedInt i = 0; i < NUM_COMP; i++) r_V[i] = 0.0; in InterpAtPoints1d() 50 for (CeedInt comp = 0; comp < NUM_COMP; comp++) { in InterpAtPoints1d() 64 template <int NUM_COMP, int NUM_POINTS, int P_1D, int Q_1D> 70 for (CeedInt comp = 0; comp < NUM_COMP; comp++) { in InterpTransposeAtPoints1d() 89 template <int NUM_COMP, int NUM_POINTS, int P_1D, int Q_1D> 95 for (CeedInt i = 0; i < NUM_COMP; i++) r_V[i] = 0.0; in GradAtPoints1d() 96 for (CeedInt comp = 0; comp < NUM_COMP; comp++) { in GradAtPoints1d() 111 template <int NUM_COMP, int NUM_POINTS, int P_1D, int Q_1D> 117 for (CeedInt comp = 0; comp < NUM_COMP; comp++) { in GradTransposeAtPoints1d() [all …]
|
| H A D | cuda-ref-operator-assemble-diagonal.h | 82 for (IndexType comp_out = 0; comp_out < NUM_COMP; comp_out++) { in __launch_bounds__() 85 for (IndexType comp_in = 0; comp_in < NUM_COMP; comp_in++) { in __launch_bounds__() 91 …assembled_qf_array[((((e_in * NUM_COMP + comp_in) * NUM_EVAL_MODES_OUT + e_out) * NUM_COMP + comp_… in __launch_bounds__() 96 … elem_diag_array[((comp_out * NUM_COMP + comp_in) * num_elem + e) * NUM_NODES + tid] += e_value; in __launch_bounds__() 105 …assembled_qf_array[((((e_in * NUM_COMP + comp_out) * NUM_EVAL_MODES_OUT + e_out) * NUM_COMP + comp… in __launch_bounds__()
|
| H A D | cuda-ref-basis-nontensor-templates.h | 15 template <int NUM_COMP, int Q_COMP, int P, int Q> 24 for (CeedInt comp = 0; comp < NUM_COMP; comp++) { in Contract() 42 template <int NUM_COMP, int Q_COMP, int P, int Q> 51 for (CeedInt comp = 0; comp < NUM_COMP; comp++) { in ContractTranspose()
|
| /libCEED/include/ceed/jit-source/hip/ |
| H A D | hip-shared-basis-tensor-templates.h | 19 template <int NUM_COMP, int P_1D, int Q_1D> 35 template <int NUM_COMP, int P_1D, int Q_1D> 51 template <int NUM_COMP, int P_1D, int Q_1D, int T_1D> 53 for (CeedInt comp = 0; comp < NUM_COMP; comp++) { in Interp1d() 54 ContractX1d<NUM_COMP, P_1D, Q_1D>(data, &r_U[comp], c_B, &r_V[comp]); in Interp1d() 61 template <int NUM_COMP, int P_1D, int Q_1D, int T_1D> 64 for (CeedInt comp = 0; comp < NUM_COMP; comp++) { in InterpTranspose1d() 65 ContractTransposeX1d<NUM_COMP, P_1D, Q_1D>(data, &r_U[comp], c_B, &r_V[comp]); in InterpTranspose1d() 72 template <int NUM_COMP, int P_1D, int Q_1D, int T_1D> 75 for (CeedInt comp = 0; comp < NUM_COMP; comp++) { in InterpCollocatedNodes1d() [all …]
|
| H A D | hip-shared-basis-tensor-flattened-templates.h | 19 template <int NUM_COMP, int P_1D, int Q_1D, int T_1D> 36 template <int NUM_COMP, int P_1D, int Q_1D, int T_1D> 53 template <int NUM_COMP, int P_1D, int Q_1D, int T_1D> 70 template <int NUM_COMP, int P_1D, int Q_1D, int T_1D> 87 template <int NUM_COMP, int P_1D, int Q_1D, int T_1D> 103 template <int NUM_COMP, int Q_1D, int T_1D> 107 for (CeedInt comp = 0; comp < NUM_COMP; comp++) { in QPack2d() 115 template <int NUM_COMP, int Q_1D, int T_1D> 119 for (CeedInt comp = 0; comp < NUM_COMP; comp++) { in QUnpack2d() 130 template <int NUM_COMP, int P_1D, int Q_1D, int T_1D> [all …]
|
| H A D | hip-shared-basis-nontensor-templates.h | 15 template <int NUM_COMP, int P_1D, int Q_1D> 31 template <int NUM_COMP, int P_1D, int Q_1D> 46 template <int NUM_COMP, int P, int Q, int T_1D> 49 for (CeedInt comp = 0; comp < NUM_COMP; comp++) { in InterpNonTensor() 50 Contract1d<NUM_COMP, P, Q>(data, &r_U[comp], c_B, &r_V[comp]); in InterpNonTensor() 57 template <int NUM_COMP, int P, int Q, int T_1D> 60 for (CeedInt comp = 0; comp < NUM_COMP; comp++) { in InterpTransposeNonTensor() 62 ContractTranspose1d<NUM_COMP, P, Q>(data, &r_U[comp], c_B, &r_V[comp]); in InterpTransposeNonTensor() 69 template <int NUM_COMP, int DIM, int P, int Q, int T_1D> 72 for (CeedInt comp = 0; comp < NUM_COMP; comp++) { in GradNonTensor() [all …]
|
| H A D | hip-gen-templates.h | 27 template <int NUM_COMP, int COMP_STRIDE, int NUM_PTS> 32 for (CeedInt comp = 0; comp < NUM_COMP; comp++) { in ReadPoint() 40 template <int NUM_COMP, int COMP_STRIDE, int NUM_PTS> 46 for (CeedInt comp = 0; comp < NUM_COMP; comp++) { in WritePoint() 59 template <int NUM_COMP, int P_1D> 72 template <int NUM_COMP, int COMP_STRIDE, int P_1D> 79 for (CeedInt comp = 0; comp < NUM_COMP; comp++) r_u[comp] = d_u[ind + COMP_STRIDE * comp]; in ReadLVecStandard1d() 86 template <int NUM_COMP, int P_1D, int STRIDES_NODE, int STRIDES_COMP, int STRIDES_ELEM> 92 for (CeedInt comp = 0; comp < NUM_COMP; comp++) r_u[comp] = d_u[ind + comp * STRIDES_COMP]; in ReadLVecStrided1d() 99 template <int NUM_COMP, int COMP_STRIDE, int P_1D> [all …]
|
| H A D | hip-shared-basis-read-write-templates.h | 27 template <int NUM_COMP, int P_1D> 34 for (CeedInt comp = 0; comp < NUM_COMP; comp++) { in ReadElementStrided1d() 43 template <int NUM_COMP, int P_1D> 50 for (CeedInt comp = 0; comp < NUM_COMP; comp++) { in WriteElementStrided1d() 56 template <int NUM_COMP, int P_1D> 63 for (CeedInt comp = 0; comp < NUM_COMP; comp++) { in SumElementStrided1d() 76 template <int NUM_COMP, int P_1D> 83 for (CeedInt comp = 0; comp < NUM_COMP; comp++) { in ReadElementStrided2d() 92 template <int NUM_COMP, int P_1D> 99 for (CeedInt comp = 0; comp < NUM_COMP; comp++) { in WriteElementStrided2d() [all …]
|
| H A D | hip-shared-basis-tensor-at-points-templates.h | 43 template <int NUM_COMP, int NUM_POINTS, int P_1D, int Q_1D> 48 for (CeedInt i = 0; i < NUM_COMP; i++) r_V[i] = 0.0; in InterpAtPoints1d() 50 for (CeedInt comp = 0; comp < NUM_COMP; comp++) { in InterpAtPoints1d() 65 template <int NUM_COMP, int NUM_POINTS, int P_1D, int Q_1D> 71 for (CeedInt comp = 0; comp < NUM_COMP; comp++) { in InterpTransposeAtPoints1d() 90 template <int NUM_COMP, int NUM_POINTS, int P_1D, int Q_1D> 96 for (CeedInt i = 0; i < NUM_COMP; i++) r_V[i] = 0.0; in GradAtPoints1d() 97 for (CeedInt comp = 0; comp < NUM_COMP; comp++) { in GradAtPoints1d() 112 template <int NUM_COMP, int NUM_POINTS, int P_1D, int Q_1D> 118 for (CeedInt comp = 0; comp < NUM_COMP; comp++) { in GradTransposeAtPoints1d() [all …]
|
| H A D | hip-ref-operator-assemble-diagonal.h | 82 for (IndexType comp_out = 0; comp_out < NUM_COMP; comp_out++) { in __launch_bounds__() 85 for (IndexType comp_in = 0; comp_in < NUM_COMP; comp_in++) { in __launch_bounds__() 91 …assembled_qf_array[((((e_in * NUM_COMP + comp_in) * NUM_EVAL_MODES_OUT + e_out) * NUM_COMP + comp_… in __launch_bounds__() 96 … elem_diag_array[((comp_out * NUM_COMP + comp_in) * num_elem + e) * NUM_NODES + tid] += e_value; in __launch_bounds__() 105 …assembled_qf_array[((((e_in * NUM_COMP + comp_out) * NUM_EVAL_MODES_OUT + e_out) * NUM_COMP + comp… in __launch_bounds__()
|
| H A D | hip-ref-basis-nontensor-templates.h | 15 template <int NUM_COMP, int Q_COMP, int P, int Q> 24 for (CeedInt comp = 0; comp < NUM_COMP; comp++) { in Contract() 42 template <int NUM_COMP, int Q_COMP, int P, int Q> 51 for (CeedInt comp = 0; comp < NUM_COMP; comp++) { in ContractTranspose()
|
| /libCEED/include/ceed/jit-source/magma/ |
| H A D | magma-common-tensor.h | 18 template <typename T, int LENGTH, int NUM_COMP> 19 …_inline__ void read_1d(const T *devptr, const int compstride, T *sBuffer[NUM_COMP], const int tx) { in read_1d() argument 21 for (int comp = 0; comp < NUM_COMP; comp++) { in read_1d() 30 template <typename T, int LENGTH, int NUM_COMP> 31 static __device__ __inline__ void write_1d(T *sBuffer[NUM_COMP], T *devptr, const int compstride, c… in write_1d() argument 33 for (int comp = 0; comp < NUM_COMP; comp++) { in write_1d() 42 template <typename T, int LENGTH, int NUM_COMP> 43 static __device__ __inline__ void sum_1d(T *sBuffer[NUM_COMP], T *devptr, const int compstride, con… in sum_1d() argument 45 for (int comp = 0; comp < NUM_COMP; comp++) { in sum_1d() 58 template <typename T, int P, int DIM_U, int NUM_COMP, int rU_SIZE, int i_DIM> [all …]
|
| H A D | magma-basis-weight-2d.h | 14 template <typename T, int DIM, int NUM_COMP, int Q, int i_DIM, int i_COMP> 15 static __device__ __inline__ void magma_weight_2d_device(const T *sTweight, T rV[DIM][NUM_COMP][Q],… in magma_weight_2d_device()
|
| H A D | magma-basis-weight-3d.h | 14 template <typename T, int DIM, int NUM_COMP, int Q, int i_DIM, int i_COMP> 15 static __device__ __inline__ void magma_weight_3d_device(const T *sTweight, T rV[DIM][NUM_COMP][Q],… in magma_weight_3d_device()
|
| H A D | magma-basis-grad-1d.h | 17 template <typename T, int DIM, int NUM_COMP, int P, int Q> 18 static __device__ __inline__ void magma_grad_1d_device(const T *sT, T *sU[NUM_COMP], T *sV[NUM_COMP… in magma_grad_1d_device() argument 29 for (int comp = 0; comp < NUM_COMP; comp++) { in magma_grad_1d_device()
|
| H A D | magma-basis-interp-1d.h | 17 template <typename T, int DIM, int NUM_COMP, int P, int Q> 18 …ice__ __inline__ void magma_interp_1d_device(const T *sT, T *sU[NUM_COMP], T *sV[NUM_COMP], const … in magma_interp_1d_device() argument 29 for (int comp = 0; comp < NUM_COMP; comp++) { in magma_interp_1d_device()
|
| H A D | magma-basis-interp-2d.h | 18 template <typename T, int DIM_U, int DIM_V, int NUM_COMP, int P, int Q, int rU_SIZE, int rV_SIZE> 19 …__ void magma_interp_2d_device(const T *sT, T rU[DIM_U][NUM_COMP][rU_SIZE], T rV[DIM_V][NUM_COMP][… in magma_interp_2d_device() 31 for (int comp = 0; comp < NUM_COMP; comp++) { in magma_interp_2d_device()
|
| H A D | magma-basis-interp-3d.h | 18 template <typename T, int DIM_U, int DIM_V, int NUM_COMP, int P, int Q, int rU_SIZE, int rV_SIZE> 19 …__ void magma_interp_3d_device(const T *sT, T rU[DIM_U][NUM_COMP][rU_SIZE], T rV[DIM_V][NUM_COMP][… in magma_interp_3d_device() 32 for (int comp = 0; comp < NUM_COMP; comp++) { in magma_interp_3d_device()
|
| /libCEED/include/ceed/jit-source/sycl/ |
| H A D | sycl-shared-basis-read-write-templates.h | 28 inline void ReadElementStrided1d(const CeedInt NUM_COMP, const CeedInt P_1D, const CeedInt num_elem… in ReadElementStrided1d() argument 37 for (CeedInt comp = 0; comp < NUM_COMP; comp++) { in ReadElementStrided1d() 46 inline void WriteElementStrided1d(const CeedInt NUM_COMP, const CeedInt P_1D, const CeedInt num_ele… in WriteElementStrided1d() argument 55 for (CeedInt comp = 0; comp < NUM_COMP; comp++) { in WriteElementStrided1d() 68 inline void ReadElementStrided2d(const CeedInt NUM_COMP, const CeedInt P_1D, const CeedInt num_elem… in ReadElementStrided2d() argument 78 for (CeedInt comp = 0; comp < NUM_COMP; comp++) { in ReadElementStrided2d() 87 inline void WriteElementStrided2d(const CeedInt NUM_COMP, const CeedInt P_1D, const CeedInt num_ele… in WriteElementStrided2d() argument 97 for (CeedInt comp = 0; comp < NUM_COMP; comp++) { in WriteElementStrided2d() 110 inline void ReadElementStrided3d(const CeedInt NUM_COMP, const CeedInt P_1D, const CeedInt num_elem… in ReadElementStrided3d() argument 121 for (CeedInt comp = 0; comp < NUM_COMP; comp++) { in ReadElementStrided3d() [all …]
|
| H A D | sycl-shared-basis-tensor-templates.h | 57 inline void Interp1d(const CeedInt NUM_COMP, const CeedInt P_1D, const CeedInt Q_1D, private const … in Interp1d() argument 59 for (CeedInt comp = 0; comp < NUM_COMP; comp++) { in Interp1d() 67 inline void InterpTranspose1d(const CeedInt NUM_COMP, const CeedInt P_1D, const CeedInt Q_1D, priva… in InterpTranspose1d() argument 69 for (CeedInt comp = 0; comp < NUM_COMP; comp++) { in InterpTranspose1d() 77 inline void Grad1d(const CeedInt NUM_COMP, const CeedInt P_1D, const CeedInt Q_1D, private const Ce… in Grad1d() argument 79 for (CeedInt comp = 0; comp < NUM_COMP; comp++) { in Grad1d() 87 inline void GradTranspose1d(const CeedInt NUM_COMP, const CeedInt P_1D, const CeedInt Q_1D, private… in GradTranspose1d() argument 89 for (CeedInt comp = 0; comp < NUM_COMP; comp++) { in GradTranspose1d() 208 inline void InterpTensor2d(const CeedInt NUM_COMP, const CeedInt P_1D, const CeedInt Q_1D, private … in InterpTensor2d() argument 212 for (CeedInt comp = 0; comp < NUM_COMP; comp++) { in InterpTensor2d() [all …]
|