trying to debug. most changes are WIP

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 <vector>\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<CeedQFunction_" << qf_name_view << ">(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));

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<SyclKernel_t>(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;

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<sycl/sycl.hpp>\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<std::string> &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;
 }
 

backends/sycl/libprtc/dynamic_library.h

@@ -5,6 +5,7 @@
 #include <string>
 
 #include "function_traits.h"
+#include <functional>
 
 namespace prtc {
 
@@ -22,7 +23,8 @@ class DynamicLibrary : public std::enable_shared_from_this<DynamicLibrary> {
 
   template <class F>
   F getFunction(const std::string& name) {
-    return function_cast<F>(getSymbol(name));
+    // return function_cast<F>(getSymbol(name));
+    return reinterpret_cast<F>(getSymbol(name));
   }
 
  private:

backends/sycl/libprtc/shell_compiler.cpp

@@ -7,7 +7,7 @@
 #include "shell_command.h"
 
 namespace prtc {
-namespace {
+// namespace {
 
 std::string concatenateFlags(const std::vector<std::string>& flags) {
   std::string all_flags{};
@@ -17,7 +17,7 @@ std::string concatenateFlags(const std::vector<std::string>& flags) {
   // Remove last space
   if (!all_flags.empty()) all_flags.pop_back();
   return all_flags;
-}
+// }
 
 }  // namespace