xref: /libCEED/include/ceed/jit-source/magma/magma-basis-weight-1d.h (revision 3e551a327d6c97f9de071b988b42ffdb7bed19a7)
1 // Copyright (c) 2017-2024, 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>
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 ////////////////////////////////////////////////////////////////////////////////
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