Merge branch 'intel:sycl' into int-buffers

Author: mateuszpn

clang/test/Driver/clang-linker-wrapper.cpp

@@ -63,13 +63,20 @@
 // CHK-CMDS-AOT-GEN: spirv-to-ir-wrapper{{.*}} -o [[FIRSTLLVMLINKIN:.*]].bc --llvm-spirv-opts --spirv-preserve-auxdata --spirv-target-env=SPV-IR --spirv-builtin-format=global
 // CHK-CMDS-AOT-GEN-NEXT: llvm-link{{.*}} --suppress-warnings [[FIRSTLLVMLINKIN]].bc -o [[FIRSTLLVMLINKOUT:.*]].bc
 // CHK-CMDS-AOT-GEN-NEXT: llvm-link{{.*}} -only-needed --suppress-warnings [[FIRSTLLVMLINKOUT]].bc {{.*}}.bc -o [[SECONDLLVMLINKOUT:.*]].bc
-// CHK-CMDS-AOT-GEN-NEXT: sycl-post-link{{.*}} SYCL_POST_LINK_OPTIONS -o [[SYCLPOSTLINKOUT:.*]].table [[SECONDLLVMLINKOUT]].bc
+// Check that target specified by -fsycl-targets is passed to sycl-post-link for filtering.
+// CHK-CMDS-AOT-GEN-NEXT: sycl-post-link{{.*}} SYCL_POST_LINK_OPTIONS -o intel_gpu_pvc,[[SYCLPOSTLINKOUT:.*]].table [[SECONDLLVMLINKOUT]].bc
 // CHK-CMDS-AOT-GEN-NEXT: llvm-spirv{{.*}} LLVM_SPIRV_OPTIONS -o {{.*}}
 // CHK-CMDS-AOT-GEN-NEXT: ocloc{{.*}} -output_no_suffix -spirv_input -device pvc{{.*}} -output {{.*}} -file {{.*}}
 // CHK-CMDS-AOT-GEN-NEXT: offload-wrapper: output: [[WRAPPEROUT:.*]].bc, input: {{.*}}, compile-opts: , link-opts:
 // CHK-CMDS-AOT-GEN-NEXT: clang{{.*}} -c -o [[LLCOUT:.*]].o [[WRAPPEROUT]].bc
 // CHK-CMDS-AOT-GEN-NEXT: "{{.*}}/ld" -- HOST_LINKER_FLAGS -dynamic-linker HOST_DYN_LIB -o a.out [[LLCOUT]].o HOST_LIB_PATH HOST_STAT_LIB {{.*}}.o
 
+// Check that when --gpu-tool-arg is specified in clang-linker-wrapper 
+// (happen when AOT device is specified via -Xsycl-target-backend '-device pvc' in clang), 
+// the target is not passed to sycl-post-link for filtering.
+// RUN: clang-linker-wrapper -sycl-embed-ir -sycl-device-libraries=%t1.devicelib.o -sycl-post-link-options="SYCL_POST_LINK_OPTIONS" -llvm-spirv-options="LLVM_SPIRV_OPTIONS" "--host-triple=x86_64-unknown-linux-gnu" "--gpu-tool-arg=-device pvc" "--linker-path=/usr/bin/ld" "--" HOST_LINKER_FLAGS "-dynamic-linker" HOST_DYN_LIB "-o" "a.out" HOST_LIB_PATH HOST_STAT_LIB %t1.o --dry-run 2>&1 | FileCheck -check-prefix=CHK-NO-CMDS-AOT-GEN %s
+// CHK-NO-CMDS-AOT-GEN-NOT: sycl-post-link{{.*}} -o intel_gpu_pv,{{.*}}
+
 /// Check for list of commands for standalone clang-linker-wrapper run for sycl (AOT for Intel CPU)
 // -------
 // Generate .o file as linker wrapper input.

clang/test/SemaSYCL/attr-add-ir-attributes-function-conflict.cpp

