1 // Copyright (c) 2017-2022, 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 #ifndef CEED_MAGMA_BASIS_WEIGHT_1D_H 11 #define CEED_MAGMA_BASIS_WEIGHT_1D_H 12 13 #include "magma-common-tensor.h" 14 15 //////////////////////////////////////////////////////////////////////////////// 16 // weight basis action -- 1D 17 template <typename T, int Q> 18 static __device__ __inline__ void magma_weight_1d_device(const T *sTweight, T *sV, const int tx) { 19 // Assumptions 20 // 1. 1D thread configuration of size Q 21 // 2. The output sV is in shared memory -- size Q 22 if (tx < Q) { 23 sV[tx] = sTweight[tx]; 24 } 25 } 26 27 //////////////////////////////////////////////////////////////////////////////// 28 extern "C" __launch_bounds__(MAGMA_BASIS_BOUNDS(BASIS_Q, MAGMA_MAXTHREADS_1D)) __global__ 29 void magma_weight_1d_kernel(const CeedScalar *dqweight1d, CeedScalar *dV, const int v_stride, const int nelem) { 30 MAGMA_DEVICE_SHARED(CeedScalar, shared_data) 31 32 const int tx = threadIdx.x; 33 const int ty = threadIdx.y; 34 const int elem_id = (blockIdx.x * blockDim.y) + ty; 35 36 if (elem_id >= nelem) return; 37 38 // global memory pointers 39 dV += elem_id * v_stride; 40 41 // shared memory pointers 42 CeedScalar *sTweight = (CeedScalar *)shared_data; 43 CeedScalar *sV = sTweight + BASIS_Q; 44 sV += ty * BASIS_Q; 45 46 // read dqweight_1d 47 if (ty == 0 && tx < BASIS_Q) { 48 sTweight[tx] = dqweight1d[tx]; 49 } 50 51 __syncthreads(); 52 magma_weight_1d_device<CeedScalar, BASIS_Q>(sTweight, sV, tx); 53 __syncthreads(); 54 55 // write V 56 dV[tx] = sV[tx]; 57 } 58 59 #endif // CEED_MAGMA_BASIS_WEIGHT_1D_H 60