| /libCEED/include/ceed/jit-source/cuda/ |
| H A D | cuda-shared-basis-tensor-flattened-templates.h | 20 … ContractX2dFlattened(SharedData_Cuda &data, const int t_id_x, const int t_id_y, const CeedScalar … 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 … ContractY2dFlattened(SharedData_Cuda &data, const int t_id_x, const int t_id_y, const CeedScalar … 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 …TransposeY2dFlattened(SharedData_Cuda &data, const int t_id_x, const int t_id_y, const CeedScalar … in ContractTransposeY2dFlattened() argument
57 data.slice[t_id_x + t_id_y * T_1D] = *U; in ContractTransposeY2dFlattened()
[all …]
|
| H A D | cuda-gen-templates.h | 186 if (data.t_id_x == target_node_x && data.t_id_y == target_node_y) { in SetEVecStandard2d_Single()
197 if (data.t_id_x < P_1D && data.t_id_y < P_1D) { in ReadLVecStandard2d()
198 const CeedInt node = data.t_id_x + data.t_id_y * P_1D; in ReadLVecStandard2d()
211 if (data.t_id_x < P_1D && data.t_id_y < P_1D) { in ReadLVecStrided2d()
212 const CeedInt node = data.t_id_x + data.t_id_y * P_1D; in ReadLVecStrided2d()
225 if (data.t_id_x < P_1D && data.t_id_y < P_1D) { in WriteLVecStandard2d()
226 const CeedInt node = data.t_id_x + data.t_id_y * P_1D; in WriteLVecStandard2d()
241 if (data.t_id_x == target_node_x && data.t_id_y == target_node_y) { in WriteLVecStandard2d_Single()
242 const CeedInt node = data.t_id_x + data.t_id_y * P_1D; in WriteLVecStandard2d_Single()
261 if (data.t_id_x < P_1D && data.t_id_y < P_1D) { in WriteLVecStandard2d_Assembly()
[all …]
|
| H A D | cuda-shared-basis-tensor-at-points-templates.h | 150 …if (data.t_id_x < Q_1D && data.t_id_y < Q_1D) data.slice[data.t_id_x + data.t_id_y * Q_1D] = r_C[c… in InterpAtPoints2d()
179 … if (data.t_id_x < Q_1D && data.t_id_y < Q_1D) data.slice[data.t_id_x + data.t_id_y * Q_1D] = 0.0; in InterpTransposeAtPoints2d()
192 const CeedInt ii = (i + data.t_id_y) % Q_1D; in InterpTransposeAtPoints2d()
197 …if (data.t_id_x < Q_1D && data.t_id_y < Q_1D) atomicAdd_block(&data.slice[jj + ii * Q_1D], chebysh… in InterpTransposeAtPoints2d()
202 …if (data.t_id_x < Q_1D && data.t_id_y < Q_1D) r_C[comp] += data.slice[data.t_id_x + data.t_id_y * … in InterpTransposeAtPoints2d()
219 …if (data.t_id_x < Q_1D && data.t_id_y < Q_1D) data.slice[data.t_id_x + data.t_id_y * Q_1D] = r_C[c… in GradAtPoints2d()
252 … if (data.t_id_x < Q_1D && data.t_id_y < Q_1D) data.slice[data.t_id_x + data.t_id_y * Q_1D] = 0.0; in GradTransposeAtPoints2d()
268 const CeedInt ii = (i + data.t_id_y) % Q_1D; in GradTransposeAtPoints2d()
273 …if (data.t_id_x < Q_1D && data.t_id_y < Q_1D) atomicAdd_block(&data.slice[jj + ii * Q_1D], chebysh… in GradTransposeAtPoints2d()
279 …if (data.t_id_x < Q_1D && data.t_id_y < Q_1D) r_C[comp] += data.slice[data.t_id_x + data.t_id_y * … in GradTransposeAtPoints2d()
[all …]
|
| H A D | cuda-shared-basis-tensor-templates.h | 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()
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()
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()
179 data.slice[data.t_id_x + data.t_id_y * T_1D] = *U; in ContractTransposeX2d()
[all …]
|
| H A D | cuda-shared-basis-read-write-templates.h | 79 if (data.t_id_x < P_1D && data.t_id_y < P_1D) { in ReadElementStrided2d()
80 const CeedInt node = data.t_id_x + data.t_id_y * P_1D; in ReadElementStrided2d()
95 if (data.t_id_x < P_1D && data.t_id_y < P_1D) { in WriteElementStrided2d()
96 const CeedInt node = data.t_id_x + data.t_id_y * P_1D; in WriteElementStrided2d()
108 if (data.t_id_x < P_1D && data.t_id_y < P_1D) { in SumElementStrided2d()
109 const CeedInt node = data.t_id_x + data.t_id_y * P_1D; in SumElementStrided2d()
128 if (data.t_id_x < P_1D && data.t_id_y < P_1D) { in ReadElementStrided3d()
130 const CeedInt node = data.t_id_x + data.t_id_y * P_1D + z * P_1D * P_1D; in ReadElementStrided3d()
146 if (data.t_id_x < P_1D && data.t_id_y < P_1D) { in WriteElementStrided3d()
148 const CeedInt node = data.t_id_x + data.t_id_y * P_1D + z * P_1D * P_1D; in WriteElementStrided3d()
[all …]
|
| H A D | cuda-shared-basis-nontensor.h | 23 data.t_id_y = threadIdx.y; in Interp()
50 data.t_id_y = threadIdx.y; in InterpTranspose()
77 data.t_id_y = threadIdx.y; in InterpTransposeAdd()
106 data.t_id_y = threadIdx.y; in Grad()
133 data.t_id_y = threadIdx.y; in GradTranspose()
160 data.t_id_y = threadIdx.y; in GradTransposeAdd()
189 data.t_id_y = threadIdx.y; in Weight()
|
| H A D | cuda-types.h | 35 CeedInt t_id_y; member
|
| H A D | cuda-shared-basis-tensor.h | 23 data.t_id_y = threadIdx.y; in Interp()
62 data.t_id_y = threadIdx.y; in InterpCollocated()
92 data.t_id_y = threadIdx.y; in InterpTranspose()
131 data.t_id_y = threadIdx.y; in InterpCollocatedTranspose()
161 data.t_id_y = threadIdx.y; in InterpTransposeAdd()
200 data.t_id_y = threadIdx.y; in InterpCollocatedTransposeAdd()
233 data.t_id_y = threadIdx.y; in Grad()
276 data.t_id_y = threadIdx.y; in GradCollocated()
316 data.t_id_y = threadIdx.y; in GradTranspose()
359 data.t_id_y = threadIdx.y; in GradCollocatedTranspose()
[all …]
|
| H A D | cuda-shared-basis-tensor-at-points.h | 29 data.t_id_y = threadIdx.y; in InterpAtPoints()
85 data.t_id_y = threadIdx.y; in InterpTransposeAtPoints()
155 data.t_id_y = threadIdx.y; in InterpTransposeAddAtPoints()
216 data.t_id_y = threadIdx.y; in GradAtPoints()
272 data.t_id_y = threadIdx.y; in GradTransposeAtPoints()
343 data.t_id_y = threadIdx.y; in GradTransposeAddAtPoints()
|
| /libCEED/include/ceed/jit-source/hip/ |
| H A D | hip-shared-basis-tensor-flattened-templates.h | 20 …d ContractX2dFlattened(SharedData_Hip &data, const int t_id_x, const int t_id_y, const CeedScalar … 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 …d ContractY2dFlattened(SharedData_Hip &data, const int t_id_x, const int t_id_y, const CeedScalar … 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 …tTransposeY2dFlattened(SharedData_Hip &data, const int t_id_x, const int t_id_y, const CeedScalar … in ContractTransposeY2dFlattened() argument
57 data.slice[t_id_x + t_id_y * T_1D] = *U; in ContractTransposeY2dFlattened()
[all …]
|
| H A D | hip-gen-templates.h | 185 if (data.t_id_x == target_node_x && data.t_id_y == target_node_y) { in SetEVecStandard2d_Single()
196 if (data.t_id_x < P_1D && data.t_id_y < P_1D) { in ReadLVecStandard2d()
197 const CeedInt node = data.t_id_x + data.t_id_y * P_1D; in ReadLVecStandard2d()
209 if (data.t_id_x < P_1D && data.t_id_y < P_1D) { in ReadLVecStrided2d()
210 const CeedInt node = data.t_id_x + data.t_id_y * P_1D; in ReadLVecStrided2d()
223 if (data.t_id_x < P_1D && data.t_id_y < P_1D) { in WriteLVecStandard2d()
224 const CeedInt node = data.t_id_x + data.t_id_y * P_1D; in WriteLVecStandard2d()
239 if (data.t_id_x == target_node_x && data.t_id_y == target_node_y) { in WriteLVecStandard2d_Single()
240 const CeedInt node = data.t_id_x + data.t_id_y * P_1D; in WriteLVecStandard2d_Single()
259 if (data.t_id_x < P_1D && data.t_id_y < P_1D) { in WriteLVecStandard2d_Assembly()
[all …]
|
| H A D | hip-shared-basis-tensor-at-points-templates.h | 151 …if (data.t_id_x < Q_1D && data.t_id_y < Q_1D) data.slice[data.t_id_x + data.t_id_y * Q_1D] = r_C[c… in InterpAtPoints2d()
180 … if (data.t_id_x < Q_1D && data.t_id_y < Q_1D) data.slice[data.t_id_x + data.t_id_y * Q_1D] = 0.0; in InterpTransposeAtPoints2d()
193 const CeedInt ii = (i + data.t_id_y) % Q_1D; in InterpTransposeAtPoints2d()
198 …if (data.t_id_x < Q_1D && data.t_id_y < Q_1D) atomicAdd(&data.slice[jj + ii * Q_1D], chebyshev_x[j… in InterpTransposeAtPoints2d()
203 …if (data.t_id_x < Q_1D && data.t_id_y < Q_1D) r_C[comp] += data.slice[data.t_id_x + data.t_id_y * … in InterpTransposeAtPoints2d()
220 …if (data.t_id_x < Q_1D && data.t_id_y < Q_1D) data.slice[data.t_id_x + data.t_id_y * Q_1D] = r_C[c… in GradAtPoints2d()
253 … if (data.t_id_x < Q_1D && data.t_id_y < Q_1D) data.slice[data.t_id_x + data.t_id_y * Q_1D] = 0.0; in GradTransposeAtPoints2d()
269 const CeedInt ii = (i + data.t_id_y) % Q_1D; in GradTransposeAtPoints2d()
274 …if (data.t_id_x < Q_1D && data.t_id_y < Q_1D) atomicAdd(&data.slice[jj + ii * Q_1D], chebyshev_x[j… in GradTransposeAtPoints2d()
280 …if (data.t_id_x < Q_1D && data.t_id_y < Q_1D) r_C[comp] += data.slice[data.t_id_x + data.t_id_y * … in GradTransposeAtPoints2d()
[all …]
|
| H A D | hip-shared-basis-tensor-templates.h | 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()
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()
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()
179 data.slice[data.t_id_x + data.t_id_y * T_1D] = *U; in ContractTransposeX2d()
[all …]
|
| H A D | hip-shared-basis-read-write-templates.h | 79 if (data.t_id_x < P_1D && data.t_id_y < P_1D) { in ReadElementStrided2d()
80 const CeedInt node = data.t_id_x + data.t_id_y * P_1D; in ReadElementStrided2d()
95 if (data.t_id_x < P_1D && data.t_id_y < P_1D) { in WriteElementStrided2d()
96 const CeedInt node = data.t_id_x + data.t_id_y * P_1D; in WriteElementStrided2d()
108 if (data.t_id_x < P_1D && data.t_id_y < P_1D) { in SumElementStrided2d()
109 const CeedInt node = data.t_id_x + data.t_id_y * P_1D; in SumElementStrided2d()
128 if (data.t_id_x < P_1D && data.t_id_y < P_1D) { in ReadElementStrided3d()
130 const CeedInt node = data.t_id_x + data.t_id_y * P_1D + z * P_1D * P_1D; in ReadElementStrided3d()
146 if (data.t_id_x < P_1D && data.t_id_y < P_1D) { in WriteElementStrided3d()
148 const CeedInt node = data.t_id_x + data.t_id_y * P_1D + z * P_1D * P_1D; in WriteElementStrided3d()
[all …]
|
| H A D | hip-shared-basis-nontensor.h | 24 data.t_id_y = threadIdx.y; in __launch_bounds__()
51 data.t_id_y = threadIdx.y; in __launch_bounds__()
78 data.t_id_y = threadIdx.y; in __launch_bounds__()
108 data.t_id_y = threadIdx.y; in __launch_bounds__()
135 data.t_id_y = threadIdx.y; in __launch_bounds__()
162 data.t_id_y = threadIdx.y; in __launch_bounds__()
192 data.t_id_y = threadIdx.y; in __launch_bounds__()
|
| H A D | hip-types.h | 35 CeedInt t_id_y; member
|
| H A D | hip-shared-basis-tensor.h | 24 data.t_id_y = threadIdx.y; in __launch_bounds__()
63 data.t_id_y = threadIdx.y; in __launch_bounds__()
93 data.t_id_y = threadIdx.y; in __launch_bounds__()
132 data.t_id_y = threadIdx.y; in __launch_bounds__()
162 data.t_id_y = threadIdx.y; in __launch_bounds__()
202 data.t_id_y = threadIdx.y; in __launch_bounds__()
235 data.t_id_y = threadIdx.y; in __launch_bounds__()
279 data.t_id_y = threadIdx.y; in __launch_bounds__()
320 data.t_id_y = threadIdx.y; in __launch_bounds__()
364 data.t_id_y = threadIdx.y; in __launch_bounds__()
[all …]
|
| H A D | hip-shared-basis-tensor-at-points.h | 30 data.t_id_y = threadIdx.y; in __launch_bounds__()
86 data.t_id_y = threadIdx.y; in __launch_bounds__()
156 data.t_id_y = threadIdx.y; in __launch_bounds__()
218 data.t_id_y = threadIdx.y; in __launch_bounds__()
274 data.t_id_y = threadIdx.y; in __launch_bounds__()
345 data.t_id_y = threadIdx.y; in __launch_bounds__()
|