From c73fe5dbaf00a9967b04662cf5d05d58ac1aa4fe Mon Sep 17 00:00:00 2001
From: Kris Rowe <kris.rowe@anl.gov>
Date: Thu, 6 Jun 2024 19:10:18 +0000
Subject: [PATCH] Save WIP.

---
 backends/sycl/ceed-sycl-compile.hpp      |  2 +-
 backends/sycl/ceed-sycl-compile.sycl.cpp | 18 +++++++++++++-----
 2 files changed, 14 insertions(+), 6 deletions(-)

diff --git a/backends/sycl/ceed-sycl-compile.hpp b/backends/sycl/ceed-sycl-compile.hpp
index d24e1ee617..e05f9e6079 100644
--- a/backends/sycl/ceed-sycl-compile.hpp
+++ b/backends/sycl/ceed-sycl-compile.hpp
@@ -15,7 +15,7 @@
 #include <libprtc/prtc.h>
 
 using SyclModule_t = std::shared_ptr<prtc::DynamicLibrary>;
-using SyclKernel_t = void*; // Revisit
+using SyclKernel_t = void*; // Revisit this
 
 CEED_INTERN int CeedBuildModule_Sycl(Ceed ceed, const std::string &kernel_source, SyclModule_t *sycl_module,
                                      const std::map<std::string, CeedInt> &constants = {});
diff --git a/backends/sycl/ceed-sycl-compile.sycl.cpp b/backends/sycl/ceed-sycl-compile.sycl.cpp
index c667298df9..401f425bc9 100644
--- a/backends/sycl/ceed-sycl-compile.sycl.cpp
+++ b/backends/sycl/ceed-sycl-compile.sycl.cpp
@@ -12,6 +12,7 @@
 #include <ceed/jit-tools.h>
 #include <level_zero/ze_api.h>
 
+#include <cstdlib>
 #include <map>
 #include <sstream>
 #include <sycl/sycl.hpp>
@@ -67,18 +68,25 @@ static inline int CeedJitGetFlags_Sycl(std::vector<std::string> &flags) {
 
 //------------------------------------------------------------------------------
 // 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 &kernel_source, std::string& output_path,
                                             const std::vector<std::string> &flags = {}) {
-  
-  // Q: Should we write the kernel source string to a file here
-  //    or handle this in a separarte function;
 
-  // Should outputfile name be generated here?
+  // Get cache path from env variable
+  const char* cache_path = std::getenv("CEED_CACHE_DIR");
+
+  // 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
+
+  // Write source string to file
+  std::string source_path;
 
   // 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(kernel_source,output_path,flags);
+  const auto [success, message] = compiler.compileAndLink(source_path,output_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;