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.
3 //
4 // SPDX-License-Identifier: BSD-2-Clause
5 //
6 // This file is part of CEED: http://github.com/ceed
7
8 /// @file
9 /// Internal header for MAGMA tensor basis weight in 1D
10 #include "magma-common-tensor.h"
11
12 ////////////////////////////////////////////////////////////////////////////////
13 // weight basis action -- 1D
14 template <typename T, int Q>
magma_weight_1d_device(const T * sTweight,T * sV,const int tx)15 static __device__ __inline__ void magma_weight_1d_device(const T *sTweight, T *sV, const int tx) {
16 // Assumptions
17 // 1. 1D thread configuration of size Q
18 // 2. The output sV is in shared memory -- size Q
19 if (tx < Q) {
20 sV[tx] = sTweight[tx];
21 }
22 }
23
24 ////////////////////////////////////////////////////////////////////////////////
__launch_bounds__(MAGMA_BASIS_BOUNDS (BASIS_Q,MAGMA_MAXTHREADS_1D))25 extern "C" __launch_bounds__(MAGMA_BASIS_BOUNDS(BASIS_Q, MAGMA_MAXTHREADS_1D)) __global__
26 void magma_weight_1d_kernel(const CeedScalar *dqweight1d, CeedScalar *dV, const int v_stride, const int nelem) {
27 MAGMA_DEVICE_SHARED(CeedScalar, shared_data)
28
29 const int tx = threadIdx.x;
30 const int ty = threadIdx.y;
31 const int elem_id = (blockIdx.x * blockDim.y) + ty;
32
33 if (elem_id >= nelem) return;
34
35 // global memory pointers
36 dV += elem_id * v_stride;
37
38 // shared memory pointers
39 CeedScalar *sTweight = (CeedScalar *)shared_data;
40 CeedScalar *sV = sTweight + BASIS_Q;
41 sV += ty * BASIS_Q;
42
43 // read dqweight_1d
44 if (ty == 0 && tx < BASIS_Q) {
45 sTweight[tx] = dqweight1d[tx];
46 }
47
48 __syncthreads();
49 magma_weight_1d_device<CeedScalar, BASIS_Q>(sTweight, sV, tx);
50 __syncthreads();
51
52 // write V
53 dV[tx] = sV[tx];
54 }
55