Lines Matching refs:data
16 inline __device__ void LoadMatrix(SharedData_Cuda &data, const CeedScalar *__restrict__ d_B, CeedSc… in LoadMatrix() argument
17 for (CeedInt i = data.t_id; i < P * Q; i += blockDim.x * blockDim.y * blockDim.z) B[i] = d_B[i]; in LoadMatrix()
28 inline __device__ void ReadPoint(SharedData_Cuda &data, const CeedInt elem, const CeedInt p, const … in ReadPoint() argument
41 inline __device__ void WritePoint(SharedData_Cuda &data, const CeedInt elem, const CeedInt p, const… in WritePoint() argument
60 inline __device__ void SetEVecStandard1d_Single(SharedData_Cuda &data, const CeedInt n, const CeedS… in SetEVecStandard1d_Single() argument
64 if (data.t_id_x == target_node) { in SetEVecStandard1d_Single()
73 inline __device__ void ReadLVecStandard1d(SharedData_Cuda &data, const CeedInt num_nodes, const Cee… in ReadLVecStandard1d() argument
75 if (data.t_id_x < P_1D) { in ReadLVecStandard1d()
76 const CeedInt node = data.t_id_x; in ReadLVecStandard1d()
87 inline __device__ void ReadLVecStrided1d(SharedData_Cuda &data, const CeedInt elem, const CeedScala… in ReadLVecStrided1d() argument
89 if (data.t_id_x < P_1D) { in ReadLVecStrided1d()
90 const CeedInt node = data.t_id_x; in ReadLVecStrided1d()
101 inline __device__ void WriteLVecStandard1d(SharedData_Cuda &data, const CeedInt num_nodes, const Ce… in WriteLVecStandard1d() argument
103 if (data.t_id_x < P_1D) { in WriteLVecStandard1d()
104 const CeedInt node = data.t_id_x; in WriteLVecStandard1d()
112 inline __device__ void WriteLVecStandard1d_Single(SharedData_Cuda &data, const CeedInt num_nodes, c… in WriteLVecStandard1d_Single() argument
118 if (data.t_id_x == target_node) { in WriteLVecStandard1d_Single()
129 inline __device__ void WriteLVecStandard1d_Assembly(SharedData_Cuda &data, const CeedInt num_nodes,… in WriteLVecStandard1d_Assembly() argument
135 if (data.t_id_x < P_1D) { in WriteLVecStandard1d_Assembly()
136 const CeedInt out_node = data.t_id_x; in WriteLVecStandard1d_Assembly()
148 inline __device__ void WriteLVecStandard1d_QFAssembly(SharedData_Cuda &data, const CeedInt num_elem… in WriteLVecStandard1d_QFAssembly() argument
150 if (data.t_id_x < Q_1D) { in WriteLVecStandard1d_QFAssembly()
151 const CeedInt ind = data.t_id_x + elem * Q_1D; in WriteLVecStandard1d_QFAssembly()
163 inline __device__ void WriteLVecStrided1d(SharedData_Cuda &data, const CeedInt elem, const CeedScal… in WriteLVecStrided1d() argument
165 if (data.t_id_x < P_1D) { in WriteLVecStrided1d()
166 const CeedInt node = data.t_id_x; in WriteLVecStrided1d()
181 inline __device__ void SetEVecStandard2d_Single(SharedData_Cuda &data, const CeedInt n, const CeedS… in SetEVecStandard2d_Single() argument
186 if (data.t_id_x == target_node_x && data.t_id_y == target_node_y) { in SetEVecStandard2d_Single()
195 inline __device__ void ReadLVecStandard2d(SharedData_Cuda &data, const CeedInt num_nodes, const Cee… in ReadLVecStandard2d() argument
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()
209 inline __device__ void ReadLVecStrided2d(SharedData_Cuda &data, const CeedInt elem, const CeedScala… in ReadLVecStrided2d() argument
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()
223 inline __device__ void WriteLVecStandard2d(SharedData_Cuda &data, const CeedInt num_nodes, const Ce… in WriteLVecStandard2d() argument
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()
234 inline __device__ void WriteLVecStandard2d_Single(SharedData_Cuda &data, const CeedInt num_nodes, c… in WriteLVecStandard2d_Single() argument
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()
253 inline __device__ void WriteLVecStandard2d_Assembly(SharedData_Cuda &data, const CeedInt num_nodes,… in WriteLVecStandard2d_Assembly() argument
261 if (data.t_id_x < P_1D && data.t_id_y < P_1D) { in WriteLVecStandard2d_Assembly()
263 const CeedInt out_node = data.t_id_x + data.t_id_y * P_1D; in WriteLVecStandard2d_Assembly()
277 inline __device__ void WriteLVecStandard2d_QFAssembly(SharedData_Cuda &data, const CeedInt num_elem… in WriteLVecStandard2d_QFAssembly() argument
279 if (data.t_id_x < Q_1D && data.t_id_y < Q_1D) { in WriteLVecStandard2d_QFAssembly()
280 const CeedInt ind = (data.t_id_x + data.t_id_y * Q_1D) + elem * Q_1D * Q_1D; in WriteLVecStandard2d_QFAssembly()
292 inline __device__ void WriteLVecStrided2d(SharedData_Cuda &data, const CeedInt elem, const CeedScal… in WriteLVecStrided2d() argument
294 if (data.t_id_x < P_1D && data.t_id_y < P_1D) { in WriteLVecStrided2d()
295 const CeedInt node = data.t_id_x + data.t_id_y * P_1D; in WriteLVecStrided2d()
310 inline __device__ void SetEVecStandard3d_Single(SharedData_Cuda &data, const CeedInt n, const CeedS… in SetEVecStandard3d_Single() argument
316 if (data.t_id_x == target_node_x && data.t_id_y == target_node_y) { in SetEVecStandard3d_Single()
325 inline __device__ void ReadLVecStandard3d(SharedData_Cuda &data, const CeedInt num_nodes, const Cee… in ReadLVecStandard3d() argument
327 if (data.t_id_x < P_1D && data.t_id_y < P_1D) { in ReadLVecStandard3d()
329 const CeedInt node = data.t_id_x + data.t_id_y * P_1D + z * P_1D * P_1D; in ReadLVecStandard3d()
341 inline __device__ void ReadLVecStrided3d(SharedData_Cuda &data, const CeedInt elem, const CeedScala… in ReadLVecStrided3d() argument
343 if (data.t_id_x < P_1D && data.t_id_y < P_1D) { in ReadLVecStrided3d()
345 const CeedInt node = data.t_id_x + data.t_id_y * P_1D + z * P_1D * P_1D; in ReadLVecStrided3d()
357 inline __device__ void ReadEVecSliceStandard3d(SharedData_Cuda &data, const CeedInt nquads, const C… in ReadEVecSliceStandard3d() argument
360 if (data.t_id_x < Q_1D && data.t_id_y < Q_1D) { in ReadEVecSliceStandard3d()
361 const CeedInt node = data.t_id_x + data.t_id_y * Q_1D + q * Q_1D * Q_1D; in ReadEVecSliceStandard3d()
372 inline __device__ void ReadEVecSliceStrided3d(SharedData_Cuda &data, const CeedInt elem, const Ceed… in ReadEVecSliceStrided3d() argument
374 if (data.t_id_x < Q_1D && data.t_id_y < Q_1D) { in ReadEVecSliceStrided3d()
375 const CeedInt node = data.t_id_x + data.t_id_y * Q_1D + q * Q_1D * Q_1D; in ReadEVecSliceStrided3d()
386 inline __device__ void WriteLVecStandard3d(SharedData_Cuda &data, const CeedInt num_nodes, const Ce… in WriteLVecStandard3d() argument
388 if (data.t_id_x < P_1D && data.t_id_y < P_1D) { in WriteLVecStandard3d()
390 const CeedInt node = data.t_id_x + data.t_id_y * P_1D + z * P_1D * P_1D; in WriteLVecStandard3d()
399 inline __device__ void WriteLVecStandard3d_Single(SharedData_Cuda &data, const CeedInt num_nodes, c… in WriteLVecStandard3d_Single() argument
407 if (data.t_id_x == target_node_x && data.t_id_y == target_node_y) { in WriteLVecStandard3d_Single()
408 const CeedInt node = data.t_id_x + data.t_id_y * P_1D + target_node_z * P_1D * P_1D; in WriteLVecStandard3d_Single()
419 inline __device__ void WriteLVecStandard3d_Assembly(SharedData_Cuda &data, const CeedInt num_nodes,… in WriteLVecStandard3d_Assembly() argument
428 if (data.t_id_x < P_1D && data.t_id_y < P_1D) { in WriteLVecStandard3d_Assembly()
431 const CeedInt out_node = data.t_id_x + data.t_id_y * P_1D + z * P_1D * P_1D; in WriteLVecStandard3d_Assembly()
446 inline __device__ void WriteLVecStandard3d_QFAssembly(SharedData_Cuda &data, const CeedInt num_elem… in WriteLVecStandard3d_QFAssembly() argument
448 if (data.t_id_x < Q_1D && data.t_id_y < Q_1D) { in WriteLVecStandard3d_QFAssembly()
450 …const CeedInt ind = (data.t_id_x + data.t_id_y * Q_1D + z * Q_1D * Q_1D) + elem * Q_1D * Q_1D * Q_… in WriteLVecStandard3d_QFAssembly()
463 inline __device__ void WriteLVecStrided3d(SharedData_Cuda &data, const CeedInt elem, const CeedScal… in WriteLVecStrided3d() argument
465 if (data.t_id_x < P_1D && data.t_id_y < P_1D) { in WriteLVecStrided3d()
467 const CeedInt node = data.t_id_x + data.t_id_y * P_1D + z * P_1D * P_1D; in WriteLVecStrided3d()
479 inline __device__ void GradColloSlice3d(SharedData_Cuda &data, const CeedInt q, const CeedScalar *_… in GradColloSlice3d() argument
481 if (data.t_id_x < Q_1D && data.t_id_y < Q_1D) { in GradColloSlice3d()
484 data.slice[data.t_id_x + data.t_id_y * T_1D] = r_U[q + comp * Q_1D]; in GradColloSlice3d()
489 … r_V[comp + 0 * NUM_COMP] += c_G[i + data.t_id_x * Q_1D] * data.slice[i + data.t_id_y * T_1D]; in GradColloSlice3d()
494 … r_V[comp + 1 * NUM_COMP] += c_G[i + data.t_id_y * Q_1D] * data.slice[data.t_id_x + i * T_1D]; in GradColloSlice3d()
509 inline __device__ void GradColloSliceTranspose3d(SharedData_Cuda &data, const CeedInt q, const Ceed… in GradColloSliceTranspose3d() argument
511 if (data.t_id_x < Q_1D && data.t_id_y < Q_1D) { in GradColloSliceTranspose3d()
514 data.slice[data.t_id_x + data.t_id_y * T_1D] = r_U[comp + 0 * NUM_COMP]; in GradColloSliceTranspose3d()
518 r_V[q + comp * Q_1D] += c_G[data.t_id_x + i * Q_1D] * data.slice[i + data.t_id_y * T_1D]; in GradColloSliceTranspose3d()
522 data.slice[data.t_id_x + data.t_id_y * T_1D] = r_U[comp + 1 * NUM_COMP]; in GradColloSliceTranspose3d()
525 r_V[q + comp * Q_1D] += c_G[data.t_id_y + i * Q_1D] * data.slice[data.t_id_x + i * T_1D]; in GradColloSliceTranspose3d()