xref: /libCEED/include/ceed/jit-source/hip/hip-shared-basis-read-write-templates.h (revision 5aed82e4fa97acf4ba24a7f10a35f5303a6798e0)
1*5aed82e4SJeremy L Thompson // Copyright (c) 2017-2024, Lawrence Livermore National Security, LLC and other CEED contributors.
29e201c85SYohann // All Rights Reserved. See the top-level LICENSE and NOTICE files for details.
39e201c85SYohann //
49e201c85SYohann // SPDX-License-Identifier: BSD-2-Clause
59e201c85SYohann //
69e201c85SYohann // This file is part of CEED:  http://github.com/ceed
79e201c85SYohann 
89e201c85SYohann /// @file
99e201c85SYohann /// Internal header for HIP shared memory basis read/write templates
1094b7b29bSJeremy L Thompson #ifndef CEED_HIP_SHARED_BASIS_READ_WRITE_TEMPLATES_H
1194b7b29bSJeremy L Thompson #define CEED_HIP_SHARED_BASIS_READ_WRITE_TEMPLATES_H
129e201c85SYohann 
139e201c85SYohann #include <ceed.h>
149e201c85SYohann 
159e201c85SYohann //------------------------------------------------------------------------------
169e201c85SYohann // Helper function: load matrices for basis actions
179e201c85SYohann //------------------------------------------------------------------------------
189e201c85SYohann template <int SIZE>
199e201c85SYohann inline __device__ void loadMatrix(const CeedScalar *d_B, CeedScalar *B) {
209e201c85SYohann   CeedInt tid = threadIdx.x + threadIdx.y * blockDim.x + threadIdx.z * blockDim.y * blockDim.x;
21672b0f2aSSebastian Grimberg 
222b730f8bSJeremy L Thompson   for (CeedInt i = tid; i < SIZE; i += blockDim.x * blockDim.y * blockDim.z) B[i] = d_B[i];
239e201c85SYohann }
249e201c85SYohann 
259e201c85SYohann //------------------------------------------------------------------------------
269e201c85SYohann // 1D
279e201c85SYohann //------------------------------------------------------------------------------
289e201c85SYohann 
299e201c85SYohann //------------------------------------------------------------------------------
309e201c85SYohann // E-vector -> single element
319e201c85SYohann //------------------------------------------------------------------------------
329e201c85SYohann template <int NUM_COMP, int P_1D>
332b730f8bSJeremy L Thompson inline __device__ void ReadElementStrided1d(SharedData_Hip &data, const CeedInt elem, const CeedInt strides_node, const CeedInt strides_comp,
342b730f8bSJeremy L Thompson                                             const CeedInt strides_elem, const CeedScalar *__restrict__ d_u, CeedScalar *r_u) {
359e201c85SYohann   if (data.t_id_x < P_1D) {
369e201c85SYohann     const CeedInt node = data.t_id_x;
379e201c85SYohann     const CeedInt ind  = node * strides_node + elem * strides_elem;
38672b0f2aSSebastian Grimberg 
399e201c85SYohann     for (CeedInt comp = 0; comp < NUM_COMP; comp++) {
409e201c85SYohann       r_u[comp] = d_u[ind + comp * strides_comp];
419e201c85SYohann     }
429e201c85SYohann   }
439e201c85SYohann }
449e201c85SYohann 
459e201c85SYohann //------------------------------------------------------------------------------
469e201c85SYohann // Single element -> E-vector
479e201c85SYohann //------------------------------------------------------------------------------
489e201c85SYohann template <int NUM_COMP, int P_1D>
492b730f8bSJeremy L Thompson inline __device__ void WriteElementStrided1d(SharedData_Hip &data, const CeedInt elem, const CeedInt strides_node, const CeedInt strides_comp,
502b730f8bSJeremy L Thompson                                              const CeedInt strides_elem, const CeedScalar *r_v, CeedScalar *d_v) {
519e201c85SYohann   if (data.t_id_x < P_1D) {
529e201c85SYohann     const CeedInt node = data.t_id_x;
539e201c85SYohann     const CeedInt ind  = node * strides_node + elem * strides_elem;
54672b0f2aSSebastian Grimberg 
559e201c85SYohann     for (CeedInt comp = 0; comp < NUM_COMP; comp++) {
569e201c85SYohann       d_v[ind + comp * strides_comp] = r_v[comp];
579e201c85SYohann     }
589e201c85SYohann   }
599e201c85SYohann }
609e201c85SYohann 
619e201c85SYohann //------------------------------------------------------------------------------
629e201c85SYohann // 2D
639e201c85SYohann //------------------------------------------------------------------------------
649e201c85SYohann 
659e201c85SYohann //------------------------------------------------------------------------------
669e201c85SYohann // E-vector -> single element
679e201c85SYohann //------------------------------------------------------------------------------
689e201c85SYohann template <int NUM_COMP, int P_1D>
692b730f8bSJeremy L Thompson inline __device__ void ReadElementStrided2d(SharedData_Hip &data, const CeedInt elem, const CeedInt strides_node, const CeedInt strides_comp,
702b730f8bSJeremy L Thompson                                             const CeedInt strides_elem, const CeedScalar *__restrict__ d_u, CeedScalar *r_u) {
719e201c85SYohann   if (data.t_id_x < P_1D && data.t_id_y < P_1D) {
729e201c85SYohann     const CeedInt node = data.t_id_x + data.t_id_y * P_1D;
739e201c85SYohann     const CeedInt ind  = node * strides_node + elem * strides_elem;
74672b0f2aSSebastian Grimberg 
759e201c85SYohann     for (CeedInt comp = 0; comp < NUM_COMP; comp++) {
769e201c85SYohann       r_u[comp] = d_u[ind + comp * strides_comp];
779e201c85SYohann     }
789e201c85SYohann   }
799e201c85SYohann }
809e201c85SYohann 
819e201c85SYohann //------------------------------------------------------------------------------
829e201c85SYohann // Single element -> E-vector
839e201c85SYohann //------------------------------------------------------------------------------
849e201c85SYohann template <int NUM_COMP, int P_1D>
852b730f8bSJeremy L Thompson inline __device__ void WriteElementStrided2d(SharedData_Hip &data, const CeedInt elem, const CeedInt strides_node, const CeedInt strides_comp,
862b730f8bSJeremy L Thompson                                              const CeedInt strides_elem, const CeedScalar *r_v, CeedScalar *d_v) {
879e201c85SYohann   if (data.t_id_x < P_1D && data.t_id_y < P_1D) {
889e201c85SYohann     const CeedInt node = data.t_id_x + data.t_id_y * P_1D;
899e201c85SYohann     const CeedInt ind  = node * strides_node + elem * strides_elem;
90672b0f2aSSebastian Grimberg 
919e201c85SYohann     for (CeedInt comp = 0; comp < NUM_COMP; comp++) {
929e201c85SYohann       d_v[ind + comp * strides_comp] = r_v[comp];
939e201c85SYohann     }
949e201c85SYohann   }
959e201c85SYohann }
969e201c85SYohann 
979e201c85SYohann //------------------------------------------------------------------------------
989e201c85SYohann // 3D
999e201c85SYohann //------------------------------------------------------------------------------
1009e201c85SYohann 
1019e201c85SYohann //------------------------------------------------------------------------------
1029e201c85SYohann // E-vector -> single element
1039e201c85SYohann //------------------------------------------------------------------------------
1049e201c85SYohann template <int NUM_COMP, int P_1D>
1052b730f8bSJeremy L Thompson inline __device__ void ReadElementStrided3d(SharedData_Hip &data, const CeedInt elem, const CeedInt strides_node, const CeedInt strides_comp,
1062b730f8bSJeremy L Thompson                                             const CeedInt strides_elem, const CeedScalar *__restrict__ d_u, CeedScalar *r_u) {
1079e201c85SYohann   if (data.t_id_x < P_1D && data.t_id_y < P_1D) {
1089e201c85SYohann     for (CeedInt z = 0; z < P_1D; z++) {
1099e201c85SYohann       const CeedInt node = data.t_id_x + data.t_id_y * P_1D + z * P_1D * P_1D;
1109e201c85SYohann       const CeedInt ind  = node * strides_node + elem * strides_elem;
111672b0f2aSSebastian Grimberg 
1129e201c85SYohann       for (CeedInt comp = 0; comp < NUM_COMP; comp++) {
1139e201c85SYohann         r_u[z + comp * P_1D] = d_u[ind + comp * strides_comp];
1149e201c85SYohann       }
1159e201c85SYohann     }
1169e201c85SYohann   }
1179e201c85SYohann }
1189e201c85SYohann 
1199e201c85SYohann //------------------------------------------------------------------------------
1209e201c85SYohann // Single element -> E-vector
1219e201c85SYohann //------------------------------------------------------------------------------
1229e201c85SYohann template <int NUM_COMP, int P_1D>
1232b730f8bSJeremy L Thompson inline __device__ void WriteElementStrided3d(SharedData_Hip &data, const CeedInt elem, const CeedInt strides_node, const CeedInt strides_comp,
1242b730f8bSJeremy L Thompson                                              const CeedInt strides_elem, const CeedScalar *r_v, CeedScalar *d_v) {
1259e201c85SYohann   if (data.t_id_x < P_1D && data.t_id_y < P_1D) {
1269e201c85SYohann     for (CeedInt z = 0; z < P_1D; z++) {
1279e201c85SYohann       const CeedInt node = data.t_id_x + data.t_id_y * P_1D + z * P_1D * P_1D;
1289e201c85SYohann       const CeedInt ind  = node * strides_node + elem * strides_elem;
129672b0f2aSSebastian Grimberg 
1309e201c85SYohann       for (CeedInt comp = 0; comp < NUM_COMP; comp++) {
1319e201c85SYohann         d_v[ind + comp * strides_comp] = r_v[z + comp * P_1D];
1329e201c85SYohann       }
1339e201c85SYohann     }
1349e201c85SYohann   }
1359e201c85SYohann }
1369e201c85SYohann 
1379e201c85SYohann //------------------------------------------------------------------------------
1389e201c85SYohann 
13994b7b29bSJeremy L Thompson #endif  // CEED_HIP_SHARED_BASIS_READ_WRITE_TEMPLATES_H
140