1 // Copyright (c) 2017-2024, 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_offsets; 38 CeedInt *h_offsets_borrowed; 39 CeedInt *h_offsets_owned; 40 CeedInt *d_offsets; 41 CeedInt *d_offsets_borrowed; 42 CeedInt *d_offsets_owned; 43 CeedInt *d_t_offsets; 44 CeedInt *d_t_indices; 45 CeedInt *d_l_vec_indices; 46 } CeedElemRestriction_Sycl; 47 48 typedef struct { 49 CeedInt dim; 50 CeedInt P_1d; 51 CeedInt Q_1d; 52 CeedInt num_comp; 53 CeedInt num_nodes; 54 CeedInt num_qpts; 55 CeedInt buf_len; 56 CeedInt op_len; 57 SyclModule_t *sycl_module; 58 CeedScalar *d_interp_1d; 59 CeedScalar *d_grad_1d; 60 CeedScalar *d_q_weight_1d; 61 } CeedBasis_Sycl; 62 63 typedef struct { 64 CeedInt dim; 65 CeedInt num_comp; 66 CeedInt num_nodes; 67 CeedInt num_qpts; 68 CeedScalar *d_interp; 69 CeedScalar *d_grad; 70 CeedScalar *d_q_weight; 71 } CeedBasisNonTensor_Sycl; 72 73 typedef struct { 74 SyclModule_t *sycl_module; 75 sycl::kernel *QFunction; 76 } CeedQFunction_Sycl; 77 78 typedef struct { 79 void *h_data; 80 void *h_data_borrowed; 81 void *h_data_owned; 82 void *d_data; 83 void *d_data_borrowed; 84 void *d_data_owned; 85 } CeedQFunctionContext_Sycl; 86 87 typedef struct { 88 CeedBasis basis_in, basis_out; 89 CeedElemRestriction diag_rstr, point_block_diag_rstr; 90 CeedVector elem_diag, point_block_elem_diag; 91 CeedInt num_e_mode_in, num_e_mode_out, num_nodes; 92 CeedInt num_qpts, num_comp; // Kernel parameters 93 CeedEvalMode *h_e_mode_in, *h_e_mode_out; 94 CeedEvalMode *d_e_mode_in, *d_e_mode_out; 95 CeedScalar *d_identity, *d_interp_in, *d_interp_out, *d_grad_in, *d_grad_out; 96 } CeedOperatorDiag_Sycl; 97 98 typedef struct { 99 CeedInt num_elem, block_size_x, block_size_y, elems_per_block; 100 CeedInt num_e_mode_in, num_e_mode_out, num_qpts, num_nodes, block_size, num_comp; // Kernel parameters 101 bool fallback; 102 CeedScalar *d_B_in, *d_B_out; 103 } CeedOperatorAssemble_Sycl; 104 105 typedef struct { 106 CeedVector *e_vecs; // E-vectors, inputs followed by outputs 107 CeedVector *q_vecs_in; // Input Q-vectors needed to apply operator 108 CeedVector *q_vecs_out; // Output Q-vectors needed to apply operator 109 CeedInt num_e_in; 110 CeedInt num_e_out; 111 CeedInt num_inputs, num_outputs; 112 CeedInt num_active_in, num_active_out; 113 CeedVector *qf_active_in; 114 CeedOperatorDiag_Sycl *diag; 115 CeedOperatorAssemble_Sycl *asmb; 116 } CeedOperator_Sycl; 117 118 CEED_INTERN int CeedVectorCreate_Sycl(CeedSize n, CeedVector vec); 119 120 CEED_INTERN int CeedBasisCreateTensorH1_Sycl(CeedInt dim, CeedInt P_1d, CeedInt Q_1d, const CeedScalar *interp_1d, const CeedScalar *grad_1d, 121 const CeedScalar *q_ref_1d, const CeedScalar *q_weight_1d, CeedBasis basis); 122 123 CEED_INTERN int CeedBasisCreateH1_Sycl(CeedElemTopology topo, CeedInt dim, CeedInt num_dof, CeedInt num_qpts, const CeedScalar *interp, 124 const CeedScalar *grad, const CeedScalar *q_ref, const CeedScalar *q_weight, CeedBasis basis); 125 126 CEED_INTERN int CeedElemRestrictionCreate_Sycl(CeedMemType mem_type, CeedCopyMode copy_mode, const CeedInt *offsets, const bool *orients, 127 const CeedInt8 *curl_orients, CeedElemRestriction r); 128 129 CEED_INTERN int CeedQFunctionCreate_Sycl(CeedQFunction qf); 130 131 CEED_INTERN int CeedQFunctionContextCreate_Sycl(CeedQFunctionContext ctx); 132 133 CEED_INTERN int CeedOperatorCreate_Sycl(CeedOperator op); 134 135 #endif // CEED_SYCL_REF_HPP 136