| /libCEED/include/ceed/jit-source/magma/ |
| H A D | magma-basis-interp-2d.h | 13 #define sT(i, j) sT[(j) * P + (i)] macro
19 static __device__ __inline__ void magma_interp_2d_device(const T *sT, T rU[DIM_U][NUM_COMP][rU_SIZE… in magma_interp_2d_device() argument
41 rTmp += rU[0][comp][i] * sT(i, j); in magma_interp_2d_device()
56 rTmp += sTmp(tx, i, sld) * sT(i, j); in magma_interp_2d_device()
86 CeedScalar *sT = (CeedScalar *)shared_data; in __launch_bounds__() local
87 CeedScalar *sTmp = sT + BASIS_P * BASIS_Q; in __launch_bounds__()
92 read_T_notrans_gm2sm<BASIS_P, BASIS_Q>(tx, dT, sT); in __launch_bounds__()
99 …magma_interp_2d_device<CeedScalar, 1, 1, BASIS_NUM_COMP, BASIS_P, BASIS_Q, BASIS_P, BASIS_Q>(sT, r… in __launch_bounds__()
127 CeedScalar *sT = (CeedScalar *)shared_data; in __launch_bounds__() local
128 CeedScalar *sTmp = sT + BASIS_Q * BASIS_P; in __launch_bounds__()
[all …]
|
| H A D | magma-basis-grad-1d.h | 13 #define sT(i, j) sT[(j) * P + (i)] macro
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
32 rv += sU[comp][i] * sT(i, tx); in magma_grad_1d_device()
59 CeedScalar *sT = (CeedScalar *)shared_data; in __launch_bounds__() local
60 CeedScalar *sW = sT + BASIS_P * BASIS_Q; in __launch_bounds__()
70 read_T_notrans_gm2sm<BASIS_P, BASIS_Q>(tx, dTgrad, sT); in __launch_bounds__()
77 magma_grad_1d_device<CeedScalar, BASIS_DIM, BASIS_NUM_COMP, BASIS_P, BASIS_Q>(sT, sU, sV, tx); in __launch_bounds__()
104 CeedScalar *sT = (CeedScalar *)shared_data; in __launch_bounds__() local
105 CeedScalar *sW = sT + BASIS_Q * BASIS_P; in __launch_bounds__()
115 read_T_trans_gm2sm<BASIS_Q, BASIS_P>(tx, dTgrad, sT); in __launch_bounds__()
[all …]
|
| H A D | magma-basis-interp-1d.h | 13 #define sT(i, j) sT[(j) * P + (i)] macro
18 static __device__ __inline__ void magma_interp_1d_device(const T *sT, T *sU[NUM_COMP], T *sV[NUM_CO… in magma_interp_1d_device() argument
32 rv += sU[comp][i] * sT(i, tx); // sT[tx * P + i]; in magma_interp_1d_device()
59 CeedScalar *sT = (CeedScalar *)shared_data; in __launch_bounds__() local
60 CeedScalar *sW = sT + BASIS_P * BASIS_Q; in __launch_bounds__()
70 read_T_notrans_gm2sm<BASIS_P, BASIS_Q>(tx, dT, sT); in __launch_bounds__()
77 magma_interp_1d_device<CeedScalar, BASIS_DIM, BASIS_NUM_COMP, BASIS_P, BASIS_Q>(sT, sU, sV, tx); in __launch_bounds__()
104 CeedScalar *sT = (CeedScalar *)shared_data; in __launch_bounds__() local
105 CeedScalar *sW = sT + BASIS_Q * BASIS_P; in __launch_bounds__()
115 read_T_trans_gm2sm<BASIS_Q, BASIS_P>(tx, dT, sT); in __launch_bounds__()
[all …]
|
| H A D | magma-basis-interp-3d.h | 13 #define sT(i, j) sT[(j) * P + (i)] macro
19 static __device__ __inline__ void magma_interp_3d_device(const T *sT, T rU[DIM_U][NUM_COMP][rU_SIZE… in magma_interp_3d_device() argument
41 rTmp[0] += rU[0][comp][i] * sT(i, j); in magma_interp_3d_device()
57 rTmp[j] += sTmp(tx_, i, sld) * sT(i, j); in magma_interp_3d_device()
84 rTmp[0] += sTmp(tx, i, sld) * sT(i, j); in magma_interp_3d_device()
114 CeedScalar *sT = (CeedScalar *)shared_data; in __launch_bounds__() local
115 CeedScalar *sTmp = sT + BASIS_P * BASIS_Q; in __launch_bounds__()
120 read_T_notrans_gm2sm<BASIS_P, BASIS_Q>(tx, dT, sT); in __launch_bounds__()
127 …magma_interp_3d_device<CeedScalar, 1, 1, BASIS_NUM_COMP, BASIS_P, BASIS_Q, BASIS_P, BASIS_Q>(sT, r… in __launch_bounds__()
155 CeedScalar *sT = (CeedScalar *)shared_data; in __launch_bounds__() local
[all …]
|
| H A D | magma-basis-grad-3d.h | 13 #define sT(i, j) sT[(j) * P + (i)] macro
62 const T *sT = (i_DIM == 0) ? sTgrad : sTinterp;
67 rTmp += rU[i_DIM_U][comp][i] * sT(i, j);
79 const T *sT = (i_DIM == 1) ? sTgrad : sTinterp;
85 rTmp += sTmp(tx_, i, sld) * sT(i, j);
97 const T *sT = (i_DIM == 2) ? sTgrad : sTinterp;
102 rTmp += sTmp(tx, i, sld) * sT(i, j);
|
| H A D | magma-basis-grad-2d.h | 13 #define sT(i, j) sT[(j) * P + (i)] macro
59 const T *sT = (i_DIM == 0) ? sTgrad : sTinterp;
64 rTmp += rU[i_DIM_U][comp][i] * sT(i, j);
75 const T *sT = (i_DIM == 1) ? sTgrad : sTinterp;
80 rTmp += sTmp(tx, i, sld) * sT(i, j);
|
| H A D | magma-common-tensor.h | 232 …device__ __inline__ void read_T_notrans_gm2sm(const int tx, const CeedScalar *dT, CeedScalar *sT) { in read_T_notrans_gm2sm() argument
235 sT[i * B + tx] = dT[i * B + tx]; in read_T_notrans_gm2sm()
246 …__device__ __inline__ void read_T_trans_gm2sm(const int tx, const CeedScalar *dT, CeedScalar *sT) { in read_T_trans_gm2sm() argument
249 sT[tx * B + i] = dT[i * J + tx]; in read_T_trans_gm2sm()
|