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 …ReadElementStrided1d(SharedData_Hip &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 …riteElementStrided1d(SharedData_Hip &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_Hip &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 …ReadElementStrided2d(SharedData_Hip &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 …riteElementStrided2d(SharedData_Hip &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_Hip &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 …ReadElementStrided3d(SharedData_Hip &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 …riteElementStrided3d(SharedData_Hip &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_Hip &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 … ReadPoint(SharedData_Hip &data, const CeedInt elem, const CeedInt p, const CeedInt points_in_elem…  in ReadPoint()
182 …const CeedInt strides_comp, const CeedInt strides_elem, const CeedScalar *__restrict__ d_u, CeedSc…  in ReadPoint()
183   const CeedInt ind = (p % NUM_PTS) * strides_point + elem * strides_elem;  in ReadPoint()
186     for (CeedInt comp = 0; comp < NUM_COMP; comp++) {  in ReadPoint()
190     for (CeedInt comp = 0; comp < NUM_COMP; comp++) {  in ReadPoint()
200 …_device__ void WritePoint(SharedData_Hip &data, const CeedInt elem, const CeedInt p, const CeedInt…  in WritePoint()
201 …const CeedInt strides_point, const CeedInt strides_comp, const CeedInt strides_elem, const CeedSca…  in WritePoint()
204     const CeedInt ind = (p % NUM_PTS) * strides_point + elem * strides_elem;  in WritePoint()
206     for (CeedInt comp = 0; comp < NUM_COMP; comp++) {  in WritePoint()