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 "cuda-shared-basis-read-write-templates.h"
13 #include "cuda-shared-basis-tensor-at-points-templates.h"
14 #include "cuda-shared-basis-tensor-templates.h"
16 //------------------------------------------------------------------------------
18 //------------------------------------------------------------------------------
20 //------------------------------------------------------------------------------
22 //------------------------------------------------------------------------------
29   data.t_id_y = threadIdx.y;  in InterpAtPoints()
31   data.t_id   = threadIdx.x + threadIdx.y * blockDim.x + threadIdx.z * blockDim.y * blockDim.x;  in InterpAtPoints()
60 …edInt point_loop_bound = (blockDim.x * blockDim.y) * ceil(1.0 * BASIS_NUM_PTS / (blockDim.x * bloc…  in InterpAtPoints()
62 …for (CeedInt i = threadIdx.x + threadIdx.y * blockDim.x; i < point_loop_bound; i += blockDim.x * b…  in InterpAtPoints()
85   data.t_id_y = threadIdx.y;  in InterpTransposeAtPoints()
87   data.t_id   = threadIdx.x + threadIdx.y * blockDim.x + threadIdx.z * blockDim.y * blockDim.x;  in InterpTransposeAtPoints()
117 …edInt point_loop_bound = (blockDim.x * blockDim.y) * ceil(1.0 * BASIS_NUM_PTS / (blockDim.x * bloc…  in InterpTransposeAtPoints()
119 …for (CeedInt i = threadIdx.x + threadIdx.y * blockDim.x; i < point_loop_bound; i += blockDim.x * b…  in InterpTransposeAtPoints()
155   data.t_id_y = threadIdx.y;  in InterpTransposeAddAtPoints()
157   data.t_id   = threadIdx.x + threadIdx.y * blockDim.x + threadIdx.z * blockDim.y * blockDim.x;  in InterpTransposeAddAtPoints()
176 …edInt point_loop_bound = (blockDim.x * blockDim.y) * ceil(1.0 * BASIS_NUM_PTS / (blockDim.x * bloc…  in InterpTransposeAddAtPoints()
178 …for (CeedInt i = threadIdx.x + threadIdx.y * blockDim.x; i < point_loop_bound; i += blockDim.x * b…  in InterpTransposeAddAtPoints()
207 //------------------------------------------------------------------------------
209 //------------------------------------------------------------------------------
216   data.t_id_y = threadIdx.y;  in GradAtPoints()
218   data.t_id   = threadIdx.x + threadIdx.y * blockDim.x + threadIdx.z * blockDim.y * blockDim.x;  in GradAtPoints()
247 …edInt point_loop_bound = (blockDim.x * blockDim.y) * ceil(1.0 * BASIS_NUM_PTS / (blockDim.x * bloc…  in GradAtPoints()
249 …for (CeedInt i = threadIdx.x + threadIdx.y * blockDim.x; i < point_loop_bound; i += blockDim.x * b…  in GradAtPoints()
272   data.t_id_y = threadIdx.y;  in GradTransposeAtPoints()
274   data.t_id   = threadIdx.x + threadIdx.y * blockDim.x + threadIdx.z * blockDim.y * blockDim.x;  in GradTransposeAtPoints()
304 …edInt point_loop_bound = (blockDim.x * blockDim.y) * ceil(1.0 * BASIS_NUM_PTS / (blockDim.x * bloc…  in GradTransposeAtPoints()
306 …for (CeedInt i = threadIdx.x + threadIdx.y * blockDim.x; i < point_loop_bound; i += blockDim.x * b…  in GradTransposeAtPoints()
343   data.t_id_y = threadIdx.y;  in GradTransposeAddAtPoints()
345   data.t_id   = threadIdx.x + threadIdx.y * blockDim.x + threadIdx.z * blockDim.y * blockDim.x;  in GradTransposeAddAtPoints()
364 …edInt point_loop_bound = (blockDim.x * blockDim.y) * ceil(1.0 * BASIS_NUM_PTS / (blockDim.x * bloc…  in GradTransposeAddAtPoints()
366 …for (CeedInt i = threadIdx.x + threadIdx.y * blockDim.x; i < point_loop_bound; i += blockDim.x * b…  in GradTransposeAddAtPoints()