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
10 #include "magma-common-tensor.h"
18 template <typename T, int DIM_U, int DIM_V, int NUM_COMP, int P, int Q, int rU_SIZE, int rV_SIZE>
19 …__device__ __inline__ void magma_interp_2d_device(const T *sT, T rU[DIM_U][NUM_COMP][rU_SIZE], T r…  in magma_interp_2d_device()
20                                                          T rTmp, T *swork) {  in magma_interp_2d_device()
32     // 1st product -- Batch P of (1xP) matrices [reg] x (PxQ) [shmem] => Batch P of (1xQ) matrices  in magma_interp_2d_device()
37       T        *sTmp    = swork + batchid * (1 * Q);  in magma_interp_2d_device()
48     // 2nd product -- Batch 1 of a (QxP) matrix [shmem] x (PxQ) [shmem] => (QxQ) matrix [reg]  in magma_interp_2d_device()
52       T        *sTmp    = swork + batchid * (Q * P);  in magma_interp_2d_device()
77 …CeedScalar rU[1][BASIS_NUM_COMP][BASIS_P] = {0.0};  // for a non-fused operator BASIS_DIM is alway…  in __launch_bounds__()
78 …CeedScalar rV[1][BASIS_NUM_COMP][BASIS_Q] = {0.0};  // for a non-fused operator BASIS_DIM is alway…  in __launch_bounds__()
90   // read T  in __launch_bounds__()
95   // read U -- there is a sync at the end of this function  in __launch_bounds__()
98   // no sync needed here -- read_U_2d already syncs at the end  in __launch_bounds__()
118 …CeedScalar rU[1][BASIS_NUM_COMP][BASIS_Q] = {0.0};  // for a non-fused operator BASIS_DIM is alway…  in __launch_bounds__()
119 …CeedScalar rV[1][BASIS_NUM_COMP][BASIS_P] = {0.0};  // for a non-fused operator BASIS_DIM is alway…  in __launch_bounds__()
131   // read T  in __launch_bounds__()
136   // read U -- there is a sync at the end of this function  in __launch_bounds__()
139   // no sync needed here -- read_U_2d already syncs at the end  in __launch_bounds__()
159 …CeedScalar rU[1][BASIS_NUM_COMP][BASIS_Q] = {0.0};  // for a non-fused operator BASIS_DIM is alway…  in __launch_bounds__()
160 …CeedScalar rV[1][BASIS_NUM_COMP][BASIS_P] = {0.0};  // for a non-fused operator BASIS_DIM is alway…  in __launch_bounds__()
172   // read T  in __launch_bounds__()
177   // read U -- there is a sync at the end of this function  in __launch_bounds__()
180   // no sync needed here -- read_U_2d already syncs at the end  in __launch_bounds__()