1*f80f4a74SSebastian Grimberg // Copyright (c) 2017-2022, Lawrence Livermore National Security, LLC and other CEED contributors. 2*f80f4a74SSebastian Grimberg // All Rights Reserved. See the top-level LICENSE and NOTICE files for details. 3*f80f4a74SSebastian Grimberg // 4*f80f4a74SSebastian Grimberg // SPDX-License-Identifier: BSD-2-Clause 5*f80f4a74SSebastian Grimberg // 6*f80f4a74SSebastian Grimberg // This file is part of CEED: http://github.com/ceed 7*f80f4a74SSebastian Grimberg 8*f80f4a74SSebastian Grimberg ////////////////////////////////////////////////////////////////////////////////////////// 9*f80f4a74SSebastian Grimberg // weight basis action -- 1D 10*f80f4a74SSebastian Grimberg template <typename T, int Q_> 11*f80f4a74SSebastian Grimberg __device__ __inline__ void magma_weight_1d_device(const T *sTweight, T *sV, const int tx) { 12*f80f4a74SSebastian Grimberg // Assumptions 13*f80f4a74SSebastian Grimberg // 1. 1D thread configuration of size Q_ 14*f80f4a74SSebastian Grimberg // 2. The output sV is in shared memory -- size 1xQ_ 15*f80f4a74SSebastian Grimberg if (tx < Q_) { 16*f80f4a74SSebastian Grimberg sV[tx] = sTweight[tx]; 17*f80f4a74SSebastian Grimberg } 18*f80f4a74SSebastian Grimberg } 19*f80f4a74SSebastian Grimberg 20*f80f4a74SSebastian Grimberg ////////////////////////////////////////////////////////////////////////////////////////// 21*f80f4a74SSebastian Grimberg extern "C" __launch_bounds__(MAGMA_BASIS_BOUNDS(Q, MAGMA_MAXTHREADS_1D)) __global__ 22*f80f4a74SSebastian Grimberg void magma_weight_1d_kernel(const CeedScalar *dqweight1d, CeedScalar *dV, const int v_stride, const int nelem) { 23*f80f4a74SSebastian Grimberg MAGMA_DEVICE_SHARED(CeedScalar, shared_data) 24*f80f4a74SSebastian Grimberg 25*f80f4a74SSebastian Grimberg const int tx = threadIdx.x; 26*f80f4a74SSebastian Grimberg const int ty = threadIdx.y; 27*f80f4a74SSebastian Grimberg const int elem_id = (blockIdx.x * blockDim.y) + ty; 28*f80f4a74SSebastian Grimberg 29*f80f4a74SSebastian Grimberg if (elem_id >= nelem) return; 30*f80f4a74SSebastian Grimberg 31*f80f4a74SSebastian Grimberg // global memory pointers 32*f80f4a74SSebastian Grimberg dV += elem_id * v_stride; 33*f80f4a74SSebastian Grimberg 34*f80f4a74SSebastian Grimberg // shared memory pointers 35*f80f4a74SSebastian Grimberg CeedScalar *sTweight = (CeedScalar *)shared_data; 36*f80f4a74SSebastian Grimberg CeedScalar *sV = sTweight + Q; 37*f80f4a74SSebastian Grimberg sV += ty * Q; 38*f80f4a74SSebastian Grimberg 39*f80f4a74SSebastian Grimberg // read dqweight_1d 40*f80f4a74SSebastian Grimberg if (ty == 0 && tx < Q) { 41*f80f4a74SSebastian Grimberg sTweight[tx] = dqweight1d[tx]; 42*f80f4a74SSebastian Grimberg } 43*f80f4a74SSebastian Grimberg 44*f80f4a74SSebastian Grimberg __syncthreads(); 45*f80f4a74SSebastian Grimberg magma_weight_1d_device<CeedScalar, Q>(sTweight, sV, tx); 46*f80f4a74SSebastian Grimberg __syncthreads(); 47*f80f4a74SSebastian Grimberg 48*f80f4a74SSebastian Grimberg // write V 49*f80f4a74SSebastian Grimberg dV[tx] = sV[tx]; 50*f80f4a74SSebastian Grimberg } 51