Lines Matching refs:data
16 inline __device__ void LoadMatrix(SharedData_Hip &data, const CeedScalar *__restrict__ d_B, CeedSca… 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_Hip &data, const CeedInt elem, const CeedInt p, const C… in ReadPoint() argument
41 inline __device__ void WritePoint(SharedData_Hip &data, const CeedInt elem, const CeedInt p, const … in WritePoint() argument
60 inline __device__ void SetEVecStandard1d_Single(SharedData_Hip &data, const CeedInt n, const CeedSc… in SetEVecStandard1d_Single() argument
64 if (data.t_id_x == target_node) { in SetEVecStandard1d_Single()
73 inline __device__ void ReadLVecStandard1d(SharedData_Hip &data, const CeedInt num_nodes, const Ceed… 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_Hip &data, const CeedInt elem, const CeedScalar… in ReadLVecStrided1d() argument
88 if (data.t_id_x < P_1D) { in ReadLVecStrided1d()
89 const CeedInt node = data.t_id_x; in ReadLVecStrided1d()
100 inline __device__ void WriteLVecStandard1d(SharedData_Hip &data, const CeedInt num_nodes, const Cee… in WriteLVecStandard1d() argument
102 if (data.t_id_x < P_1D) { in WriteLVecStandard1d()
103 const CeedInt node = data.t_id_x; in WriteLVecStandard1d()
111 inline __device__ void WriteLVecStandard1d_Single(SharedData_Hip &data, const CeedInt num_nodes, co… in WriteLVecStandard1d_Single() argument
117 if (data.t_id_x == target_node) { in WriteLVecStandard1d_Single()
128 inline __device__ void WriteLVecStandard1d_Assembly(SharedData_Hip &data, const CeedInt num_nodes, … in WriteLVecStandard1d_Assembly() argument
134 if (data.t_id_x < P_1D) { in WriteLVecStandard1d_Assembly()
135 const CeedInt out_node = data.t_id_x; in WriteLVecStandard1d_Assembly()
147 inline __device__ void WriteLVecStandard1d_QFAssembly(SharedData_Hip &data, const CeedInt num_elem,… in WriteLVecStandard1d_QFAssembly() argument
149 if (data.t_id_x < Q_1D) { in WriteLVecStandard1d_QFAssembly()
150 const CeedInt ind = data.t_id_x + elem * Q_1D; in WriteLVecStandard1d_QFAssembly()
162 inline __device__ void WriteLVecStrided1d(SharedData_Hip &data, const CeedInt elem, const CeedScala… in WriteLVecStrided1d() argument
164 if (data.t_id_x < P_1D) { in WriteLVecStrided1d()
165 const CeedInt node = data.t_id_x; in WriteLVecStrided1d()
180 inline __device__ void SetEVecStandard2d_Single(SharedData_Hip &data, const CeedInt n, const CeedSc… in SetEVecStandard2d_Single() argument
185 if (data.t_id_x == target_node_x && data.t_id_y == target_node_y) { in SetEVecStandard2d_Single()
194 inline __device__ void ReadLVecStandard2d(SharedData_Hip &data, const CeedInt num_nodes, const Ceed… in ReadLVecStandard2d() argument
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()
208 inline __device__ void ReadLVecStrided2d(SharedData_Hip &data, const CeedInt elem, const CeedScalar… in ReadLVecStrided2d() argument
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()
221 inline __device__ void WriteLVecStandard2d(SharedData_Hip &data, const CeedInt num_nodes, const Cee… in WriteLVecStandard2d() argument
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()
232 inline __device__ void WriteLVecStandard2d_Single(SharedData_Hip &data, const CeedInt num_nodes, co… in WriteLVecStandard2d_Single() argument
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()
251 inline __device__ void WriteLVecStandard2d_Assembly(SharedData_Hip &data, const CeedInt num_nodes, … in WriteLVecStandard2d_Assembly() argument
259 if (data.t_id_x < P_1D && data.t_id_y < P_1D) { in WriteLVecStandard2d_Assembly()
261 const CeedInt out_node = data.t_id_x + data.t_id_y * P_1D; in WriteLVecStandard2d_Assembly()
275 inline __device__ void WriteLVecStandard2d_QFAssembly(SharedData_Hip &data, const CeedInt num_elem,… in WriteLVecStandard2d_QFAssembly() argument
277 if (data.t_id_x < Q_1D && data.t_id_y < Q_1D) { in WriteLVecStandard2d_QFAssembly()
278 const CeedInt ind = (data.t_id_x + data.t_id_y * Q_1D) + elem * Q_1D * Q_1D; in WriteLVecStandard2d_QFAssembly()
290 inline __device__ void WriteLVecStrided2d(SharedData_Hip &data, const CeedInt elem, const CeedScala… in WriteLVecStrided2d() argument
292 if (data.t_id_x < P_1D && data.t_id_y < P_1D) { in WriteLVecStrided2d()
293 const CeedInt node = data.t_id_x + data.t_id_y * P_1D; in WriteLVecStrided2d()
308 inline __device__ void SetEVecStandard3d_Single(SharedData_Hip &data, const CeedInt n, const CeedSc… in SetEVecStandard3d_Single() argument
314 if (data.t_id_x == target_node_x && data.t_id_y == target_node_y) { in SetEVecStandard3d_Single()
323 inline __device__ void ReadLVecStandard3d(SharedData_Hip &data, const CeedInt num_nodes, const Ceed… in ReadLVecStandard3d() argument
325 if (data.t_id_x < P_1D && data.t_id_y < P_1D) { in ReadLVecStandard3d()
327 const CeedInt node = data.t_id_x + data.t_id_y * P_1D + z * P_1D * P_1D; in ReadLVecStandard3d()
339 inline __device__ void ReadLVecStrided3d(SharedData_Hip &data, const CeedInt elem, const CeedScalar… in ReadLVecStrided3d() argument
340 if (data.t_id_x < P_1D && data.t_id_y < P_1D) { in ReadLVecStrided3d()
342 const CeedInt node = data.t_id_x + data.t_id_y * P_1D + z * P_1D * P_1D; in ReadLVecStrided3d()
354 inline __device__ void ReadEVecSliceStandard3d(SharedData_Hip &data, const CeedInt nquads, const Ce… in ReadEVecSliceStandard3d() argument
357 if (data.t_id_x < Q_1D && data.t_id_y < Q_1D) { in ReadEVecSliceStandard3d()
358 const CeedInt node = data.t_id_x + data.t_id_y * Q_1D + q * Q_1D * Q_1D; in ReadEVecSliceStandard3d()
369 inline __device__ void ReadEVecSliceStrided3d(SharedData_Hip &data, const CeedInt elem, const CeedI… in ReadEVecSliceStrided3d() argument
371 if (data.t_id_x < Q_1D && data.t_id_y < Q_1D) { in ReadEVecSliceStrided3d()
372 const CeedInt node = data.t_id_x + data.t_id_y * Q_1D + q * Q_1D * Q_1D; in ReadEVecSliceStrided3d()
383 inline __device__ void WriteLVecStandard3d(SharedData_Hip &data, const CeedInt num_nodes, const Cee… in WriteLVecStandard3d() argument
385 if (data.t_id_x < P_1D && data.t_id_y < P_1D) { in WriteLVecStandard3d()
387 const CeedInt node = data.t_id_x + data.t_id_y * P_1D + z * P_1D * P_1D; in WriteLVecStandard3d()
396 inline __device__ void WriteLVecStandard3d_Single(SharedData_Hip &data, const CeedInt num_nodes, co… in WriteLVecStandard3d_Single() argument
404 if (data.t_id_x == target_node_x && data.t_id_y == target_node_y) { in WriteLVecStandard3d_Single()
405 const CeedInt node = data.t_id_x + data.t_id_y * P_1D + target_node_z * P_1D * P_1D; in WriteLVecStandard3d_Single()
416 inline __device__ void WriteLVecStandard3d_Assembly(SharedData_Hip &data, const CeedInt num_nodes, … in WriteLVecStandard3d_Assembly() argument
425 if (data.t_id_x < P_1D && data.t_id_y < P_1D) { in WriteLVecStandard3d_Assembly()
428 const CeedInt out_node = data.t_id_x + data.t_id_y * P_1D + z * P_1D * P_1D; in WriteLVecStandard3d_Assembly()
443 inline __device__ void WriteLVecStandard3d_QFAssembly(SharedData_Hip &data, const CeedInt num_elem,… in WriteLVecStandard3d_QFAssembly() argument
445 if (data.t_id_x < Q_1D && data.t_id_y < Q_1D) { in WriteLVecStandard3d_QFAssembly()
447 …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()
460 inline __device__ void WriteLVecStrided3d(SharedData_Hip &data, const CeedInt elem, const CeedScala… in WriteLVecStrided3d() argument
462 if (data.t_id_x < P_1D && data.t_id_y < P_1D) { in WriteLVecStrided3d()
464 const CeedInt node = data.t_id_x + data.t_id_y * P_1D + z * P_1D * P_1D; in WriteLVecStrided3d()
476 inline __device__ void GradColloSlice3d(SharedData_Hip &data, const CeedInt q, const CeedScalar *__… in GradColloSlice3d() argument
478 if (data.t_id_x < Q_1D && data.t_id_y < Q_1D) { in GradColloSlice3d()
481 data.slice[data.t_id_x + data.t_id_y * T_1D] = r_U[q + comp * Q_1D]; in GradColloSlice3d()
486 … 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()
491 … 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()
506 inline __device__ void GradColloSliceTranspose3d(SharedData_Hip &data, const CeedInt q, const CeedS… in GradColloSliceTranspose3d() argument
508 if (data.t_id_x < Q_1D && data.t_id_y < Q_1D) { in GradColloSliceTranspose3d()
512 data.slice[data.t_id_x + data.t_id_y * T_1D] = r_U[comp + 0 * NUM_COMP]; in GradColloSliceTranspose3d()
515 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()
519 data.slice[data.t_id_x + data.t_id_y * T_1D] = r_U[comp + 1 * NUM_COMP]; in GradColloSliceTranspose3d()
522 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()