17   for (CeedInt i = data.t_id; i < P * Q; i += blockDim.x * blockDim.y * blockDim.z) B[i] = d_B[i];  in LoadMatrix()
28 …eadElementStrided1d(SharedData_Cuda &data, const CeedInt elem, const CeedInt strides_node, const C…  in ReadElementStrided1d()
29 …                 const CeedInt strides_elem, const CeedScalar *__restrict__ d_u, CeedScalar *r_u) {  in ReadElementStrided1d()
31     const CeedInt node = data.t_id_x;  in ReadElementStrided1d()
32     const CeedInt ind  = node * strides_node + elem * strides_elem;  in ReadElementStrided1d()
34     for (CeedInt comp = 0; comp < NUM_COMP; comp++) {  in ReadElementStrided1d()
44 …iteElementStrided1d(SharedData_Cuda &data, const CeedInt elem, const CeedInt strides_node, const C…  in WriteElementStrided1d()
45 …                              const CeedInt strides_elem, const CeedScalar *r_v, CeedScalar *d_v) {  in WriteElementStrided1d()
47     const CeedInt node = data.t_id_x;  in WriteElementStrided1d()
48     const CeedInt ind  = node * strides_node + elem * strides_elem;  in WriteElementStrided1d()
50     for (CeedInt comp = 0; comp < NUM_COMP; comp++) {  in WriteElementStrided1d()
57 …SumElementStrided1d(SharedData_Cuda &data, const CeedInt elem, const CeedInt strides_node, const C…  in SumElementStrided1d()
58 …                              const CeedInt strides_elem, const CeedScalar *r_v, CeedScalar *d_v) {  in SumElementStrided1d()
60     const CeedInt node = data.t_id_x;  in SumElementStrided1d()
61     const CeedInt ind  = node * strides_node + elem * strides_elem;  in SumElementStrided1d()
63     for (CeedInt comp = 0; comp < NUM_COMP; comp++) {  in SumElementStrided1d()
77 …eadElementStrided2d(SharedData_Cuda &data, const CeedInt elem, const CeedInt strides_node, const C…  in ReadElementStrided2d()
78 …                 const CeedInt strides_elem, const CeedScalar *__restrict__ d_u, CeedScalar *r_u) {  in ReadElementStrided2d()
80     const CeedInt node = data.t_id_x + data.t_id_y * P_1D;  in ReadElementStrided2d()
81     const CeedInt ind  = node * strides_node + elem * strides_elem;  in ReadElementStrided2d()
83     for (CeedInt comp = 0; comp < NUM_COMP; comp++) {  in ReadElementStrided2d()
93 …iteElementStrided2d(SharedData_Cuda &data, const CeedInt elem, const CeedInt strides_node, const C…  in WriteElementStrided2d()
94 …                              const CeedInt strides_elem, const CeedScalar *r_v, CeedScalar *d_v) {  in WriteElementStrided2d()
96     const CeedInt node = data.t_id_x + data.t_id_y * P_1D;  in WriteElementStrided2d()
97     const CeedInt ind  = node * strides_node + elem * strides_elem;  in WriteElementStrided2d()
99     for (CeedInt comp = 0; comp < NUM_COMP; comp++) {  in WriteElementStrided2d()
106 …SumElementStrided2d(SharedData_Cuda &data, const CeedInt elem, const CeedInt strides_node, const C…  in SumElementStrided2d()
107 …                              const CeedInt strides_elem, const CeedScalar *r_v, CeedScalar *d_v) {  in SumElementStrided2d()
109     const CeedInt node = data.t_id_x + data.t_id_y * P_1D;  in SumElementStrided2d()
110     const CeedInt ind  = node * strides_node + elem * strides_elem;  in SumElementStrided2d()
112     for (CeedInt comp = 0; comp < NUM_COMP; comp++) {  in SumElementStrided2d()
126 …eadElementStrided3d(SharedData_Cuda &data, const CeedInt elem, const CeedInt strides_node, const C…  in ReadElementStrided3d()
127 …                 const CeedInt strides_elem, const CeedScalar *__restrict__ d_u, CeedScalar *r_u) {  in ReadElementStrided3d()
129     for (CeedInt z = 0; z < P_1D; z++) {  in ReadElementStrided3d()
130       const CeedInt node = data.t_id_x + data.t_id_y * P_1D + z * P_1D * P_1D;  in ReadElementStrided3d()
131       const CeedInt ind  = node * strides_node + elem * strides_elem;  in ReadElementStrided3d()
133       for (CeedInt comp = 0; comp < NUM_COMP; comp++) {  in ReadElementStrided3d()
144 …iteElementStrided3d(SharedData_Cuda &data, const CeedInt elem, const CeedInt strides_node, const C…  in WriteElementStrided3d()
145 …                              const CeedInt strides_elem, const CeedScalar *r_v, CeedScalar *d_v) {  in WriteElementStrided3d()
147     for (CeedInt z = 0; z < P_1D; z++) {  in WriteElementStrided3d()
148       const CeedInt node = data.t_id_x + data.t_id_y * P_1D + z * P_1D * P_1D;  in WriteElementStrided3d()
149       const CeedInt ind  = node * strides_node + elem * strides_elem;  in WriteElementStrided3d()
151       for (CeedInt comp = 0; comp < NUM_COMP; comp++) {  in WriteElementStrided3d()
159 …SumElementStrided3d(SharedData_Cuda &data, const CeedInt elem, const CeedInt strides_node, const C…  in SumElementStrided3d()
160 …                              const CeedInt strides_elem, const CeedScalar *r_v, CeedScalar *d_v) {  in SumElementStrided3d()
162     for (CeedInt z = 0; z < P_1D; z++) {  in SumElementStrided3d()
163       const CeedInt node = data.t_id_x + data.t_id_y * P_1D + z * P_1D * P_1D;  in SumElementStrided3d()
164       const CeedInt ind  = node * strides_node + elem * strides_elem;  in SumElementStrided3d()
166       for (CeedInt comp = 0; comp < NUM_COMP; comp++) {  in SumElementStrided3d()
181 …_device__ void ReadPoint(SharedData_Cuda &data, const CeedInt elem, const CeedInt p, const CeedInt…  in ReadPoint()
182 …               const CeedInt strides_point, const CeedInt strides_comp, const CeedInt strides_elem,  in ReadPoint()
184   const CeedInt ind = (p % NUM_PTS) * strides_point + elem * strides_elem;  in ReadPoint()
187     for (CeedInt comp = 0; comp < NUM_COMP; comp++) {  in ReadPoint()
191     for (CeedInt comp = 0; comp < NUM_COMP; comp++) {  in ReadPoint()
201 …device__ void WritePoint(SharedData_Cuda &data, const CeedInt elem, const CeedInt p, const CeedInt…  in WritePoint()
202 …const CeedInt strides_point, const CeedInt strides_comp, const CeedInt strides_elem, const CeedSca…  in WritePoint()
205     const CeedInt ind = (p % NUM_PTS) * strides_point + elem * strides_elem;  in WritePoint()
207     for (CeedInt comp = 0; comp < NUM_COMP; comp++) {  in WritePoint()