15 inline void loadMatrix(const CeedInt N, const CeedScalar *restrict d_B, CeedScalar *restrict B) {  in loadMatrix()
16   const CeedInt item_id    = get_local_linear_id();  in loadMatrix()
17   const CeedInt group_size = get_local_size(0) * get_local_size(1) * get_local_size(2);  in loadMatrix()
18   for (CeedInt i = item_id; i < N; i += group_size) B[i] = d_B[i];  in loadMatrix()
28 inline void ReadElementStrided1d(const CeedInt NUM_COMP, const CeedInt P_1D, const CeedInt num_elem…  in ReadElementStrided1d()
29 …     const CeedInt strides_comp, const CeedInt strides_elem, global const CeedScalar *restrict d_u,  in ReadElementStrided1d()
31   const CeedInt item_id_x = get_local_id(0);  in ReadElementStrided1d()
32   const CeedInt elem      = get_global_id(2);  in ReadElementStrided1d()
35     const CeedInt node = item_id_x;  in ReadElementStrided1d()
36     const CeedInt ind  = node * strides_node + elem * strides_elem;  in ReadElementStrided1d()
37     for (CeedInt comp = 0; comp < NUM_COMP; comp++) {  in ReadElementStrided1d()
46 …nline void WriteElementStrided1d(const CeedInt NUM_COMP, const CeedInt P_1D, const CeedInt num_ele…  in WriteElementStrided1d()
47 …    const CeedInt strides_comp, const CeedInt strides_elem, private const CeedScalar *restrict r_v,  in WriteElementStrided1d()
49   const CeedInt item_id_x = get_local_id(0);  in WriteElementStrided1d()
50   const CeedInt elem      = get_global_id(2);  in WriteElementStrided1d()
53     const CeedInt node = item_id_x;  in WriteElementStrided1d()
54     const CeedInt ind  = node * strides_node + elem * strides_elem;  in WriteElementStrided1d()
55     for (CeedInt comp = 0; comp < NUM_COMP; comp++) {  in WriteElementStrided1d()
68 inline void ReadElementStrided2d(const CeedInt NUM_COMP, const CeedInt P_1D, const CeedInt num_elem…  in ReadElementStrided2d()
69 …     const CeedInt strides_comp, const CeedInt strides_elem, global const CeedScalar *restrict d_u,  in ReadElementStrided2d()
71   const CeedInt item_id_x = get_local_id(0);  in ReadElementStrided2d()
72   const CeedInt item_id_y = get_local_id(1);  in ReadElementStrided2d()
73   const CeedInt elem      = get_global_id(2);  in ReadElementStrided2d()
76     const CeedInt node = item_id_x + item_id_y * P_1D;  in ReadElementStrided2d()
77     const CeedInt ind  = node * strides_node + elem * strides_elem;  in ReadElementStrided2d()
78     for (CeedInt comp = 0; comp < NUM_COMP; comp++) {  in ReadElementStrided2d()
87 …nline void WriteElementStrided2d(const CeedInt NUM_COMP, const CeedInt P_1D, const CeedInt num_ele…  in WriteElementStrided2d()
88 …    const CeedInt strides_comp, const CeedInt strides_elem, private const CeedScalar *restrict r_v,  in WriteElementStrided2d()
90   const CeedInt item_id_x = get_local_id(0);  in WriteElementStrided2d()
91   const CeedInt item_id_y = get_local_id(1);  in WriteElementStrided2d()
92   const CeedInt elem      = get_global_id(2);  in WriteElementStrided2d()
95     const CeedInt node = item_id_x + item_id_y * P_1D;  in WriteElementStrided2d()
96     const CeedInt ind  = node * strides_node + elem * strides_elem;  in WriteElementStrided2d()
97     for (CeedInt comp = 0; comp < NUM_COMP; comp++) {  in WriteElementStrided2d()
110 inline void ReadElementStrided3d(const CeedInt NUM_COMP, const CeedInt P_1D, const CeedInt num_elem…  in ReadElementStrided3d()
111 …     const CeedInt strides_comp, const CeedInt strides_elem, global const CeedScalar *restrict d_u,  in ReadElementStrided3d()
113   const CeedInt item_id_x = get_local_id(0);  in ReadElementStrided3d()
114   const CeedInt item_id_y = get_local_id(1);  in ReadElementStrided3d()
115   const CeedInt elem      = get_global_id(2);  in ReadElementStrided3d()
118     for (CeedInt z = 0; z < P_1D; z++) {  in ReadElementStrided3d()
119       const CeedInt node = item_id_x + item_id_y * P_1D + z * P_1D * P_1D;  in ReadElementStrided3d()
120       const CeedInt ind  = node * strides_node + elem * strides_elem;  in ReadElementStrided3d()
121       for (CeedInt comp = 0; comp < NUM_COMP; comp++) {  in ReadElementStrided3d()
131 …nline void WriteElementStrided3d(const CeedInt NUM_COMP, const CeedInt P_1D, const CeedInt num_ele…  in WriteElementStrided3d()
132 …    const CeedInt strides_comp, const CeedInt strides_elem, private const CeedScalar *restrict r_v,  in WriteElementStrided3d()
134   const CeedInt item_id_x = get_local_id(0);  in WriteElementStrided3d()
135   const CeedInt item_id_y = get_local_id(1);  in WriteElementStrided3d()
136   const CeedInt elem      = get_global_id(2);  in WriteElementStrided3d()
139     for (CeedInt z = 0; z < P_1D; z++) {  in WriteElementStrided3d()
140       const CeedInt node = item_id_x + item_id_y * P_1D + z * P_1D * P_1D;  in WriteElementStrided3d()
141       const CeedInt ind  = node * strides_node + elem * strides_elem;  in WriteElementStrided3d()
142       for (CeedInt comp = 0; comp < NUM_COMP; comp++) {  in WriteElementStrided3d()