diff --git a/backends/sycl-ref/ceed-sycl-ref-qfunction-load.sycl.cpp b/backends/sycl-ref/ceed-sycl-ref-qfunction-load.sycl.cpp index bc0af096fd..7654b8f265 100644 --- a/backends/sycl-ref/ceed-sycl-ref-qfunction-load.sycl.cpp +++ b/backends/sycl-ref/ceed-sycl-ref-qfunction-load.sycl.cpp @@ -86,6 +86,7 @@ extern "C" int CeedQFunctionBuildKernel_Sycl(CeedQFunction qf) { // This needs to be revisited if all qfunctions require this. // code << "__attribute__((intel_reqd_sub_group_size(" << SUB_GROUP_SIZE_QF << "))) extern \"C\" void " << kernel_name code << "#include \n\n"; + code << "class CeedQFunction_" << qf_name_view << ";\n\n"; code << "extern \"C\" void " << kernel_name << "(sycl::queue &sycl_queue, sycl::nd_range<1> kernel_range, void *ctx, CeedInt Q, Fields_Sycl *fields) {\n"; @@ -101,7 +102,7 @@ extern "C" int CeedQFunctionBuildKernel_Sycl(CeedQFunction qf) { // Output parameters code << " " - << "const CeedScalar *fields_outputs[" << num_output_fields << "];\n"; + << "CeedScalar *fields_outputs[" << num_output_fields << "];\n"; for (CeedInt i = 0; i < num_output_fields; ++i) { code << " " << "fields_outputs[" << i << "] = fields->outputs[" << i << "];\n"; @@ -115,7 +116,7 @@ extern "C" int CeedQFunctionBuildKernel_Sycl(CeedQFunction qf) { // Begin kernel function body code << " " << "sycl_queue.parallel_for(kernel_range, e, " - << "[=](sycl::id<1> id) {\n"; + << "[=](sycl::nd_item<1> item) {\n"; // Inputs code << " // Input fields\n"; @@ -139,7 +140,7 @@ extern "C" int CeedQFunctionBuildKernel_Sycl(CeedQFunction qf) { } code << "};\n\n"; - code << " const CeedInt q = id;\n\n"; + code << " const CeedInt q = item.get_global_id(0);\n\n"; code << " if(q < Q) { \n\n"; @@ -174,7 +175,9 @@ extern "C" int CeedQFunctionBuildKernel_Sycl(CeedQFunction qf) { // Compile kernel CeedCallBackend(CeedBuildModule_Sycl(ceed, code.str(), impl->sycl_module)); + std::cout << " Module built \n"; CeedCallBackend(CeedGetKernel_Sycl(ceed, impl->sycl_module, kernel_name, &impl->QFunction)); + std::cout << " Kernel retrieved \n"; // Cleanup CeedCallBackend(CeedFree(&read_write_kernel_path)); diff --git a/backends/sycl/ceed-sycl-compile.hpp b/backends/sycl/ceed-sycl-compile.hpp index a977403e67..e2e5a6e2fc 100644 --- a/backends/sycl/ceed-sycl-compile.hpp +++ b/backends/sycl/ceed-sycl-compile.hpp @@ -25,6 +25,7 @@ int CeedGetKernel_Sycl(Ceed ceed, const SyclModule_t sycl_module, const std::str try { *sycl_kernel = sycl_module->getFunction(kernel_name); } catch (const std::exception& e) { + std::cout<< "\nUnable to retrieve kernel\n"; return CeedError((ceed), CEED_ERROR_BACKEND, e.what()); } return CEED_ERROR_SUCCESS; diff --git a/backends/sycl/ceed-sycl-compile.sycl.cpp b/backends/sycl/ceed-sycl-compile.sycl.cpp index bf7e1c9e6f..8db7bfc1c9 100644 --- a/backends/sycl/ceed-sycl-compile.sycl.cpp +++ b/backends/sycl/ceed-sycl-compile.sycl.cpp @@ -33,6 +33,8 @@ static int CeedJitAddDefinitions_Sycl(Ceed ceed, const std::string &kernel_sourc const char *jit_defs_path, *jit_defs_source; const char *sycl_jith_path = "ceed/jit-source/sycl/sycl-jit.h"; + oss << "#include\n\n"; + // Prepend defined constants for (const auto &[name, value] : constants) { oss << "#define " << name << " " << value << "\n"; @@ -65,7 +67,7 @@ static int CeedJitAddDefinitions_Sycl(Ceed ceed, const std::string &kernel_sourc static inline int CeedJitGetFlags_Sycl(std::vector &flags) { // 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")}; + flags = {std::string("-fsycl"), std::string("-fno-sycl-id-queries-fit-in-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\" ")) @@ -117,6 +119,10 @@ static inline int CeedJitCompileSource_Sycl(Ceed ceed, const sycl::device &sycl_ 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"; + std::string mkdir_command = std::string("mkdir -p ") + cache_path; + prtc::ShellCommand make_dir(mkdir_command); + auto [mkdir_success, mkdir_message] = make_dir.result(); + if(!mkdir_success) return CeedError((ceed), CEED_ERROR_BACKEND, mkdir_message.c_str()); // Write source string to file std::ofstream source_file; @@ -126,9 +132,9 @@ static inline int CeedJitCompileSource_Sycl(Ceed ceed, const sycl::device &sycl_ // TODO: Get compiler-path and flags from env or some other means prtc::ShellCompiler compiler("icpx","-o","-c","-fPIC","-shared"); - const auto [build_success, message] = compiler.compileAndLink(source_file_path,object_file_path,flags); + const auto [build_success, build_message] = compiler.compileAndLink(source_file_path,object_file_path,flags); // Q: Should we always output the compiler output in verbose/debug mode? - if (!build_success) return CeedError((ceed), CEED_ERROR_BACKEND, message.c_str()); + if (!build_success) return CeedError((ceed), CEED_ERROR_BACKEND, build_message.c_str()); return CEED_ERROR_SUCCESS; } diff --git a/backends/sycl/libprtc/dynamic_library.h b/backends/sycl/libprtc/dynamic_library.h index 709cf72800..98f1d56358 100755 --- a/backends/sycl/libprtc/dynamic_library.h +++ b/backends/sycl/libprtc/dynamic_library.h @@ -5,6 +5,7 @@ #include #include "function_traits.h" +#include namespace prtc { @@ -22,7 +23,8 @@ class DynamicLibrary : public std::enable_shared_from_this { template F getFunction(const std::string& name) { - return function_cast(getSymbol(name)); + // return function_cast(getSymbol(name)); + return reinterpret_cast(getSymbol(name)); } private: diff --git a/backends/sycl/libprtc/shell_compiler.cpp b/backends/sycl/libprtc/shell_compiler.cpp index 44aa6b0ac6..ebf86ce4af 100755 --- a/backends/sycl/libprtc/shell_compiler.cpp +++ b/backends/sycl/libprtc/shell_compiler.cpp @@ -7,7 +7,7 @@ #include "shell_command.h" namespace prtc { -namespace { +// namespace { std::string concatenateFlags(const std::vector& flags) { std::string all_flags{}; @@ -17,7 +17,7 @@ std::string concatenateFlags(const std::vector& flags) { // Remove last space if (!all_flags.empty()) all_flags.pop_back(); return all_flags; -} +// } } // namespace