[SYCL] moving kernel_compiler sycl cache testing to its own test. (#16727)

cperkinsintel · web-flow · commit e8ed1342ed62 · 2025-01-23T10:42:55.000-08:00
moving kernel_compiler sycl cache testing to its own test. Also updating
some of the 'unsupported-intended' information. This also increments the
#include &lt;sycl.hpp&gt; counter, not because the new test does that (it does
not), but the kernel string it contains does and that's incorrectly
getting picked up by the no_sycl_in_hpp test.
diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_opencl.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_opencl.cpp
@@ -8,6 +8,7 @@
 
 // REQUIRES: ocloc && (opencl || level_zero)
 // UNSUPPORTED: accelerator
+// UNSUPPORTED-INTENDED: while accelerator is AoT only, this cannot run there.
 
 // -- Test the kernel_compiler with OpenCL source.
 // RUN: %{build} -o %t.out
diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp
@@ -7,12 +7,10 @@
 //===----------------------------------------------------------------------===//
 
 // REQUIRES: (opencl || level_zero)
-// UNSUPPORTED: accelerator
+// REQUIRES: aspect-usm_device_allocations
 
-// Flaky timeout on CPU. Enable when fixed.
-// Depends on SPIR-V Backend & run-time drivers version.
-// UNSUPPORTED: spirv-backend && cpu
-// UNSUPPORTED-TRACKER: CMPLRLLVM-64705
+// UNSUPPORTED: accelerator
+// UNSUPPORTED-INTENDED: while accelerator is AoT only, this cannot run there.
 
 // -- Test the kernel_compiler with SYCL source.
 // RUN: %{build} -o %t.out
@@ -23,28 +21,6 @@
 // RUN: %{run} %t.out
 // RUN: %{l0_leak_check} %{run} %t.out
 
-// -- Test again, with caching.
-// 'reading-from-cache' is just a string we pass to differentiate between the
-// two runs.
-
-// DEFINE: %{cache_vars} = %{l0_leak_check} env SYCL_CACHE_PERSISTENT=1 SYCL_CACHE_TRACE=5 SYCL_CACHE_DIR=%t/cache_dir
-// RUN: rm -rf %t/cache_dir
-// RUN:  %{cache_vars} %t.out 2>&1 |  FileCheck %s --check-prefixes=CHECK-WRITTEN-TO-CACHE
-// RUN:  %{cache_vars} %t.out reading-from-cache 2>&1 |  FileCheck %s --check-prefixes=CHECK-READ-FROM-CACHE
-
-// -- Add leak check.
-// RUN: rm -rf %t/cache_dir
-// RUN: %{l0_leak_check} %{cache_vars} %t.out 2>&1 |  FileCheck %s --check-prefixes=CHECK-WRITTEN-TO-CACHE
-// RUN: %{l0_leak_check} %{cache_vars} %t.out reading-from-cache 2>&1 |  FileCheck %s --check-prefixes=CHECK-READ-FROM-CACHE
-
-// CHECK-WRITTEN-TO-CACHE: [Persistent Cache]: enabled
-// CHECK-WRITTEN-TO-CACHE-NOT: [kernel_compiler Persistent Cache]: using cached binary
-// CHECK-WRITTEN-TO-CACHE: [kernel_compiler Persistent Cache]: binary has been cached
-
-// CHECK-READ-FROM-CACHE: [Persistent Cache]: enabled
-// CHECK-READ-FROM-CACHE-NOT: [kernel_compiler Persistent Cache]: binary has been cached
-// CHECK-READ-FROM-CACHE: [kernel_compiler Persistent Cache]: using cached binary
-
 #include <sycl/detail/core.hpp>
 #include <sycl/kernel_bundle.hpp>
 #include <sycl/usm.hpp>
@@ -149,7 +125,7 @@ void test_1(sycl::queue &Queue, sycl::kernel &Kernel, int seed) {
   sycl::free(usmPtr, Queue);
 }
 
