xref: /libCEED/include/ceed/jit-source/hip/hip-ref-qfunction.h (revision d4cc18453651bd0f94c1a2e078b2646a92dafdcc)
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 HIP backend QFunction read/write kernels
10 #include <ceed/types.h>
11 
12 //------------------------------------------------------------------------------
13 // Read from quadrature points
14 //------------------------------------------------------------------------------
15 template <int SIZE>
readQuads(const CeedInt quad,const CeedInt num_qpts,const CeedScalar * d_u,CeedScalar * r_u)16 inline __device__ void readQuads(const CeedInt quad, const CeedInt num_qpts, const CeedScalar *d_u, CeedScalar *r_u) {
17   for (CeedInt comp = 0; comp < SIZE; comp++) {
18     r_u[comp] = d_u[quad + num_qpts * comp];
19   }
20 }
21 
22 //------------------------------------------------------------------------------
23 // Write at quadrature points
24 //------------------------------------------------------------------------------
25 template <int SIZE>
writeQuads(const CeedInt quad,const CeedInt num_qpts,const CeedScalar * r_v,CeedScalar * d_v)26 inline __device__ void writeQuads(const CeedInt quad, const CeedInt num_qpts, const CeedScalar *r_v, CeedScalar *d_v) {
27   for (CeedInt comp = 0; comp < SIZE; comp++) {
28     d_v[quad + num_qpts * comp] = r_v[comp];
29   }
30 }
31