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