1*940a72f1SSebastian Grimberg // Copyright (c) 2017-2022, Lawrence Livermore National Security, LLC and other CEED contributors. 2*940a72f1SSebastian Grimberg // All Rights Reserved. See the top-level LICENSE and NOTICE files for details. 3*940a72f1SSebastian Grimberg // 4*940a72f1SSebastian Grimberg // SPDX-License-Identifier: BSD-2-Clause 5*940a72f1SSebastian Grimberg // 6*940a72f1SSebastian Grimberg // This file is part of CEED: http://github.com/ceed 7*940a72f1SSebastian Grimberg 8*940a72f1SSebastian Grimberg /// @file 9*940a72f1SSebastian Grimberg /// Internal header for MAGMA non-tensor basis weight 10*940a72f1SSebastian Grimberg #ifndef CEED_MAGMA_BASIS_WEIGHT_NONTENSOR_H 11*940a72f1SSebastian Grimberg #define CEED_MAGMA_BASIS_WEIGHT_NONTENSOR_H 12*940a72f1SSebastian Grimberg 13*940a72f1SSebastian Grimberg #include "magma-common-nontensor.h" 14*940a72f1SSebastian Grimberg 15*940a72f1SSebastian Grimberg //////////////////////////////////////////////////////////////////////////////// 16*940a72f1SSebastian Grimberg extern "C" __launch_bounds__(MAGMA_BASIS_BOUNDS(BASIS_Q, MAGMA_MAXTHREADS_1D)) __global__ 17*940a72f1SSebastian Grimberg void magma_weight_nontensor(int n, const CeedScalar *dqweight, CeedScalar *dV, int lddv) { 18*940a72f1SSebastian Grimberg MAGMA_DEVICE_SHARED(CeedScalar, shared_data); 19*940a72f1SSebastian Grimberg 20*940a72f1SSebastian Grimberg const int tx = threadIdx.x; 21*940a72f1SSebastian Grimberg const int ty = threadIdx.y; 22*940a72f1SSebastian Grimberg const int id = blockIdx.x * blockDim.y + ty; 23*940a72f1SSebastian Grimberg 24*940a72f1SSebastian Grimberg // terminate threads with no work 25*940a72f1SSebastian Grimberg if (id >= n) return; 26*940a72f1SSebastian Grimberg 27*940a72f1SSebastian Grimberg dV += id * lddv; 28*940a72f1SSebastian Grimberg 29*940a72f1SSebastian Grimberg // shared memory pointers 30*940a72f1SSebastian Grimberg CeedScalar *sqweight = (CeedScalar *)shared_data; 31*940a72f1SSebastian Grimberg CeedScalar *sV = sqweight + BASIS_Q; 32*940a72f1SSebastian Grimberg sV += ty * BASIS_Q; 33*940a72f1SSebastian Grimberg 34*940a72f1SSebastian Grimberg // read qweight 35*940a72f1SSebastian Grimberg if (ty == 0 && tx < BASIS_Q) { 36*940a72f1SSebastian Grimberg sqweight[tx] = dqweight[tx]; 37*940a72f1SSebastian Grimberg } 38*940a72f1SSebastian Grimberg __syncthreads(); 39*940a72f1SSebastian Grimberg 40*940a72f1SSebastian Grimberg if (tx < BASIS_Q) { 41*940a72f1SSebastian Grimberg sV[tx] = sqweight[tx]; 42*940a72f1SSebastian Grimberg } 43*940a72f1SSebastian Grimberg 44*940a72f1SSebastian Grimberg // write V 45*940a72f1SSebastian Grimberg dV[tx] = sV[tx]; 46*940a72f1SSebastian Grimberg } 47*940a72f1SSebastian Grimberg 48*940a72f1SSebastian Grimberg #endif // CEED_MAGMA_BASIS_WEIGHT_NONTENSOR_H 49