Skip to content

Commit

Permalink
Integrating libprtc with sycl-compiler
Browse files Browse the repository at this point in the history
  • Loading branch information
@uumesh
uumesh committed Sep 25, 2024
1 parent f43fdc8 commit 5026b11
Show file tree
Hide file tree
Showing 2 changed files with 100 additions and 80 deletions.
11 changes: 8 additions & 3 deletions backends/sycl/ceed-sycl-compile.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -12,11 +12,16 @@
#include <map>
#include <sycl/sycl.hpp>

using SyclModule_t = sycl::kernel_bundle<sycl::bundle_state::executable>;
#include <libprtc/prtc.h>

using SyclModule_t = std::shared_ptr<prtc::DynamicLibrary>;

CEED_INTERN int CeedBuildModule_Sycl(Ceed ceed, const std::string &kernel_source, SyclModule_t **sycl_module,
const std::map<std::string, CeedInt> &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 <class SyclKernel_t>
CEED_INTERN int CeedGetKernel_Sycl(Ceed ceed, const SyclModule_t *sycl_module, const std::string &kernel_name, SyclKernel_t **sycl_kernel);

template <class SyclKernel_t>
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);
169 changes: 92 additions & 77 deletions backends/sycl/ceed-sycl-compile.sycl.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -12,14 +12,15 @@
#include <ceed/jit-tools.h>
#include <level_zero/ze_api.h>

#include <cstdlib>
#include <map>
#include <string>
#include <sstream>
#include <fstream>
#include <sycl/sycl.hpp>

#include "./online_compiler.hpp"
#include "ceed-sycl-common.hpp"

using ByteVector_t = std::vector<unsigned char>;
#include "libprtc/prtc.h"

//------------------------------------------------------------------------------
// Add defined constants at the beginning of kernel source
Expand Down Expand Up @@ -61,127 +62,141 @@ static int CeedJitAddDefinitions_Sycl(Ceed ceed, const std::string &kernel_sourc
// TODO: Add architecture flags, optimization flags
//------------------------------------------------------------------------------
static inline int CeedJitGetFlags_Sycl(std::vector<std::string> &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<std::string> &flags = {}) {
sycl::ext::libceed::online_compiler<sycl::ext::libceed::source_language::opencl_c> 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 : <cache_root>/<compiler_hash>/<build_options>/<kernel_name>/<kernel_source>

// Hash kernel name and source
std::hash<std::string> 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::backend::ext_oneapi_level_zero>(sycl_context);
auto lz_device = sycl::get_native<sycl::backend::ext_oneapi_level_zero>(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_<type> only throws errors for backend mismatch--assume we have vetted this already
*sycl_module = new SyclModule_t(sycl::make_kernel_bundle<sycl::backend::ext_oneapi_level_zero, sycl::bundle_state::executable>(
{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<std::string, CeedInt> &constants) {
int CeedBuildModule_Sycl(Ceed ceed, const std::string &kernel_source, SyclModule_t *sycl_module, const std::map<std::string, CeedInt> &constants) {
Ceed_Sycl *data;
std::string jit_source;
std::string module_path;
std::vector<std::string> 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<ze_module_handle_t> 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::backend::ext_oneapi_level_zero>(*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 <class SyclKernel_t>
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<SyclKernel_t>(kernel_name);
} catch (const std::exception& e) {
return CeedError((ceed), CEED_ERROR_BACKEND, e.what());
}

*sycl_kernel = new sycl::kernel(sycl::make_kernel<sycl::backend::ext_oneapi_level_zero>(
{*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 <class SyclKernel_t>
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);
sycl::nd_range<3> kernel_range(global_range, local_range);

//-----------
// 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;
}

0 comments on commit 5026b11

Please sign in to comment.