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