xref: /libCEED/rust/libceed-sys/c-src/backends/hip-ref/ceed-hip-ref.h (revision 0d0321e0e600f17fbb9528732fcb5c1d5c63fc0f)
1*0d0321e0SJeremy L Thompson // Copyright (c) 2017-2018, Lawrence Livermore National Security, LLC.
2*0d0321e0SJeremy L Thompson // Produced at the Lawrence Livermore National Laboratory. LLNL-CODE-734707.
3*0d0321e0SJeremy L Thompson // All Rights reserved. See files LICENSE and NOTICE for details.
4*0d0321e0SJeremy L Thompson //
5*0d0321e0SJeremy L Thompson // This file is part of CEED, a collection of benchmarks, miniapps, software
6*0d0321e0SJeremy L Thompson // libraries and APIs for efficient high-order finite element and spectral
7*0d0321e0SJeremy L Thompson // element discretizations for exascale applications. For more information and
8*0d0321e0SJeremy L Thompson // source code availability see http://github.com/ceed.
9*0d0321e0SJeremy L Thompson //
10*0d0321e0SJeremy L Thompson // The CEED research is supported by the Exascale Computing Project 17-SC-20-SC,
11*0d0321e0SJeremy L Thompson // a collaborative effort of two U.S. Department of Energy organizations (Office
12*0d0321e0SJeremy L Thompson // of Science and the National Nuclear Security Administration) responsible for
13*0d0321e0SJeremy L Thompson // the planning and preparation of a capable exascale ecosystem, including
14*0d0321e0SJeremy L Thompson // software, applications, hardware, advanced system engineering and early
15*0d0321e0SJeremy L Thompson // testbed platforms, in support of the nation's exascale computing imperative.
16*0d0321e0SJeremy L Thompson 
17*0d0321e0SJeremy L Thompson #ifndef _ceed_hip_h
18*0d0321e0SJeremy L Thompson #define _ceed_hip_h
19*0d0321e0SJeremy L Thompson 
20*0d0321e0SJeremy L Thompson #include <ceed/ceed.h>
21*0d0321e0SJeremy L Thompson #include <ceed/backend.h>
22*0d0321e0SJeremy L Thompson #include <hip/hip_runtime.h>
23*0d0321e0SJeremy L Thompson #include <hipblas.h>
24*0d0321e0SJeremy L Thompson #include "../hip/ceed-hip-common.h"
25*0d0321e0SJeremy L Thompson 
26*0d0321e0SJeremy L Thompson typedef struct {
27*0d0321e0SJeremy L Thompson   CeedScalar *h_array;
28*0d0321e0SJeremy L Thompson   CeedScalar *h_array_borrowed;
29*0d0321e0SJeremy L Thompson   CeedScalar *h_array_owned;
30*0d0321e0SJeremy L Thompson   CeedScalar *d_array;
31*0d0321e0SJeremy L Thompson   CeedScalar *d_array_borrowed;
32*0d0321e0SJeremy L Thompson   CeedScalar *d_array_owned;
33*0d0321e0SJeremy L Thompson } CeedVector_Hip;
34*0d0321e0SJeremy L Thompson 
35*0d0321e0SJeremy L Thompson typedef struct {
36*0d0321e0SJeremy L Thompson   hipModule_t module;
37*0d0321e0SJeremy L Thompson   hipFunction_t noTrStrided;
38*0d0321e0SJeremy L Thompson   hipFunction_t noTrOffset;
39*0d0321e0SJeremy L Thompson   hipFunction_t trStrided;
40*0d0321e0SJeremy L Thompson   hipFunction_t trOffset;
41*0d0321e0SJeremy L Thompson   CeedInt nnodes;
42*0d0321e0SJeremy L Thompson   CeedInt *h_ind;
43*0d0321e0SJeremy L Thompson   CeedInt *h_ind_allocated;
44*0d0321e0SJeremy L Thompson   CeedInt *d_ind;
45*0d0321e0SJeremy L Thompson   CeedInt *d_ind_allocated;
46*0d0321e0SJeremy L Thompson   CeedInt *d_toffsets;
47*0d0321e0SJeremy L Thompson   CeedInt *d_tindices;
48*0d0321e0SJeremy L Thompson   CeedInt *d_lvec_indices;
49*0d0321e0SJeremy L Thompson } CeedElemRestriction_Hip;
50*0d0321e0SJeremy L Thompson 
51*0d0321e0SJeremy L Thompson // We use a struct to avoid having to memCpy the array of pointers
52*0d0321e0SJeremy L Thompson // __global__ copies by value the struct.
53*0d0321e0SJeremy L Thompson typedef struct {
54*0d0321e0SJeremy L Thompson   const CeedScalar *inputs[CEED_FIELD_MAX];
55*0d0321e0SJeremy L Thompson   CeedScalar *outputs[CEED_FIELD_MAX];
56*0d0321e0SJeremy L Thompson } Fields_Hip;
57*0d0321e0SJeremy L Thompson 
58*0d0321e0SJeremy L Thompson typedef struct {
59*0d0321e0SJeremy L Thompson   hipModule_t module;
60*0d0321e0SJeremy L Thompson   char *qFunctionName;
61*0d0321e0SJeremy L Thompson   char *qFunctionSource;
62*0d0321e0SJeremy L Thompson   hipFunction_t qFunction;
63*0d0321e0SJeremy L Thompson   Fields_Hip fields;
64*0d0321e0SJeremy L Thompson   void *d_c;
65*0d0321e0SJeremy L Thompson } CeedQFunction_Hip;
66*0d0321e0SJeremy L Thompson 
67*0d0321e0SJeremy L Thompson typedef struct {
68*0d0321e0SJeremy L Thompson   void *h_data;
69*0d0321e0SJeremy L Thompson   void *h_data_borrowed;
70*0d0321e0SJeremy L Thompson   void *h_data_owned;
71*0d0321e0SJeremy L Thompson   void *d_data;
72*0d0321e0SJeremy L Thompson   void *d_data_borrowed;
73*0d0321e0SJeremy L Thompson   void *d_data_owned;
74*0d0321e0SJeremy L Thompson } CeedQFunctionContext_Hip;
75*0d0321e0SJeremy L Thompson 
76*0d0321e0SJeremy L Thompson typedef struct {
77*0d0321e0SJeremy L Thompson   hipModule_t module;
78*0d0321e0SJeremy L Thompson   hipFunction_t interp;
79*0d0321e0SJeremy L Thompson   hipFunction_t grad;
80*0d0321e0SJeremy L Thompson   hipFunction_t weight;
81*0d0321e0SJeremy L Thompson   CeedScalar *d_interp1d;
82*0d0321e0SJeremy L Thompson   CeedScalar *d_grad1d;
83*0d0321e0SJeremy L Thompson   CeedScalar *d_qweight1d;
84*0d0321e0SJeremy L Thompson } CeedBasis_Hip;
85*0d0321e0SJeremy L Thompson 
86*0d0321e0SJeremy L Thompson typedef struct {
87*0d0321e0SJeremy L Thompson   hipModule_t module;
88*0d0321e0SJeremy L Thompson   hipFunction_t interp;
89*0d0321e0SJeremy L Thompson   hipFunction_t grad;
90*0d0321e0SJeremy L Thompson   hipFunction_t weight;
91*0d0321e0SJeremy L Thompson   CeedScalar *d_interp;
92*0d0321e0SJeremy L Thompson   CeedScalar *d_grad;
93*0d0321e0SJeremy L Thompson   CeedScalar *d_qweight;
94*0d0321e0SJeremy L Thompson } CeedBasisNonTensor_Hip;
95*0d0321e0SJeremy L Thompson 
96*0d0321e0SJeremy L Thompson typedef struct {
97*0d0321e0SJeremy L Thompson   hipModule_t module;
98*0d0321e0SJeremy L Thompson   hipFunction_t linearDiagonal;
99*0d0321e0SJeremy L Thompson   hipFunction_t linearPointBlock;
100*0d0321e0SJeremy L Thompson   CeedBasis basisin, basisout;
101*0d0321e0SJeremy L Thompson   CeedElemRestriction diagrstr, pbdiagrstr;
102*0d0321e0SJeremy L Thompson   CeedVector elemdiag, pbelemdiag;
103*0d0321e0SJeremy L Thompson   CeedInt numemodein, numemodeout, nnodes;
104*0d0321e0SJeremy L Thompson   CeedEvalMode *h_emodein, *h_emodeout;
105*0d0321e0SJeremy L Thompson   CeedEvalMode *d_emodein, *d_emodeout;
106*0d0321e0SJeremy L Thompson   CeedScalar *d_identity, *d_interpin, *d_interpout, *d_gradin, *d_gradout;
107*0d0321e0SJeremy L Thompson } CeedOperatorDiag_Hip;
108*0d0321e0SJeremy L Thompson 
109*0d0321e0SJeremy L Thompson typedef struct {
110*0d0321e0SJeremy L Thompson   CeedVector *evecs;   // E-vectors, inputs followed by outputs
111*0d0321e0SJeremy L Thompson   CeedVector *qvecsin;    // Input Q-vectors needed to apply operator
112*0d0321e0SJeremy L Thompson   CeedVector *qvecsout;   // Output Q-vectors needed to apply operator
113*0d0321e0SJeremy L Thompson   CeedInt    numein;
114*0d0321e0SJeremy L Thompson   CeedInt    numeout;
115*0d0321e0SJeremy L Thompson   CeedInt    qfnumactivein, qfnumactiveout;
116*0d0321e0SJeremy L Thompson   CeedVector *qfactivein;
117*0d0321e0SJeremy L Thompson   CeedOperatorDiag_Hip *diag;
118*0d0321e0SJeremy L Thompson } CeedOperator_Hip;
119*0d0321e0SJeremy L Thompson 
120*0d0321e0SJeremy L Thompson CEED_INTERN int CeedHipGetHipblasHandle(Ceed ceed, hipblasHandle_t *handle);
121*0d0321e0SJeremy L Thompson 
122*0d0321e0SJeremy L Thompson CEED_INTERN int CeedVectorCreate_Hip(CeedInt n, CeedVector vec);
123*0d0321e0SJeremy L Thompson 
124*0d0321e0SJeremy L Thompson CEED_INTERN int CeedElemRestrictionCreate_Hip(CeedMemType mtype,
125*0d0321e0SJeremy L Thompson     CeedCopyMode cmode, const CeedInt *indices, CeedElemRestriction r);
126*0d0321e0SJeremy L Thompson 
127*0d0321e0SJeremy L Thompson CEED_INTERN int CeedElemRestrictionCreateBlocked_Hip(const CeedMemType mtype,
128*0d0321e0SJeremy L Thompson     const CeedCopyMode cmode, const CeedInt *indices,
129*0d0321e0SJeremy L Thompson     const CeedElemRestriction res);
130*0d0321e0SJeremy L Thompson 
131*0d0321e0SJeremy L Thompson CEED_INTERN int CeedBasisApplyElems_Hip(CeedBasis basis, const CeedInt nelem,
132*0d0321e0SJeremy L Thompson                                         CeedTransposeMode tmode, CeedEvalMode emode, const CeedVector u, CeedVector v);
133*0d0321e0SJeremy L Thompson 
134*0d0321e0SJeremy L Thompson CEED_INTERN int CeedQFunctionApplyElems_Hip(CeedQFunction qf, const CeedInt Q,
135*0d0321e0SJeremy L Thompson     const CeedVector *const u, const CeedVector *v);
136*0d0321e0SJeremy L Thompson 
137*0d0321e0SJeremy L Thompson CEED_INTERN int CeedBasisCreateTensorH1_Hip(CeedInt dim, CeedInt P1d,
138*0d0321e0SJeremy L Thompson     CeedInt Q1d,
139*0d0321e0SJeremy L Thompson     const CeedScalar *interp1d,
140*0d0321e0SJeremy L Thompson     const CeedScalar *grad1d,
141*0d0321e0SJeremy L Thompson     const CeedScalar *qref1d,
142*0d0321e0SJeremy L Thompson     const CeedScalar *qweight1d,
143*0d0321e0SJeremy L Thompson     CeedBasis basis);
144*0d0321e0SJeremy L Thompson 
145*0d0321e0SJeremy L Thompson CEED_INTERN int CeedBasisCreateH1_Hip(CeedElemTopology, CeedInt, CeedInt,
146*0d0321e0SJeremy L Thompson                                       CeedInt, const CeedScalar *,
147*0d0321e0SJeremy L Thompson                                       const CeedScalar *, const CeedScalar *,
148*0d0321e0SJeremy L Thompson                                       const CeedScalar *, CeedBasis);
149*0d0321e0SJeremy L Thompson 
150*0d0321e0SJeremy L Thompson CEED_INTERN int CeedQFunctionCreate_Hip(CeedQFunction qf);
151*0d0321e0SJeremy L Thompson 
152*0d0321e0SJeremy L Thompson CEED_INTERN int CeedQFunctionContextCreate_Hip(CeedQFunctionContext ctx);
153*0d0321e0SJeremy L Thompson 
154*0d0321e0SJeremy L Thompson CEED_INTERN int CeedOperatorCreate_Hip(CeedOperator op);
155*0d0321e0SJeremy L Thompson 
156*0d0321e0SJeremy L Thompson CEED_INTERN int CeedCompositeOperatorCreate_Hip(CeedOperator op);
157*0d0321e0SJeremy L Thompson #endif
158