xref: /libCEED/include/ceed/jit-source/hip/hip-ref-qfunction.h (revision b2165e7a2e371018feedcb47974a027ed47e0487)
1a0154adeSJed Brown // Copyright (c) 2017-2022, 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 
8*b2165e7aSSebastian Grimberg /// @file
9*b2165e7aSSebastian Grimberg /// Internal header for HIP backend QFunction read/write kernels
10*b2165e7aSSebastian Grimberg #ifndef _ceed_hip_ref_qfunction_h
11*b2165e7aSSebastian Grimberg #define _ceed_hip_ref_qfunction_h
12*b2165e7aSSebastian 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 //------------------------------------------------------------------------------
36*b2165e7aSSebastian Grimberg 
37*b2165e7aSSebastian Grimberg #endif
38