1*9ba83ac0SJeremy L Thompson // Copyright (c) 2017-2026, 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 8509d4af6SJeremy L Thompson #pragma once 9bd882c8aSJames Wright 10bd882c8aSJames Wright #include <ceed/backend.h> 11bd882c8aSJames Wright #include <ceed/ceed.h> 12bd882c8aSJames Wright 13bd882c8aSJames Wright #include <sycl/sycl.hpp> 14bd882c8aSJames Wright 15bd882c8aSJames Wright #include "../sycl/ceed-sycl-common.hpp" 16bd882c8aSJames Wright #include "../sycl/ceed-sycl-compile.hpp" 17bd882c8aSJames Wright 18bd882c8aSJames Wright typedef struct { 19bd882c8aSJames Wright CeedScalar *h_array; 20bd882c8aSJames Wright CeedScalar *h_array_borrowed; 21bd882c8aSJames Wright CeedScalar *h_array_owned; 22bd882c8aSJames Wright CeedScalar *d_array; 23bd882c8aSJames Wright CeedScalar *d_array_borrowed; 24bd882c8aSJames Wright CeedScalar *d_array_owned; 25bd882c8aSJames Wright CeedScalar *reduction_norm; 26bd882c8aSJames Wright } CeedVector_Sycl; 27bd882c8aSJames Wright 28bd882c8aSJames Wright typedef struct { 29bd882c8aSJames Wright CeedInt num_nodes; 30bd882c8aSJames Wright CeedInt num_elem; 31bd882c8aSJames Wright CeedInt num_comp; 32bd882c8aSJames Wright CeedInt elem_size; 33bd882c8aSJames Wright CeedInt comp_stride; 34bd882c8aSJames Wright CeedInt strides[3]; 35f59ebe5eSJeremy L Thompson CeedInt *h_offsets; 36f59ebe5eSJeremy L Thompson CeedInt *h_offsets_borrowed; 37f59ebe5eSJeremy L Thompson CeedInt *h_offsets_owned; 38f59ebe5eSJeremy L Thompson CeedInt *d_offsets; 39f59ebe5eSJeremy L Thompson CeedInt *d_offsets_borrowed; 40f59ebe5eSJeremy L Thompson CeedInt *d_offsets_owned; 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 { 86dd64fc84SJeremy L Thompson CeedBasis basis_in, basis_out; 87dd64fc84SJeremy L Thompson CeedElemRestriction diag_rstr, point_block_diag_rstr; 88dd64fc84SJeremy L Thompson CeedVector elem_diag, point_block_elem_diag; 89681d0ea7SJeremy L Thompson CeedInt num_eval_mode_in, num_eval_mode_out, num_nodes; 90dd64fc84SJeremy L Thompson CeedInt num_qpts, num_comp; // Kernel parameters 91681d0ea7SJeremy L Thompson CeedEvalMode *h_eval_mode_in, *h_eval_mode_out; 92681d0ea7SJeremy L Thompson CeedEvalMode *d_eval_mode_in, *d_eval_mode_out; 93dd64fc84SJeremy L Thompson CeedScalar *d_identity, *d_interp_in, *d_interp_out, *d_grad_in, *d_grad_out; 94bd882c8aSJames Wright } CeedOperatorDiag_Sycl; 95bd882c8aSJames Wright 96bd882c8aSJames Wright typedef struct { 97004e4986SSebastian Grimberg CeedInt num_elem, block_size_x, block_size_y, elems_per_block; 98681d0ea7SJeremy L Thompson CeedInt num_eval_mode_in, num_eval_mode_out, num_qpts, num_nodes, block_size, num_comp; // 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 { 104dd64fc84SJeremy L Thompson CeedVector *e_vecs; // E-vectors, inputs followed by outputs 105dd64fc84SJeremy L Thompson CeedVector *q_vecs_in; // Input Q-vectors needed to apply operator 106dd64fc84SJeremy L Thompson CeedVector *q_vecs_out; // Output Q-vectors needed to apply operator 107dd64fc84SJeremy L Thompson CeedInt num_e_in; 108dd64fc84SJeremy L Thompson CeedInt num_e_out; 109dd64fc84SJeremy L Thompson CeedInt num_active_in, num_active_out; 110dd64fc84SJeremy L Thompson CeedVector *qf_active_in; 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, 118dd64fc84SJeremy L Thompson const CeedScalar *q_ref_1d, const CeedScalar *q_weight_1d, CeedBasis basis); 119bd882c8aSJames Wright 120dd64fc84SJeremy L Thompson CEED_INTERN int CeedBasisCreateH1_Sycl(CeedElemTopology topo, CeedInt dim, CeedInt num_dof, CeedInt num_qpts, const CeedScalar *interp, 121dd64fc84SJeremy L Thompson const CeedScalar *grad, const CeedScalar *q_ref, const CeedScalar *q_weight, CeedBasis basis); 12200125730SSebastian Grimberg 123f59ebe5eSJeremy L Thompson CEED_INTERN int CeedElemRestrictionCreate_Sycl(CeedMemType mem_type, CeedCopyMode copy_mode, const CeedInt *offsets, 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); 131