1 // Copyright (c) 2017-2026, Lawrence Livermore National Security, LLC and other CEED contributors.
2 // All Rights Reserved. See the top-level LICENSE and NOTICE files for details.
4 // SPDX-License-Identifier: BSD-2-Clause
12 #include "hip-shared-basis-read-write-templates.h"
13 #include "hip-shared-basis-tensor-at-points-templates.h"
14 #include "hip-shared-basis-tensor-templates.h"
16 //------------------------------------------------------------------------------
18 //------------------------------------------------------------------------------
20 //------------------------------------------------------------------------------
22 //------------------------------------------------------------------------------
30   data.t_id_y = threadIdx.y;  in __launch_bounds__()
32   data.t_id   = threadIdx.x + threadIdx.y * blockDim.x + threadIdx.z * blockDim.y * blockDim.x;  in __launch_bounds__()
61 …edInt point_loop_bound = (blockDim.x * blockDim.y) * ceil(1.0 * BASIS_NUM_PTS / (blockDim.x * bloc…  in __launch_bounds__()
63 …for (CeedInt i = threadIdx.x + threadIdx.y * blockDim.x; i < point_loop_bound; i += blockDim.x * b…  in __launch_bounds__()
86   data.t_id_y = threadIdx.y;  in __launch_bounds__()
88   data.t_id   = threadIdx.x + threadIdx.y * blockDim.x + threadIdx.z * blockDim.y * blockDim.x;  in __launch_bounds__()
118 …edInt point_loop_bound = (blockDim.x * blockDim.y) * ceil(1.0 * BASIS_NUM_PTS / (blockDim.x * bloc…  in __launch_bounds__()
120 …for (CeedInt i = threadIdx.x + threadIdx.y * blockDim.x; i < point_loop_bound; i += blockDim.x * b…  in __launch_bounds__()
156   data.t_id_y = threadIdx.y;  in __launch_bounds__()
158   data.t_id   = threadIdx.x + threadIdx.y * blockDim.x + threadIdx.z * blockDim.y * blockDim.x;  in __launch_bounds__()
177 …edInt point_loop_bound = (blockDim.x * blockDim.y) * ceil(1.0 * BASIS_NUM_PTS / (blockDim.x * bloc…  in __launch_bounds__()
179 …for (CeedInt i = threadIdx.x + threadIdx.y * blockDim.x; i < point_loop_bound; i += blockDim.x * b…  in __launch_bounds__()
208 //------------------------------------------------------------------------------
210 //------------------------------------------------------------------------------
218   data.t_id_y = threadIdx.y;  in __launch_bounds__()
220   data.t_id   = threadIdx.x + threadIdx.y * blockDim.x + threadIdx.z * blockDim.y * blockDim.x;  in __launch_bounds__()
249 …edInt point_loop_bound = (blockDim.x * blockDim.y) * ceil(1.0 * BASIS_NUM_PTS / (blockDim.x * bloc…  in __launch_bounds__()
251 …for (CeedInt i = threadIdx.x + threadIdx.y * blockDim.x; i < point_loop_bound; i += blockDim.x * b…  in __launch_bounds__()
274   data.t_id_y = threadIdx.y;  in __launch_bounds__()
276   data.t_id   = threadIdx.x + threadIdx.y * blockDim.x + threadIdx.z * blockDim.y * blockDim.x;  in __launch_bounds__()
306 …edInt point_loop_bound = (blockDim.x * blockDim.y) * ceil(1.0 * BASIS_NUM_PTS / (blockDim.x * bloc…  in __launch_bounds__()
308 …for (CeedInt i = threadIdx.x + threadIdx.y * blockDim.x; i < point_loop_bound; i += blockDim.x * b…  in __launch_bounds__()
345   data.t_id_y = threadIdx.y;  in __launch_bounds__()
347   data.t_id   = threadIdx.x + threadIdx.y * blockDim.x + threadIdx.z * blockDim.y * blockDim.x;  in __launch_bounds__()
366 …edInt point_loop_bound = (blockDim.x * blockDim.y) * ceil(1.0 * BASIS_NUM_PTS / (blockDim.x * bloc…  in __launch_bounds__()
368 …for (CeedInt i = threadIdx.x + threadIdx.y * blockDim.x; i < point_loop_bound; i += blockDim.x * b…  in __launch_bounds__()