Lines Matching refs:data
20 inline __device__ void ContractX1d(SharedData_Hip &data, const CeedScalar *U, const CeedScalar *B, … in ContractX1d() argument
22 data.slice[data.t_id_x] = *U; in ContractX1d()
25 if (data.t_id_x < Q_1D) { in ContractX1d()
27 *V += B[i + data.t_id_x * P_1D] * data.slice[i]; // Contract x direction in ContractX1d()
36 inline __device__ void ContractTransposeX1d(SharedData_Hip &data, const CeedScalar *U, const CeedSc… in ContractTransposeX1d() argument
38 data.slice[data.t_id_x] = *U; in ContractTransposeX1d()
41 if (data.t_id_x < P_1D) { in ContractTransposeX1d()
43 *V += B[data.t_id_x + i * P_1D] * data.slice[i]; // Contract x direction in ContractTransposeX1d()
52 inline __device__ void Interp1d(SharedData_Hip &data, const CeedScalar *__restrict__ r_U, const Cee… in Interp1d() argument
54 ContractX1d<NUM_COMP, P_1D, Q_1D>(data, &r_U[comp], c_B, &r_V[comp]); in Interp1d()
62 inline __device__ void InterpTranspose1d(SharedData_Hip &data, const CeedScalar *__restrict__ r_U, … in InterpTranspose1d() argument
65 ContractTransposeX1d<NUM_COMP, P_1D, Q_1D>(data, &r_U[comp], c_B, &r_V[comp]); in InterpTranspose1d()
73 inline __device__ void InterpCollocatedNodes1d(SharedData_Hip &data, const CeedScalar *__restrict__… in InterpCollocatedNodes1d() argument
84 inline __device__ void InterpTransposeCollocatedNodes1d(SharedData_Hip &data, const CeedScalar *__r… in InterpTransposeCollocatedNodes1d() argument
95 inline __device__ void Grad1d(SharedData_Hip &data, const CeedScalar *__restrict__ r_U, const CeedS… in Grad1d() argument
98 ContractX1d<NUM_COMP, P_1D, Q_1D>(data, &r_U[comp], c_G, &r_V[comp]); in Grad1d()
106 inline __device__ void GradTranspose1d(SharedData_Hip &data, const CeedScalar *__restrict__ r_U, co… in GradTranspose1d() argument
109 ContractTransposeX1d<NUM_COMP, P_1D, Q_1D>(data, &r_U[comp], c_G, &r_V[comp]); in GradTranspose1d()
117 inline __device__ void Weight1d(SharedData_Hip &data, const CeedScalar *__restrict__ q_weight_1d, C… in Weight1d() argument
118 *w = (data.t_id_x < Q_1D) ? q_weight_1d[data.t_id_x] : 0.0; in Weight1d()
129 inline __device__ void ContractX2d(SharedData_Hip &data, const CeedScalar *U, const CeedScalar *B, … in ContractX2d() argument
131 data.slice[data.t_id_x + data.t_id_y * T_1D] = *U; in ContractX2d()
134 if (data.t_id_x < Q_1D && data.t_id_y < P_1D) { in ContractX2d()
136 *V += B[i + data.t_id_x * P_1D] * data.slice[i + data.t_id_y * T_1D]; // Contract x direction in ContractX2d()
145 inline __device__ void ContractY2d(SharedData_Hip &data, const CeedScalar *U, const CeedScalar *B, … in ContractY2d() argument
147 data.slice[data.t_id_x + data.t_id_y * T_1D] = *U; in ContractY2d()
150 if (data.t_id_x < Q_1D && data.t_id_y < Q_1D) { in ContractY2d()
152 *V += B[i + data.t_id_y * P_1D] * data.slice[data.t_id_x + i * T_1D]; // Contract y direction in ContractY2d()
161 inline __device__ void ContractTransposeY2d(SharedData_Hip &data, const CeedScalar *U, const CeedSc… in ContractTransposeY2d() argument
163 data.slice[data.t_id_x + data.t_id_y * T_1D] = *U; in ContractTransposeY2d()
166 if (data.t_id_x < Q_1D && data.t_id_y < P_1D) { in ContractTransposeY2d()
168 *V += B[data.t_id_y + i * P_1D] * data.slice[data.t_id_x + i * T_1D]; // Contract y direction in ContractTransposeY2d()
177 inline __device__ void ContractTransposeX2d(SharedData_Hip &data, const CeedScalar *U, const CeedSc… in ContractTransposeX2d() argument
179 data.slice[data.t_id_x + data.t_id_y * T_1D] = *U; in ContractTransposeX2d()
182 if (data.t_id_x < P_1D && data.t_id_y < P_1D) { in ContractTransposeX2d()
184 *V += B[data.t_id_x + i * P_1D] * data.slice[i + data.t_id_y * T_1D]; // Contract x direction in ContractTransposeX2d()
193 inline __device__ void ContractTransposeAddX2d(SharedData_Hip &data, const CeedScalar *U, const Cee… in ContractTransposeAddX2d() argument
195 data.slice[data.t_id_x + data.t_id_y * T_1D] = *U; in ContractTransposeAddX2d()
197 if (data.t_id_x < P_1D && data.t_id_y < P_1D) { in ContractTransposeAddX2d()
199 *V += B[data.t_id_x + i * P_1D] * data.slice[i + data.t_id_y * T_1D]; // Contract x direction in ContractTransposeAddX2d()
208 inline __device__ void InterpTensor2d(SharedData_Hip &data, const CeedScalar *__restrict__ r_U, con… in InterpTensor2d() argument
211 ContractX2d<NUM_COMP, P_1D, Q_1D, T_1D>(data, &r_U[comp], c_B, r_t); in InterpTensor2d()
212 ContractY2d<NUM_COMP, P_1D, Q_1D, T_1D>(data, r_t, c_B, &r_V[comp]); in InterpTensor2d()
220 inline __device__ void InterpTransposeTensor2d(SharedData_Hip &data, const CeedScalar *__restrict__… in InterpTransposeTensor2d() argument
224 ContractTransposeY2d<NUM_COMP, P_1D, Q_1D, T_1D>(data, &r_U[comp], c_B, r_t); in InterpTransposeTensor2d()
225 ContractTransposeX2d<NUM_COMP, P_1D, Q_1D, T_1D>(data, r_t, c_B, &r_V[comp]); in InterpTransposeTensor2d()
233 inline __device__ void InterpTensorCollocatedNodes2d(SharedData_Hip &data, const CeedScalar *__rest… in InterpTensorCollocatedNodes2d() argument
244 inline __device__ void InterpTransposeTensorCollocatedNodes2d(SharedData_Hip &data, const CeedScala… in InterpTransposeTensorCollocatedNodes2d() argument
255 inline __device__ void GradTensor2d(SharedData_Hip &data, const CeedScalar *__restrict__ r_U, const… in GradTensor2d() argument
259 ContractX2d<NUM_COMP, P_1D, Q_1D, T_1D>(data, &r_U[comp], c_G, r_t); in GradTensor2d()
260 ContractY2d<NUM_COMP, P_1D, Q_1D, T_1D>(data, r_t, c_B, &r_V[comp + 0 * NUM_COMP]); in GradTensor2d()
261 ContractX2d<NUM_COMP, P_1D, Q_1D, T_1D>(data, &r_U[comp], c_B, r_t); in GradTensor2d()
262 ContractY2d<NUM_COMP, P_1D, Q_1D, T_1D>(data, r_t, c_G, &r_V[comp + 1 * NUM_COMP]); in GradTensor2d()
270 inline __device__ void GradTransposeTensor2d(SharedData_Hip &data, const CeedScalar *__restrict__ r… in GradTransposeTensor2d() argument
274 ContractTransposeY2d<NUM_COMP, P_1D, Q_1D, T_1D>(data, &r_U[comp + 0 * NUM_COMP], c_B, r_t); in GradTransposeTensor2d()
275 ContractTransposeX2d<NUM_COMP, P_1D, Q_1D, T_1D>(data, r_t, c_G, &r_V[comp]); in GradTransposeTensor2d()
276 ContractTransposeY2d<NUM_COMP, P_1D, Q_1D, T_1D>(data, &r_U[comp + 1 * NUM_COMP], c_G, r_t); in GradTransposeTensor2d()
277 ContractTransposeAddX2d<NUM_COMP, P_1D, Q_1D, T_1D>(data, r_t, c_B, &r_V[comp]); in GradTransposeTensor2d()
285 inline __device__ void GradTensorCollocatedNodes2d(SharedData_Hip &data, const CeedScalar *__restri… in GradTensorCollocatedNodes2d() argument
288 ContractX2d<NUM_COMP, P_1D, Q_1D, T_1D>(data, &r_U[comp], c_G, &r_V[comp + 0 * NUM_COMP]); in GradTensorCollocatedNodes2d()
289 ContractY2d<NUM_COMP, P_1D, Q_1D, T_1D>(data, &r_U[comp], c_G, &r_V[comp + 1 * NUM_COMP]); in GradTensorCollocatedNodes2d()
297 inline __device__ void GradTransposeTensorCollocatedNodes2d(SharedData_Hip &data, const CeedScalar … in GradTransposeTensorCollocatedNodes2d() argument
300 …ContractTransposeY2d<NUM_COMP, P_1D, Q_1D, T_1D>(data, &r_U[comp + 1 * NUM_COMP], c_G, &r_V[comp]); in GradTransposeTensorCollocatedNodes2d()
301 …ContractTransposeAddX2d<NUM_COMP, P_1D, Q_1D, T_1D>(data, &r_U[comp + 0 * NUM_COMP], c_G, &r_V[com… in GradTransposeTensorCollocatedNodes2d()
309 inline __device__ void WeightTensor2d(SharedData_Hip &data, const CeedScalar *__restrict__ q_weight… in WeightTensor2d() argument
310 …*w = (data.t_id_x < Q_1D && data.t_id_y < Q_1D) ? q_weight_1d[data.t_id_x] * q_weight_1d[data.t_id… in WeightTensor2d()
321 inline __device__ void ContractX3d(SharedData_Hip &data, const CeedScalar *U, const CeedScalar *B, … in ContractX3d() argument
324 r_B[i] = B[i + data.t_id_x * P_1D]; in ContractX3d()
329 data.slice[data.t_id_x + data.t_id_y * T_1D] = U[k]; in ContractX3d()
332 if (data.t_id_x < Q_1D && data.t_id_y < P_1D) { in ContractX3d()
334 V[k] += r_B[i] * data.slice[i + data.t_id_y * T_1D]; // Contract x direction in ContractX3d()
344 inline __device__ void ContractY3d(SharedData_Hip &data, const CeedScalar *U, const CeedScalar *B, … in ContractY3d() argument
347 r_B[i] = B[i + data.t_id_y * P_1D]; in ContractY3d()
352 data.slice[data.t_id_x + data.t_id_y * T_1D] = U[k]; in ContractY3d()
355 if (data.t_id_x < Q_1D && data.t_id_y < Q_1D) { in ContractY3d()
357 V[k] += r_B[i] * data.slice[data.t_id_x + i * T_1D]; // Contract y direction in ContractY3d()
367 inline __device__ void ContractZ3d(SharedData_Hip &data, const CeedScalar *U, const CeedScalar *B, … in ContractZ3d() argument
370 if (data.t_id_x < Q_1D && data.t_id_y < Q_1D) { in ContractZ3d()
382 inline __device__ void ContractTransposeZ3d(SharedData_Hip &data, const CeedScalar *U, const CeedSc… in ContractTransposeZ3d() argument
385 if (data.t_id_x < Q_1D && data.t_id_y < Q_1D) { in ContractTransposeZ3d()
397 inline __device__ void ContractTransposeY3d(SharedData_Hip &data, const CeedScalar *U, const CeedSc… in ContractTransposeY3d() argument
400 r_B[i] = B[data.t_id_y + i * P_1D]; in ContractTransposeY3d()
405 data.slice[data.t_id_x + data.t_id_y * T_1D] = U[k]; in ContractTransposeY3d()
408 if (data.t_id_x < Q_1D && data.t_id_y < P_1D) { in ContractTransposeY3d()
410 V[k] += r_B[i] * data.slice[data.t_id_x + i * T_1D]; // Contract y direction in ContractTransposeY3d()
420 inline __device__ void ContractTransposeAddY3d(SharedData_Hip &data, const CeedScalar *U, const Cee… in ContractTransposeAddY3d() argument
423 r_B[i] = B[data.t_id_y + i * P_1D]; in ContractTransposeAddY3d()
428 data.slice[data.t_id_x + data.t_id_y * T_1D] = U[k]; in ContractTransposeAddY3d()
430 if (data.t_id_x < Q_1D && data.t_id_y < P_1D) { in ContractTransposeAddY3d()
432 V[k] += r_B[i] * data.slice[data.t_id_x + i * T_1D]; // Contract y direction in ContractTransposeAddY3d()
442 inline __device__ void ContractTransposeX3d(SharedData_Hip &data, const CeedScalar *U, const CeedSc… in ContractTransposeX3d() argument
445 r_B[i] = B[data.t_id_x + i * P_1D]; in ContractTransposeX3d()
450 data.slice[data.t_id_x + data.t_id_y * T_1D] = U[k]; in ContractTransposeX3d()
453 if (data.t_id_x < P_1D && data.t_id_y < P_1D) { in ContractTransposeX3d()
455 V[k] += r_B[i] * data.slice[i + data.t_id_y * T_1D]; // Contract x direction in ContractTransposeX3d()
465 inline __device__ void ContractTransposeAddX3d(SharedData_Hip &data, const CeedScalar *U, const Cee… in ContractTransposeAddX3d() argument
468 r_B[i] = B[data.t_id_x + i * P_1D]; in ContractTransposeAddX3d()
473 data.slice[data.t_id_x + data.t_id_y * T_1D] = U[k]; in ContractTransposeAddX3d()
475 if (data.t_id_x < P_1D && data.t_id_y < P_1D) { in ContractTransposeAddX3d()
477 V[k] += r_B[i] * data.slice[i + data.t_id_y * T_1D]; // Contract x direction in ContractTransposeAddX3d()
487 inline __device__ void InterpTensor3d(SharedData_Hip &data, const CeedScalar *__restrict__ r_U, con… in InterpTensor3d() argument
491 ContractX3d<NUM_COMP, P_1D, Q_1D, T_1D>(data, &r_U[comp * P_1D], c_B, r_t1); in InterpTensor3d()
492 ContractY3d<NUM_COMP, P_1D, Q_1D, T_1D>(data, r_t1, c_B, r_t2); in InterpTensor3d()
493 ContractZ3d<NUM_COMP, P_1D, Q_1D, T_1D>(data, r_t2, c_B, &r_V[comp * Q_1D]); in InterpTensor3d()
501 inline __device__ void InterpTransposeTensor3d(SharedData_Hip &data, const CeedScalar *__restrict__… in InterpTransposeTensor3d() argument
506 ContractTransposeZ3d<NUM_COMP, P_1D, Q_1D, T_1D>(data, &r_U[comp * Q_1D], c_B, r_t1); in InterpTransposeTensor3d()
507 ContractTransposeY3d<NUM_COMP, P_1D, Q_1D, T_1D>(data, r_t1, c_B, r_t2); in InterpTransposeTensor3d()
508 ContractTransposeX3d<NUM_COMP, P_1D, Q_1D, T_1D>(data, r_t2, c_B, &r_V[comp * P_1D]); in InterpTransposeTensor3d()
516 inline __device__ void InterpTensorCollocatedNodes3d(SharedData_Hip &data, const CeedScalar *__rest… in InterpTensorCollocatedNodes3d() argument
529 inline __device__ void InterpTransposeTensorCollocatedNodes3d(SharedData_Hip &data, const CeedScala… in InterpTransposeTensorCollocatedNodes3d() argument
542 inline __device__ void GradTensor3d(SharedData_Hip &data, const CeedScalar *__restrict__ r_U, const… in GradTensor3d() argument
547 ContractX3d<NUM_COMP, P_1D, Q_1D, T_1D>(data, &r_U[comp * P_1D], c_G, r_t1); in GradTensor3d()
548 ContractY3d<NUM_COMP, P_1D, Q_1D, T_1D>(data, r_t1, c_B, r_t2); in GradTensor3d()
549 … ContractZ3d<NUM_COMP, P_1D, Q_1D, T_1D>(data, r_t2, c_B, &r_V[comp * Q_1D + 0 * NUM_COMP * Q_1D]); in GradTensor3d()
550 ContractX3d<NUM_COMP, P_1D, Q_1D, T_1D>(data, &r_U[comp * P_1D], c_B, r_t1); in GradTensor3d()
551 ContractY3d<NUM_COMP, P_1D, Q_1D, T_1D>(data, r_t1, c_G, r_t2); in GradTensor3d()
552 … ContractZ3d<NUM_COMP, P_1D, Q_1D, T_1D>(data, r_t2, c_B, &r_V[comp * Q_1D + 1 * NUM_COMP * Q_1D]); in GradTensor3d()
553 ContractX3d<NUM_COMP, P_1D, Q_1D, T_1D>(data, &r_U[comp * P_1D], c_B, r_t1); in GradTensor3d()
554 ContractY3d<NUM_COMP, P_1D, Q_1D, T_1D>(data, r_t1, c_B, r_t2); in GradTensor3d()
555 … ContractZ3d<NUM_COMP, P_1D, Q_1D, T_1D>(data, r_t2, c_G, &r_V[comp * Q_1D + 2 * NUM_COMP * Q_1D]); in GradTensor3d()
563 inline __device__ void GradTransposeTensor3d(SharedData_Hip &data, const CeedScalar *__restrict__ r… in GradTransposeTensor3d() argument
568 …ContractTransposeZ3d<NUM_COMP, P_1D, Q_1D, T_1D>(data, &r_U[comp * Q_1D + 0 * NUM_COMP * Q_1D], c_… in GradTransposeTensor3d()
569 ContractTransposeY3d<NUM_COMP, P_1D, Q_1D, T_1D>(data, r_t1, c_B, r_t2); in GradTransposeTensor3d()
570 ContractTransposeX3d<NUM_COMP, P_1D, Q_1D, T_1D>(data, r_t2, c_G, &r_V[comp * P_1D]); in GradTransposeTensor3d()
571 …ContractTransposeZ3d<NUM_COMP, P_1D, Q_1D, T_1D>(data, &r_U[comp * Q_1D + 1 * NUM_COMP * Q_1D], c_… in GradTransposeTensor3d()
572 ContractTransposeY3d<NUM_COMP, P_1D, Q_1D, T_1D>(data, r_t1, c_G, r_t2); in GradTransposeTensor3d()
573 ContractTransposeAddX3d<NUM_COMP, P_1D, Q_1D, T_1D>(data, r_t2, c_B, &r_V[comp * P_1D]); in GradTransposeTensor3d()
574 …ContractTransposeZ3d<NUM_COMP, P_1D, Q_1D, T_1D>(data, &r_U[comp * Q_1D + 2 * NUM_COMP * Q_1D], c_… in GradTransposeTensor3d()
575 ContractTransposeY3d<NUM_COMP, P_1D, Q_1D, T_1D>(data, r_t1, c_B, r_t2); in GradTransposeTensor3d()
576 ContractTransposeAddX3d<NUM_COMP, P_1D, Q_1D, T_1D>(data, r_t2, c_B, &r_V[comp * P_1D]); in GradTransposeTensor3d()
584 inline __device__ void GradTensorCollocated3d(SharedData_Hip &data, const CeedScalar *__restrict__ … in GradTensorCollocated3d() argument
589 ContractX3d<NUM_COMP, P_1D, Q_1D, T_1D>(data, &r_U[comp * P_1D], c_B, r_t1); in GradTensorCollocated3d()
590 ContractY3d<NUM_COMP, P_1D, Q_1D, T_1D>(data, r_t1, c_B, r_t2); in GradTensorCollocated3d()
591 ContractZ3d<NUM_COMP, P_1D, Q_1D, T_1D>(data, r_t2, c_B, r_t1); in GradTensorCollocated3d()
592 … ContractX3d<NUM_COMP, Q_1D, Q_1D, T_1D>(data, r_t1, c_G, &r_V[comp * Q_1D + 0 * NUM_COMP * Q_1D]); in GradTensorCollocated3d()
593 … ContractY3d<NUM_COMP, Q_1D, Q_1D, T_1D>(data, r_t1, c_G, &r_V[comp * Q_1D + 1 * NUM_COMP * Q_1D]); in GradTensorCollocated3d()
594 … ContractZ3d<NUM_COMP, Q_1D, Q_1D, T_1D>(data, r_t1, c_G, &r_V[comp * Q_1D + 2 * NUM_COMP * Q_1D]); in GradTensorCollocated3d()
602 inline __device__ void GradTransposeTensorCollocated3d(SharedData_Hip &data, const CeedScalar *__re… in GradTransposeTensorCollocated3d() argument
607 …ContractTransposeZ3d<NUM_COMP, Q_1D, Q_1D, T_1D>(data, &r_U[comp * Q_1D + 2 * NUM_COMP * Q_1D], c_… in GradTransposeTensorCollocated3d()
608 …ContractTransposeAddY3d<NUM_COMP, Q_1D, Q_1D, T_1D>(data, &r_U[comp * Q_1D + 1 * NUM_COMP * Q_1D],… in GradTransposeTensorCollocated3d()
609 …ContractTransposeAddX3d<NUM_COMP, Q_1D, Q_1D, T_1D>(data, &r_U[comp * Q_1D + 0 * NUM_COMP * Q_1D],… in GradTransposeTensorCollocated3d()
610 ContractTransposeZ3d<NUM_COMP, P_1D, Q_1D, T_1D>(data, r_t2, c_B, r_t1); in GradTransposeTensorCollocated3d()
611 ContractTransposeY3d<NUM_COMP, P_1D, Q_1D, T_1D>(data, r_t1, c_B, r_t2); in GradTransposeTensorCollocated3d()
612 ContractTransposeX3d<NUM_COMP, P_1D, Q_1D, T_1D>(data, r_t2, c_B, &r_V[comp * P_1D]); in GradTransposeTensorCollocated3d()
620 inline __device__ void GradTensorCollocatedNodes3d(SharedData_Hip &data, const CeedScalar *__restri… in GradTensorCollocatedNodes3d() argument
623 …ContractX3d<NUM_COMP, P_1D, Q_1D, T_1D>(data, &r_U[comp * P_1D], c_G, &r_V[comp * Q_1D + 0 * NUM_C… in GradTensorCollocatedNodes3d()
624 …ContractY3d<NUM_COMP, P_1D, Q_1D, T_1D>(data, &r_U[comp * P_1D], c_G, &r_V[comp * Q_1D + 1 * NUM_C… in GradTensorCollocatedNodes3d()
625 …ContractZ3d<NUM_COMP, P_1D, Q_1D, T_1D>(data, &r_U[comp * P_1D], c_G, &r_V[comp * Q_1D + 2 * NUM_C… in GradTensorCollocatedNodes3d()
633 inline __device__ void GradTransposeTensorCollocatedNodes3d(SharedData_Hip &data, const CeedScalar … in GradTransposeTensorCollocatedNodes3d() argument
636 …ContractTransposeZ3d<NUM_COMP, P_1D, Q_1D, T_1D>(data, &r_U[comp * Q_1D + 2 * NUM_COMP * Q_1D], c_… in GradTransposeTensorCollocatedNodes3d()
637 …ContractTransposeAddY3d<NUM_COMP, P_1D, Q_1D, T_1D>(data, &r_U[comp * Q_1D + 1 * NUM_COMP * Q_1D],… in GradTransposeTensorCollocatedNodes3d()
638 …ContractTransposeAddX3d<NUM_COMP, P_1D, Q_1D, T_1D>(data, &r_U[comp * Q_1D + 0 * NUM_COMP * Q_1D],… in GradTransposeTensorCollocatedNodes3d()
646 inline __device__ void WeightTensor3d(SharedData_Hip &data, const CeedScalar *__restrict__ q_weight… in WeightTensor3d() argument
647 const bool quad = (data.t_id_x < Q_1D && data.t_id_y < Q_1D); in WeightTensor3d()
648 const CeedScalar pw = quad ? q_weight_1d[data.t_id_x] * q_weight_1d[data.t_id_y] : 0.0; in WeightTensor3d()