1*5aed82e4SJeremy L Thompson // Copyright (c) 2017-2024, Lawrence Livermore National Security, LLC and other CEED contributors. 2a0154adeSJed Brown // All Rights Reserved. See the top-level LICENSE and NOTICE files for details. 3a0154adeSJed Brown // 4a0154adeSJed Brown // SPDX-License-Identifier: BSD-2-Clause 5a0154adeSJed Brown // 6a0154adeSJed Brown // This file is part of CEED: http://github.com/ceed 7a0154adeSJed Brown 8b2165e7aSSebastian Grimberg /// @file 9b2165e7aSSebastian Grimberg /// Internal header for CUDA backend QFunction read/write kernels 1094b7b29bSJeremy L Thompson #ifndef CEED_CUDA_REF_QFUNCTION_H 1194b7b29bSJeremy L Thompson #define CEED_CUDA_REF_QFUNCTION_H 12b2165e7aSSebastian Grimberg 13c9c2c079SJeremy L Thompson #include <ceed.h> 14a0154adeSJed Brown 15a0154adeSJed Brown //------------------------------------------------------------------------------ 16a0154adeSJed Brown // Read from quadrature points 17a0154adeSJed Brown //------------------------------------------------------------------------------ 189bd0a4deSJeremy L Thompson template <int SIZE> 192b730f8bSJeremy L Thompson inline __device__ void readQuads(const CeedInt quad, const CeedInt num_qpts, const CeedScalar *d_u, CeedScalar *r_u) { 20a0154adeSJed Brown for (CeedInt comp = 0; comp < SIZE; comp++) { 21a0154adeSJed Brown r_u[comp] = d_u[quad + num_qpts * comp]; 22a0154adeSJed Brown } 23a0154adeSJed Brown } 24a0154adeSJed Brown 25a0154adeSJed Brown //------------------------------------------------------------------------------ 26a0154adeSJed Brown // Write at quadrature points 27a0154adeSJed Brown //------------------------------------------------------------------------------ 28a0154adeSJed Brown template <int SIZE> 292b730f8bSJeremy L Thompson inline __device__ void writeQuads(const CeedInt quad, const CeedInt num_qpts, const CeedScalar *r_v, CeedScalar *d_v) { 30a0154adeSJed Brown for (CeedInt comp = 0; comp < SIZE; comp++) { 31a0154adeSJed Brown d_v[quad + num_qpts * comp] = r_v[comp]; 32a0154adeSJed Brown } 33a0154adeSJed Brown } 34a0154adeSJed Brown 35a0154adeSJed Brown //------------------------------------------------------------------------------ 36b2165e7aSSebastian Grimberg 3794b7b29bSJeremy L Thompson #endif // CEED_CUDA_REF_QFUNCTION_H 38