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