[SYCL] Employ cached kernel

sycl/include/CL/sycl/detail/program_impl.hpp

@@ -42,10 +42,13 @@ class program_impl {
   program_impl(const context &Context, vector_class<device> DeviceList)
       : Context(Context), Devices(DeviceList) {}
 
+  // Don't allow kernels caching for linked programs due to only compiled
+  // state of each and every program in the list and thus unknown state of
+  // caching resolution
   program_impl(vector_class<std::shared_ptr<program_impl>> ProgramList,
                string_class LinkOptions = "")
       : State(program_state::linked), LinkOptions(LinkOptions),
-        BuildOptions(LinkOptions) {
+        BuildOptions(LinkOptions), AllowKernelsCaching(false) {
     // Verify arguments
     if (ProgramList.empty()) {
       throw runtime_error("Non-empty vector of programs expected");
@@ -93,8 +96,10 @@ class program_impl {
     }
   }
 
+  // Disallow kernels caching for programs created by interoperability c-tor
   program_impl(const context &Context, RT::PiProgram Program)
-      : Program(Program), Context(Context), IsLinkable(true) {
+      : Program(Program), Context(Context), IsLinkable(true),
+        AllowKernelsCaching(false) {
 
     // TODO handle the case when cl_program build is in progress
     cl_uint NumDevices;
@@ -203,10 +208,13 @@ class program_impl {
     if (!is_host()) {
       OSModuleHandle M = OSUtil::getOSModuleHandle(AddressInThisModule);
       // If there are no build options, program can be safely cached
-      if (BuildOptions.empty()) {
-        Program = ProgramManager::getInstance().getBuiltOpenCLProgram(M, Context);
+      if (is_cacheable_with_build_options(BuildOptions)) {
+        Program =
+            ProgramManager::getInstance().getBuiltOpenCLProgram(M, Context);
         PI_CALL(RT::piProgramRetain, Program);
       } else {
+        AllowKernelsCaching = false;
+
         create_cl_program_with_il(M);
         build(BuildOptions);
       }
@@ -217,6 +225,9 @@ class program_impl {
   void build_with_source(string_class KernelSource,
                          string_class BuildOptions = "") {
     throw_if_state_is_not(program_state::none);
+
+    AllowKernelsCaching = false;
+
     // TODO should it throw if it's host?
     if (!is_host()) {
       create_cl_program_with_source(KernelSource);
@@ -231,11 +242,10 @@ class program_impl {
       check_device_feature_support<
           info::device::is_linker_available>(Devices);
       vector_class<RT::PiDevice> Devices(get_pi_devices());
-      RT::PiResult Err;
-      Err = PI_CALL_RESULT(RT::piProgramLink,
-                           detail::getSyclObjImpl(Context)->getHandleRef(),
-                           Devices.size(), Devices.data(), LinkOptions.c_str(),
-                           1, &Program, nullptr, nullptr, &Program);
+      RT::PiResult Err = PI_CALL_RESULT(
+            RT::piProgramLink, detail::getSyclObjImpl(Context)->getHandleRef(),
+            Devices.size(), Devices.data(), LinkOptions.c_str(), 1, &Program,
+            nullptr, nullptr, &Program);
       RT::piCheckThrow<compile_program_error>(Err);
       this->LinkOptions = LinkOptions;
       BuildOptions = LinkOptions;
@@ -411,16 +421,33 @@ class program_impl {
     return false;
   }
 
+  bool is_cacheable() const {
+    return is_cacheable_with_build_options(BuildOptions) && AllowKernelsCaching;
+  }
+
+  static bool
+  is_cacheable_with_build_options(const string_class &BuildOptions) {
+    return BuildOptions.empty();
+  }
+
   RT::PiKernel get_pi_kernel(const string_class &KernelName) const {
     RT::PiKernel Kernel;
-    RT::PiResult Err;
-    Err = PI_CALL_RESULT(RT::piKernelCreate, Program, KernelName.c_str(),
-                         &Kernel);
-    if (Err == PI_RESULT_INVALID_KERNEL_NAME) {
-      throw invalid_object_error(
-          "This instance of program does not contain the kernel requested");
+
+    if (is_cacheable()) {
+      OSModuleHandle M = OSUtil::getOSModuleHandle(AddressInThisModule);
+
+      Kernel = ProgramManager::getInstance().getOrCreateKernel(M, Context,
+                                                               KernelName);
+    } else {
+      RT::PiResult Err = PI_CALL_RESULT(RT::piKernelCreate, Program,
+                                        KernelName.c_str(), &Kernel);
+      if (Err == PI_RESULT_INVALID_KERNEL_NAME) {
+        throw invalid_object_error(
+            "This instance of program does not contain the kernel requested");
+      }
+      RT::piCheckResult(Err);
     }
-    RT::piCheckResult(Err);
+
     return Kernel;
   }
 
@@ -454,6 +481,11 @@ class program_impl {
   string_class CompileOptions;
   string_class LinkOptions;
   string_class BuildOptions;
+
+  // Only allow kernel caching for programs constructed with context only (or
+  // device list and context) and built with build_with_kernel_type with
+  // default build options
+  bool AllowKernelsCaching = true;
 };
 
 template <>

sycl/source/detail/program_manager/program_manager.cpp

@@ -118,6 +118,8 @@ RT::PiKernel ProgramManager::getOrCreateKernel(OSModuleHandle M,
   RT::PiKernel &Kernel = KernelsCache[KernelName];
   if (!Kernel) {
     PI_CALL(RT::piKernelCreate, Program, KernelName.c_str(), &Kernel);
+    // TODO need some user-friendly error/exception
+    // instead of currently obscure one
   }
   return Kernel;
 }

sycl/test/kernel-and-program/cache.cpp

@@ -0,0 +1,176 @@
+// RUN: %clangxx -fsycl %s -o %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+//==------------- kernel_cache.cpp - SYCL kernel/program test --------------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include <CL/sycl.hpp>
+
+class Functor {
+public:
+  void operator()(cl::sycl::item<1> Item) { (void)Item; }
+};
+
+struct TestContext {
+  int Data;
+  cl::sycl::queue Queue;
+  cl::sycl::buffer<int, 1> Buf;
+
+  TestContext() : Data(0), Buf(&Data, cl::sycl::range<1>(1)) {}
+
+  cl::sycl::program
+  getProgram(const cl::sycl::string_class &BuildOptions = "") {
+    cl::sycl::program Prog(Queue.get_context());
+
+    Prog.build_with_kernel_type<class SingleTask>(BuildOptions);
+
+    assert(Prog.get_state() == cl::sycl::program_state::linked &&
+           "Linked state was expected");
+
+    assert(Prog.has_kernel<class SingleTask>() &&
+           "Expecting SingleTask kernel exists");
+
+    return std::move(Prog);
+  }
+
+  cl::sycl::program getCompiledProgram() {
+    cl::sycl::program Prog(Queue.get_context());
+
+    Prog.compile_with_kernel_type<class SingleTask>();
+
+    assert(Prog.get_state() == cl::sycl::program_state::compiled &&
+           "Compiled state was expected");
+
+    return std::move(Prog);
+  }
+
+  cl::sycl::kernel getKernel(cl::sycl::program &Prog) {
+    auto Kernel = Prog.get_kernel<class SingleTask>();
+
+    Queue.submit([&](cl::sycl::handler &CGH) {
+      auto acc = Buf.get_access<cl::sycl::access::mode::read_write>(CGH);
+      CGH.single_task<class SingleTask>(Kernel, [=]() { acc[0] = acc[0] + 1; });
+    });
+
+    return std::move(Kernel);
+  }
+};
+
+namespace pi = cl::sycl::detail::pi;
+namespace RT = cl::sycl::RT;
+
+static void testProgramCachePositive() {
+  TestContext TestCtx;
+
+  auto Prog = TestCtx.getProgram();
+
+  auto *CLProg = cl::sycl::detail::getSyclObjImpl(Prog)->getHandleRef();
+
+  auto *Ctx = cl::sycl::detail::getRawSyclObjImpl(Prog.get_context());
+
+  assert(Ctx->getCachedPrograms().size() == 1 &&
+         "Expecting only a single element in program cache");
+  assert(Ctx->getCachedPrograms().begin()->second ==
+             pi::cast<pi_program>(CLProg) &&
+         "Invalid data in programs cache");
+}
+
+static void testProgramCacheNegativeCustomBuildOptions() {
+  TestContext TestCtx;
+
+  auto Prog = TestCtx.getProgram("-g");
+
+  auto *Ctx = cl::sycl::detail::getRawSyclObjImpl(Prog.get_context());
+
+  assert(Ctx->getCachedPrograms().size() == 0 &&
+         "Expecting empty program cache");
+}
+
+static void testKernelCachePositive() {
+  TestContext TestCtx;
+
+  auto Prog = TestCtx.getProgram();
+  auto Kernel = TestCtx.getKernel(Prog);
+
+  if (!TestCtx.Queue.is_host()) {
+    auto *CLProg = cl::sycl::detail::getSyclObjImpl(Prog)->getHandleRef();
+    auto *CLKernel = cl::sycl::detail::getSyclObjImpl(Kernel)->getHandleRef();
+
+    auto *Ctx = cl::sycl::detail::getRawSyclObjImpl(Prog.get_context());
+
+    assert(Ctx->getCachedKernels().size() == 1 &&
+           "Expecting only a single element in kernels cache");
+    assert(Ctx->getCachedKernels().begin()->first ==
+               pi::cast<pi_program>(CLProg) &&
+           "Invalid program key in kernels cache");
+    assert(Ctx->getCachedKernels().begin()->second.size() == 1 &&
+           "Expecting only a single kernel for the program");
+    assert(Ctx->getCachedKernels().begin()->second.begin()->second ==
+               pi::cast<pi_kernel>(CLKernel) &&
+           "Invalid data in kernels cache");
+  }
+}
+
+void testKernelCacheNegativeLinkedProgram() {
+  TestContext TestCtx;
+
+  auto Prog1 = TestCtx.getCompiledProgram();
+  auto Prog2 = TestCtx.getCompiledProgram();
+
+  auto LinkedProg = cl::sycl::program({Prog1, Prog2});
+
+  auto Kernel = TestCtx.getKernel(LinkedProg);
+
+  if (!TestCtx.Queue.is_host()) {
+    auto *Ctx = cl::sycl::detail::getRawSyclObjImpl(LinkedProg.get_context());
+
+    assert(Ctx->getCachedKernels().size() == 0 &&
+           "Unexpected data in kernels cache");
+  }
+}
+
+void testKernelCacheNegativeOCLProgram() {
+  TestContext TestCtx;
+
+  auto SyclProg = TestCtx.getProgram();
+
+  auto OclProg = cl::sycl::program(TestCtx.Queue.get_context(), SyclProg.get());
+
+  auto Kernel = TestCtx.getKernel(OclProg);
+
+  if (!TestCtx.Queue.is_host()) {
+    auto *Ctx = cl::sycl::detail::getRawSyclObjImpl(OclProg.get_context());
+
+    assert(Ctx->getCachedKernels().size() == 0 &&
+           "Unexpected data in kernels cache");
+  }
+}
+
+void testKernelCacheNegativeCustomBuildOptions() {
+  TestContext TestCtx;
+
+  auto Prog = TestCtx.getProgram("-g");
+  auto Kernel = TestCtx.getKernel(Prog);
+
+  if (!TestCtx.Queue.is_host()) {
+    auto *Ctx = cl::sycl::detail::getRawSyclObjImpl(Prog.get_context());
+    assert(Ctx->getCachedKernels().size() == 0 &&
+           "Unexpected data in kernels cache");
+  }
+}
+
+int main() {
+  testProgramCachePositive();
+  testProgramCacheNegativeCustomBuildOptions();
+
+  testKernelCachePositive();
+  testKernelCacheNegativeLinkedProgram();
+  testKernelCacheNegativeOCLProgram();
+  testKernelCacheNegativeCustomBuildOptions();
+
+  return 0;
+}