| /libCEED/include/ceed/jit-source/sycl/ |
| H A D | sycl-shared-basis-tensor.h | 26 CeedScalar r_U[BASIS_NUM_COMP * (BASIS_DIM > 2 ? BASIS_P_1D : 1)]; in Interp() local
37 …1d(BASIS_NUM_COMP, BASIS_P_1D, num_elem, 1, BASIS_NUM_NODES * num_elem, BASIS_NUM_NODES, d_U, r_U); in Interp()
38 Interp1d(BASIS_NUM_COMP, BASIS_P_1D, BASIS_Q_1D, r_U, s_B, r_V, elem_scratch); in Interp()
42 …2d(BASIS_NUM_COMP, BASIS_P_1D, num_elem, 1, BASIS_NUM_NODES * num_elem, BASIS_NUM_NODES, d_U, r_U); in Interp()
43 InterpTensor2d(BASIS_NUM_COMP, BASIS_P_1D, BASIS_Q_1D, r_U, s_B, r_V, elem_scratch); in Interp()
47 …3d(BASIS_NUM_COMP, BASIS_P_1D, num_elem, 1, BASIS_NUM_NODES * num_elem, BASIS_NUM_NODES, d_U, r_U); in Interp()
48 InterpTensor3d(BASIS_NUM_COMP, BASIS_P_1D, BASIS_Q_1D, r_U, s_B, r_V, elem_scratch); in Interp()
60 CeedScalar r_U[BASIS_NUM_COMP * (BASIS_DIM > 2 ? BASIS_Q_1D : 1)]; in InterpTranspose() local
71 …ed1d(BASIS_NUM_COMP, BASIS_Q_1D, num_elem, 1, BASIS_NUM_QPTS * num_elem, BASIS_NUM_QPTS, d_U, r_U); in InterpTranspose()
72 InterpTranspose1d(BASIS_NUM_COMP, BASIS_P_1D, BASIS_Q_1D, r_U, s_B, r_V, elem_scratch); in InterpTranspose()
[all …]
|
| H A D | sycl-shared-basis-tensor-templates.h | 57 …t CeedInt NUM_COMP, const CeedInt P_1D, const CeedInt Q_1D, private const CeedScalar *restrict r_U, in Interp1d() argument
60 ContractX1d(P_1D, Q_1D, r_U + comp, s_B, r_V + comp, scratch); in Interp1d()
67 …t CeedInt NUM_COMP, const CeedInt P_1D, const CeedInt Q_1D, private const CeedScalar *restrict r_U, in InterpTranspose1d() argument
70 ContractTransposeX1d(P_1D, Q_1D, r_U + comp, s_B, r_V + comp, scratch); in InterpTranspose1d()
77 …t CeedInt NUM_COMP, const CeedInt P_1D, const CeedInt Q_1D, private const CeedScalar *restrict r_U, in Grad1d() argument
80 ContractX1d(P_1D, Q_1D, r_U + comp, s_G, r_V + comp, scratch); in Grad1d()
87 …t CeedInt NUM_COMP, const CeedInt P_1D, const CeedInt Q_1D, private const CeedScalar *restrict r_U, in GradTranspose1d() argument
90 ContractTransposeX1d(P_1D, Q_1D, r_U + comp, s_G, r_V + comp, scratch); in GradTranspose1d()
208 …t CeedInt NUM_COMP, const CeedInt P_1D, const CeedInt Q_1D, private const CeedScalar *restrict r_U, in InterpTensor2d() argument
213 ContractX2d(P_1D, Q_1D, r_U + comp, s_B, r_t, scratch); in InterpTensor2d()
[all …]
|
| H A D | sycl-gen-templates.h | 285 …onst CeedInt num_comp, const CeedInt Q_1D, const CeedInt q, const private CeedScalar *restrict r_U, in gradCollo3d() argument
292 scratch[item_id_x + item_id_y * T_1D] = r_U[q + comp * Q_1D]; in gradCollo3d()
309 …for (CeedInt i = 0; i < Q_1D; ++i) r_V[comp + 2 * num_comp] += s_G[i + q * Q_1D] * r_U[i + comp * … in gradCollo3d()
319 …onst CeedInt num_comp, const CeedInt Q_1D, const CeedInt q, const private CeedScalar *restrict r_U, in gradColloTranspose3d() argument
327 scratch[item_id_x + item_id_y * T_1D] = r_U[comp + 0 * num_comp]; in gradColloTranspose3d()
339 scratch[item_id_x + item_id_y * T_1D] = r_U[comp + 1 * num_comp]; in gradColloTranspose3d()
352 …r_V[i + comp * Q_1D] += s_G[i + q * Q_1D] * r_U[comp + 2 * num_comp]; // PARTIAL contract z direc… in gradColloTranspose3d()
|
| /libCEED/include/ceed/jit-source/cuda/ |
| H A D | cuda-shared-basis-tensor.h | 28 CeedScalar r_U[BASIS_NUM_COMP * (BASIS_DIM > 2 ? BASIS_P_1D : 1)]; in Interp() local
39 …tStrided1d<BASIS_NUM_COMP, BASIS_P_1D>(data, elem, 1, BASIS_P_1D * num_elem, BASIS_P_1D, d_U, r_U); in Interp()
40 Interp1d<BASIS_NUM_COMP, BASIS_P_1D, BASIS_Q_1D, BASIS_T_1D>(data, r_U, s_B, r_V); in Interp()
43 … BASIS_P_1D>(data, elem, 1, BASIS_P_1D * BASIS_P_1D * num_elem, BASIS_P_1D * BASIS_P_1D, d_U, r_U); in Interp()
44 InterpTensor2d<BASIS_NUM_COMP, BASIS_P_1D, BASIS_Q_1D, BASIS_T_1D>(data, r_U, s_B, r_V); in Interp()
48 … BASIS_P_1D * BASIS_P_1D * BASIS_P_1D, d_U, r_U); in Interp()
49 InterpTensor3d<BASIS_NUM_COMP, BASIS_P_1D, BASIS_Q_1D, BASIS_T_1D>(data, r_U, s_B, r_V); in Interp()
67 CeedScalar r_U[BASIS_NUM_COMP * (BASIS_DIM > 2 ? BASIS_P_1D : 1)]; in InterpCollocated() local
72 …tStrided1d<BASIS_NUM_COMP, BASIS_P_1D>(data, elem, 1, BASIS_P_1D * num_elem, BASIS_P_1D, d_U, r_U); in InterpCollocated()
73 …tStrided1d<BASIS_NUM_COMP, BASIS_Q_1D>(data, elem, 1, BASIS_Q_1D * num_elem, BASIS_Q_1D, r_U, d_V); in InterpCollocated()
[all …]
|
| H A D | cuda-shared-basis-tensor-templates.h | 52 inline __device__ void Interp1d(SharedData_Cuda &data, const CeedScalar *__restrict__ r_U, const Ce… 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_Cuda &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 …erpCollocatedNodes1d(SharedData_Cuda &data, const CeedScalar *__restrict__ r_U, const CeedScalar *… in InterpCollocatedNodes1d() argument
76 r_V[comp] = r_U[comp]; in InterpCollocatedNodes1d()
84 …oseCollocatedNodes1d(SharedData_Cuda &data, const CeedScalar *__restrict__ r_U, const CeedScalar *… in InterpTransposeCollocatedNodes1d() argument
87 r_V[comp] = r_U[comp]; in InterpTransposeCollocatedNodes1d()
95 inline __device__ void Grad1d(SharedData_Cuda &data, const CeedScalar *__restrict__ r_U, const Ceed… in Grad1d() argument
98 ContractX1d<NUM_COMP, P_1D, Q_1D>(data, &r_U[comp], c_G, &r_V[comp]); in Grad1d()
[all …]
|
| H A D | cuda-shared-basis-tensor-flattened-templates.h | 131 inline __device__ void InterpTensor2dFlattened(SharedData_Cuda &data, CeedScalar *__restrict__ r_U,… in InterpTensor2dFlattened() argument
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()
142 if (P_1D != T_1D) QPack2d<NUM_COMP, P_1D, T_1D>(data, t_id_x, t_id_y, r_U); in InterpTensor2dFlattened()
150 …TransposeTensor2dFlattened(SharedData_Cuda &data, CeedScalar *__restrict__ r_U, const CeedScalar *… in InterpTransposeTensor2dFlattened() argument
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()
168 …CollocatedNodes2dFlattened(SharedData_Cuda &data, CeedScalar *__restrict__ r_U, const CeedScalar *… in InterpTensorCollocatedNodes2dFlattened() argument
172 if (P_1D != T_1D) QUnpack2d<NUM_COMP, P_1D, T_1D>(data, t_id_x, t_id_y, r_U); in InterpTensorCollocatedNodes2dFlattened()
174 r_V[comp] = r_U[comp]; in InterpTensorCollocatedNodes2dFlattened()
[all …]
|
| H A D | cuda-shared-basis-nontensor.h | 28 CeedScalar r_U[BASIS_NUM_COMP]; in Interp() local
38 …eadElementStrided1d<BASIS_NUM_COMP, BASIS_P>(data, elem, 1, BASIS_P * num_elem, BASIS_P, d_U, r_U); in Interp()
39 InterpNonTensor<BASIS_NUM_COMP, BASIS_P, BASIS_Q, BASIS_T_1D>(data, r_U, s_B, r_V); in Interp()
55 CeedScalar r_U[BASIS_NUM_COMP]; in InterpTranspose() local
65 …eadElementStrided1d<BASIS_NUM_COMP, BASIS_Q>(data, elem, 1, BASIS_Q * num_elem, BASIS_Q, d_U, r_U); in InterpTranspose()
66 InterpTransposeNonTensor<BASIS_NUM_COMP, BASIS_P, BASIS_Q, BASIS_T_1D>(data, r_U, s_B, r_V); in InterpTranspose()
82 CeedScalar r_U[BASIS_NUM_COMP]; in InterpTransposeAdd() local
92 …eadElementStrided1d<BASIS_NUM_COMP, BASIS_Q>(data, elem, 1, BASIS_Q * num_elem, BASIS_Q, d_U, r_U); in InterpTransposeAdd()
93 InterpTransposeNonTensor<BASIS_NUM_COMP, BASIS_P, BASIS_Q, BASIS_T_1D>(data, r_U, s_B, r_V); in InterpTransposeAdd()
111 CeedScalar r_U[BASIS_NUM_COMP]; in Grad() local
[all …]
|
| H A D | cuda-shared-basis-nontensor-templates.h | 47 inline __device__ void InterpNonTensor(SharedData_Cuda &data, const CeedScalar *__restrict__ r_U, c… in InterpNonTensor() argument
50 Contract1d<NUM_COMP, P, Q>(data, &r_U[comp], c_B, &r_V[comp]); in InterpNonTensor()
58 …rpTransposeNonTensor(SharedData_Cuda &data, const CeedScalar *__restrict__ r_U, const CeedScalar *… in InterpTransposeNonTensor() argument
62 ContractTranspose1d<NUM_COMP, P, Q>(data, &r_U[comp], c_B, &r_V[comp]); in InterpTransposeNonTensor()
70 inline __device__ void GradNonTensor(SharedData_Cuda &data, const CeedScalar *__restrict__ r_U, con… in GradNonTensor() argument
73 Contract1d<NUM_COMP, P, Q>(data, &r_U[comp], &c_G[dim * P * Q], &r_V[comp + dim * NUM_COMP]); in GradNonTensor()
82 …adTransposeNonTensor(SharedData_Cuda &data, const CeedScalar *__restrict__ r_U, const CeedScalar *… in GradTransposeNonTensor() argument
87 …ContractTranspose1d<NUM_COMP, P, Q>(data, &r_U[comp + dim * NUM_COMP], &c_G[dim * P * Q], &r_V[com… in GradTransposeNonTensor()
|
| H A D | cuda-shared-basis-tensor-at-points.h | 35 CeedScalar r_U[BASIS_NUM_COMP * (BASIS_DIM > 2 ? BASIS_P_1D : 1)]; in InterpAtPoints() local
48 …tStrided1d<BASIS_NUM_COMP, BASIS_P_1D>(data, elem, 1, BASIS_P_1D * num_elem, BASIS_P_1D, d_U, r_U); in InterpAtPoints()
49 Interp1d<BASIS_NUM_COMP, BASIS_P_1D, BASIS_Q_1D, BASIS_T_1D>(data, r_U, s_B, r_C); in InterpAtPoints()
51 … BASIS_P_1D>(data, elem, 1, BASIS_P_1D * BASIS_P_1D * num_elem, BASIS_P_1D * BASIS_P_1D, d_U, r_U); in InterpAtPoints()
52 InterpTensor2d<BASIS_NUM_COMP, BASIS_P_1D, BASIS_Q_1D, BASIS_T_1D>(data, r_U, s_B, r_C); in InterpAtPoints()
55 … BASIS_P_1D * BASIS_P_1D * BASIS_P_1D, d_U, r_U); in InterpAtPoints()
56 InterpTensor3d<BASIS_NUM_COMP, BASIS_P_1D, BASIS_Q_1D, BASIS_T_1D>(data, r_U, s_B, r_C); in InterpAtPoints()
91 CeedScalar r_U[BASIS_NUM_COMP]; in InterpTransposeAtPoints() local
123 …M_PTS>(data, elem, i, points_per_elem[elem], 1, num_elem * BASIS_NUM_PTS, BASIS_NUM_PTS, d_U, r_U); in InterpTransposeAtPoints()
125 …TransposeAtPoints1d<BASIS_NUM_COMP, BASIS_NUM_PTS, BASIS_P_1D, BASIS_Q_1D>(data, i, r_U, r_X, r_C); in InterpTransposeAtPoints()
[all …]
|
| H A D | cuda-shared-basis-tensor-at-points-templates.h | 65 …s1d(SharedData_Cuda &data, const CeedInt p, const CeedScalar *__restrict__ r_U, const CeedScalar *… in InterpTransposeAtPoints1d() argument
77 ….slice[comp * Q_1D + (i + data.t_id_x) % Q_1D], chebyshev_x[(i + data.t_id_x) % Q_1D] * r_U[comp]); in InterpTransposeAtPoints1d()
112 …s1d(SharedData_Cuda &data, const CeedInt p, const CeedScalar *__restrict__ r_U, const CeedScalar *… in GradTransposeAtPoints1d() argument
124 ….slice[comp * Q_1D + (i + data.t_id_x) % Q_1D], chebyshev_x[(i + data.t_id_x) % Q_1D] * r_U[comp]); in GradTransposeAtPoints1d()
172 …s2d(SharedData_Cuda &data, const CeedInt p, const CeedScalar *__restrict__ r_U, const CeedScalar *… in InterpTransposeAtPoints2d() argument
183 const CeedScalar r_u = p < NUM_POINTS ? r_U[comp] : 0.0; in InterpTransposeAtPoints2d()
245 …s2d(SharedData_Cuda &data, const CeedInt p, const CeedScalar *__restrict__ r_U, const CeedScalar *… in GradTransposeAtPoints2d() argument
258 const CeedScalar r_u = p < NUM_POINTS ? r_U[comp + dim * NUM_COMP] : 0.0; in GradTransposeAtPoints2d()
328 …s3d(SharedData_Cuda &data, const CeedInt p, const CeedScalar *__restrict__ r_U, const CeedScalar *… in InterpTransposeAtPoints3d() argument
344 const CeedScalar r_u = p < NUM_POINTS ? r_U[comp] : 0.0; in InterpTransposeAtPoints3d()
[all …]
|
| H A D | cuda-gen-templates.h | 479 …e3d(SharedData_Cuda &data, const CeedInt q, const CeedScalar *__restrict__ r_U, const CeedScalar *… in GradColloSlice3d() argument
484 data.slice[data.t_id_x + data.t_id_y * T_1D] = r_U[q + comp * Q_1D]; in GradColloSlice3d()
499 r_V[comp + 2 * NUM_COMP] += c_G[i + q * Q_1D] * r_U[i + comp * Q_1D]; in GradColloSlice3d()
509 …e3d(SharedData_Cuda &data, const CeedInt q, const CeedScalar *__restrict__ r_U, const CeedScalar *… in GradColloSliceTranspose3d() argument
514 data.slice[data.t_id_x + data.t_id_y * T_1D] = r_U[comp + 0 * NUM_COMP]; in GradColloSliceTranspose3d()
522 data.slice[data.t_id_x + data.t_id_y * T_1D] = r_U[comp + 1 * NUM_COMP]; in GradColloSliceTranspose3d()
529 r_V[i + comp * Q_1D] += c_G[i + q * Q_1D] * r_U[comp + 2 * NUM_COMP]; in GradColloSliceTranspose3d()
|
| /libCEED/include/ceed/jit-source/hip/ |
| H A D | hip-shared-basis-tensor.h | 29 CeedScalar r_U[BASIS_NUM_COMP * (BASIS_DIM > 2 ? BASIS_P_1D : 1)]; in __launch_bounds__() local
40 …tStrided1d<BASIS_NUM_COMP, BASIS_P_1D>(data, elem, 1, BASIS_P_1D * num_elem, BASIS_P_1D, d_U, r_U); in __launch_bounds__()
41 Interp1d<BASIS_NUM_COMP, BASIS_P_1D, BASIS_Q_1D, BASIS_T_1D>(data, r_U, s_B, r_V); in __launch_bounds__()
44 … BASIS_P_1D>(data, elem, 1, BASIS_P_1D * BASIS_P_1D * num_elem, BASIS_P_1D * BASIS_P_1D, d_U, r_U); in __launch_bounds__()
45 InterpTensor2d<BASIS_NUM_COMP, BASIS_P_1D, BASIS_Q_1D, BASIS_T_1D>(data, r_U, s_B, r_V); in __launch_bounds__()
49 … BASIS_P_1D * BASIS_P_1D * BASIS_P_1D, d_U, r_U); in __launch_bounds__()
50 InterpTensor3d<BASIS_NUM_COMP, BASIS_P_1D, BASIS_Q_1D, BASIS_T_1D>(data, r_U, s_B, r_V); in __launch_bounds__()
68 CeedScalar r_U[BASIS_NUM_COMP * (BASIS_DIM > 2 ? BASIS_P_1D : 1)]; in __launch_bounds__() local
73 …tStrided1d<BASIS_NUM_COMP, BASIS_P_1D>(data, elem, 1, BASIS_P_1D * num_elem, BASIS_P_1D, d_U, r_U); in __launch_bounds__()
74 …tStrided1d<BASIS_NUM_COMP, BASIS_Q_1D>(data, elem, 1, BASIS_Q_1D * num_elem, BASIS_Q_1D, r_U, d_V); in __launch_bounds__()
[all …]
|
| H A D | hip-shared-basis-tensor-templates.h | 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 …terpCollocatedNodes1d(SharedData_Hip &data, const CeedScalar *__restrict__ r_U, const CeedScalar *… in InterpCollocatedNodes1d() argument
76 r_V[comp] = r_U[comp]; in InterpCollocatedNodes1d()
84 …poseCollocatedNodes1d(SharedData_Hip &data, const CeedScalar *__restrict__ r_U, const CeedScalar *… in InterpTransposeCollocatedNodes1d() argument
87 r_V[comp] = r_U[comp]; in InterpTransposeCollocatedNodes1d()
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()
[all …]
|
| H A D | hip-shared-basis-tensor-flattened-templates.h | 131 inline __device__ void InterpTensor2dFlattened(SharedData_Hip &data, CeedScalar *__restrict__ r_U, … in InterpTensor2dFlattened() argument
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()
142 if (P_1D != T_1D) QPack2d<NUM_COMP, P_1D, T_1D>(data, t_id_x, t_id_y, r_U); in InterpTensor2dFlattened()
150 …pTransposeTensor2dFlattened(SharedData_Hip &data, CeedScalar *__restrict__ r_U, const CeedScalar *… in InterpTransposeTensor2dFlattened() argument
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()
168 …rCollocatedNodes2dFlattened(SharedData_Hip &data, CeedScalar *__restrict__ r_U, const CeedScalar *… in InterpTensorCollocatedNodes2dFlattened() argument
172 if (P_1D != T_1D) QUnpack2d<NUM_COMP, P_1D, T_1D>(data, t_id_x, t_id_y, r_U); in InterpTensorCollocatedNodes2dFlattened()
174 r_V[comp] = r_U[comp]; in InterpTensorCollocatedNodes2dFlattened()
[all …]
|
| H A D | hip-shared-basis-nontensor.h | 29 CeedScalar r_U[BASIS_NUM_COMP]; in __launch_bounds__() local
39 …eadElementStrided1d<BASIS_NUM_COMP, BASIS_P>(data, elem, 1, BASIS_P * num_elem, BASIS_P, d_U, r_U); in __launch_bounds__()
40 InterpNonTensor<BASIS_NUM_COMP, BASIS_P, BASIS_Q, BASIS_T_1D>(data, r_U, s_B, r_V); in __launch_bounds__()
56 CeedScalar r_U[BASIS_NUM_COMP]; in __launch_bounds__() local
66 …eadElementStrided1d<BASIS_NUM_COMP, BASIS_Q>(data, elem, 1, BASIS_Q * num_elem, BASIS_Q, d_U, r_U); in __launch_bounds__()
67 InterpTransposeNonTensor<BASIS_NUM_COMP, BASIS_P, BASIS_Q, BASIS_T_1D>(data, r_U, s_B, r_V); in __launch_bounds__()
83 CeedScalar r_U[BASIS_NUM_COMP]; in __launch_bounds__() local
93 …eadElementStrided1d<BASIS_NUM_COMP, BASIS_Q>(data, elem, 1, BASIS_Q * num_elem, BASIS_Q, d_U, r_U); in __launch_bounds__()
94 InterpTransposeNonTensor<BASIS_NUM_COMP, BASIS_P, BASIS_Q, BASIS_T_1D>(data, r_U, s_B, r_V); in __launch_bounds__()
113 CeedScalar r_U[BASIS_NUM_COMP]; in __launch_bounds__() local
[all …]
|
| H A D | hip-shared-basis-nontensor-templates.h | 47 inline __device__ void InterpNonTensor(SharedData_Hip &data, const CeedScalar *__restrict__ r_U, co… in InterpNonTensor() argument
50 Contract1d<NUM_COMP, P, Q>(data, &r_U[comp], c_B, &r_V[comp]); in InterpNonTensor()
58 …erpTransposeNonTensor(SharedData_Hip &data, const CeedScalar *__restrict__ r_U, const CeedScalar *… in InterpTransposeNonTensor() argument
62 ContractTranspose1d<NUM_COMP, P, Q>(data, &r_U[comp], c_B, &r_V[comp]); in InterpTransposeNonTensor()
70 inline __device__ void GradNonTensor(SharedData_Hip &data, const CeedScalar *__restrict__ r_U, cons… in GradNonTensor() argument
73 Contract1d<NUM_COMP, P, Q>(data, &r_U[comp], &c_G[dim * P * Q], &r_V[comp + dim * NUM_COMP]); in GradNonTensor()
82 …radTransposeNonTensor(SharedData_Hip &data, const CeedScalar *__restrict__ r_U, const CeedScalar *… in GradTransposeNonTensor() argument
87 …ContractTranspose1d<NUM_COMP, P, Q>(data, &r_U[comp + dim * NUM_COMP], &c_G[dim * P * Q], &r_V[com… in GradTransposeNonTensor()
|
| H A D | hip-shared-basis-tensor-at-points.h | 36 CeedScalar r_U[BASIS_NUM_COMP * (BASIS_DIM > 2 ? BASIS_P_1D : 1)]; in __launch_bounds__() local
49 …tStrided1d<BASIS_NUM_COMP, BASIS_P_1D>(data, elem, 1, BASIS_P_1D * num_elem, BASIS_P_1D, d_U, r_U); in __launch_bounds__()
50 Interp1d<BASIS_NUM_COMP, BASIS_P_1D, BASIS_Q_1D, BASIS_T_1D>(data, r_U, s_B, r_C); in __launch_bounds__()
52 … BASIS_P_1D>(data, elem, 1, BASIS_P_1D * BASIS_P_1D * num_elem, BASIS_P_1D * BASIS_P_1D, d_U, r_U); in __launch_bounds__()
53 InterpTensor2d<BASIS_NUM_COMP, BASIS_P_1D, BASIS_Q_1D, BASIS_T_1D>(data, r_U, s_B, r_C); in __launch_bounds__()
56 … BASIS_P_1D * BASIS_P_1D * BASIS_P_1D, d_U, r_U); in __launch_bounds__()
57 InterpTensor3d<BASIS_NUM_COMP, BASIS_P_1D, BASIS_Q_1D, BASIS_T_1D>(data, r_U, s_B, r_C); in __launch_bounds__()
92 CeedScalar r_U[BASIS_NUM_COMP]; in __launch_bounds__() local
124 …M_PTS>(data, elem, i, points_per_elem[elem], 1, num_elem * BASIS_NUM_PTS, BASIS_NUM_PTS, d_U, r_U); in __launch_bounds__()
126 …TransposeAtPoints1d<BASIS_NUM_COMP, BASIS_NUM_PTS, BASIS_P_1D, BASIS_Q_1D>(data, i, r_U, r_X, r_C); in __launch_bounds__()
[all …]
|
| H A D | hip-shared-basis-tensor-at-points-templates.h | 66 …nterpTransposeAtPoints1d(SharedData_Hip &data, const CeedInt p, const CeedScalar *__restrict__ r_U, in InterpTransposeAtPoints1d() argument
78 ….slice[comp * Q_1D + (i + data.t_id_x) % Q_1D], chebyshev_x[(i + data.t_id_x) % Q_1D] * r_U[comp]); in InterpTransposeAtPoints1d()
113 … GradTransposeAtPoints1d(SharedData_Hip &data, const CeedInt p, const CeedScalar *__restrict__ r_U, in GradTransposeAtPoints1d() argument
125 ….slice[comp * Q_1D + (i + data.t_id_x) % Q_1D], chebyshev_x[(i + data.t_id_x) % Q_1D] * r_U[comp]); in GradTransposeAtPoints1d()
173 …nterpTransposeAtPoints2d(SharedData_Hip &data, const CeedInt p, const CeedScalar *__restrict__ r_U, in InterpTransposeAtPoints2d() argument
184 const CeedScalar r_u = p < NUM_POINTS ? r_U[comp] : 0.0; in InterpTransposeAtPoints2d()
246 … GradTransposeAtPoints2d(SharedData_Hip &data, const CeedInt p, const CeedScalar *__restrict__ r_U, in GradTransposeAtPoints2d() argument
259 const CeedScalar r_u = p < NUM_POINTS ? r_U[comp + dim * NUM_COMP] : 0.0; in GradTransposeAtPoints2d()
329 …nterpTransposeAtPoints3d(SharedData_Hip &data, const CeedInt p, const CeedScalar *__restrict__ r_U, in InterpTransposeAtPoints3d() argument
345 const CeedScalar r_u = p < NUM_POINTS ? r_U[comp] : 0.0; in InterpTransposeAtPoints3d()
[all …]
|
| H A D | hip-gen-templates.h | 476 …ce3d(SharedData_Hip &data, const CeedInt q, const CeedScalar *__restrict__ r_U, const CeedScalar *… in GradColloSlice3d() argument
481 data.slice[data.t_id_x + data.t_id_y * T_1D] = r_U[q + comp * Q_1D]; in GradColloSlice3d()
496 r_V[comp + 2 * NUM_COMP] += c_G[i + q * Q_1D] * r_U[i + comp * Q_1D]; in GradColloSlice3d()
506 …se3d(SharedData_Hip &data, const CeedInt q, const CeedScalar *__restrict__ r_U, const CeedScalar *… in GradColloSliceTranspose3d() argument
512 data.slice[data.t_id_x + data.t_id_y * T_1D] = r_U[comp + 0 * NUM_COMP]; in GradColloSliceTranspose3d()
519 data.slice[data.t_id_x + data.t_id_y * T_1D] = r_U[comp + 1 * NUM_COMP]; in GradColloSliceTranspose3d()
526 r_V[i + comp * Q_1D] += c_G[i + q * Q_1D] * r_U[comp + 2 * NUM_COMP]; in GradColloSliceTranspose3d()
|