From 5026b11100529d608b3aa37392bce2335ad6a76c Mon Sep 17 00:00:00 2001 From: Umesh Unnikrishnan Date: Wed, 25 Sep 2024 15:26:35 +0000 Subject: [PATCH] Integrating libprtc with sycl-compiler --- backends/sycl/ceed-sycl-compile.hpp | 11 +- backends/sycl/ceed-sycl-compile.sycl.cpp | 169 ++++++++++++----------- 2 files changed, 100 insertions(+), 80 deletions(-) diff --git a/backends/sycl/ceed-sycl-compile.hpp b/backends/sycl/ceed-sycl-compile.hpp index 67db04f294..7b30a4e818 100644 --- a/backends/sycl/ceed-sycl-compile.hpp +++ b/backends/sycl/ceed-sycl-compile.hpp @@ -12,11 +12,16 @@ #include #include -using SyclModule_t = sycl::kernel_bundle; +#include + +using SyclModule_t = std::shared_ptr; CEED_INTERN int CeedBuildModule_Sycl(Ceed ceed, const std::string &kernel_source, SyclModule_t **sycl_module, const std::map &constants = {}); -CEED_INTERN int CeedGetKernel_Sycl(Ceed ceed, const SyclModule_t *sycl_module, const std::string &kernel_name, sycl::kernel **sycl_kernel); -CEED_INTERN int CeedRunKernelDimSharedSycl(Ceed ceed, sycl::kernel *kernel, const int grid_size, const int block_size_x, const int block_size_y, +template +CEED_INTERN int CeedGetKernel_Sycl(Ceed ceed, const SyclModule_t *sycl_module, const std::string &kernel_name, SyclKernel_t **sycl_kernel); + +template +CEED_INTERN int CeedRunKernelDimSharedSycl(Ceed ceed, SyclKernel_t *kernel, const int grid_size, const int block_size_x, const int block_size_y, const int block_size_z, const int shared_mem_size, void **args); diff --git a/backends/sycl/ceed-sycl-compile.sycl.cpp b/backends/sycl/ceed-sycl-compile.sycl.cpp index 9dc0177401..28a7411309 100644 --- a/backends/sycl/ceed-sycl-compile.sycl.cpp +++ b/backends/sycl/ceed-sycl-compile.sycl.cpp @@ -12,14 +12,15 @@ #include #include +#include #include +#include #include +#include #include -#include "./online_compiler.hpp" #include "ceed-sycl-common.hpp" - -using ByteVector_t = std::vector; +#include "libprtc/prtc.h" //------------------------------------------------------------------------------ // Add defined constants at the beginning of kernel source @@ -61,111 +62,125 @@ static int CeedJitAddDefinitions_Sycl(Ceed ceed, const std::string &kernel_sourc // TODO: Add architecture flags, optimization flags //------------------------------------------------------------------------------ static inline int CeedJitGetFlags_Sycl(std::vector &flags) { - flags = {std::string("-cl-std=CL3.0"), std::string("-Dint32_t=int")}; + + // flags = {std::string("-cl-std=CL3.0"), std::string("-Dint32_t=int")}; + flags = {std::string("-fsycl"), std::string("-fno-sycl-id-queries-fit-in-int"), std::string("-Dint32_t=int")}; + // TODO : Add AOT flags and other optimization flags + // flags.push_back(std::string("-O3")); + // flags.push_back(std::string("-fsycl-targets=spir64_gen -Xsycl-target-backend \"-device pvc\" ")) + return CEED_ERROR_SUCCESS; } //------------------------------------------------------------------------------ -// Compile an OpenCL source to SPIR-V using Intel's online compiler extension +// Compile sycl source to a shared library +// TODO: Check if source, module, etc. already exists //------------------------------------------------------------------------------ -static inline int CeedJitCompileSource_Sycl(Ceed ceed, const sycl::device &sycl_device, const std::string &opencl_source, ByteVector_t &il_binary, +static inline int CeedJitCompileSource_Sycl(Ceed ceed, const sycl::device &sycl_device, const std::string &kernel_source, std::string& output_path, const std::vector &flags = {}) { - sycl::ext::libceed::online_compiler compiler(sycl_device); - try { - il_binary = compiler.compile(opencl_source, flags); - } catch (sycl::ext::libceed::online_compile_error &e) { - return CeedError((ceed), CEED_ERROR_BACKEND, e.what()); + // Get cache path from env variable + std::string cache_root; + // TODO: Add default directory to current working directory + if(std::getenv()"CEED_CACHE_DIR")) { + cache_root = std::string(std::getenv("CEED_CACHE_DIR")) + "/.ceed/cache"; + } else { + cache_root = std::string(std::getenv("PWD")) + "/.ceed/cache"; } + + // Generate kernel hash + // E.g., see https://intel.github.io/llvm-docs/design/KernelProgramCache.html + // An example of directory structure can be found here: + // https://intel.github.io/llvm-docs/design/KernelProgramCache.html#persistent-cache-storage-structure + + // UU: Plan to use cache storage structure as : //// + + // Hash kernel name and source + std::hash string_hash; + // size_t kernel_name_hash = string_hash(get_kernel_name) ! UU : TODO LATER + size_t kernel_source_hash = string_hash(kernel_source); + + // Hash compilation flags + std::sort(flags.begin(), flags.end()); + std::string all_flags = prtc::concatenateFlags(flags); + size_t build_options_hash = string_hash(all_flags); + + // Hash compiler version + prtc::ShellCommand command("icpx --version"); + const auto [success, compiler_version] = command.result(); + if (!success) return CeedError((ceed), CEED_ERROR_BACKEND, compiler_version); + size_t compiler_hash = string_hash(compiler_version); + + // Determine file paths for source and binaries based on hashes + std::string cache_path = cache_root + "/" + std::to_string(compiler_hash) + "/" + std::to_string(build_options_hash) + "/" + std::to_string(kernel_source_hash) + "/"; + std::string source_file_path = cache_path + "source.cpp"; + std::string object_file_path = cache_path + "binary.so"; + + // Write source string to file + std::ofstream source_file; + source_file.open(source_file_path); + source_file << kernel_source; + source_file.close(); + + // TODO: Get compiler-path and flags from env or some other means + prtc::ShellCompiler compiler("icpx","-o","-c","-fPIC","-shared"); + const auto [success, message] = compiler.compileAndLink(source_file_path,object_file_path,flags); + // Q: Should we always output the compiler output in verbose/debug mode? + if (!success) return CeedError((ceed), CEED_ERROR_BACKEND, message); return CEED_ERROR_SUCCESS; } // ------------------------------------------------------------------------------ // Load (compile) SPIR-V source and wrap in sycl kernel_bundle // ------------------------------------------------------------------------------ -static int CeedLoadModule_Sycl(Ceed ceed, const sycl::context &sycl_context, const sycl::device &sycl_device, const ByteVector_t &il_binary, - SyclModule_t **sycl_module) { - auto lz_context = sycl::get_native(sycl_context); - auto lz_device = sycl::get_native(sycl_device); - - ze_module_desc_t lz_mod_desc = {ZE_STRUCTURE_TYPE_MODULE_DESC, - nullptr, // extension specific structs - ZE_MODULE_FORMAT_IL_SPIRV, - il_binary.size(), - il_binary.data(), - " -ze-opt-large-register-file", // flags - nullptr}; // specialization constants - - ze_module_handle_t lz_module; - ze_module_build_log_handle_t lz_log; - ze_result_t lz_err = zeModuleCreate(lz_context, lz_device, &lz_mod_desc, &lz_module, &lz_log); - - if (ZE_RESULT_SUCCESS != lz_err) { - size_t log_size = 0; - char *log_message; - - zeModuleBuildLogGetString(lz_log, &log_size, nullptr); - - CeedCallBackend(CeedCalloc(log_size, &log_message)); - zeModuleBuildLogGetString(lz_log, &log_size, log_message); - - return CeedError(ceed, CEED_ERROR_BACKEND, "Failed to compile Level Zero module:\n%s", log_message); +static int CeedLoadModule_Sycl(Ceed ceed, const sycl::context &sycl_context, const sycl::device &sycl_device, const std::string& path, + SyclModule_t *sycl_module) { + try { + *sycl_module = prtc::DynamicLibrary::open(path); + } catch (const std::exception& e) { + return CeedError((ceed), CEED_ERROR_BACKEND, e.what()); } - - // sycl make_ only throws errors for backend mismatch--assume we have vetted this already - *sycl_module = new SyclModule_t(sycl::make_kernel_bundle( - {lz_module, sycl::ext::oneapi::level_zero::ownership::transfer}, sycl_context)); return CEED_ERROR_SUCCESS; } // ------------------------------------------------------------------------------ -// Compile kernel source to an executable `sycl::kernel_bundle` +// Compile kernel source to a shared library // ------------------------------------------------------------------------------ -int CeedBuildModule_Sycl(Ceed ceed, const std::string &kernel_source, SyclModule_t **sycl_module, const std::map &constants) { +int CeedBuildModule_Sycl(Ceed ceed, const std::string &kernel_source, SyclModule_t *sycl_module, const std::map &constants) { Ceed_Sycl *data; std::string jit_source; + std::string module_path; std::vector flags; - ByteVector_t il_binary; CeedCallBackend(CeedGetData(ceed, &data)); CeedCallBackend(CeedJitAddDefinitions_Sycl(ceed, kernel_source, jit_source, constants)); CeedCallBackend(CeedJitGetFlags_Sycl(flags)); - CeedCallBackend(CeedJitCompileSource_Sycl(ceed, data->sycl_device, jit_source, il_binary, flags)); - CeedCallBackend(CeedLoadModule_Sycl(ceed, data->sycl_context, data->sycl_device, il_binary, sycl_module)); + + CeedCallBackend(CeedJitCompileSource_Sycl(ceed, data->sycl_device, jit_source, module_path, flags)); + + CeedCallBackend(CeedLoadModule_Sycl(ceed, data->sycl_context, data->sycl_device, module_path, sycl_module)); + return CEED_ERROR_SUCCESS; } // ------------------------------------------------------------------------------ -// Get a sycl kernel from an existing kernel_bundle -// -// TODO: Error handle lz calls +// Get a sycl kernel from an existing module // ------------------------------------------------------------------------------ -int CeedGetKernel_Sycl(Ceed ceed, const SyclModule_t *sycl_module, const std::string &kernel_name, sycl::kernel **sycl_kernel) { - Ceed_Sycl *data; - - CeedCallBackend(CeedGetData(ceed, &data)); - - // sycl::get_native returns std::vector for lz backend - // https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/supported/sycl_ext_oneapi_backend_level_zero.md - ze_module_handle_t lz_module = sycl::get_native(*sycl_module).front(); - - ze_kernel_desc_t lz_kernel_desc = {ZE_STRUCTURE_TYPE_KERNEL_DESC, nullptr, 0, kernel_name.c_str()}; - ze_kernel_handle_t lz_kernel; - ze_result_t lz_err = zeKernelCreate(lz_module, &lz_kernel_desc, &lz_kernel); - - if (ZE_RESULT_SUCCESS != lz_err) { - return CeedError(ceed, CEED_ERROR_BACKEND, "Failed to retrieve kernel from Level Zero module"); +template +int CeedGetKernel_Sycl(Ceed ceed, const SyclModule_t sycl_module, const std::string &kernel_name, SyclKernel_t *sycl_kernel) { + try { + *sycl_kernel = sycl_module->getFunction(kernel_name); + } catch (const std::exception& e) { + return CeedError((ceed), CEED_ERROR_BACKEND, e.what()); } - - *sycl_kernel = new sycl::kernel(sycl::make_kernel( - {*sycl_module, lz_kernel, sycl::ext::oneapi::level_zero::ownership::transfer}, data->sycl_context)); return CEED_ERROR_SUCCESS; } //------------------------------------------------------------------------------ // Run SYCL kernel for spatial dimension with shared memory //------------------------------------------------------------------------------ -int CeedRunKernelDimSharedSycl(Ceed ceed, sycl::kernel *kernel, const int grid_size, const int block_size_x, const int block_size_y, +template +int CeedRunKernelDimSharedSycl(Ceed ceed, SyclKernel_t *kernel, const int grid_size, const int block_size_x, const int block_size_y, const int block_size_z, const int shared_mem_size, void **kernel_args) { sycl::range<3> local_range(block_size_z, block_size_y, block_size_x); sycl::range<3> global_range(grid_size * block_size_z, block_size_y, block_size_x); @@ -173,15 +188,15 @@ int CeedRunKernelDimSharedSycl(Ceed ceed, sycl::kernel *kernel, const int grid_s //----------- // Order queue - Ceed_Sycl *ceed_Sycl; + // Ceed_Sycl *ceed_Sycl; - CeedCallBackend(CeedGetData(ceed, &ceed_Sycl)); - sycl::event e = ceed_Sycl->sycl_queue.ext_oneapi_submit_barrier(); + // CeedCallBackend(CeedGetData(ceed, &ceed_Sycl)); + // sycl::event e = ceed_Sycl->sycl_queue.ext_oneapi_submit_barrier(); - ceed_Sycl->sycl_queue.submit([&](sycl::handler &cgh) { - cgh.depends_on(e); - cgh.set_args(*kernel_args); - cgh.parallel_for(kernel_range, *kernel); - }); + // ceed_Sycl->sycl_queue.submit([&](sycl::handler &cgh) { + // cgh.depends_on(e); + // cgh.set_args(*kernel_args); + // cgh.parallel_for(kernel_range, *kernel); + // }); return CEED_ERROR_SUCCESS; }