xref: /libCEED/rust/libceed-sys/c-src/backends/hip-ref/ceed-hip-ref.h (revision 51475c7c4e99ed8faf0a644c51f7b001cf768463)
13d8e8822SJeremy L Thompson // Copyright (c) 2017-2022, Lawrence Livermore National Security, LLC and other CEED contributors.
23d8e8822SJeremy L Thompson // All Rights Reserved. See the top-level LICENSE and NOTICE files for details.
30d0321e0SJeremy L Thompson //
43d8e8822SJeremy L Thompson // SPDX-License-Identifier: BSD-2-Clause
50d0321e0SJeremy L Thompson //
63d8e8822SJeremy 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 
1149aac155SJeremy L Thompson #include <ceed.h>
120d0321e0SJeremy L Thompson #include <ceed/backend.h>
1349aac155SJeremy L Thompson #include <ceed/jit-source/hip/hip-types.h>
140d0321e0SJeremy L Thompson #include <hip/hip_runtime.h>
152b730f8bSJeremy L Thompson 
160d0321e0SJeremy L Thompson #include "../hip/ceed-hip-common.h"
170d0321e0SJeremy L Thompson 
180d0321e0SJeremy L Thompson typedef struct {
190d0321e0SJeremy L Thompson   CeedScalar *h_array;
200d0321e0SJeremy L Thompson   CeedScalar *h_array_borrowed;
210d0321e0SJeremy L Thompson   CeedScalar *h_array_owned;
220d0321e0SJeremy L Thompson   CeedScalar *d_array;
230d0321e0SJeremy L Thompson   CeedScalar *d_array_borrowed;
240d0321e0SJeremy L Thompson   CeedScalar *d_array_owned;
250d0321e0SJeremy L Thompson } CeedVector_Hip;
260d0321e0SJeremy L Thompson 
270d0321e0SJeremy L Thompson typedef struct {
280d0321e0SJeremy L Thompson   hipModule_t   module;
29437930d1SJeremy L Thompson   hipFunction_t StridedTranspose;
30437930d1SJeremy L Thompson   hipFunction_t StridedNoTranspose;
31437930d1SJeremy L Thompson   hipFunction_t OffsetTranspose;
32437930d1SJeremy L Thompson   hipFunction_t OffsetNoTranspose;
33437930d1SJeremy L Thompson   CeedInt       num_nodes;
340d0321e0SJeremy L Thompson   CeedInt      *h_ind;
350d0321e0SJeremy L Thompson   CeedInt      *h_ind_allocated;
360d0321e0SJeremy L Thompson   CeedInt      *d_ind;
370d0321e0SJeremy L Thompson   CeedInt      *d_ind_allocated;
38437930d1SJeremy L Thompson   CeedInt      *d_t_offsets;
39437930d1SJeremy L Thompson   CeedInt      *d_t_indices;
40437930d1SJeremy L Thompson   CeedInt      *d_l_vec_indices;
410d0321e0SJeremy L Thompson } CeedElemRestriction_Hip;
420d0321e0SJeremy L Thompson 
43437930d1SJeremy L Thompson typedef struct {
44437930d1SJeremy L Thompson   hipModule_t   module;
45437930d1SJeremy L Thompson   hipFunction_t Interp;
46437930d1SJeremy L Thompson   hipFunction_t Grad;
47437930d1SJeremy L Thompson   hipFunction_t Weight;
48437930d1SJeremy L Thompson   CeedScalar   *d_interp_1d;
49437930d1SJeremy L Thompson   CeedScalar   *d_grad_1d;
50437930d1SJeremy L Thompson   CeedScalar   *d_q_weight_1d;
51437930d1SJeremy L Thompson } CeedBasis_Hip;
52437930d1SJeremy L Thompson 
53437930d1SJeremy L Thompson typedef struct {
54437930d1SJeremy L Thompson   hipModule_t   module;
55437930d1SJeremy L Thompson   hipFunction_t Interp;
56437930d1SJeremy L Thompson   hipFunction_t Grad;
57437930d1SJeremy L Thompson   hipFunction_t Weight;
58437930d1SJeremy L Thompson   CeedScalar   *d_interp;
59437930d1SJeremy L Thompson   CeedScalar   *d_grad;
60437930d1SJeremy L Thompson   CeedScalar   *d_q_weight;
61437930d1SJeremy L Thompson } CeedBasisNonTensor_Hip;
62437930d1SJeremy L Thompson 
630d0321e0SJeremy L Thompson typedef struct {
640d0321e0SJeremy L Thompson   hipModule_t   module;
65437930d1SJeremy L Thompson   char         *qfunction_name;
66437930d1SJeremy L Thompson   char         *qfunction_source;
67437930d1SJeremy L Thompson   hipFunction_t QFunction;
680d0321e0SJeremy L Thompson   Fields_Hip    fields;
690d0321e0SJeremy L Thompson   void         *d_c;
700d0321e0SJeremy L Thompson } CeedQFunction_Hip;
710d0321e0SJeremy L Thompson 
720d0321e0SJeremy L Thompson typedef struct {
730d0321e0SJeremy L Thompson   void *h_data;
740d0321e0SJeremy L Thompson   void *h_data_borrowed;
750d0321e0SJeremy L Thompson   void *h_data_owned;
760d0321e0SJeremy L Thompson   void *d_data;
770d0321e0SJeremy L Thompson   void *d_data_borrowed;
780d0321e0SJeremy L Thompson   void *d_data_owned;
790d0321e0SJeremy L Thompson } CeedQFunctionContext_Hip;
800d0321e0SJeremy L Thompson 
810d0321e0SJeremy L Thompson typedef struct {
820d0321e0SJeremy L Thompson   hipModule_t         module;
830d0321e0SJeremy L Thompson   hipFunction_t       linearDiagonal;
840d0321e0SJeremy L Thompson   hipFunction_t       linearPointBlock;
850d0321e0SJeremy L Thompson   CeedBasis           basisin, basisout;
860d0321e0SJeremy L Thompson   CeedElemRestriction diagrstr, pbdiagrstr;
870d0321e0SJeremy L Thompson   CeedVector          elemdiag, pbelemdiag;
880d0321e0SJeremy L Thompson   CeedInt             numemodein, numemodeout, nnodes;
890d0321e0SJeremy L Thompson   CeedEvalMode       *h_emodein, *h_emodeout;
900d0321e0SJeremy L Thompson   CeedEvalMode       *d_emodein, *d_emodeout;
910d0321e0SJeremy L Thompson   CeedScalar         *d_identity, *d_interpin, *d_interpout, *d_gradin, *d_gradout;
920d0321e0SJeremy L Thompson } CeedOperatorDiag_Hip;
930d0321e0SJeremy L Thompson 
940d0321e0SJeremy L Thompson typedef struct {
95a835093fSnbeams   hipModule_t   module;
96a835093fSnbeams   hipFunction_t linearAssemble;
9759ad764aSnbeams   CeedInt       nelem, block_size_x, block_size_y, elemsPerBlock;
98a835093fSnbeams   CeedScalar   *d_B_in, *d_B_out;
99a835093fSnbeams } CeedOperatorAssemble_Hip;
100a835093fSnbeams 
101a835093fSnbeams typedef struct {
1020d0321e0SJeremy L Thompson   CeedVector               *evecs;     // E-vectors, inputs followed by outputs
1030d0321e0SJeremy L Thompson   CeedVector               *qvecsin;   // Input Q-vectors needed to apply operator
1040d0321e0SJeremy L Thompson   CeedVector               *qvecsout;  // Output Q-vectors needed to apply operator
1050d0321e0SJeremy L Thompson   CeedInt                   numein;
1060d0321e0SJeremy L Thompson   CeedInt                   numeout;
1070d0321e0SJeremy L Thompson   CeedInt                   qfnumactivein, qfnumactiveout;
1080d0321e0SJeremy L Thompson   CeedVector               *qfactivein;
1090d0321e0SJeremy L Thompson   CeedOperatorDiag_Hip     *diag;
110a835093fSnbeams   CeedOperatorAssemble_Hip *asmb;
1110d0321e0SJeremy L Thompson } CeedOperator_Hip;
1120d0321e0SJeremy L Thompson 
1130d0321e0SJeremy L Thompson CEED_INTERN int CeedHipGetHipblasHandle(Ceed ceed, hipblasHandle_t *handle);
1140d0321e0SJeremy L Thompson 
1151f9221feSJeremy L Thompson CEED_INTERN int CeedVectorCreate_Hip(CeedSize n, CeedVector vec);
1160d0321e0SJeremy L Thompson 
1172b730f8bSJeremy L Thompson CEED_INTERN int CeedElemRestrictionCreate_Hip(CeedMemType mtype, CeedCopyMode cmode, const CeedInt *indices, CeedElemRestriction r);
1180d0321e0SJeremy L Thompson 
119*51475c7cSJeremy L Thompson CEED_INTERN int CeedElemRestrictionCreateBlocked_Hip(CeedMemType mtype, CeedCopyMode cmode, const CeedInt *indices, CeedElemRestriction res);
1200d0321e0SJeremy L Thompson 
121*51475c7cSJeremy L Thompson CEED_INTERN int CeedBasisApplyElems_Hip(CeedBasis basis, CeedInt nelem, CeedTransposeMode tmode, CeedEvalMode emode, const CeedVector u,
1222b730f8bSJeremy L Thompson                                         CeedVector v);
1230d0321e0SJeremy L Thompson 
124*51475c7cSJeremy L Thompson CEED_INTERN int CeedQFunctionApplyElems_Hip(CeedQFunction qf, CeedInt Q, const CeedVector *const u, const CeedVector *v);
1250d0321e0SJeremy L Thompson 
1266574a04fSJeremy L Thompson CEED_INTERN int CeedBasisCreateTensorH1_Hip(CeedInt dim, CeedInt P_1d, CeedInt Q_1d, const CeedScalar *interp_1d, const CeedScalar *grad_1d,
1276574a04fSJeremy L Thompson                                             const CeedScalar *q_ref_1d, const CeedScalar *q_weight_1d, CeedBasis basis);
1280d0321e0SJeremy L Thompson 
129*51475c7cSJeremy L Thompson CEED_INTERN int CeedBasisCreateH1_Hip(CeedElemTopology topo, CeedInt dim, CeedInt num_nodes, CeedInt num_qpts, const CeedScalar *interp,
130*51475c7cSJeremy L Thompson                                       const CeedScalar *grad, const CeedScalar *q_ref, const CeedScalar *q_weight, CeedBasis basis);
1310d0321e0SJeremy L Thompson 
1320d0321e0SJeremy L Thompson CEED_INTERN int CeedQFunctionCreate_Hip(CeedQFunction qf);
1330d0321e0SJeremy L Thompson 
1340d0321e0SJeremy L Thompson CEED_INTERN int CeedQFunctionContextCreate_Hip(CeedQFunctionContext ctx);
1350d0321e0SJeremy L Thompson 
1360d0321e0SJeremy L Thompson CEED_INTERN int CeedOperatorCreate_Hip(CeedOperator op);
1370d0321e0SJeremy L Thompson 
1380d0321e0SJeremy L Thompson #endif
139