Lines Matching refs:t_id_x
20 inline __device__ void ContractX2dFlattened(SharedData_Hip &data, const int t_id_x, const int t_id_… in ContractX2dFlattened() argument
23 data.slice[t_id_x + t_id_y * T_1D] = *U; in ContractX2dFlattened()
26 if (t_id_x < Q_1D && t_id_y < P_1D) { in ContractX2dFlattened()
28 *V += B[i + t_id_x * P_1D] * data.slice[i + t_id_y * T_1D]; // Contract x direction in ContractX2dFlattened()
37 inline __device__ void ContractY2dFlattened(SharedData_Hip &data, const int t_id_x, const int t_id_… in ContractY2dFlattened() argument
40 data.slice[t_id_x + t_id_y * T_1D] = *U; in ContractY2dFlattened()
43 if (t_id_x < Q_1D && t_id_y < Q_1D) { in ContractY2dFlattened()
45 *V += B[i + t_id_y * P_1D] * data.slice[t_id_x + i * T_1D]; // Contract y direction in ContractY2dFlattened()
54 inline __device__ void ContractTransposeY2dFlattened(SharedData_Hip &data, const int t_id_x, const … in ContractTransposeY2dFlattened() argument
57 data.slice[t_id_x + t_id_y * T_1D] = *U; in ContractTransposeY2dFlattened()
60 if (t_id_x < Q_1D && t_id_y < P_1D) { in ContractTransposeY2dFlattened()
62 *V += B[t_id_y + i * P_1D] * data.slice[t_id_x + i * T_1D]; // Contract y direction in ContractTransposeY2dFlattened()
71 inline __device__ void ContractTransposeX2dFlattened(SharedData_Hip &data, const int t_id_x, const … in ContractTransposeX2dFlattened() argument
74 data.slice[t_id_x + t_id_y * T_1D] = *U; in ContractTransposeX2dFlattened()
77 if (t_id_x < P_1D && t_id_y < P_1D) { in ContractTransposeX2dFlattened()
79 *V += B[t_id_x + i * P_1D] * data.slice[i + t_id_y * T_1D]; // Contract x direction in ContractTransposeX2dFlattened()
88 inline __device__ void ContractTransposeAddX2dFlattened(SharedData_Hip &data, const int t_id_x, con… in ContractTransposeAddX2dFlattened() argument
91 data.slice[t_id_x + t_id_y * T_1D] = *U; in ContractTransposeAddX2dFlattened()
93 if (t_id_x < P_1D && t_id_y < P_1D) { in ContractTransposeAddX2dFlattened()
95 *V += B[t_id_x + i * P_1D] * data.slice[i + t_id_y * T_1D]; // Contract x direction in ContractTransposeAddX2dFlattened()
104 inline __device__ void QPack2d(SharedData_Hip &data, const int t_id_x, const int t_id_y, CeedScalar… in QPack2d() argument
105 const CeedInt new_t_id_x = data.t_id_x % Q_1D, new_t_id_y = data.t_id_x / Q_1D; in QPack2d()
109 if (t_id_x < Q_1D && t_id_y < Q_1D) data.slice[t_id_x + t_id_y * T_1D] = U[comp]; in QPack2d()
111 U[comp] = data.t_id_x < (Q_1D * Q_1D) ? data.slice[new_t_id_x + new_t_id_y * T_1D] : 0.0; in QPack2d()
116 inline __device__ void QUnpack2d(SharedData_Hip &data, const int t_id_x, const int t_id_y, CeedScal… in QUnpack2d() argument
117 const CeedInt old_t_id_x = data.t_id_x % Q_1D, old_t_id_y = data.t_id_x / Q_1D; in QUnpack2d()
121 if (data.t_id_x < (Q_1D * Q_1D)) data.slice[old_t_id_x + old_t_id_y * T_1D] = U[comp]; in QUnpack2d()
123 U[comp] = (t_id_x < Q_1D && t_id_y < Q_1D) ? data.slice[t_id_x + t_id_y * T_1D] : 0.0; in QUnpack2d()
133 const int t_id_x = data.t_id_x % T_1D, t_id_y = data.t_id_x / T_1D; in InterpTensor2dFlattened() local
136 if (P_1D != T_1D) QUnpack2d<NUM_COMP, P_1D, T_1D>(data, t_id_x, t_id_y, r_U); in InterpTensor2dFlattened()
138 ContractX2dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, &r_U[comp], c_B, r_t); in InterpTensor2dFlattened()
139 ContractY2dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, r_t, c_B, &r_V[comp]); in InterpTensor2dFlattened()
142 if (P_1D != T_1D) QPack2d<NUM_COMP, P_1D, T_1D>(data, t_id_x, t_id_y, r_U); in InterpTensor2dFlattened()
143 if (Q_1D != T_1D) QPack2d<NUM_COMP, Q_1D, T_1D>(data, t_id_x, t_id_y, r_V); in InterpTensor2dFlattened()
152 const int t_id_x = data.t_id_x % T_1D, t_id_y = data.t_id_x / T_1D; in InterpTransposeTensor2dFlattened() local
155 if (Q_1D != T_1D) QUnpack2d<NUM_COMP, Q_1D, T_1D>(data, t_id_x, t_id_y, r_U); in InterpTransposeTensor2dFlattened()
157 …ContractTransposeY2dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, &r_U[comp], c_B, r… in InterpTransposeTensor2dFlattened()
158 …ContractTransposeX2dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, r_t, c_B, &r_V[com… in InterpTransposeTensor2dFlattened()
161 if (P_1D != T_1D) QPack2d<NUM_COMP, P_1D, T_1D>(data, t_id_x, t_id_y, r_V); in InterpTransposeTensor2dFlattened()
170 const int t_id_x = data.t_id_x % T_1D, t_id_y = data.t_id_x / T_1D; in InterpTensorCollocatedNodes2dFlattened() local
172 if (P_1D != T_1D) QUnpack2d<NUM_COMP, P_1D, T_1D>(data, t_id_x, t_id_y, r_U); in InterpTensorCollocatedNodes2dFlattened()
177 if (P_1D != T_1D) QPack2d<NUM_COMP, P_1D, T_1D>(data, t_id_x, t_id_y, r_U); in InterpTensorCollocatedNodes2dFlattened()
178 if (Q_1D != T_1D) QPack2d<NUM_COMP, Q_1D, T_1D>(data, t_id_x, t_id_y, r_V); in InterpTensorCollocatedNodes2dFlattened()
187 const int t_id_x = data.t_id_x % T_1D, t_id_y = data.t_id_x / T_1D; in InterpTransposeTensorCollocatedNodes2dFlattened() local
189 if (Q_1D != T_1D) QUnpack2d<NUM_COMP, Q_1D, T_1D>(data, t_id_x, t_id_y, r_U); in InterpTransposeTensorCollocatedNodes2dFlattened()
194 if (P_1D != T_1D) QPack2d<NUM_COMP, P_1D, T_1D>(data, t_id_x, t_id_y, r_V); in InterpTransposeTensorCollocatedNodes2dFlattened()
203 const int t_id_x = data.t_id_x % T_1D, t_id_y = data.t_id_x / T_1D; in GradTensor2dFlattened() local
206 if (P_1D != T_1D) QUnpack2d<NUM_COMP, P_1D, T_1D>(data, t_id_x, t_id_y, r_U); in GradTensor2dFlattened()
208 ContractX2dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, &r_U[comp], c_G, r_t); in GradTensor2dFlattened()
209 …ContractY2dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, r_t, c_B, &r_V[comp + 0 * N… in GradTensor2dFlattened()
210 ContractX2dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, &r_U[comp], c_B, r_t); in GradTensor2dFlattened()
211 …ContractY2dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, r_t, c_G, &r_V[comp + 1 * N… in GradTensor2dFlattened()
214 if (P_1D != T_1D) QPack2d<NUM_COMP, P_1D, T_1D>(data, t_id_x, t_id_y, r_U); in GradTensor2dFlattened()
215 if (Q_1D != T_1D) QPack2d<NUM_COMP * 2, Q_1D, T_1D>(data, t_id_x, t_id_y, r_V); in GradTensor2dFlattened()
224 const int t_id_x = data.t_id_x % T_1D, t_id_y = data.t_id_x / T_1D; in GradTransposeTensor2dFlattened() local
227 if (Q_1D != T_1D) QUnpack2d<NUM_COMP * 2, Q_1D, T_1D>(data, t_id_x, t_id_y, r_U); in GradTransposeTensor2dFlattened()
229 …ContractTransposeY2dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, &r_U[comp + 0 * NU… in GradTransposeTensor2dFlattened()
230 …ContractTransposeX2dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, r_t, c_G, &r_V[com… in GradTransposeTensor2dFlattened()
231 …ContractTransposeY2dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, &r_U[comp + 1 * NU… in GradTransposeTensor2dFlattened()
232 …ContractTransposeAddX2dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, r_t, c_B, &r_V[… in GradTransposeTensor2dFlattened()
235 if (P_1D != T_1D) QPack2d<NUM_COMP, P_1D, T_1D>(data, t_id_x, t_id_y, r_V); in GradTransposeTensor2dFlattened()
244 const int t_id_x = data.t_id_x % T_1D, t_id_y = data.t_id_x / T_1D; in GradTensorCollocatedNodes2dFlattened() local
246 if (P_1D != T_1D) QUnpack2d<NUM_COMP, P_1D, T_1D>(data, t_id_x, t_id_y, r_U); in GradTensorCollocatedNodes2dFlattened()
248 …ContractX2dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, &r_U[comp], c_G, &r_V[comp … in GradTensorCollocatedNodes2dFlattened()
249 …ContractY2dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, &r_U[comp], c_G, &r_V[comp … in GradTensorCollocatedNodes2dFlattened()
252 if (P_1D != T_1D) QPack2d<NUM_COMP, P_1D, T_1D>(data, t_id_x, t_id_y, r_U); in GradTensorCollocatedNodes2dFlattened()
253 if (Q_1D != T_1D) QPack2d<NUM_COMP * 2, Q_1D, T_1D>(data, t_id_x, t_id_y, r_V); in GradTensorCollocatedNodes2dFlattened()
262 const int t_id_x = data.t_id_x % T_1D, t_id_y = data.t_id_x / T_1D; in GradTransposeTensorCollocatedNodes2dFlattened() local
264 if (Q_1D != T_1D) QUnpack2d<NUM_COMP * 2, Q_1D, T_1D>(data, t_id_x, t_id_y, r_U); in GradTransposeTensorCollocatedNodes2dFlattened()
266 …ContractTransposeY2dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, &r_U[comp + 1 * NU… in GradTransposeTensorCollocatedNodes2dFlattened()
267 …ContractTransposeAddX2dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, &r_U[comp + 0 *… in GradTransposeTensorCollocatedNodes2dFlattened()
270 if (P_1D != T_1D) QPack2d<NUM_COMP, P_1D, T_1D>(data, t_id_x, t_id_y, r_V); in GradTransposeTensorCollocatedNodes2dFlattened()
278 const int t_id_x = data.t_id_x % Q_1D, t_id_y = data.t_id_x / Q_1D; in WeightTensor2dFlattened() local
280 *w = (t_id_x < Q_1D && t_id_y < Q_1D) ? q_weight_1d[t_id_x] * q_weight_1d[t_id_y] : 0.0; in WeightTensor2dFlattened()
291 inline __device__ void ContractX3dFlattened(SharedData_Hip &data, const int t_id_x, const int t_id_… in ContractX3dFlattened() argument
294 data.slice[t_id_x + t_id_y * T_1D + t_id_z * T_1D * T_1D] = *U; in ContractX3dFlattened()
297 if (t_id_x < Q_1D && t_id_y < P_1D && t_id_z < P_1D) { in ContractX3dFlattened()
299 …*V += B[i + t_id_x * P_1D] * data.slice[i + t_id_y * T_1D + t_id_z * T_1D * T_1D]; // Contract x … in ContractX3dFlattened()
308 inline __device__ void ContractY3dFlattened(SharedData_Hip &data, const int t_id_x, const int t_id_… in ContractY3dFlattened() argument
311 data.slice[t_id_x + t_id_y * T_1D + t_id_z * T_1D * T_1D] = *U; in ContractY3dFlattened()
314 if (t_id_x < Q_1D && t_id_y < Q_1D && t_id_z < P_1D) { in ContractY3dFlattened()
316 …*V += B[i + t_id_y * P_1D] * data.slice[t_id_x + i * T_1D + t_id_z * T_1D * T_1D]; // Contract y … in ContractY3dFlattened()
325 inline __device__ void ContractZ3dFlattened(SharedData_Hip &data, const int t_id_x, const int t_id_… in ContractZ3dFlattened() argument
328 data.slice[t_id_x + t_id_y * T_1D + t_id_z * T_1D * T_1D] = *U; in ContractZ3dFlattened()
331 if (t_id_x < Q_1D && t_id_y < Q_1D && t_id_z < Q_1D) { in ContractZ3dFlattened()
333 …*V += B[i + t_id_z * P_1D] * data.slice[t_id_x + t_id_y * T_1D + i * T_1D * T_1D]; // Contract z … in ContractZ3dFlattened()
342 inline __device__ void ContractTransposeZ3dFlattened(SharedData_Hip &data, const int t_id_x, const … in ContractTransposeZ3dFlattened() argument
345 data.slice[t_id_x + t_id_y * T_1D + t_id_z * T_1D * T_1D] = *U; in ContractTransposeZ3dFlattened()
348 if (t_id_x < Q_1D && t_id_y < Q_1D && t_id_z < P_1D) { in ContractTransposeZ3dFlattened()
350 …*V += B[t_id_z + i * P_1D] * data.slice[t_id_x + t_id_y * T_1D + i * T_1D * T_1D]; // Contract z … in ContractTransposeZ3dFlattened()
359 inline __device__ void ContractTransposeAddZ3dFlattened(SharedData_Hip &data, const int t_id_x, con… in ContractTransposeAddZ3dFlattened() argument
362 data.slice[t_id_x + t_id_y * T_1D + t_id_z * T_1D * T_1D] = *U; in ContractTransposeAddZ3dFlattened()
364 if (t_id_x < Q_1D && t_id_y < Q_1D && t_id_z < P_1D) { in ContractTransposeAddZ3dFlattened()
366 …*V += B[t_id_z + i * P_1D] * data.slice[t_id_x + t_id_y * T_1D + i * T_1D * T_1D]; // Contract z … in ContractTransposeAddZ3dFlattened()
375 inline __device__ void ContractTransposeY3dFlattened(SharedData_Hip &data, const int t_id_x, const … in ContractTransposeY3dFlattened() argument
378 data.slice[t_id_x + t_id_y * T_1D + t_id_z * T_1D * T_1D] = *U; in ContractTransposeY3dFlattened()
381 if (t_id_x < Q_1D && t_id_y < P_1D && t_id_z < P_1D) { in ContractTransposeY3dFlattened()
383 …*V += B[t_id_y + i * P_1D] * data.slice[t_id_x + i * T_1D + t_id_z * T_1D * T_1D]; // Contract y … in ContractTransposeY3dFlattened()
392 inline __device__ void ContractTransposeAddY3dFlattened(SharedData_Hip &data, const int t_id_x, con… in ContractTransposeAddY3dFlattened() argument
395 data.slice[t_id_x + t_id_y * T_1D + t_id_z * T_1D * T_1D] = *U; in ContractTransposeAddY3dFlattened()
397 if (t_id_x < Q_1D && t_id_y < P_1D && t_id_z < P_1D) { in ContractTransposeAddY3dFlattened()
399 …*V += B[t_id_y + i * P_1D] * data.slice[t_id_x + i * T_1D + t_id_z * T_1D * T_1D]; // Contract y … in ContractTransposeAddY3dFlattened()
408 inline __device__ void ContractTransposeX3dFlattened(SharedData_Hip &data, const int t_id_x, const … in ContractTransposeX3dFlattened() argument
411 data.slice[t_id_x + t_id_y * T_1D + t_id_z * T_1D * T_1D] = *U; in ContractTransposeX3dFlattened()
414 if (t_id_x < P_1D && t_id_y < P_1D && t_id_z < P_1D) { in ContractTransposeX3dFlattened()
416 …*V += B[t_id_x + i * P_1D] * data.slice[i + t_id_y * T_1D + t_id_z * T_1D * T_1D]; // Contract x … in ContractTransposeX3dFlattened()
425 inline __device__ void ContractTransposeAddX3dFlattened(SharedData_Hip &data, const int t_id_x, con… in ContractTransposeAddX3dFlattened() argument
428 data.slice[t_id_x + t_id_y * T_1D + t_id_z * T_1D * T_1D] = *U; in ContractTransposeAddX3dFlattened()
430 if (t_id_x < P_1D && t_id_y < P_1D && t_id_z < P_1D) { in ContractTransposeAddX3dFlattened()
432 …*V += B[t_id_x + i * P_1D] * data.slice[i + t_id_y * T_1D + t_id_z * T_1D * T_1D]; // Contract x … in ContractTransposeAddX3dFlattened()
441 inline __device__ void QPack3d(SharedData_Hip &data, const int t_id_x, const int t_id_y, const int … in QPack3d() argument
442 …const CeedInt new_t_id_x = data.t_id_x % Q_1D, new_t_id_y = (data.t_id_x / Q_1D) % Q_1D, new_t_id_… in QPack3d()
446 …if (t_id_x < Q_1D && t_id_y < Q_1D) data.slice[t_id_x + t_id_y * T_1D + t_id_z * T_1D * T_1D] = U[… in QPack3d()
448 …U[comp] = data.t_id_x < (Q_1D * Q_1D * Q_1D) ? data.slice[new_t_id_x + new_t_id_y * T_1D + new_t_i… in QPack3d()
453 inline __device__ void QUnpack3d(SharedData_Hip &data, const int t_id_x, const int t_id_y, const in… in QUnpack3d() argument
454 …const CeedInt old_t_id_x = data.t_id_x % Q_1D, old_t_id_y = (data.t_id_x / Q_1D) % Q_1D, old_t_id_… in QUnpack3d()
458 …if (data.t_id_x < Q_1D * Q_1D * Q_1D) data.slice[old_t_id_x + old_t_id_y * T_1D + old_t_id_z * T_1… in QUnpack3d()
460 …U[comp] = (t_id_x < Q_1D && t_id_y < Q_1D) ? data.slice[t_id_x + t_id_y * T_1D + t_id_z * T_1D * T… in QUnpack3d()
470 …const CeedInt t_id_x = data.t_id_x % T_1D, t_id_y = (data.t_id_x / T_1D) % T_1D, t_id_z = data.t_i… in InterpTensor3dFlattened() local
473 if (P_1D != T_1D) QUnpack3d<NUM_COMP, P_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, r_U); in InterpTensor3dFlattened()
475 …ContractX3dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, &r_U[comp], c_B, r_… in InterpTensor3dFlattened()
476 ContractY3dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, r_t1, c_B, r_t2); in InterpTensor3dFlattened()
477 …ContractZ3dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, r_t2, c_B, &r_V[com… in InterpTensor3dFlattened()
480 if (P_1D != T_1D) QPack3d<NUM_COMP, P_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, r_U); in InterpTensor3dFlattened()
481 if (Q_1D != T_1D) QPack3d<NUM_COMP, Q_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, r_V); in InterpTensor3dFlattened()
490 …const CeedInt t_id_x = data.t_id_x % T_1D, t_id_y = (data.t_id_x / T_1D) % T_1D, t_id_z = data.t_i… in InterpTransposeTensor3dFlattened() local
493 if (Q_1D != T_1D) QUnpack3d<NUM_COMP, Q_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, r_U); in InterpTransposeTensor3dFlattened()
495 …ContractTransposeZ3dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, &r_U[comp]… in InterpTransposeTensor3dFlattened()
496 …ContractTransposeY3dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, r_t1, c_B,… in InterpTransposeTensor3dFlattened()
497 …ContractTransposeX3dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, r_t2, c_B,… in InterpTransposeTensor3dFlattened()
500 if (P_1D != T_1D) QPack3d<NUM_COMP, P_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, r_V); in InterpTransposeTensor3dFlattened()
509 …const CeedInt t_id_x = data.t_id_x % T_1D, t_id_y = (data.t_id_x / T_1D) % T_1D, t_id_z = data.t_i… in InterpTensorCollocatedNodes3dFlattened() local
511 if (P_1D != T_1D) QUnpack3d<NUM_COMP, P_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, r_U); in InterpTensorCollocatedNodes3dFlattened()
516 if (P_1D != T_1D) QPack3d<NUM_COMP, P_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, r_U); in InterpTensorCollocatedNodes3dFlattened()
517 if (Q_1D != T_1D) QPack3d<NUM_COMP, Q_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, r_V); in InterpTensorCollocatedNodes3dFlattened()
526 …const CeedInt t_id_x = data.t_id_x % T_1D, t_id_y = (data.t_id_x / T_1D) % T_1D, t_id_z = data.t_i… in InterpTransposeTensorCollocatedNodes3dFlattened() local
528 if (Q_1D != T_1D) QUnpack3d<NUM_COMP, Q_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, r_U); in InterpTransposeTensorCollocatedNodes3dFlattened()
533 if (P_1D != T_1D) QPack3d<NUM_COMP, P_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, r_V); in InterpTransposeTensorCollocatedNodes3dFlattened()
542 …const CeedInt t_id_x = data.t_id_x % T_1D, t_id_y = (data.t_id_x / T_1D) % T_1D, t_id_z = data.t_i… in GradTensor3dFlattened() local
545 if (P_1D != T_1D) QUnpack3d<NUM_COMP, P_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, r_U); in GradTensor3dFlattened()
547 …ContractX3dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, &r_U[comp], c_G, r_… in GradTensor3dFlattened()
548 ContractY3dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, r_t1, c_B, r_t2); in GradTensor3dFlattened()
549 …ContractZ3dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, r_t2, c_B, &r_V[com… in GradTensor3dFlattened()
550 …ContractX3dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, &r_U[comp], c_B, r_… in GradTensor3dFlattened()
551 ContractY3dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, r_t1, c_G, r_t2); in GradTensor3dFlattened()
552 …ContractZ3dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, r_t2, c_B, &r_V[com… in GradTensor3dFlattened()
553 …ContractX3dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, &r_U[comp], c_B, r_… in GradTensor3dFlattened()
554 ContractY3dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, r_t1, c_B, r_t2); in GradTensor3dFlattened()
555 …ContractZ3dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, r_t2, c_G, &r_V[com… in GradTensor3dFlattened()
558 if (P_1D != T_1D) QPack3d<NUM_COMP, P_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, r_U); in GradTensor3dFlattened()
559 if (Q_1D != T_1D) QPack3d<NUM_COMP * 3, Q_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, r_V); in GradTensor3dFlattened()
568 …const CeedInt t_id_x = data.t_id_x % T_1D, t_id_y = (data.t_id_x / T_1D) % T_1D, t_id_z = data.t_i… in GradTransposeTensor3dFlattened() local
571 if (Q_1D != T_1D) QUnpack3d<NUM_COMP * 3, Q_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, r_U); in GradTransposeTensor3dFlattened()
573 …ContractTransposeZ3dFlattened<NUM_COMP, t_id_x, t_id_y, t_id_z, P_1D, Q_1D, T_1D>(data, &r_U[comp … in GradTransposeTensor3dFlattened()
574 …ContractTransposeY3dFlattened<NUM_COMP, t_id_x, t_id_y, t_id_z, P_1D, Q_1D, T_1D>(data, r_t1, c_B,… in GradTransposeTensor3dFlattened()
575 …ContractTransposeX3dFlattened<NUM_COMP, t_id_x, t_id_y, t_id_z, P_1D, Q_1D, T_1D>(data, r_t2, c_G,… in GradTransposeTensor3dFlattened()
576 …ContractTransposeZ3dFlattened<NUM_COMP, t_id_x, t_id_y, t_id_z, P_1D, Q_1D, T_1D>(data, &r_U[comp … in GradTransposeTensor3dFlattened()
577 …ContractTransposeY3dFlattened<NUM_COMP, t_id_x, t_id_y, t_id_z, P_1D, Q_1D, T_1D>(data, r_t1, c_G,… in GradTransposeTensor3dFlattened()
578 …ContractTransposeAddX3dFlattened<NUM_COMP, t_id_x, t_id_y, t_id_z, P_1D, Q_1D, T_1D>(data, r_t2, c… in GradTransposeTensor3dFlattened()
579 …ContractTransposeZ3dFlattened<NUM_COMP, t_id_x, t_id_y, t_id_z, P_1D, Q_1D, T_1D>(data, &r_U[comp … in GradTransposeTensor3dFlattened()
580 …ContractTransposeY3dFlattened<NUM_COMP, t_id_x, t_id_y, t_id_z, P_1D, Q_1D, T_1D>(data, r_t1, c_B,… in GradTransposeTensor3dFlattened()
581 …ContractTransposeAddX3dFlattened<NUM_COMP, t_id_x, t_id_y, t_id_z, P_1D, Q_1D, T_1D>(data, r_t2, c… in GradTransposeTensor3dFlattened()
584 if (P_1D != T_1D) QPack3d<NUM_COMP, P_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, r_V); in GradTransposeTensor3dFlattened()
593 …const CeedInt t_id_x = data.t_id_x % T_1D, t_id_y = (data.t_id_x / T_1D) % T_1D, t_id_z = data.t_i… in GradTensorCollocated3dFlattened() local
596 if (P_1D != T_1D) QUnpack3d<NUM_COMP, P_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, r_U); in GradTensorCollocated3dFlattened()
598 …ContractX3dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, &r_U[comp], c_B, r_… in GradTensorCollocated3dFlattened()
599 ContractY3dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, r_t1, c_B, r_t2); in GradTensorCollocated3dFlattened()
600 ContractZ3dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, r_t2, c_B, r_t1); in GradTensorCollocated3dFlattened()
601 …ContractX3dFlattened<NUM_COMP, Q_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, r_t1, c_G, &r_V[com… in GradTensorCollocated3dFlattened()
602 …ContractY3dFlattened<NUM_COMP, Q_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, r_t1, c_G, &r_V[com… in GradTensorCollocated3dFlattened()
603 …ContractZ3dFlattened<NUM_COMP, Q_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, r_t1, c_G, &r_V[com… in GradTensorCollocated3dFlattened()
606 if (P_1D != T_1D) QPack3d<NUM_COMP, P_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, r_U); in GradTensorCollocated3dFlattened()
607 if (Q_1D != T_1D) QPack3d<NUM_COMP * 3, Q_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, r_V); in GradTensorCollocated3dFlattened()
616 …const CeedInt t_id_x = data.t_id_x % T_1D, t_id_y = (data.t_id_x / T_1D) % T_1D, t_id_z = data.t_i… in GradTransposeTensorCollocated3dFlattened() local
619 if (Q_1D != T_1D) QUnpack3d<NUM_COMP * 3, Q_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, r_U); in GradTransposeTensorCollocated3dFlattened()
621 …ContractTransposeZ3dFlattened<NUM_COMP, Q_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, &r_U[comp … in GradTransposeTensorCollocated3dFlattened()
622 …ContractTransposeAddY3dFlattened<NUM_COMP, Q_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, &r_U[co… in GradTransposeTensorCollocated3dFlattened()
623 …ContractTransposeAddX3dFlattened<NUM_COMP, Q_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, &r_U[co… in GradTransposeTensorCollocated3dFlattened()
624 …ContractTransposeZ3dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, r_t2, c_B,… in GradTransposeTensorCollocated3dFlattened()
625 …ContractTransposeY3dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, r_t1, c_B,… in GradTransposeTensorCollocated3dFlattened()
626 …ContractTransposeX3dFlattened<NUM_COMP, P_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, r_t2, c_B,… in GradTransposeTensorCollocated3dFlattened()
629 if (P_1D != T_1D) QPack3d<NUM_COMP, P_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, r_V); in GradTransposeTensorCollocated3dFlattened()
638 …const CeedInt t_id_x = data.t_id_x % T_1D, t_id_y = (data.t_id_x / T_1D) % T_1D, t_id_z = data.t_i… in GradTensorCollocatedNodes3dFlattened() local
640 if (P_1D != T_1D) QUnpack3d<NUM_COMP, P_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, r_U); in GradTensorCollocatedNodes3dFlattened()
642 …ContractX3dFlattened<NUM_COMP, Q_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, r_U[comp], c_G, &r_… in GradTensorCollocatedNodes3dFlattened()
643 …ContractY3dFlattened<NUM_COMP, Q_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, r_U[comp], c_G, &r_… in GradTensorCollocatedNodes3dFlattened()
644 …ContractZ3dFlattened<NUM_COMP, Q_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, r_U[comp], c_G, &r_… in GradTensorCollocatedNodes3dFlattened()
647 if (P_1D != T_1D) QPack3d<NUM_COMP, P_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, r_U); in GradTensorCollocatedNodes3dFlattened()
648 if (Q_1D != T_1D) QPack3d<NUM_COMP * 3, Q_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, r_V); in GradTensorCollocatedNodes3dFlattened()
657 …const CeedInt t_id_x = data.t_id_x % T_1D, t_id_y = (data.t_id_x / T_1D) % T_1D, t_id_z = data.t_i… in GradTransposeTensorCollocatedNodes3dFlattened() local
659 if (Q_1D != T_1D) QUnpack3d<NUM_COMP * 3, Q_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, r_U); in GradTransposeTensorCollocatedNodes3dFlattened()
661 …ContractTransposeZ3dFlattened<NUM_COMP, Q_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, &r_U[comp … in GradTransposeTensorCollocatedNodes3dFlattened()
662 …ContractTransposeAddY3dFlattened<NUM_COMP, Q_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, &r_U[co… in GradTransposeTensorCollocatedNodes3dFlattened()
663 …ContractTransposeAddX3dFlattened<NUM_COMP, Q_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, &r_U[co… in GradTransposeTensorCollocatedNodes3dFlattened()
666 if (P_1D != T_1D) QPack3d<NUM_COMP, P_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, r_V); in GradTransposeTensorCollocatedNodes3dFlattened()
674 …const CeedInt t_id_x = data.t_id_x % Q_1D, t_id_y = (data.t_id_x / Q_1D) % Q_1D, t_id_z = data.t_i… in WeightTensor3dFlattened() local
676 …*w = (t_id_x < Q_1D && t_id_y < Q_1D && t_id_z < Q_1D) ? q_weight_1d[t_id_x] * q_weight_1d[t_id_y]… in WeightTensor3dFlattened()