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_REF_HPP 10 #define CEED_SYCL_REF_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 basis_in, basis_out; 87 CeedElemRestriction diag_rstr, point_block_diag_rstr; 88 CeedVector elem_diag, point_block_elem_diag; 89 CeedInt num_e_mode_in, num_e_mode_out, num_nodes; 90 CeedInt num_qpts, num_comp; // Kernel parameters 91 CeedEvalMode *h_e_mode_in, *h_e_mode_out; 92 CeedEvalMode *d_e_mode_in, *d_e_mode_out; 93 CeedScalar *d_identity, *d_interp_in, *d_interp_out, *d_grad_in, *d_grad_out; 94 } CeedOperatorDiag_Sycl; 95 96 typedef struct { 97 CeedInt num_elem, block_size_x, block_size_y, elems_per_block; 98 CeedInt num_e_mode_in, num_e_mode_out, num_qpts, num_nodes, block_size, num_comp; // Kernel parameters 99 bool fallback; 100 CeedScalar *d_B_in, *d_B_out; 101 } CeedOperatorAssemble_Sycl; 102 103 typedef struct { 104 CeedVector *e_vecs; // E-vectors, inputs followed by outputs 105 CeedVector *q_vecs_in; // Input Q-vectors needed to apply operator 106 CeedVector *q_vecs_out; // Output Q-vectors needed to apply operator 107 CeedInt num_e_in; 108 CeedInt num_e_out; 109 CeedInt num_inputs, num_outputs; 110 CeedInt num_active_in, num_active_out; 111 CeedVector *qf_active_in; 112 CeedOperatorDiag_Sycl *diag; 113 CeedOperatorAssemble_Sycl *asmb; 114 } CeedOperator_Sycl; 115 116 CEED_INTERN int CeedVectorCreate_Sycl(CeedSize n, CeedVector vec); 117 118 CEED_INTERN int CeedBasisCreateTensorH1_Sycl(CeedInt dim, CeedInt P_1d, CeedInt Q_1d, const CeedScalar *interp_1d, const CeedScalar *grad_1d, 119 const CeedScalar *q_ref_1d, const CeedScalar *q_weight_1d, CeedBasis basis); 120 121 CEED_INTERN int CeedBasisCreateH1_Sycl(CeedElemTopology topo, CeedInt dim, CeedInt num_dof, CeedInt num_qpts, const CeedScalar *interp, 122 const CeedScalar *grad, const CeedScalar *q_ref, const CeedScalar *q_weight, CeedBasis basis); 123 124 CEED_INTERN int CeedElemRestrictionCreate_Sycl(CeedMemType mem_type, CeedCopyMode copy_mode, const CeedInt *indices, const bool *orients, 125 const CeedInt8 *curl_orients, CeedElemRestriction r); 126 127 CEED_INTERN int CeedQFunctionCreate_Sycl(CeedQFunction qf); 128 129 CEED_INTERN int CeedQFunctionContextCreate_Sycl(CeedQFunctionContext ctx); 130 131 CEED_INTERN int CeedOperatorCreate_Sycl(CeedOperator op); 132 133 #endif // CEED_SYCL_REF_HPP 134