Lines Matching refs:data
20 inline __device__ void ContractX2dFlattened(SharedData_Cuda &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()
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_Cuda &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()
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_Cuda &data, const int t_id_x, const… in ContractTransposeY2dFlattened() argument
57 data.slice[t_id_x + t_id_y * T_1D] = *U; 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_Cuda &data, const int t_id_x, const… in ContractTransposeX2dFlattened() argument
74 data.slice[t_id_x + t_id_y * T_1D] = *U; 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_Cuda &data, const int t_id_x, co… in ContractTransposeAddX2dFlattened() argument
91 data.slice[t_id_x + t_id_y * T_1D] = *U; 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_Cuda &data, const int t_id_x, const int t_id_y, CeedScala… 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_Cuda &data, const int t_id_x, const int t_id_y, CeedSca… 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()
131 inline __device__ void InterpTensor2dFlattened(SharedData_Cuda &data, CeedScalar *__restrict__ r_U,… in InterpTensor2dFlattened() argument
133 const int t_id_x = data.t_id_x % T_1D, t_id_y = data.t_id_x / T_1D; in InterpTensor2dFlattened()
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()
150 inline __device__ void InterpTransposeTensor2dFlattened(SharedData_Cuda &data, CeedScalar *__restri… in InterpTransposeTensor2dFlattened() argument
152 const int t_id_x = data.t_id_x % T_1D, t_id_y = data.t_id_x / T_1D; in InterpTransposeTensor2dFlattened()
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()
168 inline __device__ void InterpTensorCollocatedNodes2dFlattened(SharedData_Cuda &data, CeedScalar *__… in InterpTensorCollocatedNodes2dFlattened() argument
170 const int t_id_x = data.t_id_x % T_1D, t_id_y = data.t_id_x / T_1D; in InterpTensorCollocatedNodes2dFlattened()
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()
185 inline __device__ void InterpTransposeTensorCollocatedNodes2dFlattened(SharedData_Cuda &data, CeedS… in InterpTransposeTensorCollocatedNodes2dFlattened() argument
187 const int t_id_x = data.t_id_x % T_1D, t_id_y = data.t_id_x / T_1D; in InterpTransposeTensorCollocatedNodes2dFlattened()
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()
201 inline __device__ void GradTensor2dFlattened(SharedData_Cuda &data, CeedScalar *__restrict__ r_U, c… in GradTensor2dFlattened() argument
203 const int t_id_x = data.t_id_x % T_1D, t_id_y = data.t_id_x / T_1D; in GradTensor2dFlattened()
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()
222 inline __device__ void GradTransposeTensor2dFlattened(SharedData_Cuda &data, CeedScalar *__restrict… in GradTransposeTensor2dFlattened() argument
224 const int t_id_x = data.t_id_x % T_1D, t_id_y = data.t_id_x / T_1D; in GradTransposeTensor2dFlattened()
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()
242 inline __device__ void GradTensorCollocatedNodes2dFlattened(SharedData_Cuda &data, CeedScalar *__re… in GradTensorCollocatedNodes2dFlattened() argument
244 const int t_id_x = data.t_id_x % T_1D, t_id_y = data.t_id_x / T_1D; in GradTensorCollocatedNodes2dFlattened()
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()
260 inline __device__ void GradTransposeTensorCollocatedNodes2dFlattened(SharedData_Cuda &data, CeedSca… in GradTransposeTensorCollocatedNodes2dFlattened() argument
262 const int t_id_x = data.t_id_x % T_1D, t_id_y = data.t_id_x / T_1D; in GradTransposeTensorCollocatedNodes2dFlattened()
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()
277 inline __device__ void WeightTensor2dFlattened(SharedData_Cuda &data, const CeedScalar *__restrict_… in WeightTensor2dFlattened() argument
278 const int t_id_x = data.t_id_x % Q_1D, t_id_y = data.t_id_x / Q_1D; in WeightTensor2dFlattened()
291 inline __device__ void ContractX3dFlattened(SharedData_Cuda &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()
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_Cuda &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()
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_Cuda &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()
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_Cuda &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()
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_Cuda &data, const int t_id_x, co… in ContractTransposeAddZ3dFlattened() argument
362 data.slice[t_id_x + t_id_y * T_1D + t_id_z * T_1D * T_1D] = *U; 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_Cuda &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()
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_Cuda &data, const int t_id_x, co… in ContractTransposeAddY3dFlattened() argument
395 data.slice[t_id_x + t_id_y * T_1D + t_id_z * T_1D * T_1D] = *U; 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_Cuda &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()
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_Cuda &data, const int t_id_x, co… in ContractTransposeAddX3dFlattened() argument
428 data.slice[t_id_x + t_id_y * T_1D + t_id_z * T_1D * T_1D] = *U; 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_Cuda &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_Cuda &data, const int t_id_x, const int t_id_y, const i… 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()
468 inline __device__ void InterpTensor3dFlattened(SharedData_Cuda &data, CeedScalar *__restrict__ r_U,… in InterpTensor3dFlattened() argument
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()
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()
488 inline __device__ void InterpTransposeTensor3dFlattened(SharedData_Cuda &data, CeedScalar *__restri… in InterpTransposeTensor3dFlattened() argument
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()
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()
507 inline __device__ void InterpTensorCollocatedNodes3dFlattened(SharedData_Cuda &data, CeedScalar *__… in InterpTensorCollocatedNodes3dFlattened() argument
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()
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()
524 inline __device__ void InterpTransposeTensorCollocatedNodes3dFlattened(SharedData_Cuda &data, CeedS… in InterpTransposeTensorCollocatedNodes3dFlattened() argument
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()
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()
540 inline __device__ void GradTensor3dFlattened(SharedData_Cuda &data, CeedScalar *__restrict__ r_U, c… in GradTensor3dFlattened() argument
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()
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()
566 inline __device__ void GradTransposeTensor3dFlattened(SharedData_Cuda &data, CeedScalar *__restrict… in GradTransposeTensor3dFlattened() argument
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()
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()
591 inline __device__ void GradTensorCollocated3dFlattened(SharedData_Cuda &data, CeedScalar *__restric… in GradTensorCollocated3dFlattened() argument
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()
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()
614 inline __device__ void GradTransposeTensor3dFlattened(SharedData_Cuda &data, CeedScalar *__restrict… in GradTransposeTensor3dFlattened() argument
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 GradTransposeTensor3dFlattened()
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 GradTransposeTensor3dFlattened()
621 …ContractTransposeZ3dFlattened<NUM_COMP, t_id_x, t_id_y, t_id_z, P_1D, Q_1D, T_1D>(data, &r_U[comp … in GradTransposeTensor3dFlattened()
622 …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()
623 …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()
624 …ContractTransposeZ3dFlattened<NUM_COMP, t_id_x, t_id_y, t_id_z, P_1D, Q_1D, T_1D>(data, &r_U[comp … in GradTransposeTensor3dFlattened()
625 …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()
626 …ContractTransposeAddX3dFlattened<NUM_COMP, t_id_x, t_id_y, t_id_z, P_1D, Q_1D, T_1D>(data, r_t2, c… in GradTransposeTensor3dFlattened()
627 …ContractTransposeZ3dFlattened<NUM_COMP, t_id_x, t_id_y, t_id_z, P_1D, Q_1D, T_1D>(data, &r_U[comp … in GradTransposeTensor3dFlattened()
628 …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()
629 …ContractTransposeAddX3dFlattened<NUM_COMP, t_id_x, t_id_y, t_id_z, P_1D, Q_1D, T_1D>(data, r_t2, c… in GradTransposeTensor3dFlattened()
632 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()
639 inline __device__ void GradTensorCollocatedNodes3dFlattened(SharedData_Cuda &data, CeedScalar *__re… in GradTensorCollocatedNodes3dFlattened() argument
641 …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()
643 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()
645 …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()
646 …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()
647 …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()
650 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()
651 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()
658 inline __device__ void GradTransposeTensorCollocatedNodes3dFlattened(SharedData_Cuda &data, CeedSca… in GradTransposeTensorCollocatedNodes3dFlattened() argument
660 …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()
662 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()
664 …ContractTransposeZ3dFlattened<NUM_COMP, Q_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, &r_U[comp … in GradTransposeTensorCollocatedNodes3dFlattened()
665 …ContractTransposeAddY3dFlattened<NUM_COMP, Q_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, &r_U[co… in GradTransposeTensorCollocatedNodes3dFlattened()
666 …ContractTransposeAddX3dFlattened<NUM_COMP, Q_1D, Q_1D, T_1D>(data, t_id_x, t_id_y, t_id_z, &r_U[co… in GradTransposeTensorCollocatedNodes3dFlattened()
669 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()
676 inline __device__ void WeightTensor3dFlattened(SharedData_Cuda &data, const CeedScalar *__restrict_… in WeightTensor3dFlattened() argument
677 …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()