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 …__device__ void ReadPoint(SharedData_Hip &data, const CeedInt elem, const CeedInt p, const CeedInt…  in ReadPoint()
29 …        const CeedInt *__restrict__ indices, const CeedScalar *__restrict__ d_u, CeedScalar *r_u) {  in ReadPoint()
30   const CeedInt ind = indices[p + elem * NUM_PTS];  in ReadPoint()
32   for (CeedInt comp = 0; comp < NUM_COMP; comp++) {  in ReadPoint()
41 …_device__ void WritePoint(SharedData_Hip &data, const CeedInt elem, const CeedInt p, const CeedInt…  in WritePoint()
42 …        const CeedInt *__restrict__ indices, const CeedScalar *__restrict__ r_u, CeedScalar *d_u) {  in WritePoint()
44     const CeedInt ind = indices[p + elem * NUM_PTS];  in WritePoint()
46     for (CeedInt comp = 0; comp < NUM_COMP; comp++) {  in WritePoint()
60 inline __device__ void SetEVecStandard1d_Single(SharedData_Hip &data, const CeedInt n, const CeedSc…  in SetEVecStandard1d_Single()
61   const CeedInt target_comp = n / P_1D;  in SetEVecStandard1d_Single()
62   const CeedInt target_node = n % P_1D;  in SetEVecStandard1d_Single()
73 …id ReadLVecStandard1d(SharedData_Hip &data, const CeedInt num_nodes, const CeedInt elem, const Cee…  in ReadLVecStandard1d()
76     const CeedInt node = data.t_id_x;  in ReadLVecStandard1d()
77     const CeedInt ind  = indices[node + elem * P_1D];  in ReadLVecStandard1d()
79     for (CeedInt comp = 0; comp < NUM_COMP; comp++) r_u[comp] = d_u[ind + COMP_STRIDE * comp];  in ReadLVecStandard1d()
87 inline __device__ void ReadLVecStrided1d(SharedData_Hip &data, const CeedInt elem, const CeedScalar…  in ReadLVecStrided1d()
89     const CeedInt node = data.t_id_x;  in ReadLVecStrided1d()
90     const CeedInt ind  = node * STRIDES_NODE + elem * STRIDES_ELEM;  in ReadLVecStrided1d()
92     for (CeedInt comp = 0; comp < NUM_COMP; comp++) r_u[comp] = d_u[ind + comp * STRIDES_COMP];  in ReadLVecStrided1d()
100 …d WriteLVecStandard1d(SharedData_Hip &data, const CeedInt num_nodes, const CeedInt elem, const Cee…  in WriteLVecStandard1d()
103     const CeedInt node = data.t_id_x;  in WriteLVecStandard1d()
104     const CeedInt ind  = indices[node + elem * P_1D];  in WriteLVecStandard1d()
106 …for (CeedInt comp = 0; comp < NUM_COMP; comp++) atomicAdd(&d_v[ind + COMP_STRIDE * comp], r_v[comp…  in WriteLVecStandard1d()
111 …LVecStandard1d_Single(SharedData_Hip &data, const CeedInt num_nodes, const CeedInt elem, const Cee…  in WriteLVecStandard1d_Single()
112 …                           const CeedInt *__restrict__ indices, const CeedScalar *__restrict__ r_v,  in WriteLVecStandard1d_Single()
114   const CeedInt target_comp = n / P_1D;  in WriteLVecStandard1d_Single()
115   const CeedInt target_node = n % P_1D;  in WriteLVecStandard1d_Single()
118     const CeedInt ind = indices[target_node + elem * P_1D];  in WriteLVecStandard1d_Single()
128 …ecStandard1d_Assembly(SharedData_Hip &data, const CeedInt num_nodes, const CeedInt elem, const Cee…  in WriteLVecStandard1d_Assembly()
130   const CeedInt in_comp    = in / P_1D;  in WriteLVecStandard1d_Assembly()
131   const CeedInt in_node    = in % P_1D;  in WriteLVecStandard1d_Assembly()
132   const CeedInt e_vec_size = P_1D * NUM_COMP;  in WriteLVecStandard1d_Assembly()
135     const CeedInt out_node = data.t_id_x;  in WriteLVecStandard1d_Assembly()
137     for (CeedInt comp = 0; comp < NUM_COMP; comp++) {  in WriteLVecStandard1d_Assembly()
147 …cStandard1d_QFAssembly(SharedData_Hip &data, const CeedInt num_elem, const CeedInt elem, const Cee…  in WriteLVecStandard1d_QFAssembly()
148 …   const CeedInt output_offset, const CeedScalar *__restrict__ r_v, CeedScalar *__restrict__ d_v) {  in WriteLVecStandard1d_QFAssembly()
150     const CeedInt ind = data.t_id_x + elem * Q_1D;  in WriteLVecStandard1d_QFAssembly()
152     for (CeedInt comp = 0; comp < NUM_COMP_FIELD; comp++) {  in WriteLVecStandard1d_QFAssembly()
162 inline __device__ void WriteLVecStrided1d(SharedData_Hip &data, const CeedInt elem, const CeedScala…  in WriteLVecStrided1d()
165     const CeedInt node = data.t_id_x;  in WriteLVecStrided1d()
166     const CeedInt ind  = node * STRIDES_NODE + elem * STRIDES_ELEM;  in WriteLVecStrided1d()
168     for (CeedInt comp = 0; comp < NUM_COMP; comp++) d_v[ind + comp * STRIDES_COMP] += r_v[comp];  in WriteLVecStrided1d()
180 inline __device__ void SetEVecStandard2d_Single(SharedData_Hip &data, const CeedInt n, const CeedSc…  in SetEVecStandard2d_Single()
181   const CeedInt target_comp   = n / (P_1D * P_1D);  in SetEVecStandard2d_Single()
182   const CeedInt target_node_x = n % P_1D;  in SetEVecStandard2d_Single()
183   const CeedInt target_node_y = (n % (P_1D * P_1D)) / P_1D;  in SetEVecStandard2d_Single()
194 …id ReadLVecStandard2d(SharedData_Hip &data, const CeedInt num_nodes, const CeedInt elem, const Cee…  in ReadLVecStandard2d()
197     const CeedInt node = data.t_id_x + data.t_id_y * P_1D;  in ReadLVecStandard2d()
198     const CeedInt ind  = indices[node + elem * P_1D * P_1D];  in ReadLVecStandard2d()
200     for (CeedInt comp = 0; comp < NUM_COMP; comp++) r_u[comp] = d_u[ind + COMP_STRIDE * comp];  in ReadLVecStandard2d()
208 inline __device__ void ReadLVecStrided2d(SharedData_Hip &data, const CeedInt elem, const CeedScalar…  in ReadLVecStrided2d()
210     const CeedInt node = data.t_id_x + data.t_id_y * P_1D;  in ReadLVecStrided2d()
211     const CeedInt ind  = node * STRIDES_NODE + elem * STRIDES_ELEM;  in ReadLVecStrided2d()
213     for (CeedInt comp = 0; comp < NUM_COMP; comp++) r_u[comp] = d_u[ind + comp * STRIDES_COMP];  in ReadLVecStrided2d()
221 …d WriteLVecStandard2d(SharedData_Hip &data, const CeedInt num_nodes, const CeedInt elem, const Cee…  in WriteLVecStandard2d()
224     const CeedInt node = data.t_id_x + data.t_id_y * P_1D;  in WriteLVecStandard2d()
225     const CeedInt ind  = indices[node + elem * P_1D * P_1D];  in WriteLVecStandard2d()
227 …for (CeedInt comp = 0; comp < NUM_COMP; comp++) atomicAdd(&d_v[ind + COMP_STRIDE * comp], r_v[comp…  in WriteLVecStandard2d()
232 …LVecStandard2d_Single(SharedData_Hip &data, const CeedInt num_nodes, const CeedInt elem, const Cee…  in WriteLVecStandard2d_Single()
233 …                           const CeedInt *__restrict__ indices, const CeedScalar *__restrict__ r_v,  in WriteLVecStandard2d_Single()
235   const CeedInt target_comp   = n / (P_1D * P_1D);  in WriteLVecStandard2d_Single()
236   const CeedInt target_node_x = n % P_1D;  in WriteLVecStandard2d_Single()
237   const CeedInt target_node_y = (n % (P_1D * P_1D)) / P_1D;  in WriteLVecStandard2d_Single()
240     const CeedInt node = data.t_id_x + data.t_id_y * P_1D;  in WriteLVecStandard2d_Single()
241     const CeedInt ind  = indices[node + elem * P_1D * P_1D];  in WriteLVecStandard2d_Single()
251 …ecStandard2d_Assembly(SharedData_Hip &data, const CeedInt num_nodes, const CeedInt elem, const Cee…  in WriteLVecStandard2d_Assembly()
253   const CeedInt elem_size  = P_1D * P_1D;  in WriteLVecStandard2d_Assembly()
254   const CeedInt in_comp    = in / elem_size;  in WriteLVecStandard2d_Assembly()
255   const CeedInt in_node_x  = in % P_1D;  in WriteLVecStandard2d_Assembly()
256   const CeedInt in_node_y  = (in % elem_size) / P_1D;  in WriteLVecStandard2d_Assembly()
257   const CeedInt e_vec_size = elem_size * NUM_COMP;  in WriteLVecStandard2d_Assembly()
260     const CeedInt in_node  = in_node_x + in_node_y * P_1D;  in WriteLVecStandard2d_Assembly()
261     const CeedInt out_node = data.t_id_x + data.t_id_y * P_1D;  in WriteLVecStandard2d_Assembly()
263     for (CeedInt comp = 0; comp < NUM_COMP; comp++) {  in WriteLVecStandard2d_Assembly()
264 …const CeedInt index = (in_comp * NUM_COMP + comp) * elem_size * elem_size + out_node * elem_size +…  in WriteLVecStandard2d_Assembly()
275 …cStandard2d_QFAssembly(SharedData_Hip &data, const CeedInt num_elem, const CeedInt elem, const Cee…  in WriteLVecStandard2d_QFAssembly()
276 …   const CeedInt output_offset, const CeedScalar *__restrict__ r_v, CeedScalar *__restrict__ d_v) {  in WriteLVecStandard2d_QFAssembly()
278     const CeedInt ind = (data.t_id_x + data.t_id_y * Q_1D) + elem * Q_1D * Q_1D;  in WriteLVecStandard2d_QFAssembly()
280     for (CeedInt comp = 0; comp < NUM_COMP_FIELD; comp++) {  in WriteLVecStandard2d_QFAssembly()
290 inline __device__ void WriteLVecStrided2d(SharedData_Hip &data, const CeedInt elem, const CeedScala…  in WriteLVecStrided2d()
293     const CeedInt node = data.t_id_x + data.t_id_y * P_1D;  in WriteLVecStrided2d()
294     const CeedInt ind  = node * STRIDES_NODE + elem * STRIDES_ELEM;  in WriteLVecStrided2d()
296     for (CeedInt comp = 0; comp < NUM_COMP; comp++) d_v[ind + comp * STRIDES_COMP] += r_v[comp];  in WriteLVecStrided2d()
308 inline __device__ void SetEVecStandard3d_Single(SharedData_Hip &data, const CeedInt n, const CeedSc…  in SetEVecStandard3d_Single()
309   const CeedInt target_comp   = n / (P_1D * P_1D * P_1D);  in SetEVecStandard3d_Single()
310   const CeedInt target_node_x = n % P_1D;  in SetEVecStandard3d_Single()
311   const CeedInt target_node_y = ((n % (P_1D * P_1D * P_1D)) / P_1D) % P_1D;  in SetEVecStandard3d_Single()
312   const CeedInt target_node_z = (n % (P_1D * P_1D * P_1D)) / (P_1D * P_1D);  in SetEVecStandard3d_Single()
323 …id ReadLVecStandard3d(SharedData_Hip &data, const CeedInt num_nodes, const CeedInt elem, const Cee…  in ReadLVecStandard3d()
326     for (CeedInt z = 0; z < P_1D; z++) {  in ReadLVecStandard3d()
327       const CeedInt node = data.t_id_x + data.t_id_y * P_1D + z * P_1D * P_1D;  in ReadLVecStandard3d()
328       const CeedInt ind  = indices[node + elem * P_1D * P_1D * P_1D];  in ReadLVecStandard3d()
330 …for (CeedInt comp = 0; comp < NUM_COMP; comp++) r_u[z + comp * P_1D] = d_u[ind + COMP_STRIDE * com…  in ReadLVecStandard3d()
339 inline __device__ void ReadLVecStrided3d(SharedData_Hip &data, const CeedInt elem, const CeedScalar…  in ReadLVecStrided3d()
341     for (CeedInt z = 0; z < P_1D; z++) {  in ReadLVecStrided3d()
342       const CeedInt node = data.t_id_x + data.t_id_y * P_1D + z * P_1D * P_1D;  in ReadLVecStrided3d()
343       const CeedInt ind  = node * STRIDES_NODE + elem * STRIDES_ELEM;  in ReadLVecStrided3d()
345 …for (CeedInt comp = 0; comp < NUM_COMP; comp++) r_u[z + comp * P_1D] = d_u[ind + comp * STRIDES_CO…  in ReadLVecStrided3d()
354 …ReadEVecSliceStandard3d(SharedData_Hip &data, const CeedInt nquads, const CeedInt elem, const Ceed…  in ReadEVecSliceStandard3d()
355 …                           const CeedInt *__restrict__ indices, const CeedScalar *__restrict__ d_u,  in ReadEVecSliceStandard3d()
358     const CeedInt node = data.t_id_x + data.t_id_y * Q_1D + q * Q_1D * Q_1D;  in ReadEVecSliceStandard3d()
359     const CeedInt ind  = indices[node + elem * Q_1D * Q_1D * Q_1D];  in ReadEVecSliceStandard3d()
361     for (CeedInt comp = 0; comp < NUM_COMP; comp++) r_u[comp] = d_u[ind + COMP_STRIDE * comp];  in ReadEVecSliceStandard3d()
369 …vice__ void ReadEVecSliceStrided3d(SharedData_Hip &data, const CeedInt elem, const CeedInt q, cons…  in ReadEVecSliceStrided3d()
372     const CeedInt node = data.t_id_x + data.t_id_y * Q_1D + q * Q_1D * Q_1D;  in ReadEVecSliceStrided3d()
373     const CeedInt ind  = node * STRIDES_NODE + elem * STRIDES_ELEM;  in ReadEVecSliceStrided3d()
375     for (CeedInt comp = 0; comp < NUM_COMP; comp++) r_u[comp] = d_u[ind + comp * STRIDES_COMP];  in ReadEVecSliceStrided3d()
383 …d WriteLVecStandard3d(SharedData_Hip &data, const CeedInt num_nodes, const CeedInt elem, const Cee…  in WriteLVecStandard3d()
386     for (CeedInt z = 0; z < P_1D; z++) {  in WriteLVecStandard3d()
387       const CeedInt node = data.t_id_x + data.t_id_y * P_1D + z * P_1D * P_1D;  in WriteLVecStandard3d()
388       const CeedInt ind  = indices[node + elem * P_1D * P_1D * P_1D];  in WriteLVecStandard3d()
390 …for (CeedInt comp = 0; comp < NUM_COMP; comp++) atomicAdd(&d_v[ind + COMP_STRIDE * comp], r_v[z + …  in WriteLVecStandard3d()
396 …LVecStandard3d_Single(SharedData_Hip &data, const CeedInt num_nodes, const CeedInt elem, const Cee…  in WriteLVecStandard3d_Single()
397 …                           const CeedInt *__restrict__ indices, const CeedScalar *__restrict__ r_v,  in WriteLVecStandard3d_Single()
399   const CeedInt target_comp   = n / (P_1D * P_1D * P_1D);  in WriteLVecStandard3d_Single()
400   const CeedInt target_node_x = n % P_1D;  in WriteLVecStandard3d_Single()
401   const CeedInt target_node_y = ((n % (P_1D * P_1D * P_1D)) / P_1D) % P_1D;  in WriteLVecStandard3d_Single()
402   const CeedInt target_node_z = (n % (P_1D * P_1D * P_1D)) / (P_1D * P_1D);  in WriteLVecStandard3d_Single()
405     const CeedInt node = data.t_id_x + data.t_id_y * P_1D + target_node_z * P_1D * P_1D;  in WriteLVecStandard3d_Single()
406     const CeedInt ind  = indices[node + elem * P_1D * P_1D * P_1D];  in WriteLVecStandard3d_Single()
416 …ecStandard3d_Assembly(SharedData_Hip &data, const CeedInt num_nodes, const CeedInt elem, const Cee…  in WriteLVecStandard3d_Assembly()
418   const CeedInt elem_size  = P_1D * P_1D * P_1D;  in WriteLVecStandard3d_Assembly()
419   const CeedInt in_comp    = in / elem_size;  in WriteLVecStandard3d_Assembly()
420   const CeedInt in_node_x  = in % P_1D;  in WriteLVecStandard3d_Assembly()
421   const CeedInt in_node_y  = (in % (P_1D * P_1D)) / P_1D;  in WriteLVecStandard3d_Assembly()
422   const CeedInt in_node_z  = (in % elem_size) / (P_1D * P_1D);  in WriteLVecStandard3d_Assembly()
423   const CeedInt e_vec_size = elem_size * NUM_COMP;  in WriteLVecStandard3d_Assembly()
426     const CeedInt in_node = in_node_x + in_node_y * P_1D + in_node_z * P_1D * P_1D;  in WriteLVecStandard3d_Assembly()
427     for (CeedInt z = 0; z < P_1D; z++) {  in WriteLVecStandard3d_Assembly()
428       const CeedInt out_node = data.t_id_x + data.t_id_y * P_1D + z * P_1D * P_1D;  in WriteLVecStandard3d_Assembly()
430       for (CeedInt comp = 0; comp < NUM_COMP; comp++) {  in WriteLVecStandard3d_Assembly()
431 …const CeedInt index = (in_comp * NUM_COMP + comp) * elem_size * elem_size + out_node * elem_size +…  in WriteLVecStandard3d_Assembly()
443 …cStandard3d_QFAssembly(SharedData_Hip &data, const CeedInt num_elem, const CeedInt elem, const Cee…  in WriteLVecStandard3d_QFAssembly()
444 …   const CeedInt output_offset, const CeedScalar *__restrict__ r_v, CeedScalar *__restrict__ d_v) {  in WriteLVecStandard3d_QFAssembly()
446     for (CeedInt z = 0; z < Q_1D; z++) {  in WriteLVecStandard3d_QFAssembly()
447 …const CeedInt ind = (data.t_id_x + data.t_id_y * Q_1D + z * Q_1D * Q_1D) + elem * Q_1D * Q_1D * Q_…  in WriteLVecStandard3d_QFAssembly()
449       for (CeedInt comp = 0; comp < NUM_COMP_FIELD; comp++) {  in WriteLVecStandard3d_QFAssembly()
460 inline __device__ void WriteLVecStrided3d(SharedData_Hip &data, const CeedInt elem, const CeedScala…  in WriteLVecStrided3d()
463     for (CeedInt z = 0; z < P_1D; z++) {  in WriteLVecStrided3d()
464       const CeedInt node = data.t_id_x + data.t_id_y * P_1D + z * P_1D * P_1D;  in WriteLVecStrided3d()
465       const CeedInt ind  = node * STRIDES_NODE + elem * STRIDES_ELEM;  in WriteLVecStrided3d()
467 …for (CeedInt comp = 0; comp < NUM_COMP; comp++) d_v[ind + comp * STRIDES_COMP] += r_v[z + comp * P…  in WriteLVecStrided3d()
476 inline __device__ void GradColloSlice3d(SharedData_Hip &data, const CeedInt q, const CeedScalar *__…  in GradColloSlice3d()
479     for (CeedInt comp = 0; comp < NUM_COMP; comp++) {  in GradColloSlice3d()
485       for (CeedInt i = 0; i < Q_1D; i++) {  in GradColloSlice3d()
490       for (CeedInt i = 0; i < Q_1D; i++) {  in GradColloSlice3d()
495       for (CeedInt i = 0; i < Q_1D; i++) {  in GradColloSlice3d()
506 inline __device__ void GradColloSliceTranspose3d(SharedData_Hip &data, const CeedInt q, const CeedS…  in GradColloSliceTranspose3d()
509     for (CeedInt comp = 0; comp < NUM_COMP; comp++) {  in GradColloSliceTranspose3d()
514       for (CeedInt i = 0; i < Q_1D; i++) {  in GradColloSliceTranspose3d()
521       for (CeedInt i = 0; i < Q_1D; i++) {  in GradColloSliceTranspose3d()
525       for (CeedInt i = 0; i < Q_1D; i++) {  in GradColloSliceTranspose3d()