-void test_build_and_run(bool readingFromCache) {
+void test_build_and_run() {
   namespace syclex = sycl::ext::oneapi::experimental;
   using source_kb = sycl::kernel_bundle<sycl::bundle_state::ext_oneapi_source>;
   using exe_kb = sycl::kernel_bundle<sycl::bundle_state::executable>;
@@ -192,12 +168,8 @@ void test_build_and_run(bool readingFromCache) {
       syclex::properties{syclex::build_options{flags}, syclex::save_log{&log},
                          syclex::registered_kernel_names{"ff_templated<int>"}});
 
-  // If the kernel was restored from cache, there will not have been
-  // any warning issued by the compilation of the kernel.
-  if (!readingFromCache) {
-    assert(log.find("warning: 'this_nd_item<1>' is deprecated") !=
-           std::string::npos);
-  }
+  assert(log.find("warning: 'this_nd_item<1>' is deprecated") !=
+         std::string::npos);
 
   // clang-format off
 
@@ -311,23 +283,8 @@ void test_esimd() {
 }
 
 int main(int argc, char *argv[]) {
-  namespace syclex = sycl::ext::oneapi::experimental;
-  bool readingFromCache = false;
-
-  // Check if the argument is present
-  if (argc > 1) {
-    std::string argument(argv[1]);
-    if (argument == "reading-from-cache") {
-      readingFromCache = true;
-    } else if (argument == "available") {
-      sycl::device d;
-      bool avail = d.ext_oneapi_can_compile(syclex::source_language::sycl);
-      return avail;
-    }
-  }
-
 #ifdef SYCL_EXT_ONEAPI_KERNEL_COMPILER
-  test_build_and_run(readingFromCache);
+  test_build_and_run();
   test_error();
   test_esimd();
 #else
diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp
@@ -7,7 +7,10 @@
 //===----------------------------------------------------------------------===//
 
 // REQUIRES: (opencl || level_zero)
+// REQUIRES: aspect-usm_device_allocations
+
 // UNSUPPORTED: accelerator
+// UNSUPPORTED-INTENDED: while accelerator is AoT only, this cannot run there.
 
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out 1
diff --git a/sycl/test-e2e/KernelCompiler/multi_device.cpp b/sycl/test-e2e/KernelCompiler/multi_device.cpp
@@ -1,5 +1,6 @@
 // REQUIRES: (opencl || level_zero) && ocloc
 // UNSUPPORTED: accelerator
+// UNSUPPORTED-INTENDED: while accelerator is AoT only, this cannot run there.
 
 // RUN: %{build} -o %t.out
 // RUN: env NEOReadDebugKeys=1 CreateMultipleRootDevices=3 %{run} %t.out
diff --git a/sycl/test-e2e/KernelCompiler/sycl_and_cache.cpp b/sycl/test-e2e/KernelCompiler/sycl_and_cache.cpp
@@ -0,0 +1,123 @@
+//==- sycl_and_cache.cpp - cache works with kernel_compiler sycl  ----------==//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+// REQUIRES: (opencl || level_zero)
+// REQUIRES: aspect-usm_device_allocations
+
+// UNSUPPORTED: accelerator
+// UNSUPPORTED-INTENDED: while accelerator is AoT only, this cannot run there.
+
+// -- Test the kernel_compiler with SYCL source.
+// RUN: %{build} -o %t.out
+
+// -- Run with caching.
+
+// DEFINE: %{cache_vars} = %{l0_leak_check} env SYCL_CACHE_PERSISTENT=1 SYCL_CACHE_TRACE=5 SYCL_CACHE_DIR=%t/cache_dir
+// RUN: rm -rf %t/cache_dir
+// RUN:  %{cache_vars} %t.out 2>&1 |  FileCheck %s --check-prefixes=CHECK-WRITTEN-TO-CACHE
+// RUN:  %{cache_vars} %t.out 2>&1 |  FileCheck %s --check-prefixes=CHECK-READ-FROM-CACHE
+
+// CHECK-WRITTEN-TO-CACHE: [Persistent Cache]: enabled
+// CHECK-WRITTEN-TO-CACHE-NOT: [kernel_compiler Persistent Cache]: using cached binary
+// CHECK-WRITTEN-TO-CACHE: [kernel_compiler Persistent Cache]: binary has been cached
+
+// CHECK-READ-FROM-CACHE: [Persistent Cache]: enabled
+// CHECK-READ-FROM-CACHE-NOT: [kernel_compiler Persistent Cache]: binary has been cached
+// CHECK-READ-FROM-CACHE: [kernel_compiler Persistent Cache]: using cached binary
+
+#include <sycl/detail/core.hpp>
+#include <sycl/kernel_bundle.hpp>
+#include <sycl/usm.hpp>
+
+// TODO: remove SYCL_EXTERNAL once it is no longer needed.
+auto constexpr SYCLSource = R"===(
+#include <sycl/sycl.hpp>
+
+int AddEm(int a, int b){
+    return a + b + 5;
+}
+
+// use extern "C" to avoid name mangling
+extern "C" SYCL_EXTERNAL SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((sycl::ext::oneapi::experimental::nd_range_kernel<1>))
+void ff_cp(int *ptr) {
+
+  // intentionally using deprecated routine, as opposed to this_work_item::get_nd_item<1>()
+  sycl::nd_item<1> Item = sycl::ext::oneapi::experimental::this_nd_item<1>();
+
+  sycl::id<1> GId = Item.get_global_id();
+  ptr[GId.get(0)] = AddEm(GId.get(0), 37);
+}
+)===";
+
+void test_1(sycl::queue &Queue, sycl::kernel &Kernel, int seed) {
+  constexpr int Range = 10;
+  int *usmPtr = sycl::malloc_shared<int>(Range, Queue);
+  int start = 3;
+
+  sycl::nd_range<1> R1{{Range}, {1}};
+
+  bool Passa = true;
+
+  memset(usmPtr, 0, Range * sizeof(int));
+  Queue.submit([&](sycl::handler &Handler) {
+    Handler.set_arg(0, usmPtr);
+    Handler.parallel_for(R1, Kernel);
+  });
+  Queue.wait();
+
+  for (int i = 0; i < Range; i++) {
+    std::cout << usmPtr[i] << "=" << (i + seed) << " ";
+    assert(usmPtr[i] == i + seed);
+  }
+  std::cout << std::endl;
+
+  sycl::free(usmPtr, Queue);
+}
+
+void test_build_and_run() {
+  namespace syclex = sycl::ext::oneapi::experimental;
+  using source_kb = sycl::kernel_bundle<sycl::bundle_state::ext_oneapi_source>;
+  using exe_kb = sycl::kernel_bundle<sycl::bundle_state::executable>;
+
+  sycl::queue q;
+  sycl::context ctx = q.get_context();
+
+  bool ok =
+      q.get_device().ext_oneapi_can_compile(syclex::source_language::sycl);
+  if (!ok) {
+    std::cout << "Apparently this device does not support SYCL source "
+                 "kernel bundle extension: "
+              << q.get_device().get_info<sycl::info::device::name>()
+              << std::endl;
+    return;
+  }
+
+  // Create from source.
+  source_kb kbSrc = syclex::create_kernel_bundle_from_source(
+      ctx, syclex::source_language::sycl, SYCLSource);
+
+  // Compilation of empty prop list, no devices.
+  exe_kb kbExe = syclex::build(kbSrc);
+
+  // extern "C" was used, so the name "ff_cp" is not mangled and can be used
+  // directly.
+  sycl::kernel k = kbExe.ext_oneapi_get_kernel("ff_cp");
+
+  // Test the kernels.
+  test_1(q, k, 37 + 5); // ff_cp seeds 37. AddEm will add 5 more.
+}
+
+int main(int argc, char *argv[]) {
+
+#ifdef SYCL_EXT_ONEAPI_KERNEL_COMPILER
+  test_build_and_run();
+#else
+  static_assert(false, "Kernel Compiler feature test macro undefined");
+#endif
+  return 0;
+}
diff --git a/sycl/test-e2e/KernelCompiler/sycl_device_flags.cpp b/sycl/test-e2e/KernelCompiler/sycl_device_flags.cpp
@@ -7,9 +7,10 @@
 //===----------------------------------------------------------------------===//
 
 // REQUIRES: level_zero
-// UNSUPPORTED: windows
+// REQUIRES: aspect-usm_device_allocations
 
-// IGC shader dump not available on Windows.
+// UNSUPPORTED: windows
+// UNSUPPORTED-INTENDED:  IGC shader dump not available on Windows.
 
 // RUN: %{build} -o %t.out
 // RUN: env IGC_DumpToCustomDir=%T.dump IGC_ShaderDumpEnable=1 NEO_CACHE_PERSISTENT=0 %{run} %t.out %T.dump/
diff --git a/sycl/test/e2e_test_requirements/no-unsupported-without-info.cpp b/sycl/test/e2e_test_requirements/no-unsupported-without-info.cpp
@@ -54,7 +54,7 @@
 // tests to match the required format and in that case you should just update
 // (i.e. reduce) the number and the list below.
 //
-// NUMBER-OF-UNSUPPORTED-WITHOUT-INFO: 414
+// NUMBER-OF-UNSUPPORTED-WITHOUT-INFO: 409
 //
 // List of improperly UNSUPPORTED tests.
 // Remove the CHECK once the test has been properly UNSUPPORTED.
@@ -283,11 +283,6 @@
 // CHECK-NEXT: KernelAndProgram/spec_constants_after_link.cpp
 // CHECK-NEXT: KernelAndProgram/spec_constants_after_link.cpp
 // CHECK-NEXT: KernelAndProgram/undefined-symbol.cpp
-// CHECK-NEXT: KernelCompiler/kernel_compiler_opencl.cpp
-// CHECK-NEXT: KernelCompiler/kernel_compiler_sycl.cpp
-// CHECK-NEXT: KernelCompiler/kernel_compiler_sycl_jit.cpp
-// CHECK-NEXT: KernelCompiler/multi_device.cpp
-// CHECK-NEXT: KernelCompiler/sycl_device_flags.cpp
 // CHECK-NEXT: LLVMIntrinsicLowering/bitreverse.cpp
 // CHECK-NEXT: LLVMIntrinsicLowering/sub_byte_bitreverse.cpp
 // CHECK-NEXT: Matrix/SG32/element_wise_abc.cpp
diff --git a/sycl/test/e2e_test_requirements/no_sycl_hpp_in_e2e_tests.cpp b/sycl/test/e2e_test_requirements/no_sycl_hpp_in_e2e_tests.cpp
@@ -6,7 +6,7 @@
 // CHECK-DAG: README.md
 // CHECK-DAG: lit.cfg.py
 //
-// CHECK-NUM-MATCHES: 5
+// CHECK-NUM-MATCHES: 6
 //
 // This test verifies that `<sycl/sycl.hpp>` isn't used in E2E tests. Instead,
 // fine-grained includes should used, see
