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