@@ -112,6 +112,11 @@ int main() {
   NonemptyWrapper.kernel_single_task<class NEK9>([]() [[sycl::device_has(sycl::aspect::cpu)]] {});
   // expected-warning@+1 {{kernel has both attribute 'sycl::device_has' and kernel properties; conflicting properties are ignored}}
   NonemptyWrapper.kernel_single_task<class NEK10>([]() [[sycl::device_has(sycl::aspect::cpu, sycl::aspect::gpu)]] {});
+  // expected-warning@+4 {{kernel has both attribute 'sycl::reqd_work_group_size' and kernel properties; conflicting properties are ignored}}
+  // expected-warning@+3 {{kernel has both attribute 'sycl::work_group_size_hint' and kernel properties; conflicting properties are ignored}}
+  // expected-warning@+2 {{kernel has both attribute 'sycl::reqd_sub_group_size' and kernel properties; conflicting properties are ignored}}
+  // expected-warning@+1 {{kernel has both attribute 'sycl::device_has' and kernel properties; conflicting properties are ignored}}
+  NonemptyWrapper.kernel_single_task<class NEK11>([]() [[sycl::reqd_work_group_size(1)]] [[sycl::work_group_size_hint(1)]] [[sycl::reqd_sub_group_size(1)]] [[sycl::device_has(sycl::aspect::cpu, sycl::aspect::gpu)]] {});
 
   // expected-warning@+1 {{kernel has both attribute 'sycl::reqd_work_group_size' and kernel properties; conflicting properties are ignored}}
   NonemptyWrapperWithImplicit.kernel_single_task<class NEKWI1>([]() [[sycl::reqd_work_group_size(1)]] {});
@@ -133,6 +138,11 @@ int main() {
   NonemptyWrapperWithImplicit.kernel_single_task<class NEKWI9>([]() [[sycl::device_has(sycl::aspect::cpu)]] {});
   // expected-warning@+1 {{kernel has both attribute 'sycl::device_has' and kernel properties; conflicting properties are ignored}}
   NonemptyWrapperWithImplicit.kernel_single_task<class NEKWI10>([]() [[sycl::device_has(sycl::aspect::cpu, sycl::aspect::gpu)]] {});
+  // expected-warning@+4 {{kernel has both attribute 'sycl::reqd_work_group_size' and kernel properties; conflicting properties are ignored}}
+  // expected-warning@+3 {{kernel has both attribute 'sycl::work_group_size_hint' and kernel properties; conflicting properties are ignored}}
+  // expected-warning@+2 {{kernel has both attribute 'sycl::reqd_sub_group_size' and kernel properties; conflicting properties are ignored}}
+  // expected-warning@+1 {{kernel has both attribute 'sycl::device_has' and kernel properties; conflicting properties are ignored}}
+  NonemptyWrapperWithImplicit.kernel_single_task<class NEKWI11>([]() [[sycl::reqd_work_group_size(1)]] [[sycl::work_group_size_hint(1)]] [[sycl::reqd_sub_group_size(1)]] [[sycl::device_has(sycl::aspect::cpu, sycl::aspect::gpu)]] {});
 
   // expected-warning@+1 {{kernel has both attribute 'sycl::reqd_work_group_size' and kernel properties; conflicting properties are ignored}}
   NonemptyWrapperWithFilter.kernel_single_task<class NEKWF1>([]() [[sycl::reqd_work_group_size(1)]] {});
@@ -154,6 +164,11 @@ int main() {
   NonemptyWrapperWithFilter.kernel_single_task<class NEKWF9>([]() [[sycl::device_has(sycl::aspect::cpu)]] {});
   // expected-warning@+1 {{kernel has both attribute 'sycl::device_has' and kernel properties; conflicting properties are ignored}}
   NonemptyWrapperWithFilter.kernel_single_task<class NEKWF10>([]() [[sycl::device_has(sycl::aspect::cpu, sycl::aspect::gpu)]] {});
+  // expected-warning@+4 {{kernel has both attribute 'sycl::reqd_work_group_size' and kernel properties; conflicting properties are ignored}}
+  // expected-warning@+3 {{kernel has both attribute 'sycl::work_group_size_hint' and kernel properties; conflicting properties are ignored}}
+  // expected-warning@+2 {{kernel has both attribute 'sycl::reqd_sub_group_size' and kernel properties; conflicting properties are ignored}}
+  // expected-warning@+1 {{kernel has both attribute 'sycl::device_has' and kernel properties; conflicting properties are ignored}}
+  NonemptyWrapperWithFilter.kernel_single_task<class NEKWF11>([]() [[sycl::reqd_work_group_size(1)]] [[sycl::work_group_size_hint(1)]] [[sycl::reqd_sub_group_size(1)]] [[sycl::device_has(sycl::aspect::cpu, sycl::aspect::gpu)]] {});
 
   // expected-warning@+1 {{kernel has both attribute 'sycl::reqd_work_group_size' and kernel properties; conflicting properties are ignored}}
   NonemptyWrapperWithImplicitAndFilter.kernel_single_task<class NEKWIF1>([]() [[sycl::reqd_work_group_size(1)]] {});
@@ -175,4 +190,9 @@ int main() {
   NonemptyWrapperWithImplicitAndFilter.kernel_single_task<class NEKWIF9>([]() [[sycl::device_has(sycl::aspect::cpu)]] {});
   // expected-warning@+1 {{kernel has both attribute 'sycl::device_has' and kernel properties; conflicting properties are ignored}}
   NonemptyWrapperWithImplicitAndFilter.kernel_single_task<class NEKWIF10>([]() [[sycl::device_has(sycl::aspect::cpu, sycl::aspect::gpu)]] {});
+  // expected-warning@+4 {{kernel has both attribute 'sycl::reqd_work_group_size' and kernel properties; conflicting properties are ignored}}
+  // expected-warning@+3 {{kernel has both attribute 'sycl::work_group_size_hint' and kernel properties; conflicting properties are ignored}}
+  // expected-warning@+2 {{kernel has both attribute 'sycl::reqd_sub_group_size' and kernel properties; conflicting properties are ignored}}
+  // expected-warning@+1 {{kernel has both attribute 'sycl::device_has' and kernel properties; conflicting properties are ignored}}
+  NonemptyWrapperWithImplicitAndFilter.kernel_single_task<class NEKWIF11>([]() [[sycl::reqd_work_group_size(1)]] [[sycl::work_group_size_hint(1)]] [[sycl::reqd_sub_group_size(1)]] [[sycl::device_has(sycl::aspect::cpu, sycl::aspect::gpu)]] {});
 }

clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp

@@ -727,10 +727,24 @@ runSYCLPostLinkTool(ArrayRef<StringRef> InputFiles, const ArgList &Args) {
       createOutputFile(sys::path::filename(ExecutableName), "table");
   if (!TempFileOrErr)
     return TempFileOrErr.takeError();
+  std::string OutputPathWithArch = TempFileOrErr->str();
+
+  // Enable the driver to invoke sycl-post-link with the device architecture
+  // when Intel GPU targets are passed in -fsycl-targets.
+  // OPT_gpu_tool_arg_EQ is checked to ensure the device architecture is not
+  // passed through -Xsycl-target-backend=spir64_gen "-device <arch>" format
+  const llvm::Triple Triple(Args.getLastArgValue(OPT_triple_EQ));
+  StringRef Arch = Args.getLastArgValue(OPT_arch_EQ);
+  StringRef IsGPUTool = Args.getLastArgValue(OPT_gpu_tool_arg_EQ);
+
+  if (Triple.getSubArch() == llvm::Triple::SPIRSubArch_gen && !Arch.empty() &&
+      IsGPUTool.empty() && Arch != "*")
+    OutputPathWithArch = "intel_gpu_" + Arch.str() + "," + OutputPathWithArch;
+  else if (Triple.getSubArch() == llvm::Triple::SPIRSubArch_x86_64)
+    OutputPathWithArch = "spir64_x86_64," + OutputPathWithArch;
 
   SmallVector<StringRef, 8> CmdArgs;
   CmdArgs.push_back(*SYCLPostLinkPath);
-  const llvm::Triple Triple(Args.getLastArgValue(OPT_triple_EQ));
   Arg *SYCLDeviceLibLoc = Args.getLastArg(OPT_sycl_device_library_location_EQ);
   if (SYCLDeviceLibLoc && !Triple.isSPIRAOT()) {
     std::string SYCLDeviceLibLocParam = SYCLDeviceLibLoc->getValue();
@@ -748,7 +762,7 @@ runSYCLPostLinkTool(ArrayRef<StringRef> InputFiles, const ArgList &Args) {
   SYCLPostLinkOptions.split(CmdArgs, " ", /* MaxSplit = */ -1,
                             /* KeepEmpty = */ false);
   CmdArgs.push_back("-o");
-  CmdArgs.push_back(*TempFileOrErr);
+  CmdArgs.push_back(Args.MakeArgString(OutputPathWithArch));
   for (auto &File : InputFiles)
     CmdArgs.push_back(File);
   if (Error Err = executeCommands(*SYCLPostLinkPath, CmdArgs))

devops/actions/run-tests/benchmark/action.yml

@@ -231,6 +231,10 @@ runs:
       WORKDIR="$(realpath ./llvm_test_workdir)"
       if [ -n "$WORKDIR" ] && [ -d "$WORKDIR" ] && [[ "$WORKDIR" == *llvm_test_workdir* ]]; then rm -rf "$WORKDIR" ; fi
 
+      # Clean up potentially existing, old summary files
+      [ -f "github_summary_exe.md" ] && rm github_summary_exe.md
+      [ -f "github_summary_reg.md" ] && rm github_summary_reg.md
+
       numactl --cpunodebind "$NUMA_NODE" --membind "$NUMA_NODE" \
       ./devops/scripts/benchmarks/main.py "$WORKDIR" \
         --sycl "$(realpath ./toolchain)" \
@@ -243,6 +247,7 @@ runs:
         --preset "$PRESET" \
         --timestamp-override "$SAVE_TIMESTAMP" \
         --detect-version sycl,compute_runtime \
+        --produce-github-summary \
         ${{ inputs.exit_on_failure == 'true' && '--exit-on-failure --iterations 1' || '' }}
       # TODO: add back: "--flamegraph inclusive" once works properly
 
@@ -273,8 +278,9 @@ runs:
     if: always()
     shell: bash
     run: |
-      # Cache changes and upload github summary
-      [ -f "github_summary.md" ] && cat github_summary.md >> $GITHUB_STEP_SUMMARY
+      # Cache changes and upload github summaries
+      [ -f "github_summary_exe.md" ] && cat github_summary_exe.md >> $GITHUB_STEP_SUMMARY
+      [ -f "github_summary_reg.md" ] && cat github_summary_reg.md >> $GITHUB_STEP_SUMMARY
 
       cd "./llvm-ci-perf-results"
       git add .

devops/scripts/benchmarks/benches/compute.py

@@ -61,8 +61,8 @@ def git_url(self) -> str:
         return "https://github.com/intel/compute-benchmarks.git"
 
     def git_hash(self) -> str:
-        # Nov 7, 2025
-        return "d985da634fc1a9416ca0bd067cfb9886b02d0211"
+        # Nov 17, 2025
+        return "932ae79f7cca7e156285fc10a59610927c769e89"
 
     def setup(self) -> None:
         if options.sycl is None:
@@ -860,6 +860,7 @@ def _bin_args(self, run_trace: TracingType = TracingType.NONE) -> list[str]:
             "--multiplier=1",
             "--vectorSize=1",
             "--lws=256",
+            "--prefetch=0",
         ]
 
 

devops/scripts/benchmarks/compare.py

@@ -357,7 +357,8 @@ def to_hist(
     parser_avg.add_argument(
         "--produce-github-summary",
         action="store_true",
-        help=f"Create a summary file '{options.github_summary_filename}' for Github workflow summaries.",
+        help=f"Produce regression summary for Github workflow, in file '{options.github_summary_regression_filename}'.",
+        default=False,
     )
 
     args = parser.parse_args()
@@ -473,14 +474,16 @@ def print_regression(entry: dict, is_warning: bool = False):
 
             if not args.dry_run:
                 if args.produce_github_summary:
-                    with open(options.github_summary_filename, "w") as f:
+                    with open(options.github_summary_regression_filename, "w") as f:
                         f.write("\n".join(gh_summary))
                 exit(1)  # Exit 1 to trigger Github test failure
 
         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_filename, "w") as f:
+            with open(options.github_summary_regression_filename, "w") as f:
                 f.write("\n".join(gh_summary))
 
     else: