1 // Copyright (c) 2017-2025, Lawrence Livermore National Security, LLC and other CEED contributors. 2 // All Rights Reserved. See the top-level LICENSE and NOTICE files for details. 3 // 4 // SPDX-License-Identifier: BSD-2-Clause 5 // 6 // This file is part of CEED: http://github.com/ceed 7 8 #include "ceed-sycl-compile.hpp" 9 10 #include <ceed/backend.h> 11 #include <ceed/ceed.h> 12 #include <ceed/jit-tools.h> 13 #include <level_zero/ze_api.h> 14 15 #include <map> 16 #include <sstream> 17 #include <sycl/sycl.hpp> 18 19 #include "./online_compiler.hpp" 20 #include "ceed-sycl-common.hpp" 21 22 using ByteVector_t = std::vector<unsigned char>; 23 24 //------------------------------------------------------------------------------ 25 // Add defined constants at the beginning of kernel source 26 //------------------------------------------------------------------------------ 27 static int CeedJitAddDefinitions_Sycl(Ceed ceed, const std::string &kernel_source, std::string &jit_source, 28 const std::map<std::string, CeedInt> &constants = {}) { 29 std::ostringstream oss; 30 31 const char *jit_defs_path, *jit_defs_source; 32 const char *sycl_jith_path = "ceed/jit-source/sycl/sycl-jit.h"; 33 34 // Prepend defined constants 35 for (const auto &[name, value] : constants) { 36 oss << "#define " << name << " " << value << "\n"; 37 } 38 39 // libCeed definitions for Sycl Backends 40 CeedCallBackend(CeedGetJitAbsolutePath(ceed, sycl_jith_path, &jit_defs_path)); 41 { 42 char *source; 43 44 CeedCallBackend(CeedLoadSourceToBuffer(ceed, jit_defs_path, &source)); 45 jit_defs_source = source; 46 } 47 48 oss << jit_defs_source << "\n"; 49 50 CeedCallBackend(CeedFree(&jit_defs_path)); 51 CeedCallBackend(CeedFree(&jit_defs_source)); 52 53 // Append kernel_source 54 oss << "\n" << kernel_source; 55 56 jit_source = oss.str(); 57 return CEED_ERROR_SUCCESS; 58 } 59 60 //------------------------------------------------------------------------------ 61 // TODO: Add architecture flags, optimization flags 62 //------------------------------------------------------------------------------ 63 static inline int CeedJitGetFlags_Sycl(std::vector<std::string> &flags) { 64 flags = {std::string("-cl-std=CL3.0"), std::string("-Dint32_t=int"), std::string("-DCEED_RUNNING_JIT_PASS=1")}; 65 return CEED_ERROR_SUCCESS; 66 } 67 68 //------------------------------------------------------------------------------ 69 // Compile an OpenCL source to SPIR-V using Intel's online compiler extension 70 //------------------------------------------------------------------------------ 71 static inline int CeedJitCompileSource_Sycl(Ceed ceed, const sycl::device &sycl_device, const std::string &opencl_source, ByteVector_t &il_binary, 72 const std::vector<std::string> &flags = {}) { 73 sycl::ext::libceed::online_compiler<sycl::ext::libceed::source_language::opencl_c> compiler(sycl_device); 74 75 try { 76 il_binary = compiler.compile(opencl_source, flags); 77 } catch (sycl::ext::libceed::online_compile_error &e) { 78 return CeedError((ceed), CEED_ERROR_BACKEND, e.what()); 79 } 80 return CEED_ERROR_SUCCESS; 81 } 82 83 // ------------------------------------------------------------------------------ 84 // Load (compile) SPIR-V source and wrap in sycl kernel_bundle 85 // ------------------------------------------------------------------------------ 86 static int CeedLoadModule_Sycl(Ceed ceed, const sycl::context &sycl_context, const sycl::device &sycl_device, const ByteVector_t &il_binary, 87 SyclModule_t **sycl_module) { 88 auto lz_context = sycl::get_native<sycl::backend::ext_oneapi_level_zero>(sycl_context); 89 auto lz_device = sycl::get_native<sycl::backend::ext_oneapi_level_zero>(sycl_device); 90 91 ze_module_desc_t lz_mod_desc = {ZE_STRUCTURE_TYPE_MODULE_DESC, 92 nullptr, // extension specific structs 93 ZE_MODULE_FORMAT_IL_SPIRV, 94 il_binary.size(), 95 il_binary.data(), 96 " -ze-opt-large-register-file", // flags 97 nullptr}; // specialization constants 98 99 ze_module_handle_t lz_module; 100 ze_module_build_log_handle_t lz_log; 101 ze_result_t lz_err = zeModuleCreate(lz_context, lz_device, &lz_mod_desc, &lz_module, &lz_log); 102 103 if (ZE_RESULT_SUCCESS != lz_err) { 104 size_t log_size = 0; 105 char *log_message; 106 107 zeModuleBuildLogGetString(lz_log, &log_size, nullptr); 108 109 CeedCallBackend(CeedCalloc(log_size, &log_message)); 110 zeModuleBuildLogGetString(lz_log, &log_size, log_message); 111 112 return CeedError(ceed, CEED_ERROR_BACKEND, "Failed to compile Level Zero module:\n%s", log_message); 113 } 114 115 // sycl make_<type> only throws errors for backend mismatch--assume we have vetted this already 116 *sycl_module = new SyclModule_t(sycl::make_kernel_bundle<sycl::backend::ext_oneapi_level_zero, sycl::bundle_state::executable>( 117 {lz_module, sycl::ext::oneapi::level_zero::ownership::transfer}, sycl_context)); 118 return CEED_ERROR_SUCCESS; 119 } 120 121 // ------------------------------------------------------------------------------ 122 // Compile kernel source to an executable `sycl::kernel_bundle` 123 // ------------------------------------------------------------------------------ 124 int CeedBuildModule_Sycl(Ceed ceed, const std::string &kernel_source, SyclModule_t **sycl_module, const std::map<std::string, CeedInt> &constants) { 125 Ceed_Sycl *data; 126 std::string jit_source; 127 std::vector<std::string> flags; 128 ByteVector_t il_binary; 129 130 CeedCallBackend(CeedGetData(ceed, &data)); 131 CeedCallBackend(CeedJitAddDefinitions_Sycl(ceed, kernel_source, jit_source, constants)); 132 CeedCallBackend(CeedJitGetFlags_Sycl(flags)); 133 CeedCallBackend(CeedJitCompileSource_Sycl(ceed, data->sycl_device, jit_source, il_binary, flags)); 134 CeedCallBackend(CeedLoadModule_Sycl(ceed, data->sycl_context, data->sycl_device, il_binary, sycl_module)); 135 return CEED_ERROR_SUCCESS; 136 } 137 138 // ------------------------------------------------------------------------------ 139 // Get a sycl kernel from an existing kernel_bundle 140 // 141 // TODO: Error handle lz calls 142 // ------------------------------------------------------------------------------ 143 int CeedGetKernel_Sycl(Ceed ceed, const SyclModule_t *sycl_module, const std::string &kernel_name, sycl::kernel **sycl_kernel) { 144 Ceed_Sycl *data; 145 146 CeedCallBackend(CeedGetData(ceed, &data)); 147 148 // sycl::get_native returns std::vector<ze_module_handle_t> for lz backend 149 // https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/supported/sycl_ext_oneapi_backend_level_zero.md 150 ze_module_handle_t lz_module = sycl::get_native<sycl::backend::ext_oneapi_level_zero>(*sycl_module).front(); 151 152 ze_kernel_desc_t lz_kernel_desc = {ZE_STRUCTURE_TYPE_KERNEL_DESC, nullptr, 0, kernel_name.c_str()}; 153 ze_kernel_handle_t lz_kernel; 154 ze_result_t lz_err = zeKernelCreate(lz_module, &lz_kernel_desc, &lz_kernel); 155 156 if (ZE_RESULT_SUCCESS != lz_err) { 157 return CeedError(ceed, CEED_ERROR_BACKEND, "Failed to retrieve kernel from Level Zero module"); 158 } 159 160 *sycl_kernel = new sycl::kernel(sycl::make_kernel<sycl::backend::ext_oneapi_level_zero>( 161 {*sycl_module, lz_kernel, sycl::ext::oneapi::level_zero::ownership::transfer}, data->sycl_context)); 162 return CEED_ERROR_SUCCESS; 163 } 164 165 //------------------------------------------------------------------------------ 166 // Run SYCL kernel for spatial dimension with shared memory 167 //------------------------------------------------------------------------------ 168 int CeedRunKernelDimSharedSycl(Ceed ceed, sycl::kernel *kernel, const int grid_size, const int block_size_x, const int block_size_y, 169 const int block_size_z, const int shared_mem_size, void **kernel_args) { 170 sycl::range<3> local_range(block_size_z, block_size_y, block_size_x); 171 sycl::range<3> global_range(grid_size * block_size_z, block_size_y, block_size_x); 172 sycl::nd_range<3> kernel_range(global_range, local_range); 173 174 //----------- 175 // Order queue 176 Ceed_Sycl *ceed_Sycl; 177 178 CeedCallBackend(CeedGetData(ceed, &ceed_Sycl)); 179 sycl::event e = ceed_Sycl->sycl_queue.ext_oneapi_submit_barrier(); 180 181 ceed_Sycl->sycl_queue.submit([&](sycl::handler &cgh) { 182 cgh.depends_on(e); 183 cgh.set_args(*kernel_args); 184 cgh.parallel_for(kernel_range, *kernel); 185 }); 186 return CEED_ERROR_SUCCESS; 187 } 188