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 non-tensor basis weight 10 #include "magma-common-nontensor.h" 11 12 //////////////////////////////////////////////////////////////////////////////// 13 extern "C" __launch_bounds__(MAGMA_BASIS_BOUNDS(BASIS_Q, MAGMA_MAXTHREADS_1D)) __global__ 14 void magma_weight_nontensor(int n, const CeedScalar *__restrict__ dqweight, CeedScalar *__restrict__ dV) { 15 MAGMA_DEVICE_SHARED(CeedScalar, shared_data); 16 17 const int tx = threadIdx.x; 18 const int ty = threadIdx.y; 19 const int id = blockIdx.x * blockDim.y + ty; 20 21 // terminate threads with no work 22 if (id >= n) return; 23 24 dV += id * BASIS_Q; 25 26 // shared memory pointers 27 CeedScalar *sqweight = (CeedScalar *)shared_data; 28 CeedScalar *sV = sqweight + BASIS_Q; 29 sV += ty * BASIS_Q; 30 31 // read qweight 32 if (ty == 0 && tx < BASIS_Q) { 33 sqweight[tx] = dqweight[tx]; 34 } 35 __syncthreads(); 36 37 if (tx < BASIS_Q) { 38 sV[tx] = sqweight[tx]; 39 } 40 41 // write V 42 dV[tx] = sV[tx]; 43 } 44