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