1bd882c8aSJames Wright // Copyright (c) 2017-2022, Lawrence Livermore National Security, LLC and other 2bd882c8aSJames Wright // CEED contributors. All Rights Reserved. See the top-level LICENSE and NOTICE 3bd882c8aSJames Wright // files for details. 4bd882c8aSJames Wright // 5bd882c8aSJames Wright // SPDX-License-Identifier: BSD-2-Clause 6bd882c8aSJames Wright // 7bd882c8aSJames Wright // This file is part of CEED: http://github.com/ceed 8bd882c8aSJames Wright 9*17b5e52fSSebastian Grimberg #ifndef CEED_SYCL_REF_HPP 10*17b5e52fSSebastian Grimberg #define CEED_SYCL_REF_HPP 11bd882c8aSJames Wright 12bd882c8aSJames Wright #include <ceed/backend.h> 13bd882c8aSJames Wright #include <ceed/ceed.h> 14bd882c8aSJames Wright 15bd882c8aSJames Wright #include <sycl/sycl.hpp> 16bd882c8aSJames Wright 17bd882c8aSJames Wright #include "../sycl/ceed-sycl-common.hpp" 18bd882c8aSJames Wright #include "../sycl/ceed-sycl-compile.hpp" 19bd882c8aSJames Wright 20bd882c8aSJames Wright typedef struct { 21bd882c8aSJames Wright CeedScalar *h_array; 22bd882c8aSJames Wright CeedScalar *h_array_borrowed; 23bd882c8aSJames Wright CeedScalar *h_array_owned; 24bd882c8aSJames Wright CeedScalar *d_array; 25bd882c8aSJames Wright CeedScalar *d_array_borrowed; 26bd882c8aSJames Wright CeedScalar *d_array_owned; 27bd882c8aSJames Wright CeedScalar *reduction_norm; 28bd882c8aSJames Wright } CeedVector_Sycl; 29bd882c8aSJames Wright 30bd882c8aSJames Wright typedef struct { 31bd882c8aSJames Wright CeedInt num_nodes; 32bd882c8aSJames Wright CeedInt num_elem; 33bd882c8aSJames Wright CeedInt num_comp; 34bd882c8aSJames Wright CeedInt elem_size; 35bd882c8aSJames Wright CeedInt comp_stride; 36bd882c8aSJames Wright CeedInt strides[3]; 37bd882c8aSJames Wright CeedInt *h_ind; 38bd882c8aSJames Wright CeedInt *h_ind_allocated; 39bd882c8aSJames Wright CeedInt *d_ind; 40bd882c8aSJames Wright CeedInt *d_ind_allocated; 41bd882c8aSJames Wright CeedInt *d_t_offsets; 42bd882c8aSJames Wright CeedInt *d_t_indices; 43bd882c8aSJames Wright CeedInt *d_l_vec_indices; 44bd882c8aSJames Wright } CeedElemRestriction_Sycl; 45bd882c8aSJames Wright 46bd882c8aSJames Wright typedef struct { 47bd882c8aSJames Wright CeedInt dim; 48bd882c8aSJames Wright CeedInt P_1d; 49bd882c8aSJames Wright CeedInt Q_1d; 50bd882c8aSJames Wright CeedInt num_comp; 51bd882c8aSJames Wright CeedInt num_nodes; 52bd882c8aSJames Wright CeedInt num_qpts; 53bd882c8aSJames Wright CeedInt buf_len; 54bd882c8aSJames Wright CeedInt op_len; 55bd882c8aSJames Wright SyclModule_t *sycl_module; 56bd882c8aSJames Wright CeedScalar *d_interp_1d; 57bd882c8aSJames Wright CeedScalar *d_grad_1d; 58bd882c8aSJames Wright CeedScalar *d_q_weight_1d; 59bd882c8aSJames Wright } CeedBasis_Sycl; 60bd882c8aSJames Wright 61bd882c8aSJames Wright typedef struct { 62bd882c8aSJames Wright CeedInt dim; 63bd882c8aSJames Wright CeedInt num_comp; 64bd882c8aSJames Wright CeedInt num_nodes; 65bd882c8aSJames Wright CeedInt num_qpts; 66bd882c8aSJames Wright CeedScalar *d_interp; 67bd882c8aSJames Wright CeedScalar *d_grad; 68bd882c8aSJames Wright CeedScalar *d_q_weight; 69bd882c8aSJames Wright } CeedBasisNonTensor_Sycl; 70bd882c8aSJames Wright 71bd882c8aSJames Wright typedef struct { 72bd882c8aSJames Wright SyclModule_t *sycl_module; 73bd882c8aSJames Wright sycl::kernel *QFunction; 74bd882c8aSJames Wright } CeedQFunction_Sycl; 75bd882c8aSJames Wright 76bd882c8aSJames Wright typedef struct { 77bd882c8aSJames Wright void *h_data; 78bd882c8aSJames Wright void *h_data_borrowed; 79bd882c8aSJames Wright void *h_data_owned; 80bd882c8aSJames Wright void *d_data; 81bd882c8aSJames Wright void *d_data_borrowed; 82bd882c8aSJames Wright void *d_data_owned; 83bd882c8aSJames Wright } CeedQFunctionContext_Sycl; 84bd882c8aSJames Wright 85bd882c8aSJames Wright typedef struct { 86bd882c8aSJames Wright CeedBasis basisin, basisout; 87bd882c8aSJames Wright CeedElemRestriction diagrstr, pbdiagrstr; 88bd882c8aSJames Wright CeedVector elemdiag, pbelemdiag; 89bd882c8aSJames Wright CeedInt numemodein, numemodeout, nnodes; 90bd882c8aSJames Wright CeedInt nqpts, ncomp; // Kernel parameters 91bd882c8aSJames Wright CeedEvalMode *h_emodein, *h_emodeout; 92bd882c8aSJames Wright CeedEvalMode *d_emodein, *d_emodeout; 93bd882c8aSJames Wright CeedScalar *d_identity, *d_interpin, *d_interpout, *d_gradin, *d_gradout; 94bd882c8aSJames Wright } CeedOperatorDiag_Sycl; 95bd882c8aSJames Wright 96bd882c8aSJames Wright typedef struct { 97bd882c8aSJames Wright CeedInt nelem, block_size_x, block_size_y, elemsPerBlock; 98bd882c8aSJames Wright CeedInt numemodein, numemodeout, nqpts, nnodes, block_size, ncomp; // Kernel parameters 99bd882c8aSJames Wright bool fallback; 100bd882c8aSJames Wright CeedScalar *d_B_in, *d_B_out; 101bd882c8aSJames Wright } CeedOperatorAssemble_Sycl; 102bd882c8aSJames Wright 103bd882c8aSJames Wright typedef struct { 104bd882c8aSJames Wright CeedVector *evecs; // E-vectors, inputs followed by outputs 105bd882c8aSJames Wright CeedVector *qvecsin; // Input Q-vectors needed to apply operator 106bd882c8aSJames Wright CeedVector *qvecsout; // Output Q-vectors needed to apply operator 107bd882c8aSJames Wright CeedInt numein; 108bd882c8aSJames Wright CeedInt numeout; 109bd882c8aSJames Wright CeedInt qfnumactivein, qfnumactiveout; 110bd882c8aSJames Wright CeedVector *qfactivein; 111bd882c8aSJames Wright CeedOperatorDiag_Sycl *diag; 112bd882c8aSJames Wright CeedOperatorAssemble_Sycl *asmb; 113bd882c8aSJames Wright } CeedOperator_Sycl; 114bd882c8aSJames Wright 115bd882c8aSJames Wright CEED_INTERN int CeedVectorCreate_Sycl(CeedSize n, CeedVector vec); 116bd882c8aSJames Wright 117bd882c8aSJames Wright CEED_INTERN int CeedBasisCreateTensorH1_Sycl(CeedInt dim, CeedInt P_1d, CeedInt Q_1d, const CeedScalar *interp_1d, const CeedScalar *grad_1d, 118bd882c8aSJames Wright const CeedScalar *qref_1d, const CeedScalar *qweight_1d, CeedBasis basis); 119bd882c8aSJames Wright 12000125730SSebastian Grimberg CEED_INTERN int CeedBasisCreateH1_Sycl(CeedElemTopology topo, CeedInt dim, CeedInt ndof, CeedInt nqpts, const CeedScalar *interp, 12100125730SSebastian Grimberg const CeedScalar *grad, const CeedScalar *qref, const CeedScalar *qweight, CeedBasis basis); 12200125730SSebastian Grimberg 12300125730SSebastian Grimberg CEED_INTERN int CeedElemRestrictionCreate_Sycl(CeedMemType mem_type, CeedCopyMode copy_mode, const CeedInt *indices, const bool *orients, 12400125730SSebastian Grimberg const CeedInt8 *curl_orients, CeedElemRestriction r); 125bd882c8aSJames Wright 126bd882c8aSJames Wright CEED_INTERN int CeedQFunctionCreate_Sycl(CeedQFunction qf); 127bd882c8aSJames Wright 128bd882c8aSJames Wright CEED_INTERN int CeedQFunctionContextCreate_Sycl(CeedQFunctionContext ctx); 129bd882c8aSJames Wright 130bd882c8aSJames Wright CEED_INTERN int CeedOperatorCreate_Sycl(CeedOperator op); 131bd882c8aSJames Wright 132*17b5e52fSSebastian Grimberg #endif // CEED_SYCL_REF_HPP 133