1 // Copyright (c) 2017-2022, 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 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 CeedCallBackend(CeedLoadSourceToBuffer(ceed, jit_defs_path, &jit_defs_source)); 42 43 oss << jit_defs_source << "\n"; 44 45 CeedCallBackend(CeedFree(&jit_defs_path)); 46 CeedCallBackend(CeedFree(&jit_defs_source)); 47 48 // Append kernel_source 49 oss << "\n" << kernel_source; 50 51 jit_source = oss.str(); 52 return CEED_ERROR_SUCCESS; 53 } 54 55 //------------------------------------------------------------------------------ 56 // TODO: Add architecture flags, optimization flags 57 //------------------------------------------------------------------------------ 58 static inline int CeedJitGetFlags_Sycl(std::vector<std::string> &flags) { 59 flags = {std::string("-cl-std=CL3.0"), std::string("-Dint32_t=int")}; 60 return CEED_ERROR_SUCCESS; 61 } 62 63 //------------------------------------------------------------------------------ 64 // Compile an OpenCL source to SPIR-V using Intel's online compiler extension 65 //------------------------------------------------------------------------------ 66 static inline int CeedJitCompileSource_Sycl(Ceed ceed, const sycl::device &sycl_device, const std::string &opencl_source, ByteVector_t &il_binary, 67 const std::vector<std::string> &flags = {}) { 68 sycl::ext::libceed::online_compiler<sycl::ext::libceed::source_language::opencl_c> compiler(sycl_device); 69 70 try { 71 il_binary = compiler.compile(opencl_source, flags); 72 } catch (sycl::ext::libceed::online_compile_error &e) { 73 return CeedError((ceed), CEED_ERROR_BACKEND, e.what()); 74 } 75 return CEED_ERROR_SUCCESS; 76 } 77 78 // ------------------------------------------------------------------------------ 79 // Load (compile) SPIR-V source and wrap in sycl kernel_bundle 80 // ------------------------------------------------------------------------------ 81 static int CeedLoadModule_Sycl(Ceed ceed, const sycl::context &sycl_context, const sycl::device &sycl_device, const ByteVector_t &il_binary, 82 SyclModule_t **sycl_module) { 83 auto lz_context = sycl::get_native<sycl::backend::ext_oneapi_level_zero>(sycl_context); 84 auto lz_device = sycl::get_native<sycl::backend::ext_oneapi_level_zero>(sycl_device); 85 86 ze_module_desc_t lz_mod_desc = {ZE_STRUCTURE_TYPE_MODULE_DESC, 87 nullptr, // extension specific structs 88 ZE_MODULE_FORMAT_IL_SPIRV, 89 il_binary.size(), 90 il_binary.data(), 91 " -ze-opt-large-register-file", // flags 92 nullptr}; // specialization constants 93 94 ze_module_handle_t lz_module; 95 ze_module_build_log_handle_t lz_log; 96 ze_result_t lz_err = zeModuleCreate(lz_context, lz_device, &lz_mod_desc, &lz_module, &lz_log); 97 98 if (ZE_RESULT_SUCCESS != lz_err) { 99 size_t log_size = 0; 100 char *log_message; 101 102 zeModuleBuildLogGetString(lz_log, &log_size, nullptr); 103 104 CeedCall(CeedCalloc(log_size, &log_message)); 105 zeModuleBuildLogGetString(lz_log, &log_size, log_message); 106 107 return CeedError(ceed, CEED_ERROR_BACKEND, "Failed to compile Level Zero module:\n%s", log_message); 108 } 109 110 // sycl make_<type> only throws errors for backend mismatch--assume we have vetted this already 111 *sycl_module = new SyclModule_t(sycl::make_kernel_bundle<sycl::backend::ext_oneapi_level_zero, sycl::bundle_state::executable>( 112 {lz_module, sycl::ext::oneapi::level_zero::ownership::transfer}, sycl_context)); 113 return CEED_ERROR_SUCCESS; 114 } 115 116 // ------------------------------------------------------------------------------ 117 // Compile kernel source to an executable `sycl::kernel_bundle` 118 // ------------------------------------------------------------------------------ 119 int CeedBuildModule_Sycl(Ceed ceed, const std::string &kernel_source, SyclModule_t **sycl_module, const std::map<std::string, CeedInt> &constants) { 120 Ceed_Sycl *data; 121 std::string jit_source; 122 std::vector<std::string> flags; 123 ByteVector_t il_binary; 124 125 CeedCallBackend(CeedGetData(ceed, &data)); 126 CeedCallBackend(CeedJitAddDefinitions_Sycl(ceed, kernel_source, jit_source, constants)); 127 CeedCallBackend(CeedJitGetFlags_Sycl(flags)); 128 CeedCallBackend(CeedJitCompileSource_Sycl(ceed, data->sycl_device, jit_source, il_binary, flags)); 129 CeedCallBackend(CeedLoadModule_Sycl(ceed, data->sycl_context, data->sycl_device, il_binary, sycl_module)); 130 return CEED_ERROR_SUCCESS; 131 } 132 133 // ------------------------------------------------------------------------------ 134 // Get a sycl kernel from an existing kernel_bundle 135 // 136 // TODO: Error handle lz calls 137 // ------------------------------------------------------------------------------ 138 int CeedGetKernel_Sycl(Ceed ceed, const SyclModule_t *sycl_module, const std::string &kernel_name, sycl::kernel **sycl_kernel) { 139 Ceed_Sycl *data; 140 141 CeedCallBackend(CeedGetData(ceed, &data)); 142 143 // sycl::get_native returns std::vector<ze_module_handle_t> for lz backend 144 // https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/supported/sycl_ext_oneapi_backend_level_zero.md 145 ze_module_handle_t lz_module = sycl::get_native<sycl::backend::ext_oneapi_level_zero>(*sycl_module).front(); 146 147 ze_kernel_desc_t lz_kernel_desc = {ZE_STRUCTURE_TYPE_KERNEL_DESC, nullptr, 0, kernel_name.c_str()}; 148 ze_kernel_handle_t lz_kernel; 149 ze_result_t lz_err = zeKernelCreate(lz_module, &lz_kernel_desc, &lz_kernel); 150 151 if (ZE_RESULT_SUCCESS != lz_err) { 152 return CeedError(ceed, CEED_ERROR_BACKEND, "Failed to retrieve kernel from Level Zero module"); 153 } 154 155 *sycl_kernel = new sycl::kernel(sycl::make_kernel<sycl::backend::ext_oneapi_level_zero>( 156 {*sycl_module, lz_kernel, sycl::ext::oneapi::level_zero::ownership::transfer}, data->sycl_context)); 157 return CEED_ERROR_SUCCESS; 158 } 159 160 //------------------------------------------------------------------------------ 161 // Run SYCL kernel for spatial dimension with shared memory 162 //------------------------------------------------------------------------------ 163 int CeedRunKernelDimSharedSycl(Ceed ceed, sycl::kernel *kernel, const int grid_size, const int block_size_x, const int block_size_y, 164 const int block_size_z, const int shared_mem_size, void **kernel_args) { 165 sycl::range<3> local_range(block_size_z, block_size_y, block_size_x); 166 sycl::range<3> global_range(grid_size * block_size_z, block_size_y, block_size_x); 167 sycl::nd_range<3> kernel_range(global_range, local_range); 168 169 //----------- 170 // Order queue 171 Ceed_Sycl *ceed_Sycl; 172 173 CeedCallBackend(CeedGetData(ceed, &ceed_Sycl)); 174 sycl::event e = ceed_Sycl->sycl_queue.ext_oneapi_submit_barrier(); 175 176 ceed_Sycl->sycl_queue.submit([&](sycl::handler &cgh) { 177 cgh.depends_on(e); 178 cgh.set_args(*kernel_args); 179 cgh.parallel_for(kernel_range, *kernel); 180 }); 181 return CEED_ERROR_SUCCESS; 182 } 183