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