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 // Prepend defined constants 32 for (const auto &[name, value] : constants) { 33 oss << "#define " << name << " " << value << "\n"; 34 } 35 36 // libCeed definitions for Sycl Backends 37 char *jit_defs_path, *jit_defs_source; 38 const char *sycl_jith_path = "ceed/jit-source/sycl/sycl-jit.h"; 39 CeedCallBackend(CeedGetJitAbsolutePath(ceed, sycl_jith_path, &jit_defs_path)); 40 CeedCallBackend(CeedLoadSourceToBuffer(ceed, jit_defs_path, &jit_defs_source)); 41 42 oss << jit_defs_source << "\n"; 43 44 CeedCallBackend(CeedFree(&jit_defs_path)); 45 CeedCallBackend(CeedFree(&jit_defs_source)); 46 47 // Append kernel_source 48 oss << "\n" << kernel_source; 49 50 jit_source = oss.str(); 51 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 zeModuleBuildLogGetString(lz_log, &log_size, nullptr); 101 102 char *log_message; 103 CeedCall(CeedCalloc(log_size, &log_message)); 104 zeModuleBuildLogGetString(lz_log, &log_size, log_message); 105 106 return CeedError(ceed, CEED_ERROR_BACKEND, "Failed to compile Level Zero module:\n%s", log_message); 107 } 108 109 // sycl make_<type> only throws errors for backend mismatch--assume we have vetted this already 110 *sycl_module = new SyclModule_t(sycl::make_kernel_bundle<sycl::backend::ext_oneapi_level_zero, sycl::bundle_state::executable>( 111 {lz_module, sycl::ext::oneapi::level_zero::ownership::transfer}, sycl_context)); 112 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 CeedCallBackend(CeedGetData(ceed, &data)); 122 123 std::string jit_source; 124 CeedCallBackend(CeedJitAddDefinitions_Sycl(ceed, kernel_source, jit_source, constants)); 125 126 std::vector<std::string> flags; 127 CeedCallBackend(CeedJitGetFlags_Sycl(flags)); 128 129 ByteVector_t il_binary; 130 CeedCallBackend(CeedJitCompileSource_Sycl(ceed, data->sycl_device, jit_source, il_binary, flags)); 131 132 CeedCallBackend(CeedLoadModule_Sycl(ceed, data->sycl_context, data->sycl_device, il_binary, sycl_module)); 133 134 return CEED_ERROR_SUCCESS; 135 } 136 137 // ------------------------------------------------------------------------------ 138 // Get a sycl kernel from an existing kernel_bundle 139 // 140 // TODO: Error handle lz calls 141 // ------------------------------------------------------------------------------ 142 int CeedGetKernel_Sycl(Ceed ceed, const SyclModule_t *sycl_module, const std::string &kernel_name, sycl::kernel **sycl_kernel) { 143 Ceed_Sycl *data; 144 CeedCallBackend(CeedGetData(ceed, &data)); 145 146 // sycl::get_native returns std::vector<ze_module_handle_t> for lz backend 147 // https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/supported/sycl_ext_oneapi_backend_level_zero.md 148 ze_module_handle_t lz_module = sycl::get_native<sycl::backend::ext_oneapi_level_zero>(*sycl_module).front(); 149 150 ze_kernel_desc_t lz_kernel_desc = {ZE_STRUCTURE_TYPE_KERNEL_DESC, nullptr, 0, kernel_name.c_str()}; 151 ze_kernel_handle_t lz_kernel; 152 ze_result_t lz_err = zeKernelCreate(lz_module, &lz_kernel_desc, &lz_kernel); 153 154 if (ZE_RESULT_SUCCESS != lz_err) { 155 return CeedError(ceed, CEED_ERROR_BACKEND, "Failed to retrieve kernel from Level Zero module"); 156 } 157 158 *sycl_kernel = new sycl::kernel(sycl::make_kernel<sycl::backend::ext_oneapi_level_zero>( 159 {*sycl_module, lz_kernel, sycl::ext::oneapi::level_zero::ownership::transfer}, data->sycl_context)); 160 161 return CEED_ERROR_SUCCESS; 162 } 163 164 //------------------------------------------------------------------------------ 165 // Run SYCL kernel for spatial dimension with shared memory 166 //------------------------------------------------------------------------------ 167 int CeedRunKernelDimSharedSycl(Ceed ceed, sycl::kernel *kernel, const int grid_size, const int block_size_x, const int block_size_y, 168 const int block_size_z, const int shared_mem_size, void **kernel_args) { 169 sycl::range<3> local_range(block_size_z, block_size_y, block_size_x); 170 sycl::range<3> global_range(grid_size * block_size_z, block_size_y, block_size_x); 171 sycl::nd_range<3> kernel_range(global_range, local_range); 172 173 //----------- 174 // Order queue 175 Ceed_Sycl *ceed_Sycl; 176 CeedCallBackend(CeedGetData(ceed, &ceed_Sycl)); 177 sycl::event e = ceed_Sycl->sycl_queue.ext_oneapi_submit_barrier(); 178 179 ceed_Sycl->sycl_queue.submit([&](sycl::handler &cgh) { 180 cgh.depends_on(e); 181 cgh.set_args(*kernel_args); 182 cgh.parallel_for(kernel_range, *kernel); 183 }); 184 185 return CEED_ERROR_SUCCESS; 186 } 187