Home
last modified time | relevance | path

Searched refs:t_id_x (Results 1 – 20 of 20) sorted by relevance

/libCEED/include/ceed/jit-source/cuda/
H A Dcuda-shared-basis-tensor-flattened-templates.h20 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()
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_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()
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_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()
[all …]
H A Dcuda-shared-basis-tensor-at-points-templates.h52 if (data.t_id_x < Q_1D) data.slice[data.t_id_x] = r_C[comp]; in InterpAtPoints1d()
72 if (data.t_id_x < Q_1D) data.slice[data.t_id_x] = 0.0; in InterpTransposeAtPoints1d()
77 …atomicAdd_block(&data.slice[comp * Q_1D + (i + data.t_id_x) % Q_1D], chebyshev_x[(i + data.t_id_x)… in InterpTransposeAtPoints1d()
82 if (data.t_id_x < Q_1D) r_C[comp] += data.slice[data.t_id_x]; in InterpTransposeAtPoints1d()
99 if (data.t_id_x < Q_1D) data.slice[data.t_id_x] = r_C[comp]; in GradAtPoints1d()
119 if (data.t_id_x < Q_1D) data.slice[data.t_id_x] = 0.0; in GradTransposeAtPoints1d()
124 …atomicAdd_block(&data.slice[comp * Q_1D + (i + data.t_id_x) % Q_1D], chebyshev_x[(i + data.t_id_x)… in GradTransposeAtPoints1d()
129 if (data.t_id_x < Q_1D) r_C[comp] += data.slice[data.t_id_x]; in GradTransposeAtPoints1d()
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()
[all …]
H A Dcuda-gen-templates.h64 if (data.t_id_x == target_node) { in SetEVecStandard1d_Single()
75 if (data.t_id_x < P_1D) { in ReadLVecStandard1d()
76 const CeedInt node = data.t_id_x; in ReadLVecStandard1d()
89 if (data.t_id_x < P_1D) { in ReadLVecStrided1d()
90 const CeedInt node = data.t_id_x; in ReadLVecStrided1d()
103 if (data.t_id_x < P_1D) { in WriteLVecStandard1d()
104 const CeedInt node = data.t_id_x; in WriteLVecStandard1d()
118 if (data.t_id_x == target_node) { in WriteLVecStandard1d_Single()
135 if (data.t_id_x < P_1D) { in WriteLVecStandard1d_Assembly()
136 const CeedInt out_node = data.t_id_x; in WriteLVecStandard1d_Assembly()
[all …]
H A Dcuda-shared-basis-read-write-templates.h30 if (data.t_id_x < P_1D) { in ReadElementStrided1d()
31 const CeedInt node = data.t_id_x; in ReadElementStrided1d()
46 if (data.t_id_x < P_1D) { in WriteElementStrided1d()
47 const CeedInt node = data.t_id_x; in WriteElementStrided1d()
59 if (data.t_id_x < P_1D) { in SumElementStrided1d()
60 const CeedInt node = data.t_id_x; in SumElementStrided1d()
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()
[all …]
H A Dcuda-shared-basis-tensor-templates.h22 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()
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()
118 *w = (data.t_id_x < Q_1D) ? q_weight_1d[data.t_id_x] : 0.0; in Weight1d()
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()
[all …]
H A Dcuda-shared-basis-nontensor-templates.h17 data.slice[data.t_id_x] = *U; in Contract1d()
20 if (data.t_id_x < Q_1D) { in Contract1d()
22 *V += B[i + data.t_id_x * P_1D] * data.slice[i]; // Contract x direction in Contract1d()
33 data.slice[data.t_id_x] = *U; in ContractTranspose1d()
35 if (data.t_id_x < P_1D) { in ContractTranspose1d()
37 *V += B[data.t_id_x + i * P_1D] * data.slice[i]; // Contract x direction in ContractTranspose1d()
97 *w = (data.t_id_x < Q) ? q_weight[data.t_id_x] : 0.0; in WeightNonTensor()
H A Dcuda-shared-basis-nontensor.h22 data.t_id_x = threadIdx.x; in Interp()
49 data.t_id_x = threadIdx.x; in InterpTranspose()
76 data.t_id_x = threadIdx.x; in InterpTransposeAdd()
105 data.t_id_x = threadIdx.x; in Grad()
132 data.t_id_x = threadIdx.x; in GradTranspose()
159 data.t_id_x = threadIdx.x; in GradTransposeAdd()
188 data.t_id_x = threadIdx.x; in Weight()
H A Dcuda-types.h34 CeedInt t_id_x; member
H A Dcuda-shared-basis-tensor.h22 data.t_id_x = threadIdx.x; in Interp()
61 data.t_id_x = threadIdx.x; in InterpCollocated()
91 data.t_id_x = threadIdx.x; in InterpTranspose()
130 data.t_id_x = threadIdx.x; in InterpCollocatedTranspose()
160 data.t_id_x = threadIdx.x; in InterpTransposeAdd()
199 data.t_id_x = threadIdx.x; in InterpCollocatedTransposeAdd()
232 data.t_id_x = threadIdx.x; in Grad()
275 data.t_id_x = threadIdx.x; in GradCollocated()
315 data.t_id_x = threadIdx.x; in GradTranspose()
358 data.t_id_x = threadIdx.x; in GradCollocatedTranspose()
[all …]
H A Dcuda-shared-basis-tensor-at-points.h28 data.t_id_x = threadIdx.x; in InterpAtPoints()
84 data.t_id_x = threadIdx.x; in InterpTransposeAtPoints()
154 data.t_id_x = threadIdx.x; in InterpTransposeAddAtPoints()
215 data.t_id_x = threadIdx.x; in GradAtPoints()
271 data.t_id_x = threadIdx.x; in GradTransposeAtPoints()
342 data.t_id_x = threadIdx.x; in GradTransposeAddAtPoints()
/libCEED/include/ceed/jit-source/hip/
H A Dhip-shared-basis-tensor-flattened-templates.h20 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()
[all …]
H A Dhip-shared-basis-tensor-at-points-templates.h53 if (data.t_id_x < Q_1D) data.slice[data.t_id_x] = r_C[comp]; in InterpAtPoints1d()
73 if (data.t_id_x < Q_1D) data.slice[data.t_id_x] = 0.0; in InterpTransposeAtPoints1d()
78 …atomicAdd(&data.slice[comp * Q_1D + (i + data.t_id_x) % Q_1D], chebyshev_x[(i + data.t_id_x) % Q_1… in InterpTransposeAtPoints1d()
83 if (data.t_id_x < Q_1D) r_C[comp] += data.slice[data.t_id_x]; in InterpTransposeAtPoints1d()
100 if (data.t_id_x < Q_1D) data.slice[data.t_id_x] = r_C[comp]; in GradAtPoints1d()
120 if (data.t_id_x < Q_1D) data.slice[data.t_id_x] = 0.0; in GradTransposeAtPoints1d()
125 …atomicAdd(&data.slice[comp * Q_1D + (i + data.t_id_x) % Q_1D], chebyshev_x[(i + data.t_id_x) % Q_1… in GradTransposeAtPoints1d()
130 if (data.t_id_x < Q_1D) r_C[comp] += data.slice[data.t_id_x]; in GradTransposeAtPoints1d()
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()
[all …]
H A Dhip-gen-templates.h64 if (data.t_id_x == target_node) { in SetEVecStandard1d_Single()
75 if (data.t_id_x < P_1D) { in ReadLVecStandard1d()
76 const CeedInt node = data.t_id_x; in ReadLVecStandard1d()
88 if (data.t_id_x < P_1D) { in ReadLVecStrided1d()
89 const CeedInt node = data.t_id_x; in ReadLVecStrided1d()
102 if (data.t_id_x < P_1D) { in WriteLVecStandard1d()
103 const CeedInt node = data.t_id_x; in WriteLVecStandard1d()
117 if (data.t_id_x == target_node) { in WriteLVecStandard1d_Single()
134 if (data.t_id_x < P_1D) { in WriteLVecStandard1d_Assembly()
135 const CeedInt out_node = data.t_id_x; in WriteLVecStandard1d_Assembly()
[all …]
H A Dhip-shared-basis-read-write-templates.h30 if (data.t_id_x < P_1D) { in ReadElementStrided1d()
31 const CeedInt node = data.t_id_x; in ReadElementStrided1d()
46 if (data.t_id_x < P_1D) { in WriteElementStrided1d()
47 const CeedInt node = data.t_id_x; in WriteElementStrided1d()
59 if (data.t_id_x < P_1D) { in SumElementStrided1d()
60 const CeedInt node = data.t_id_x; in SumElementStrided1d()
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()
[all …]
H A Dhip-shared-basis-tensor-templates.h22 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()
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()
118 *w = (data.t_id_x < Q_1D) ? q_weight_1d[data.t_id_x] : 0.0; in Weight1d()
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()
[all …]
H A Dhip-shared-basis-nontensor-templates.h17 data.slice[data.t_id_x] = *U; in Contract1d()
20 if (data.t_id_x < Q_1D) { in Contract1d()
22 *V += B[i + data.t_id_x * P_1D] * data.slice[i]; // Contract x direction in Contract1d()
33 data.slice[data.t_id_x] = *U; in ContractTranspose1d()
35 if (data.t_id_x < P_1D) { in ContractTranspose1d()
37 *V += B[data.t_id_x + i * P_1D] * data.slice[i]; // Contract x direction in ContractTranspose1d()
97 *w = (data.t_id_x < Q) ? q_weight[data.t_id_x] : 0.0; in WeightNonTensor()
H A Dhip-shared-basis-nontensor.h23 data.t_id_x = threadIdx.x; in __launch_bounds__()
50 data.t_id_x = threadIdx.x; in __launch_bounds__()
77 data.t_id_x = threadIdx.x; in __launch_bounds__()
107 data.t_id_x = threadIdx.x; in __launch_bounds__()
134 data.t_id_x = threadIdx.x; in __launch_bounds__()
161 data.t_id_x = threadIdx.x; in __launch_bounds__()
191 data.t_id_x = threadIdx.x; in __launch_bounds__()
H A Dhip-types.h34 CeedInt t_id_x; member
H A Dhip-shared-basis-tensor.h23 data.t_id_x = threadIdx.x; in __launch_bounds__()
62 data.t_id_x = threadIdx.x; in __launch_bounds__()
92 data.t_id_x = threadIdx.x; in __launch_bounds__()
131 data.t_id_x = threadIdx.x; in __launch_bounds__()
161 data.t_id_x = threadIdx.x; in __launch_bounds__()
201 data.t_id_x = threadIdx.x; in __launch_bounds__()
234 data.t_id_x = threadIdx.x; in __launch_bounds__()
278 data.t_id_x = threadIdx.x; in __launch_bounds__()
319 data.t_id_x = threadIdx.x; in __launch_bounds__()
363 data.t_id_x = threadIdx.x; in __launch_bounds__()
[all …]
H A Dhip-shared-basis-tensor-at-points.h29 data.t_id_x = threadIdx.x; in __launch_bounds__()
85 data.t_id_x = threadIdx.x; in __launch_bounds__()
155 data.t_id_x = threadIdx.x; in __launch_bounds__()
217 data.t_id_x = threadIdx.x; in __launch_bounds__()
273 data.t_id_x = threadIdx.x; in __launch_bounds__()
344 data.t_id_x = threadIdx.x; in __launch_bounds__()