xref: /libCEED/backends/cuda-ref/ceed-cuda-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_cuda_h
90d0321e0SJeremy L Thompson #define _ceed_cuda_h
100d0321e0SJeremy L Thompson 
110d0321e0SJeremy L Thompson #include <ceed/ceed.h>
120d0321e0SJeremy L Thompson #include <ceed/backend.h>
130d0321e0SJeremy L Thompson #include <cuda.h>
140d0321e0SJeremy L Thompson #include "../cuda/ceed-cuda-common.h"
150d0321e0SJeremy L Thompson 
160d0321e0SJeremy L Thompson typedef struct {
170d0321e0SJeremy L Thompson   CeedScalar *h_array;
180d0321e0SJeremy L Thompson   CeedScalar *h_array_borrowed;
190d0321e0SJeremy L Thompson   CeedScalar *h_array_owned;
200d0321e0SJeremy L Thompson   CeedScalar *d_array;
210d0321e0SJeremy L Thompson   CeedScalar *d_array_borrowed;
220d0321e0SJeremy L Thompson   CeedScalar *d_array_owned;
230d0321e0SJeremy L Thompson } CeedVector_Cuda;
240d0321e0SJeremy L Thompson 
250d0321e0SJeremy L Thompson typedef struct {
260d0321e0SJeremy L Thompson   CUmodule module;
27437930d1SJeremy L Thompson   CUfunction StridedTranspose;
28437930d1SJeremy L Thompson   CUfunction StridedNoTranspose;
29437930d1SJeremy L Thompson   CUfunction OffsetTranspose;
30437930d1SJeremy L Thompson   CUfunction OffsetNoTranspose;
31437930d1SJeremy L Thompson   CeedInt num_nodes;
320d0321e0SJeremy L Thompson   CeedInt *h_ind;
330d0321e0SJeremy L Thompson   CeedInt *h_ind_allocated;
340d0321e0SJeremy L Thompson   CeedInt *d_ind;
350d0321e0SJeremy L Thompson   CeedInt *d_ind_allocated;
36437930d1SJeremy L Thompson   CeedInt *d_t_offsets;
37437930d1SJeremy L Thompson   CeedInt *d_t_indices;
38437930d1SJeremy L Thompson   CeedInt *d_l_vec_indices;
390d0321e0SJeremy L Thompson } CeedElemRestriction_Cuda;
400d0321e0SJeremy L Thompson 
41437930d1SJeremy L Thompson typedef struct {
42437930d1SJeremy L Thompson   CUmodule module;
43437930d1SJeremy L Thompson   CUfunction Interp;
44437930d1SJeremy L Thompson   CUfunction Grad;
45437930d1SJeremy L Thompson   CUfunction Weight;
46437930d1SJeremy L Thompson   CeedScalar *d_interp_1d;
47437930d1SJeremy L Thompson   CeedScalar *d_grad_1d;
48437930d1SJeremy L Thompson   CeedScalar *d_q_weight_1d;
49437930d1SJeremy L Thompson } CeedBasis_Cuda;
50437930d1SJeremy L Thompson 
51437930d1SJeremy L Thompson typedef struct {
52437930d1SJeremy L Thompson   CUmodule module;
53437930d1SJeremy L Thompson   CUfunction Interp;
54437930d1SJeremy L Thompson   CUfunction Grad;
55437930d1SJeremy L Thompson   CUfunction Weight;
56437930d1SJeremy L Thompson   CeedScalar *d_interp;
57437930d1SJeremy L Thompson   CeedScalar *d_grad;
58437930d1SJeremy L Thompson   CeedScalar *d_q_weight;
59437930d1SJeremy L Thompson } CeedBasisNonTensor_Cuda;
60437930d1SJeremy L Thompson 
610d0321e0SJeremy L Thompson // We use a struct to avoid having to memCpy the array of pointers
620d0321e0SJeremy L Thompson // __global__ copies by value the struct.
630d0321e0SJeremy L Thompson typedef struct {
640d0321e0SJeremy L Thompson   const CeedScalar *inputs[CEED_FIELD_MAX];
650d0321e0SJeremy L Thompson   CeedScalar *outputs[CEED_FIELD_MAX];
660d0321e0SJeremy L Thompson } Fields_Cuda;
670d0321e0SJeremy L Thompson 
680d0321e0SJeremy L Thompson typedef struct {
690d0321e0SJeremy L Thompson   CUmodule module;
70437930d1SJeremy L Thompson   char *qfunction_name;
71437930d1SJeremy L Thompson   char *qfunction_source;
72437930d1SJeremy L Thompson   CUfunction QFunction;
730d0321e0SJeremy L Thompson   Fields_Cuda fields;
740d0321e0SJeremy L Thompson   void *d_c;
750d0321e0SJeremy L Thompson } CeedQFunction_Cuda;
760d0321e0SJeremy L Thompson 
770d0321e0SJeremy L Thompson typedef struct {
780d0321e0SJeremy L Thompson   void *h_data;
790d0321e0SJeremy L Thompson   void *h_data_borrowed;
800d0321e0SJeremy L Thompson   void *h_data_owned;
810d0321e0SJeremy L Thompson   void *d_data;
820d0321e0SJeremy L Thompson   void *d_data_borrowed;
830d0321e0SJeremy L Thompson   void *d_data_owned;
840d0321e0SJeremy L Thompson } CeedQFunctionContext_Cuda;
850d0321e0SJeremy L Thompson 
860d0321e0SJeremy L Thompson typedef struct {
870d0321e0SJeremy L Thompson   CUmodule module;
880d0321e0SJeremy L Thompson   CUfunction linearDiagonal;
890d0321e0SJeremy L Thompson   CUfunction linearPointBlock;
900d0321e0SJeremy L Thompson   CeedBasis basisin, basisout;
910d0321e0SJeremy L Thompson   CeedElemRestriction diagrstr, pbdiagrstr;
920d0321e0SJeremy L Thompson   CeedVector elemdiag, pbelemdiag;
930d0321e0SJeremy L Thompson   CeedInt numemodein, numemodeout, nnodes;
940d0321e0SJeremy L Thompson   CeedEvalMode *h_emodein, *h_emodeout;
950d0321e0SJeremy L Thompson   CeedEvalMode *d_emodein, *d_emodeout;
960d0321e0SJeremy L Thompson   CeedScalar *d_identity, *d_interpin, *d_interpout, *d_gradin, *d_gradout;
970d0321e0SJeremy L Thompson } CeedOperatorDiag_Cuda;
980d0321e0SJeremy L Thompson 
990d0321e0SJeremy L Thompson typedef struct {
1000d0321e0SJeremy L Thompson   CeedVector *evecs;   // E-vectors, inputs followed by outputs
1010d0321e0SJeremy L Thompson   CeedVector *qvecsin;    // Input Q-vectors needed to apply operator
1020d0321e0SJeremy L Thompson   CeedVector *qvecsout;   // Output Q-vectors needed to apply operator
1030d0321e0SJeremy L Thompson   CeedInt    numein;
1040d0321e0SJeremy L Thompson   CeedInt    numeout;
1050d0321e0SJeremy L Thompson   CeedInt    qfnumactivein, qfnumactiveout;
1060d0321e0SJeremy L Thompson   CeedVector *qfactivein;
1070d0321e0SJeremy L Thompson   CeedOperatorDiag_Cuda *diag;
1080d0321e0SJeremy L Thompson } CeedOperator_Cuda;
1090d0321e0SJeremy L Thompson 
1100d0321e0SJeremy L Thompson CEED_INTERN int CeedCudaGetCublasHandle(Ceed ceed, cublasHandle_t *handle);
1110d0321e0SJeremy L Thompson 
1121f9221feSJeremy L Thompson CEED_INTERN int CeedVectorCreate_Cuda(CeedSize n, CeedVector vec);
1130d0321e0SJeremy L Thompson 
114437930d1SJeremy L Thompson CEED_INTERN int CeedElemRestrictionCreate_Cuda(CeedMemType mem_type,
115437930d1SJeremy L Thompson     CeedCopyMode copy_mode, const CeedInt *indices, CeedElemRestriction r);
1160d0321e0SJeremy L Thompson 
117437930d1SJeremy L Thompson CEED_INTERN int CeedElemRestrictionCreateBlocked_Cuda(const CeedMemType
118437930d1SJeremy L Thompson     mem_type,
119437930d1SJeremy L Thompson     const CeedCopyMode copy_mode, const CeedInt *indices,
1200d0321e0SJeremy L Thompson     const CeedElemRestriction res);
1210d0321e0SJeremy L Thompson 
122437930d1SJeremy L Thompson CEED_INTERN int CeedBasisApplyElems_Cuda(CeedBasis basis,
123437930d1SJeremy L Thompson     const CeedInt num_elem,
124437930d1SJeremy L Thompson     CeedTransposeMode t_mode, CeedEvalMode eval_mode, const CeedVector u,
125437930d1SJeremy L Thompson     CeedVector v);
1260d0321e0SJeremy L Thompson 
1270d0321e0SJeremy L Thompson CEED_INTERN int CeedQFunctionApplyElems_Cuda(CeedQFunction qf, const CeedInt Q,
1280d0321e0SJeremy L Thompson     const CeedVector *const u, const CeedVector *v);
1290d0321e0SJeremy L Thompson 
130437930d1SJeremy L Thompson CEED_INTERN int CeedBasisCreateTensorH1_Cuda(CeedInt dim, CeedInt P_1d,
131437930d1SJeremy L Thompson     CeedInt Q_1d, const CeedScalar *interp_1d, const CeedScalar *grad_1d,
132437930d1SJeremy L Thompson     const CeedScalar *qref_1d, const CeedScalar *qweight_1d, CeedBasis basis);
1330d0321e0SJeremy L Thompson 
1340d0321e0SJeremy L Thompson CEED_INTERN int CeedBasisCreateH1_Cuda(CeedElemTopology, CeedInt, CeedInt,
1350d0321e0SJeremy L Thompson                                        CeedInt, const CeedScalar *,
1360d0321e0SJeremy L Thompson                                        const CeedScalar *, const CeedScalar *,
1370d0321e0SJeremy L Thompson                                        const CeedScalar *, CeedBasis);
1380d0321e0SJeremy L Thompson 
1390d0321e0SJeremy L Thompson CEED_INTERN int CeedQFunctionCreate_Cuda(CeedQFunction qf);
1400d0321e0SJeremy L Thompson 
1410d0321e0SJeremy L Thompson CEED_INTERN int CeedQFunctionContextCreate_Cuda(CeedQFunctionContext ctx);
1420d0321e0SJeremy L Thompson 
1430d0321e0SJeremy L Thompson CEED_INTERN int CeedOperatorCreate_Cuda(CeedOperator op);
1440d0321e0SJeremy L Thompson 
1450d0321e0SJeremy L Thompson CEED_INTERN int CeedCompositeOperatorCreate_Cuda(CeedOperator op);
1460d0321e0SJeremy L Thompson #endif
147