Fixed the cache to store binaries instead of OpenCL programs

5 files changed, 48 insertions, 27 deletions

diff --git a/include/clblast.h b/include/clblast.h
index 57fca119..e473adbe 100644
--- a/include/clblast.h
+++ b/include/clblast.h
@@ -556,7 +556,7 @@ StatusCode Trsm(const Layout layout, const Side side, const Triangle triangle, c
 
 // CLBlast stores binaries of compiled kernels into a cache in case the same kernel is used later on
 // for the same device. This cache can be cleared to free up system memory or in case of debugging.
-StatusCode ClearCompiledProgramCache();
+StatusCode ClearCache();
 
 // =================================================================================================
 
diff --git a/include/clblast_c.h b/include/clblast_c.h
index e23f0305..45e50cff 100644
--- a/include/clblast_c.h
+++ b/include/clblast_c.h
@@ -1074,7 +1074,7 @@ StatusCode PUBLIC_API CLBlastZtrsm(const Layout layout, const Side side, const T
 
 // CLBlast stores binaries of compiled kernels into a cache in case the same kernel is used later on
 // for the same device. This cache can be cleared to free up system memory or in case of debugging.
-StatusCode PUBLIC_API CLBlastClearCompiledProgramCache();
+StatusCode PUBLIC_API CLBlastClearCache();
 
 // =================================================================================================
 
diff --git a/include/internal/cache.h b/include/internal/cache.h
index 44fad68d..fa33b78f 100644
--- a/include/internal/cache.h
+++ b/include/internal/cache.h
@@ -24,9 +24,9 @@ namespace clblast {
 namespace cache {
 // =================================================================================================
 
-// The cache of compiled OpenCL programs, along with some meta-data
-struct ProgramCache {
-  Program program;
+// The cache of compiled OpenCL binaries, along with some meta-data
+struct BinaryCache {
+  std::string binary;
   std::string device_name;
   Precision precision;
   std::string routine_name_;
@@ -41,28 +41,28 @@ struct ProgramCache {
 };
 
 // The actual cache, implemented as a vector of the above data-type, and its mutex
-static std::vector<ProgramCache> program_cache_;
-static std::mutex program_cache_mutex_;
+static std::vector<BinaryCache> binary_cache_;
+static std::mutex binary_cache_mutex_;
 
 // =================================================================================================
 
-// Stores the compiled program in the cache
-void StoreProgramToCache(const Program& program, const std::string &device_name,
-                         const Precision &precision, const std::string &routine_name);
+// Stores the compiled binary in the cache
+void StoreBinaryToCache(const std::string& binary, const std::string &device_name,
+                        const Precision &precision, const std::string &routine_name);
 
-// Queries the cache and retrieves a matching program. Assumes that the match is available, throws
+// Queries the cache and retrieves a matching binary. Assumes that the match is available, throws
 // otherwise.
-const Program& GetProgramFromCache(const std::string &device_name, const Precision &precision,
-                                   const std::string &routine_name);
+const std::string& GetBinaryFromCache(const std::string &device_name, const Precision &precision,
+                                      const std::string &routine_name);
 
 // Queries the cache to see whether or not the compiled kernel is already there
-bool ProgramIsInCache(const std::string &device_name, const Precision &precision,
-                      const std::string &routine_name);
+bool BinaryIsInCache(const std::string &device_name, const Precision &precision,
+                     const std::string &routine_name);
 
 // =================================================================================================
 
-// Clears the cache of stored program binaries
-StatusCode ClearCompiledProgramCache();
+// Clears the cache of stored binaries
+StatusCode ClearCache();
 
 // =================================================================================================
 } // namespace cache
diff --git a/include/internal/clpp11.h b/include/internal/clpp11.h
index 543d423a..b865ab1e 100644
--- a/include/internal/clpp11.h
+++ b/include/internal/clpp11.h
@@ -283,7 +283,7 @@ class Program {
  public:
   // Note that there is no constructor based on the regular OpenCL data-type because of extra state
 
-  // Regular constructor with memory management
+  // Source-based constructor with memory management
   explicit Program(const Context &context, std::string source):
       program_(new cl_program, [](cl_program* p) { CheckError(clReleaseProgram(*p)); delete p; }),
       length_(source.length()),
@@ -294,6 +294,22 @@ class Program {
     CheckError(status);
   }
 
+  // Binary-based constructor with memory management
+  explicit Program(const Device &device, const Context &context, const std::string& binary):
+      program_(new cl_program, [](cl_program* p) { CheckError(clReleaseProgram(*p)); delete p; }),
+      length_(binary.length()),
+      source_(binary),
+      source_ptr_(&source_[0]) {
+    auto status1 = CL_SUCCESS;
+    auto status2 = CL_SUCCESS;
+    const cl_device_id dev = device();
+    *program_ = clCreateProgramWithBinary(context(), 1, &dev, &length_,
+                                          reinterpret_cast<const unsigned char**>(&source_ptr_),
+                                          &status1, &status2);
+    CheckError(status1);
+    CheckError(status2);
+  }
+
   // Compiles the device program and returns whether or not there where any warnings/errors
   BuildStatus Build(const Device &device, std::vector<std::string> &options) {
     auto options_string = std::accumulate(options.begin(), options.end(), std::string{" "});
@@ -322,7 +338,7 @@ class Program {
     return result;
   }
 
-  // Retrieves an intermediate representation of the compiled program
+  // Retrieves a binary or an intermediate representation of the compiled program
   std::string GetIR() const {
     auto bytes = size_t{0};
     CheckError(clGetProgramInfo(*program_, CL_PROGRAM_BINARY_SIZES, sizeof(size_t), &bytes, nullptr));
@@ -338,7 +354,7 @@ class Program {
  private:
   std::shared_ptr<cl_program> program_;
   size_t length_;
-  std::string source_;
+  std::string source_; // Note: the source can also be a binary or IR
   const char* source_ptr_;
 };
 
@@ -633,7 +649,8 @@ class Kernel {
     // Launches the kernel while waiting for other events
     CheckError(clEnqueueNDRangeKernel(queue(), *kernel_, static_cast<cl_uint>(global.size()),
                                       nullptr, global.data(), local.data(),
-                                      waitForEventsPlain.size(), waitForEventsPlain.data(),
+                                      static_cast<cl_uint>(waitForEventsPlain.size()),
+                                      waitForEventsPlain.data(),
                                       event));
   }
 
diff --git a/include/internal/routine.h b/include/internal/routine.h
index 013769d8..32be6012 100644
--- a/include/internal/routine.h
+++ b/include/internal/routine.h
@@ -84,19 +84,23 @@ class Routine {
                                     const bool upper = false, const bool lower = false,
                                     const bool diagonal_imag_zero = false);
 
-  // Stores a newly compiled program into the cache
-  void StoreProgramToCache(const Program& program) const {
-    return cache::StoreProgramToCache(program, device_name_, precision_, routine_name_);
+  // Stores a newly compiled binary into the cache
+  void StoreBinaryToCache(const std::string& binary) const {
+    return cache::StoreBinaryToCache(binary, device_name_, precision_, routine_name_);
   }
 
   // Queries the cache and retrieve either a matching program or a boolean whether a match exists.
   // The first assumes that the program is available in the cache and will throw an exception
   // otherwise.
-  const Program& GetProgramFromCache() const {
-    return cache::GetProgramFromCache(device_name_, precision_, routine_name_);
+  Program GetProgramFromCache() const {
+    auto& binary = cache::GetBinaryFromCache(device_name_, precision_, routine_name_);
+    auto program = Program(device_, context_, binary);
+    auto options = std::vector<std::string>();
+    program.Build(device_, options);
+    return program;
   }
   bool ProgramIsInCache() const {
-    return cache::ProgramIsInCache(device_name_, precision_, routine_name_);
+    return cache::BinaryIsInCache(device_name_, precision_, routine_name_);
   }
 
   // Non-static variable for the precision. Note that the same variable (but static) might exist in