Merge from 'sycl' to 'sycl-web' (15 commits)

Author: iclsrc

.github/workflows/sycl-ur-perf-benchmarking.yml

@@ -24,6 +24,7 @@ on:
           - Full
           - SYCL
           - Minimal
+          - Core
           - Normal
           - Test
           - Gromacs

clang/lib/CodeGen/BackendUtil.cpp

@@ -1141,6 +1141,9 @@ void EmitAssemblyHelper::RunOptimizationPipeline(
         MPM.addPass(SYCLPropagateJointMatrixUsagePass());
         // Lowers static/dynamic local memory builtin calls.
         MPM.addPass(SYCLLowerWGLocalMemoryPass());
+        // Compile-time properties pass must create standard metadata as early
+        // as possible to make them available for other passes.
+        MPM.addPass(CompileTimePropertiesPass());
       });
     else if (LangOpts.SYCLIsHost && !LangOpts.SYCLESIMDBuildHostCode)
       PB.registerPipelineStartEPCallback(
@@ -1303,9 +1306,6 @@ void EmitAssemblyHelper::RunOptimizationPipeline(
         MPM.addPass(SPIRITTAnnotationsPass());
       }
 
-      // Process properties and annotations
-      MPM.addPass(CompileTimePropertiesPass());
-
       // Record SYCL aspect names (this should come after propagating aspects
       // and before cleaning up metadata)
       MPM.addPass(RecordSYCLAspectNamesPass());

clang/test/CodeGenSYCL/kernel-early-optimization-pipeline.cpp

@@ -9,6 +9,7 @@
 // CHECK: SYCLPropagateAspectsUsagePass
 // CHECK: SYCLPropagateJointMatrixUsagePass
 // CHECK: SYCLLowerWGLocalMemoryPass
+// CHECK: CompileTimePropertiesPass
 // CHECK: InferFunctionAttrsPass
 // CHECK: AlwaysInlinerPass
 // CHECK: ModuleInlinerWrapperPass
@@ -17,7 +18,6 @@
 // CHECK: SYCLMutatePrintfAddrspacePass
 // CHECK: SYCLPropagateAspectsUsagePass
 // CHECK: SYCLAddOptLevelAttributePass
-// CHECK: CompileTimePropertiesPass
 // CHECK: RecordSYCLAspectNamesPass
 // CHECK: CleanupSYCLMetadataPass
 //

…ng/test/Driver/sycl-offload-new-driver.c …/test/Driver/sycl-offload-new-driver.cppclang/test/Driver/sycl-offload-new-driver.c renamed to clang/test/Driver/sycl-offload-new-driver.cpp clang/test/Driver/sycl-offload-new-driver.c renamed to clang/test/Driver/sycl-offload-new-driver.cpp

@@ -2,7 +2,7 @@
 /// Verify --offload-new-driver option phases
 // RUN:  %clang --target=x86_64-unknown-linux-gnu -fsycl -fsycl-targets=nvptx64-nvidia-cuda,spir64 --offload-new-driver -ccc-print-phases %s 2>&1 \
 // RUN:   | FileCheck -check-prefix=OFFLOAD-NEW-DRIVER %s
-// OFFLOAD-NEW-DRIVER: 0: input, "[[INPUT:.+\.c]]", c++, (host-sycl)
+// OFFLOAD-NEW-DRIVER: 0: input, "[[INPUT:.+\.cpp]]", c++, (host-sycl)
 // OFFLOAD-NEW_DRIVER: 1: preprocessor, {0}, c++-cpp-output, (host-sycl)
 // OFFLOAD-NEW_DRIVER: 2: compiler, {1}, ir, (host-sycl)
 // OFFLOAD-NEW_DRIVER: 3: input, "[[INPUT]]", c++, (device-sycl)
@@ -123,7 +123,7 @@
 // RUN:          -fsycl-targets=intel_gpu_dg1,intel_gpu_pvc \
 // RUN:          --offload-new-driver -ccc-print-phases %s 2>&1 \
 // RUN:  | FileCheck -check-prefix=MULT_TARG_PHASES %s
-// MULT_TARG_PHASES: 0: input, "[[INPUT:.+\.c]]", c++, (host-sycl)
+// MULT_TARG_PHASES: 0: input, "[[INPUT:.+\.cpp]]", c++, (host-sycl)
 // MULT_TARG_PHASES: 1: preprocessor, {0}, c++-cpp-output, (host-sycl)
 // MULT_TARG_PHASES: 2: compiler, {1}, ir, (host-sycl)
 // MULT_TARG_PHASES: 3: input, "[[INPUT]]", c++, (device-sycl, dg1)

devops/scripts/benchmarks/README.md

@@ -115,6 +115,7 @@ The available benchmarks options are:
 * `Full` (BenchDNN, Compute, Gromacs, llama, SYCL, Velocity and UMF benchmarks)
 * `SYCL` (Compute, llama, SYCL, Velocity)
 * `Minimal` (Compute)
+* `Core` (Compute: SubmitKernel)
 * `Normal` (BenchDNN, Compute, Gromacs, llama, Velocity)
 * `Gromacs` (Gromacs)
 * `OneDNN` (BenchDNN)

devops/scripts/benchmarks/benches/compute.py

@@ -353,6 +353,41 @@ def createRrBench(variant_name: str, **kwargs):
         return benches
 
 
+class ComputeBenchCoreSuite(ComputeBench):
+    """
+    A suite for core compute benchmarks scenarios for quick runs.
+    """
+
+    def name(self) -> str:
+        return "Compute Benchmarks Core"
+
+    def benchmarks(self) -> list[Benchmark]:
+        core_benches = []
+        submit_kernel_params = product(
+            list(RUNTIMES),
+            [0, 1],  # in_order_queue
+            [0, 1],  # measure_completion
+            [0, 1],  # use_events
+        )
+        for (
+            runtime,
+            in_order_queue,
+            measure_completion,
+            use_events,
+        ) in submit_kernel_params:
+            core_benches.append(
+                SubmitKernel(
+                    self,
+                    runtime,
+                    in_order_queue,
+                    measure_completion,
+                    use_events,
+                    KernelExecTime=1,
+                )
+            )
+        return core_benches
+
+
 class ComputeBenchmark(Benchmark):
     def __init__(
         self,

devops/scripts/benchmarks/compare.py

@@ -411,7 +411,7 @@ def print_regression(entry: dict, is_warning: bool = False):
             log_func(f"-- Delta: {entry['delta']}")
             log_func("")
             if args.produce_github_summary:
-                gh_summary.append(f"#### {entry['name']}:")
+                gh_summary.append(f"##### {entry['name']}:")
                 gh_summary.append(
                     f"- Historic {entry['avg_type']}: {entry['hist_avg']}"
                 )
@@ -427,12 +427,16 @@ def print_regression(entry: dict, is_warning: bool = False):
                 )
                 gh_summary.append("")
 
+        if args.produce_github_summary:
+            gh_summary.append("")
+            gh_summary.append("### Regressions and Improvements")
+
         if improvements:
             log.info("#")
             log.info("# Improvements:")
             log.info("#")
             if args.produce_github_summary:
-                gh_summary.append(f"### Improvements")
+                gh_summary.append(f"#### Improvements")
                 gh_summary.append(
                     f"<details><summary>{len(improvements)} improved tests:</summary>"
                 )
@@ -444,12 +448,16 @@ def print_regression(entry: dict, is_warning: bool = False):
                 gh_summary.append("")
         if regressions_ignored:
             log.info("#")
-            log.info("# Regressions (filtered out by --regression-filter):")
+            log.info(
+                f"# Regressions Ignored (filtered out by --regression-filter: {filter_type_capitalized})"
+            )
             log.info("#")
             if args.produce_github_summary:
-                gh_summary.append(f"### Non-{filter_type_capitalized} Regressions")
                 gh_summary.append(
-                    f"<details><summary>{len(regressions_ignored)} non-{args.regression_filter_type} regressions:</summary>"
+                    f"#### Regressions Ignored (filtered out by --regression-filter: {filter_type_capitalized})"
+                )
+                gh_summary.append(
+                    f"<details><summary>{len(regressions_ignored)} non-'{args.regression_filter_type}' regressions:</summary>"
                 )
                 gh_summary.append("")
             for test in regressions_ignored:
@@ -462,7 +470,7 @@ def print_regression(entry: dict, is_warning: bool = False):
             log.warning("# Regressions:")
             log.warning("#")
             if args.produce_github_summary:
-                gh_summary.append(f"### {filter_type_capitalized} Regressions")
+                gh_summary.append(f"#### {filter_type_capitalized} Regressions")
                 gh_summary.append(
                     f"{len(regressions_of_concern)} {args.regression_filter_type} regressions. These regressions warrant a CI failure:"
                 )
@@ -480,8 +488,6 @@ def print_regression(entry: dict, is_warning: bool = False):
 
         log.info("No unexpected regressions found!")
         if args.produce_github_summary:
-            gh_summary.append("")
-            gh_summary.append("### Regressions")
             gh_summary.append("No unexpected regressions found!")
             with open(options.github_summary_regression_filename, "w") as f:
                 f.write("\n".join(gh_summary))

devops/scripts/benchmarks/main.py

@@ -269,6 +269,7 @@ def main(directory, additional_env_vars, compare_names, filter, execution_stats)
 
     suites = [
         ComputeBench(),
+        ComputeBenchCoreSuite(),
         VelocityBench(),
         SyclBench(),
         LlamaCppBench(),

devops/scripts/benchmarks/presets.py

@@ -26,6 +26,9 @@
     "Minimal": [
         "Compute Benchmarks",
     ],
+    "Core": [
+        "Compute Benchmarks Core",
+    ],
     "Normal": [
         "BenchDNN",
         "Compute Benchmarks",

libdevice/nativecpu_utils.cpp

@@ -95,10 +95,14 @@ DefGenericCastToPtrExpl(ToGlobal, OCL_GLOBAL);
   DefSubgroupBlockINTEL_vt(Type, v8)
 
 namespace ncpu_types {
+template <typename DataT, int NumElements>
+using native_vector_t =
+    sycl::detail::ConvertToOpenCLType_t<sycl::vec<DataT, NumElements>>;
+
 template <class T> struct vtypes {
-  using v2 = typename sycl::vec<T, 2>::vector_t;
-  using v4 = typename sycl::vec<T, 4>::vector_t;
-  using v8 = typename sycl::vec<T, 8>::vector_t;
+  using v2 = native_vector_t<T, 2>;
+  using v4 = native_vector_t<T, 4>;
+  using v8 = native_vector_t<T, 8>;
 };
 } // namespace ncpu_types
 
@@ -224,15 +228,15 @@ DefineLogicalGroupOp(bool, bool, i1);
   }                                                                            \
                                                                                \
   DEVICE_EXTERNAL Type __spirv_GroupBroadcast(                                 \
-      int32_t g, Type v, sycl::vec<IDType, 2>::vector_t l) noexcept {          \
+      int32_t g, Type v, ncpu_types::native_vector_t<IDType, 2> l) noexcept {  \
     if (__spv::Scope::Flag::Subgroup == g)                                     \
       return __mux_sub_group_broadcast_##Sfx(v, l[0]);                         \
     else                                                                       \
       return __mux_work_group_broadcast_##Sfx(0, v, l[0], l[1], 0);            \
   }                                                                            \
                                                                                \
   DEVICE_EXTERNAL Type __spirv_GroupBroadcast(                                 \
-      int32_t g, Type v, sycl::vec<IDType, 3>::vector_t l) noexcept {          \
+      int32_t g, Type v, ncpu_types::native_vector_t<IDType, 3> l) noexcept {  \
     if (__spv::Scope::Flag::Subgroup == g)                                     \
       return __mux_sub_group_broadcast_##Sfx(v, l[0]);                         \
     else                                                                       \
@@ -310,8 +314,8 @@ DefShuffleINTEL_All(float, f32, float);
 DefShuffleINTEL_All(_Float16, f16, _Float16);
 
 #define DefineShuffleVec(T, N, Sfx, MuxType)                                   \
-  using vt##T##N = sycl::vec<T, N>::vector_t;                                  \
-  using vt##MuxType##N = sycl::vec<MuxType, N>::vector_t;                      \
+  using vt##T##N = ncpu_types::native_vector_t<T, N>;                          \
+  using vt##MuxType##N = ncpu_types::native_vector_t<MuxType, N>;              \
   DefShuffleINTEL_All(vt##T##N, v##N##Sfx, vt##MuxType##N)
 
 #define DefineShuffleVec2to16(Type, Sfx, MuxType)                              \