19 …_inline__ void read_1d(const T *devptr, const int compstride, T *sBuffer[NUM_COMP], const int tx) {  in read_1d()  argument
20   if (tx < LENGTH) {  in read_1d()
22       sBuffer[comp][tx] = devptr[comp * compstride + tx];  in read_1d()
31 …e__ __inline__ void write_1d(T *sBuffer[NUM_COMP], T *devptr, const int compstride, const int tx) {  in write_1d()  argument
32   if (tx < LENGTH) {  in write_1d()
34       devptr[comp * compstride + tx] = sBuffer[comp][tx];  in write_1d()
43 …ice__ __inline__ void sum_1d(T *sBuffer[NUM_COMP], T *devptr, const int compstride, const int tx) {  in sum_1d()  argument
44   if (tx < LENGTH) {  in sum_1d()
46       devptr[comp * compstride + tx] += sBuffer[comp][tx];  in sum_1d()
59 …ad_U_2d(const T *dU, const int compstride, T rU[DIM_U][NUM_COMP][rU_SIZE], T *sTmp, const int tx) {  in read_U_2d()  argument
72     if (tx < P) {  in read_U_2d()
74         sTmp[i * P + tx] = dU[comp * compstride + i * P + tx];  in read_U_2d()
79     if (tx < P) {  in read_U_2d()
81         rU[i_DIM][comp][i] = sTmp[tx * P + i];  in read_U_2d()
95 …_ void read_V_2d(const T *dV, const int compstride, T rV[DIM_V][NUM_COMP][rV_SIZE], const int tx) {  in read_V_2d()  argument
96   if (tx < Q) {  in read_V_2d()
99         rV[i_DIM][comp][j] = dV[comp * compstride + j * Q + tx];  in read_V_2d()
112 …line__ void write_V_2d(T *dV, const int compstride, T rV[DIM_V][NUM_COMP][rV_SIZE], const int tx) {  in write_V_2d()  argument
113   if (tx < Q) {  in write_V_2d()
116         dV[comp * compstride + j * Q + tx] = rV[i_DIM][comp][j];  in write_V_2d()
129 …inline__ void sum_V_2d(T *dV, const int compstride, T rV[DIM_V][NUM_COMP][rV_SIZE], const int tx) {  in sum_V_2d()  argument
130   if (tx < Q) {  in sum_V_2d()
133         dV[comp * compstride + j * Q + tx] += rV[i_DIM][comp][j];  in sum_V_2d()
147 …ad_U_3d(const T *dU, const int compstride, T rU[DIM_U][NUM_COMP][rU_SIZE], T *sTmp, const int tx) {  in read_U_3d()  argument
160     if (tx < P * P) {  in read_U_3d()
162         sTmp[i * P * P + tx] = dU[comp * compstride + i * P * P + tx];  in read_U_3d()
167     if (tx < P * P) {  in read_U_3d()
169         rU[i_DIM][comp][i] = sTmp[tx * P + i];  in read_U_3d()
183 …_ void read_V_3d(const T *dV, const int compstride, T rV[DIM_V][NUM_COMP][rV_SIZE], const int tx) {  in read_V_3d()  argument
184   if (tx < Q * Q) {  in read_V_3d()
187         rV[i_DIM][comp][j] = dV[comp * compstride + j * (Q * Q) + tx];  in read_V_3d()
200 …line__ void write_V_3d(T *dV, const int compstride, T rV[DIM_V][NUM_COMP][rV_SIZE], const int tx) {  in write_V_3d()  argument
201   if (tx < (Q * Q)) {  in write_V_3d()
204         dV[comp * compstride + j * (Q * Q) + tx] = rV[i_DIM][comp][j];  in write_V_3d()
217 …inline__ void sum_V_3d(T *dV, const int compstride, T rV[DIM_V][NUM_COMP][rV_SIZE], const int tx) {  in sum_V_3d()  argument
218   if (tx < (Q * Q)) {  in sum_V_3d()
221         dV[comp * compstride + j * (Q * Q) + tx] += rV[i_DIM][comp][j];  in sum_V_3d()
232 static __device__ __inline__ void read_T_notrans_gm2sm(const int tx, const CeedScalar *dT, CeedScal…  in read_T_notrans_gm2sm()  argument
233   if (tx < B) {  in read_T_notrans_gm2sm()
235       sT[i * B + tx] = dT[i * B + tx];  in read_T_notrans_gm2sm()
246 static __device__ __inline__ void read_T_trans_gm2sm(const int tx, const CeedScalar *dT, CeedScalar…  in read_T_trans_gm2sm()  argument
247   if (tx < J) {  in read_T_trans_gm2sm()
249       sT[tx * B + i] = dT[i * J + tx];  in read_T_trans_gm2sm()