xref: /libCEED/backends/hip-ref/ceed-hip-ref.h (revision 3d8e882215d238700cdceb37404f76ca7fa24eaa)
1*3d8e8822SJeremy L Thompson // Copyright (c) 2017-2022, Lawrence Livermore National Security, LLC and other CEED contributors.
2*3d8e8822SJeremy L Thompson // All Rights Reserved. See the top-level LICENSE and NOTICE files for details.
30d0321e0SJeremy L Thompson //
4*3d8e8822SJeremy L Thompson // SPDX-License-Identifier: BSD-2-Clause
50d0321e0SJeremy L Thompson //
6*3d8e8822SJeremy L Thompson // This file is part of CEED:  http://github.com/ceed
70d0321e0SJeremy L Thompson 
80d0321e0SJeremy L Thompson #ifndef _ceed_hip_h
90d0321e0SJeremy L Thompson #define _ceed_hip_h
100d0321e0SJeremy L Thompson 
110d0321e0SJeremy L Thompson #include <ceed/ceed.h>
120d0321e0SJeremy L Thompson #include <ceed/backend.h>
130d0321e0SJeremy L Thompson #include <hip/hip_runtime.h>
140d0321e0SJeremy L Thompson #include <hipblas.h>
150d0321e0SJeremy L Thompson #include "../hip/ceed-hip-common.h"
160d0321e0SJeremy L Thompson 
170d0321e0SJeremy L Thompson typedef struct {
180d0321e0SJeremy L Thompson   CeedScalar *h_array;
190d0321e0SJeremy L Thompson   CeedScalar *h_array_borrowed;
200d0321e0SJeremy L Thompson   CeedScalar *h_array_owned;
210d0321e0SJeremy L Thompson   CeedScalar *d_array;
220d0321e0SJeremy L Thompson   CeedScalar *d_array_borrowed;
230d0321e0SJeremy L Thompson   CeedScalar *d_array_owned;
240d0321e0SJeremy L Thompson } CeedVector_Hip;
250d0321e0SJeremy L Thompson 
260d0321e0SJeremy L Thompson typedef struct {
270d0321e0SJeremy L Thompson   hipModule_t module;
28437930d1SJeremy L Thompson   hipFunction_t StridedTranspose;
29437930d1SJeremy L Thompson   hipFunction_t StridedNoTranspose;
30437930d1SJeremy L Thompson   hipFunction_t OffsetTranspose;
31437930d1SJeremy L Thompson   hipFunction_t OffsetNoTranspose;
32437930d1SJeremy L Thompson   CeedInt num_nodes;
330d0321e0SJeremy L Thompson   CeedInt *h_ind;
340d0321e0SJeremy L Thompson   CeedInt *h_ind_allocated;
350d0321e0SJeremy L Thompson   CeedInt *d_ind;
360d0321e0SJeremy L Thompson   CeedInt *d_ind_allocated;
37437930d1SJeremy L Thompson   CeedInt *d_t_offsets;
38437930d1SJeremy L Thompson   CeedInt *d_t_indices;
39437930d1SJeremy L Thompson   CeedInt *d_l_vec_indices;
400d0321e0SJeremy L Thompson } CeedElemRestriction_Hip;
410d0321e0SJeremy L Thompson 
42437930d1SJeremy L Thompson typedef struct {
43437930d1SJeremy L Thompson   hipModule_t module;
44437930d1SJeremy L Thompson   hipFunction_t Interp;
45437930d1SJeremy L Thompson   hipFunction_t Grad;
46437930d1SJeremy L Thompson   hipFunction_t Weight;
47437930d1SJeremy L Thompson   CeedScalar *d_interp_1d;
48437930d1SJeremy L Thompson   CeedScalar *d_grad_1d;
49437930d1SJeremy L Thompson   CeedScalar *d_q_weight_1d;
50437930d1SJeremy L Thompson } CeedBasis_Hip;
51437930d1SJeremy L Thompson 
52437930d1SJeremy L Thompson typedef struct {
53437930d1SJeremy L Thompson   hipModule_t module;
54437930d1SJeremy L Thompson   hipFunction_t Interp;
55437930d1SJeremy L Thompson   hipFunction_t Grad;
56437930d1SJeremy L Thompson   hipFunction_t Weight;
57437930d1SJeremy L Thompson   CeedScalar *d_interp;
58437930d1SJeremy L Thompson   CeedScalar *d_grad;
59437930d1SJeremy L Thompson   CeedScalar *d_q_weight;
60437930d1SJeremy L Thompson } CeedBasisNonTensor_Hip;
61437930d1SJeremy L Thompson 
620d0321e0SJeremy L Thompson // We use a struct to avoid having to memCpy the array of pointers
630d0321e0SJeremy L Thompson // __global__ copies by value the struct.
640d0321e0SJeremy L Thompson typedef struct {
650d0321e0SJeremy L Thompson   const CeedScalar *inputs[CEED_FIELD_MAX];
660d0321e0SJeremy L Thompson   CeedScalar *outputs[CEED_FIELD_MAX];
670d0321e0SJeremy L Thompson } Fields_Hip;
680d0321e0SJeremy L Thompson 
690d0321e0SJeremy L Thompson typedef struct {
700d0321e0SJeremy L Thompson   hipModule_t module;
71437930d1SJeremy L Thompson   char *qfunction_name;
72437930d1SJeremy L Thompson   char *qfunction_source;
73437930d1SJeremy L Thompson   hipFunction_t QFunction;
740d0321e0SJeremy L Thompson   Fields_Hip fields;
750d0321e0SJeremy L Thompson   void *d_c;
760d0321e0SJeremy L Thompson } CeedQFunction_Hip;
770d0321e0SJeremy L Thompson 
780d0321e0SJeremy L Thompson typedef struct {
790d0321e0SJeremy L Thompson   void *h_data;
800d0321e0SJeremy L Thompson   void *h_data_borrowed;
810d0321e0SJeremy L Thompson   void *h_data_owned;
820d0321e0SJeremy L Thompson   void *d_data;
830d0321e0SJeremy L Thompson   void *d_data_borrowed;
840d0321e0SJeremy L Thompson   void *d_data_owned;
850d0321e0SJeremy L Thompson } CeedQFunctionContext_Hip;
860d0321e0SJeremy L Thompson 
870d0321e0SJeremy L Thompson typedef struct {
880d0321e0SJeremy L Thompson   hipModule_t module;
890d0321e0SJeremy L Thompson   hipFunction_t linearDiagonal;
900d0321e0SJeremy L Thompson   hipFunction_t linearPointBlock;
910d0321e0SJeremy L Thompson   CeedBasis basisin, basisout;
920d0321e0SJeremy L Thompson   CeedElemRestriction diagrstr, pbdiagrstr;
930d0321e0SJeremy L Thompson   CeedVector elemdiag, pbelemdiag;
940d0321e0SJeremy L Thompson   CeedInt numemodein, numemodeout, nnodes;
950d0321e0SJeremy L Thompson   CeedEvalMode *h_emodein, *h_emodeout;
960d0321e0SJeremy L Thompson   CeedEvalMode *d_emodein, *d_emodeout;
970d0321e0SJeremy L Thompson   CeedScalar *d_identity, *d_interpin, *d_interpout, *d_gradin, *d_gradout;
980d0321e0SJeremy L Thompson } CeedOperatorDiag_Hip;
990d0321e0SJeremy L Thompson 
1000d0321e0SJeremy L Thompson typedef struct {
1010d0321e0SJeremy L Thompson   CeedVector *evecs;   // E-vectors, inputs followed by outputs
1020d0321e0SJeremy L Thompson   CeedVector *qvecsin;    // Input Q-vectors needed to apply operator
1030d0321e0SJeremy L Thompson   CeedVector *qvecsout;   // Output Q-vectors needed to apply operator
1040d0321e0SJeremy L Thompson   CeedInt    numein;
1050d0321e0SJeremy L Thompson   CeedInt    numeout;
1060d0321e0SJeremy L Thompson   CeedInt    qfnumactivein, qfnumactiveout;
1070d0321e0SJeremy L Thompson   CeedVector *qfactivein;
1080d0321e0SJeremy L Thompson   CeedOperatorDiag_Hip *diag;
1090d0321e0SJeremy L Thompson } CeedOperator_Hip;
1100d0321e0SJeremy L Thompson 
1110d0321e0SJeremy L Thompson CEED_INTERN int CeedHipGetHipblasHandle(Ceed ceed, hipblasHandle_t *handle);
1120d0321e0SJeremy L Thompson 
1131f9221feSJeremy L Thompson CEED_INTERN int CeedVectorCreate_Hip(CeedSize n, CeedVector vec);
1140d0321e0SJeremy L Thompson 
1150d0321e0SJeremy L Thompson CEED_INTERN int CeedElemRestrictionCreate_Hip(CeedMemType mtype,
1160d0321e0SJeremy L Thompson     CeedCopyMode cmode, const CeedInt *indices, CeedElemRestriction r);
1170d0321e0SJeremy L Thompson 
1180d0321e0SJeremy L Thompson CEED_INTERN int CeedElemRestrictionCreateBlocked_Hip(const CeedMemType mtype,
1190d0321e0SJeremy L Thompson     const CeedCopyMode cmode, const CeedInt *indices,
1200d0321e0SJeremy L Thompson     const CeedElemRestriction res);
1210d0321e0SJeremy L Thompson 
1220d0321e0SJeremy L Thompson CEED_INTERN int CeedBasisApplyElems_Hip(CeedBasis basis, const CeedInt nelem,
1230d0321e0SJeremy L Thompson                                         CeedTransposeMode tmode, CeedEvalMode emode, const CeedVector u, CeedVector v);
1240d0321e0SJeremy L Thompson 
1250d0321e0SJeremy L Thompson CEED_INTERN int CeedQFunctionApplyElems_Hip(CeedQFunction qf, const CeedInt Q,
1260d0321e0SJeremy L Thompson     const CeedVector *const u, const CeedVector *v);
1270d0321e0SJeremy L Thompson 
1280d0321e0SJeremy L Thompson CEED_INTERN int CeedBasisCreateTensorH1_Hip(CeedInt dim, CeedInt P1d,
1290d0321e0SJeremy L Thompson     CeedInt Q1d,
1300d0321e0SJeremy L Thompson     const CeedScalar *interp1d,
1310d0321e0SJeremy L Thompson     const CeedScalar *grad1d,
1320d0321e0SJeremy L Thompson     const CeedScalar *qref1d,
1330d0321e0SJeremy L Thompson     const CeedScalar *qweight1d,
1340d0321e0SJeremy L Thompson     CeedBasis basis);
1350d0321e0SJeremy L Thompson 
1360d0321e0SJeremy L Thompson CEED_INTERN int CeedBasisCreateH1_Hip(CeedElemTopology, CeedInt, CeedInt,
1370d0321e0SJeremy L Thompson                                       CeedInt, const CeedScalar *,
1380d0321e0SJeremy L Thompson                                       const CeedScalar *, const CeedScalar *,
1390d0321e0SJeremy L Thompson                                       const CeedScalar *, CeedBasis);
1400d0321e0SJeremy L Thompson 
1410d0321e0SJeremy L Thompson CEED_INTERN int CeedQFunctionCreate_Hip(CeedQFunction qf);
1420d0321e0SJeremy L Thompson 
1430d0321e0SJeremy L Thompson CEED_INTERN int CeedQFunctionContextCreate_Hip(CeedQFunctionContext ctx);
1440d0321e0SJeremy L Thompson 
1450d0321e0SJeremy L Thompson CEED_INTERN int CeedOperatorCreate_Hip(CeedOperator op);
1460d0321e0SJeremy L Thompson 
1470d0321e0SJeremy L Thompson CEED_INTERN int CeedCompositeOperatorCreate_Hip(CeedOperator op);
1480d0321e0SJeremy L Thompson #endif
149