[SYCL] Fix regression with program building

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

@@ -44,13 +44,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
+  // Kernels caching for linked programs won't be allowed 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), AllowKernelsCaching(false) {
+        BuildOptions(LinkOptions) {
     // Verify arguments
     if (ProgramList.empty()) {
       throw runtime_error("Non-empty vector of programs expected");
@@ -96,10 +96,9 @@ class program_impl {
     }
   }
 
-  // Disallow kernels caching for programs created by interoperability c-tor
+  // Kernel caching for programs created by interoperability c-tor isn't allowed
   program_impl(const context &Context, RT::PiProgram Program)
-      : Program(Program), Context(Context), IsLinkable(true),
-        AllowKernelsCaching(false) {
+      : Program(Program), Context(Context), IsLinkable(true) {
 
     // TODO handle the case when cl_program build is in progress
     cl_uint NumDevices;
@@ -210,13 +209,12 @@ class program_impl {
     if (!is_host()) {
       OSModuleHandle M = OSUtil::getOSModuleHandle(AddressInThisModule);
       // If there are no build options, program can be safely cached
-      if (is_cacheable_with_build_options(BuildOptions)) {
+      if (is_cacheable_with_options(BuildOptions)) {
+        IsProgramAndKernelCachingAllowed = true;
         Program =
             ProgramManager::getInstance().getBuiltOpenCLProgram(M, Context);
         PI_CALL(piProgramRetain)(Program);
       } else {
-        AllowKernelsCaching = false;
-
         create_cl_program_with_il(M);
         build(BuildOptions);
       }
@@ -227,9 +225,6 @@ 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);
@@ -425,12 +420,12 @@ class program_impl {
   }
 
   bool is_cacheable() const {
-    return is_cacheable_with_build_options(BuildOptions) && AllowKernelsCaching;
+    return IsProgramAndKernelCachingAllowed;
   }
 
   static bool
-  is_cacheable_with_build_options(const string_class &BuildOptions) {
-    return BuildOptions.empty();
+  is_cacheable_with_options(const string_class &Options) {
+    return Options.empty();
   }
 
   RT::PiKernel get_pi_kernel(const string_class &KernelName) const {
@@ -488,7 +483,7 @@ class program_impl {
   // 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;
+  bool IsProgramAndKernelCachingAllowed = false;
 };
 
 template <>

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

@@ -10,6 +10,13 @@
 
 #include <CL/sycl.hpp>
 
+namespace pi = cl::sycl::detail::pi;
+namespace RT = cl::sycl::RT;
+
+#define KERNEL_NAME_SRC "kernel_source"
+#define TEST_SOURCE "kernel void " KERNEL_NAME_SRC "(global int* a) " \
+  "{ a[get_global_id(0)] += 1; }\n"
+
 class Functor {
 public:
   void operator()(cl::sycl::item<1> Item) { (void)Item; }
@@ -37,6 +44,21 @@ struct TestContext {
     return std::move(Prog);
   }
 
+  cl::sycl::program
+  getProgramWSource(const cl::sycl::string_class &BuildOptions = "") {
+    cl::sycl::program Prog(Queue.get_context());
+
+    Prog.build_with_source(TEST_SOURCE, 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());
 
@@ -48,6 +70,43 @@ struct TestContext {
     return std::move(Prog);
   }
 
+  cl::sycl::program
+  getCompiledAndLinkedProgram(const cl::sycl::string_class &CompileOptions = "",
+                              const cl::sycl::string_class &LinkOptions = "") {
+    cl::sycl::program Prog(Queue.get_context());
+
+    Prog.compile_with_kernel_type<class SingleTask>(CompileOptions);
+
+    assert(Prog.get_state() == cl::sycl::program_state::compiled &&
+           "Compiled state was expected");
+
+    Prog.link(LinkOptions);
+
+    assert(Prog.get_state() == cl::sycl::program_state::linked &&
+           "Linked state was expected");
+
+    return std::move(Prog);
+  }
+
+  cl::sycl::program
+  getCompiledAndLinkedProgramWSource(
+      const cl::sycl::string_class &CompileOptions = "",
+      const cl::sycl::string_class &LinkOptions = "") {
+    cl::sycl::program Prog(Queue.get_context());
+
+    Prog.compile_with_source(TEST_SOURCE, CompileOptions);
+
+    assert(Prog.get_state() == cl::sycl::program_state::compiled &&
+           "Compiled state was expected");
+
+    Prog.link(LinkOptions);
+
+    assert(Prog.get_state() == cl::sycl::program_state::linked &&
+           "Linked state was expected");
+
+    return std::move(Prog);
+  }
+
   cl::sycl::kernel getKernel(cl::sycl::program &Prog) {
     auto Kernel = Prog.get_kernel<class SingleTask>();
 
@@ -58,10 +117,13 @@ struct TestContext {
 
     return std::move(Kernel);
   }
-};
 
-namespace pi = cl::sycl::detail::pi;
-namespace RT = cl::sycl::RT;
+  cl::sycl::kernel getKernelWSource(cl::sycl::program &Prog) {
+    auto Kernel = Prog.get_kernel(KERNEL_NAME_SRC);
+
+    return std::move(Kernel);
+  }
+};
 
 static void testProgramCachePositive() {
   TestContext TestCtx;
@@ -90,6 +152,86 @@ static void testProgramCacheNegativeCustomBuildOptions() {
          "Expecting empty program cache");
 }
 
+static void testProgramCacheNegativeCompileLinkCustomOpts() {
+  TestContext TestCtx;
+
+  {
+    auto Prog = TestCtx.getCompiledAndLinkedProgram();
+
+    auto *Ctx = cl::sycl::detail::getRawSyclObjImpl(Prog.get_context());
+
+    assert(Ctx->getCachedPrograms().size() == 0 &&
+           "Expecting empty program cache");
+  }
+
+  {
+    auto Prog = TestCtx.getCompiledAndLinkedProgram("-g", "-cl-no-signed-zeroes");
+
+    auto *Ctx = cl::sycl::detail::getRawSyclObjImpl(Prog.get_context());
+
+    assert(Ctx->getCachedPrograms().size() == 0 &&
+           "Expecting empty program cache");
+  }
+
+  {
+    auto Prog = TestCtx.getCompiledAndLinkedProgram("", "-cl-no-signed-zeroes");
+
+    auto *Ctx = cl::sycl::detail::getRawSyclObjImpl(Prog.get_context());
+
+    assert(Ctx->getCachedPrograms().size() == 0 &&
+           "Expecting empty program cache");
+  }
+
+  {
+    auto Prog = TestCtx.getCompiledAndLinkedProgram("-g", "");
+
+    auto *Ctx = cl::sycl::detail::getRawSyclObjImpl(Prog.get_context());
+
+    assert(Ctx->getCachedPrograms().size() == 0 &&
+           "Expecting empty program cache");
+  }
+}
+
+static void testProgramCacheNegativeCompileLinkSource() {
+  TestContext TestCtx;
+
+  {
+    auto Prog = TestCtx.getCompiledAndLinkedProgramWSource();
+
+    auto *Ctx = cl::sycl::detail::getRawSyclObjImpl(Prog.get_context());
+
+    assert(Ctx->getCachedPrograms().size() == 0 &&
+           "Expecting empty program cache");
+  }
+
+  {
+    auto Prog = TestCtx.getCompiledAndLinkedProgramWSource("-g", "-cl-no-signed-zeroes");
+
+    auto *Ctx = cl::sycl::detail::getRawSyclObjImpl(Prog.get_context());
+
+    assert(Ctx->getCachedPrograms().size() == 0 &&
+           "Expecting empty program cache");
+  }
+
+  {
+    auto Prog = TestCtx.getCompiledAndLinkedProgramWSource("", "-cl-no-signed-zeroes");
+
+    auto *Ctx = cl::sycl::detail::getRawSyclObjImpl(Prog.get_context());
+
+    assert(Ctx->getCachedPrograms().size() == 0 &&
+           "Expecting empty program cache");
+  }
+
+  {
+    auto Prog = TestCtx.getCompiledAndLinkedProgramWSource("-g", "");
+
+    auto *Ctx = cl::sycl::detail::getRawSyclObjImpl(Prog.get_context());
+
+    assert(Ctx->getCachedPrograms().size() == 0 &&
+           "Expecting empty program cache");
+  }
+}
+
 static void testKernelCachePositive() {
   TestContext TestCtx;
 
@@ -163,14 +305,115 @@ void testKernelCacheNegativeCustomBuildOptions() {
   }
 }
 
+void testKernelCacheNegativeCompileLink() {
+  TestContext TestCtx;
+
+  {
+    auto Prog = TestCtx.getCompiledAndLinkedProgram();
+    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");
+    }
+  }
+
+  {
+    TestContext TestCtx1;
+    auto Prog = TestCtx1.getCompiledAndLinkedProgram("-g", "-cl-no-signed-zeroes");
+    auto Kernel = TestCtx1.getKernel(Prog);
+
+    if (!TestCtx1.Queue.is_host()) {
+      auto *Ctx = cl::sycl::detail::getRawSyclObjImpl(Prog.get_context());
+      assert(Ctx->getCachedKernels().size() == 0 &&
+             "Unexpected data in kernels cache");
+    }
+  }
+
+  {
+    auto Prog = TestCtx.getCompiledAndLinkedProgram("-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");
+    }
+  }
+
+  {
+    auto Prog = TestCtx.getCompiledAndLinkedProgram("", "-cl-no-signed-zeroes");
+    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");
+    }
+  }
+}
+
+void testKernelCacheNegativeCompileLinkSource() {
+  TestContext TestCtx;
+
+  {
+    auto Prog = TestCtx.getCompiledAndLinkedProgramWSource();
+    auto Kernel = TestCtx.getKernelWSource(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");
+    }
+  }
+
+  {
+    auto Prog = TestCtx.getCompiledAndLinkedProgramWSource("-g", "-cl-no-signed-zeroes");
+    auto Kernel = TestCtx.getKernelWSource(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");
+    }
+  }
+
+  {
+    auto Prog = TestCtx.getCompiledAndLinkedProgramWSource("-g", "");
+    auto Kernel = TestCtx.getKernelWSource(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");
+    }
+  }
+
+  {
+    auto Prog = TestCtx.getCompiledAndLinkedProgramWSource("", "-cl-no-signed-zeroes");
+    auto Kernel = TestCtx.getKernelWSource(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();
+  testProgramCacheNegativeCompileLinkCustomOpts();
+  testProgramCacheNegativeCompileLinkSource();
 
   testKernelCachePositive();
   testKernelCacheNegativeLinkedProgram();
   testKernelCacheNegativeOCLProgram();
   testKernelCacheNegativeCustomBuildOptions();
+  testKernelCacheNegativeCompileLink();
+  testKernelCacheNegativeCompileLinkSource();
 
   return 0;
 }