From 98012b288741fcd8c99ddbf2ef63a9703195df90 Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Fri, 9 Oct 2026 13:47:38 -0700 Subject: [PATCH 01/22] [SYCL][E2E] Select Intel GPU AOT targets via gpu_aot_target E2E tests compiled for Intel GPUs ahead of time with '-fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen "-device *"'. That spelling is being deprecated in favor of '-fsycl-targets=intel_gpu_', which has no "all devices" equivalent of '-device *', so the tests need to be told which architecture to build for. Replace the gpu_aot_target_opts lit parameter (GPU_AOT_TARGET_OPTS in CMake), which only carried the '-device' value for spir64_gen, with gpu_aot_target (GPU_AOT_TARGET). It is a -fsycl-targets value exposed to tests as %{gpu_aot_target}. When unset, it defaults to the Intel GPU architectures detected via sycl-ls, or to intel_gpu_pvc when none are detected (e.g. build-only mode). This mirrors how amd_arch is handled for HIP. The spir_gen build target now compiles for %{gpu_aot_target}. Tests that pinned a specific device keep it, spelled as intel_gpu_. --- sycl/test-e2e/AOT/early_aot.cpp | 4 +-- sycl/test-e2e/AOT/gpu.cpp | 2 +- sycl/test-e2e/AOT/multiple-devices.cpp | 6 ++--- sycl/test-e2e/AddressSanitizer/lit.local.cfg | 3 +-- .../BFloat16/bfloat16_example_aot.cpp | 4 +-- .../BFloat16/bfloat16_example_aot_cpu.cpp | 4 +-- .../BFloat16/bfloat16_example_aot_gpu.cpp | 4 +-- sycl/test-e2e/Basic/build_log.cpp | 2 +- sycl/test-e2e/DeviceCodeSplit/aot-gpu.cpp | 4 +-- .../DeviceImageDependencies/dynamic_aot.cpp | 2 +- .../DeviceImageDependencies/objects_aot.cpp | 5 ++-- sycl/test-e2e/ESIMD/aot_mixed.cpp | 4 +-- .../KernelAndProgram/test_cache_jit_aot.cpp | 4 +-- sycl/test-e2e/MemorySanitizer/lit.local.cfg | 3 +-- sycl/test-e2e/NewOffloadDriver/aot-gpu.cpp | 4 +-- .../NewOffloadDriver/aot-multiple-devices.cpp | 6 ++--- .../is_compatible_several_targets.cpp | 2 +- .../is_compatible_spir64_gen.cpp | 2 +- .../ProgramManager/uneven_kernel_split.cpp | 6 ++--- sycl/test-e2e/README.md | 27 ++++++++++--------- sycl/test-e2e/SYCLBIN/aot_object_load.cpp | 8 +++--- .../SYCLBIN/aot_object_no_symbols.cpp | 7 +++-- .../SYCLBIN/basic_executable_bmg_aot.cpp | 2 +- .../SYCLBIN/dae_executable_bmg_aot.cpp | 2 +- .../SYCLBIN/dg_executable_bmg_aot.cpp | 2 +- sycl/test-e2e/SYCLBIN/link_object_aot.cpp | 6 ++--- sycl/test-e2e/SYCLBIN/lit.local.cfg | 6 ++--- ...nal_kernel_features_executable_bmg_aot.cpp | 2 +- sycl/test-e2e/SYCLBIN/simple_kernel_aot.cpp | 10 +++---- .../SpecConstants/2020/image_selection.cpp | 10 +++---- .../2020/native_specialization_constant.cpp | 2 +- .../SpecConstants/2020/non_native/gpu.cpp | 2 +- .../non_native/spec_const_alignment_test.cpp | 2 +- sycl/test-e2e/ThreadSanitizer/lit.local.cfg | 3 +-- sycl/test-e2e/format.py | 14 ++++++---- sycl/test-e2e/lit.cfg.py | 17 ++++++++---- sycl/test-e2e/lit.site.cfg.py.in | 2 +- 37 files changed, 97 insertions(+), 98 deletions(-) diff --git a/sycl/test-e2e/AOT/early_aot.cpp b/sycl/test-e2e/AOT/early_aot.cpp index 6727d91a53f69..4c86bbeccf0a0 100644 --- a/sycl/test-e2e/AOT/early_aot.cpp +++ b/sycl/test-e2e/AOT/early_aot.cpp @@ -6,8 +6,8 @@ // at the linking step. Old Offload Model requires it at the compilation step. // Build the early AOT device binaries -// RUN: %clangxx -fsycl -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen %gpu_aot_target_opts %if !new-offload-model %{ -fno-sycl-rdc %} -c -DADD_CPP %s -o %t_add.o -// RUN: %clangxx -fsycl -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen %gpu_aot_target_opts %if !new-offload-model %{ -fno-sycl-rdc %} -c -DSUB_CPP %s -o %t_sub.o +// RUN: %clangxx -fsycl -fsycl-targets=%{gpu_aot_target} %if !new-offload-model %{ -fno-sycl-rdc %} -c -DADD_CPP %s -o %t_add.o +// RUN: %clangxx -fsycl -fsycl-targets=%{gpu_aot_target} %if !new-offload-model %{ -fno-sycl-rdc %} -c -DSUB_CPP %s -o %t_sub.o // RUN: %clangxx -fsycl %if new-offload-model %{ -fno-sycl-rdc %} -DMAIN_CPP %s %t_add.o %t_sub.o -o %t.out // RUN: %{run} %t.out diff --git a/sycl/test-e2e/AOT/gpu.cpp b/sycl/test-e2e/AOT/gpu.cpp index 2572b59eadf7f..c8060c7c0ab99 100644 --- a/sycl/test-e2e/AOT/gpu.cpp +++ b/sycl/test-e2e/AOT/gpu.cpp @@ -8,5 +8,5 @@ // REQUIRES: ocloc, gpu, target-spir // -// RUN: %clangxx -fsycl -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen %gpu_aot_target_opts %S/Inputs/aot.cpp -o %t.out +// RUN: %clangxx -fsycl -fsycl-targets=%{gpu_aot_target} %S/Inputs/aot.cpp -o %t.out // RUN: %{run} %t.out diff --git a/sycl/test-e2e/AOT/multiple-devices.cpp b/sycl/test-e2e/AOT/multiple-devices.cpp index 61de08b4041cc..af44e27d8c2a5 100644 --- a/sycl/test-e2e/AOT/multiple-devices.cpp +++ b/sycl/test-e2e/AOT/multiple-devices.cpp @@ -28,17 +28,17 @@ // REQUIRES: opencl-aot, ocloc, any-device-is-cpu, any-device-is-gpu, target-spir, opencl-cpu-rt // Produce a fat object for all targets (generic SPIR-V, CPU, GPU) -// RUN: %clangxx -fsycl -fsycl-targets=spir64,spir64_x86_64,spir64_gen %S/Inputs/aot.cpp -c -o %t.o +// RUN: %clangxx -fsycl -fsycl-targets=spir64,spir64_x86_64,%{gpu_aot_target} %S/Inputs/aot.cpp -c -o %t.o // Verify that AOT compilation occurs for the targets (GPU, CPU) // specified in the link command line. Note that generic SPIR-V compilation // is enabled by default even when AOT GPU or CPU targets are specified. -// RUN: %{run-aux} %clangxx -fsycl -fsycl-targets=spir64_x86_64,spir64_gen -Xsycl-target-backend=spir64_gen %gpu_aot_target_opts -v %t.o -o %t_cpu_gpu.out -Wno-unused-command-line-argument 2>&1 | FileCheck %s --check-prefix=CHECK-AOT-CPU --check-prefix=CHECK-AOT-GPU --check-prefix=CHECK-GENERIC +// RUN: %{run-aux} %clangxx -fsycl -fsycl-targets=spir64_x86_64,%{gpu_aot_target} -v %t.o -o %t_cpu_gpu.out -Wno-unused-command-line-argument 2>&1 | FileCheck %s --check-prefix=CHECK-AOT-CPU --check-prefix=CHECK-AOT-GPU --check-prefix=CHECK-GENERIC // RUN: %{run} %t_cpu_gpu.out // Verify that AOT compilation occurs for the targets (GPU) specified in // the link command line. -// RUN: %clangxx -fsycl -fsycl-targets=spir64,spir64_gen -Xsycl-target-backend=spir64_gen %gpu_aot_target_opts -v %t.o -o %t_spv_gpu.out -Wno-unused-command-line-argument 2>&1 | FileCheck %s --check-prefix=CHECK-GENERIC --check-prefix=CHECK-AOT-GPU --check-prefix=CHECK-NO-AOT-CPU +// RUN: %clangxx -fsycl -fsycl-targets=spir64,%{gpu_aot_target} -v %t.o -o %t_spv_gpu.out -Wno-unused-command-line-argument 2>&1 | FileCheck %s --check-prefix=CHECK-GENERIC --check-prefix=CHECK-AOT-GPU --check-prefix=CHECK-NO-AOT-CPU // Check that execution on AOT-compatible devices is unaffected // RUN: env ONEAPI_DEVICE_SELECTOR="*:gpu" %{run-unfiltered-devices} %t_spv_gpu.out diff --git a/sycl/test-e2e/AddressSanitizer/lit.local.cfg b/sycl/test-e2e/AddressSanitizer/lit.local.cfg index a9a68083cd3c3..d1e227f89ce90 100644 --- a/sycl/test-e2e/AddressSanitizer/lit.local.cfg +++ b/sycl/test-e2e/AddressSanitizer/lit.local.cfg @@ -16,8 +16,7 @@ config.substitutions.append( config.substitutions.append( ( "%device_asan_aot_gpu_flags", - "-Xarch_device -fsanitize=address -fsycl-targets=spir64_gen " - "-Xsycl-target-backend=spir64_gen '-device pvc,dg2'", + "-Xarch_device -fsanitize=address -fsycl-targets=intel_gpu_pvc,intel_gpu_dg2", ) ) config.substitutions.append( diff --git a/sycl/test-e2e/BFloat16/bfloat16_example_aot.cpp b/sycl/test-e2e/BFloat16/bfloat16_example_aot.cpp index 6d1e1e217bbbd..2d3de11d6f75e 100644 --- a/sycl/test-e2e/BFloat16/bfloat16_example_aot.cpp +++ b/sycl/test-e2e/BFloat16/bfloat16_example_aot.cpp @@ -8,11 +8,11 @@ // RUN: %clangxx -fsycl -fsycl-targets=spir64 %s -o %t.out // RUN: %{run} %t.out -// RUN: %clangxx -fsycl -fsycl-targets=spir64,spir64_gen -Xsycl-target-backend=spir64_gen "-device dg2" %s -o %t.out +// RUN: %clangxx -fsycl -fsycl-targets=spir64,intel_gpu_dg2 %s -o %t.out // RUN: %{run} %t.out // CPU AOT targets host isa, so we compile on the run system instead. -// RUN: %{run-aux} %clangxx -fsycl -fsycl-targets=spir64_x86_64,spir64_gen -Xsycl-target-backend=spir64_gen "-device dg2" %s -o %t.out +// RUN: %{run-aux} %clangxx -fsycl -fsycl-targets=spir64_x86_64,intel_gpu_dg2 %s -o %t.out // RUN: %{run} %t.out #include "bfloat16_example.hpp" diff --git a/sycl/test-e2e/BFloat16/bfloat16_example_aot_cpu.cpp b/sycl/test-e2e/BFloat16/bfloat16_example_aot_cpu.cpp index edf317f5b57e9..f110d125703ef 100644 --- a/sycl/test-e2e/BFloat16/bfloat16_example_aot_cpu.cpp +++ b/sycl/test-e2e/BFloat16/bfloat16_example_aot_cpu.cpp @@ -5,11 +5,11 @@ // REQUIRES: opencl-aot, ocloc, gpu-intel-dg2, any-device-is-cpu, target-spir -// RUN: %clangxx -fsycl -fsycl-targets=spir64,spir64_gen -Xsycl-target-backend=spir64_gen "-device dg2" %s -o %t.out +// RUN: %clangxx -fsycl -fsycl-targets=spir64,intel_gpu_dg2 %s -o %t.out // RUN: %if cpu %{ %{run} %t.out %} // CPU AOT targets host isa, so we compile on the run system instead. -// RUN: %{run-aux} %clangxx -fsycl -fsycl-targets=spir64_x86_64,spir64_gen -Xsycl-target-backend=spir64_gen "-device dg2" %s -o %t.out +// RUN: %{run-aux} %clangxx -fsycl -fsycl-targets=spir64_x86_64,intel_gpu_dg2 %s -o %t.out // RUN: %if cpu %{ %{run} %t.out %} #include "bfloat16_example.hpp" diff --git a/sycl/test-e2e/BFloat16/bfloat16_example_aot_gpu.cpp b/sycl/test-e2e/BFloat16/bfloat16_example_aot_gpu.cpp index 6b1dfe8f64631..846b9b4cc99ba 100644 --- a/sycl/test-e2e/BFloat16/bfloat16_example_aot_gpu.cpp +++ b/sycl/test-e2e/BFloat16/bfloat16_example_aot_gpu.cpp @@ -5,10 +5,10 @@ // REQUIRES: opencl-aot, ocloc, gpu-intel-dg2, any-device-is-gpu, target-spir -// RUN: %clangxx -fsycl -fsycl-targets=spir64_gen -Xsycl-target-backend "-device dg2" %s -o %t.out +// RUN: %clangxx -fsycl -fsycl-targets=intel_gpu_dg2 %s -o %t.out // RUN: %if gpu %{%{run} %t.out %} -// RUN: %clangxx -fsycl -fsycl-targets=spir64_gen -Xsycl-target-backend "-device *" %s -o %t.out +// RUN: %clangxx -fsycl -fsycl-targets=%{gpu_aot_target} %s -o %t.out // RUN: %if gpu %{%{run} %t.out %} #include "bfloat16_example.hpp" diff --git a/sycl/test-e2e/Basic/build_log.cpp b/sycl/test-e2e/Basic/build_log.cpp index 0dc3f4fb8019c..359584f1e777a 100644 --- a/sycl/test-e2e/Basic/build_log.cpp +++ b/sycl/test-e2e/Basic/build_log.cpp @@ -2,7 +2,7 @@ // UNSUPPORTED: arch-intel_gpu_dg2 // UNSUPPORTED-INTENDED: see https://github.com/intel/llvm/pull/20643 -// RUN: %clangxx -fsycl -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen "-device dg2" %s -o %t.out +// RUN: %clangxx -fsycl -fsycl-targets=intel_gpu_dg2 %s -o %t.out // RUN: env SYCL_RT_WARNING_LEVEL=2 %{run} %t.out 2>&1 | FileCheck %s #include diff --git a/sycl/test-e2e/DeviceCodeSplit/aot-gpu.cpp b/sycl/test-e2e/DeviceCodeSplit/aot-gpu.cpp index 2582be8a498fe..b2c76fef3b071 100644 --- a/sycl/test-e2e/DeviceCodeSplit/aot-gpu.cpp +++ b/sycl/test-e2e/DeviceCodeSplit/aot-gpu.cpp @@ -1,9 +1,7 @@ // REQUIRES: ocloc, gpu, target-spir, !gpu-intel-gen12 // // RUN: %clangxx -fsycl -fsycl-device-code-split=per_source \ -// RUN: -fsycl-targets=spir64_gen \ -// RUN: -Xsycl-target-backend=spir64_gen \ -// RUN: "-device dg2" -I %S/Inputs -o %t.out \ +// RUN: -fsycl-targets=intel_gpu_dg2 -I %S/Inputs -o %t.out \ // RUN: %S/split-per-source-main.cpp \ // RUN: %S/Inputs/split-per-source-second-file.cpp \ // RUN: -fsycl-dead-args-optimization diff --git a/sycl/test-e2e/DeviceImageDependencies/dynamic_aot.cpp b/sycl/test-e2e/DeviceImageDependencies/dynamic_aot.cpp index 504d61bb9efc0..498929050e4b4 100644 --- a/sycl/test-e2e/DeviceImageDependencies/dynamic_aot.cpp +++ b/sycl/test-e2e/DeviceImageDependencies/dynamic_aot.cpp @@ -2,7 +2,7 @@ // REQUIRES: ocloc, gpu, target-spir -// DEFINE: %{aot_options} = -fsycl -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen %gpu_aot_target_opts -DUSE_AOT +// DEFINE: %{aot_options} = -fsycl -fsycl-targets=%{gpu_aot_target} -DUSE_AOT // DEFINE: %{dynamic_lib_options} = %{aot_options} %fPIC %shared_lib -fsycl-allow-device-image-dependencies -ftarget-export-symbols -I %S/Inputs %if windows %{-DMAKE_DLL %} // DEFINE: %{dynamic_lib_suffix} = %if windows %{dll%} %else %{so%} diff --git a/sycl/test-e2e/DeviceImageDependencies/objects_aot.cpp b/sycl/test-e2e/DeviceImageDependencies/objects_aot.cpp index cc2ee41c8f9e5..aec81f397bc98 100644 --- a/sycl/test-e2e/DeviceImageDependencies/objects_aot.cpp +++ b/sycl/test-e2e/DeviceImageDependencies/objects_aot.cpp @@ -8,12 +8,11 @@ // UNSUPPORTED: new-offload-model // UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/23302 -// DEFINE: %{aot_options} = -fsycl -fsycl-targets=spir64_gen -DUSE_AOT -// DEFINE: %{aot_backend_options} = -Xsycl-target-backend=spir64_gen %gpu_aot_target_opts +// DEFINE: %{aot_options} = -fsycl -fsycl-targets=%{gpu_aot_target} -DUSE_AOT // RUN: %clangxx %{aot_options} %S/Inputs/a.cpp -I %S/Inputs -c -o %t_a.o // RUN: %clangxx %{aot_options} %S/Inputs/b.cpp -I %S/Inputs -c -o %t_b.o // RUN: %clangxx %{aot_options} %S/Inputs/c.cpp -I %S/Inputs -c -o %t_c.o // RUN: %clangxx %{aot_options} %S/Inputs/d.cpp -I %S/Inputs -c -o %t_d.o -// RUN: %clangxx %{aot_options} %{aot_backend_options} -fsycl-device-code-split=per_kernel -fsycl-allow-device-image-dependencies -ftarget-export-symbols %t_a.o %t_b.o %t_c.o %t_d.o %S/Inputs/basic.cpp -o %t.out +// RUN: %clangxx %{aot_options} -fsycl-device-code-split=per_kernel -fsycl-allow-device-image-dependencies -ftarget-export-symbols %t_a.o %t_b.o %t_c.o %t_d.o %S/Inputs/basic.cpp -o %t.out // RUN: %{run} %t.out diff --git a/sycl/test-e2e/ESIMD/aot_mixed.cpp b/sycl/test-e2e/ESIMD/aot_mixed.cpp index 0b678400af5cd..07ea7677f18f6 100644 --- a/sycl/test-e2e/ESIMD/aot_mixed.cpp +++ b/sycl/test-e2e/ESIMD/aot_mixed.cpp @@ -7,9 +7,9 @@ //===----------------------------------------------------------------------===// // TODO: Enable on other GPUs once internal ticket is fixed // REQUIRES: ocloc && gpu-intel-dg2 -// RUN: %clangxx -fsycl -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen "-device dg2" -o %t.sycl.out -DENABLE_SYCL=0 %s +// RUN: %clangxx -fsycl -fsycl-targets=intel_gpu_dg2 -o %t.sycl.out -DENABLE_SYCL=0 %s // RUN: %{run} %t.sycl.out -// RUN: %clangxx -fsycl -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen "-device dg2" -o %t.out %s +// RUN: %clangxx -fsycl -fsycl-targets=intel_gpu_dg2 -o %t.out %s // RUN: %{run} %t.out // This test checks the following ESIMD ahead-of-time compilation scenarios: diff --git a/sycl/test-e2e/KernelAndProgram/test_cache_jit_aot.cpp b/sycl/test-e2e/KernelAndProgram/test_cache_jit_aot.cpp index bece8886570c6..24a6b92cb12cc 100644 --- a/sycl/test-e2e/KernelAndProgram/test_cache_jit_aot.cpp +++ b/sycl/test-e2e/KernelAndProgram/test_cache_jit_aot.cpp @@ -27,7 +27,7 @@ // RUN: %{cache_vars} %{run-unfiltered-devices} %t.out 2>&1 | FileCheck %s --check-prefixes RESULT1 // ****************************** // -// REDEFINE: %{build_cmd} = %clangxx -fsycl -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen "-device acm-g10" %s +// REDEFINE: %{build_cmd} = %clangxx -fsycl -fsycl-targets=intel_gpu_acm_g10 %s // ****************************** // Check the logs first. // RUN: %{run-aux} %{build_cmd} -DVALUE=1 -o %t.out @@ -47,7 +47,7 @@ // RUN: %{cache_vars} %{run-unfiltered-devices} %t.out 2>&1 | FileCheck %s --check-prefixes RESULT1 // ****************************** // -// REDEFINE: %{build_cmd} = %clangxx -fsycl -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen "-device acm-g10,acm-g11" %s +// REDEFINE: %{build_cmd} = %clangxx -fsycl -fsycl-targets=intel_gpu_acm_g10,intel_gpu_acm_g11 %s // ****************************** // Check the logs first. // RUN: %{run-aux} %{build_cmd} -DVALUE=1 -o %t.out diff --git a/sycl/test-e2e/MemorySanitizer/lit.local.cfg b/sycl/test-e2e/MemorySanitizer/lit.local.cfg index b27814baeea78..48d52b892a4b9 100644 --- a/sycl/test-e2e/MemorySanitizer/lit.local.cfg +++ b/sycl/test-e2e/MemorySanitizer/lit.local.cfg @@ -11,8 +11,7 @@ config.substitutions.append( config.substitutions.append( ( "%device_msan_aot_gpu_flags", - "-Xarch_device -fsanitize=memory -fsycl-targets=spir64_gen " - "-Xsycl-target-backend=spir64_gen '-device pvc'", + "-Xarch_device -fsanitize=memory -fsycl-targets=intel_gpu_pvc", ) ) diff --git a/sycl/test-e2e/NewOffloadDriver/aot-gpu.cpp b/sycl/test-e2e/NewOffloadDriver/aot-gpu.cpp index 99dab3ed39d52..729688a747eba 100644 --- a/sycl/test-e2e/NewOffloadDriver/aot-gpu.cpp +++ b/sycl/test-e2e/NewOffloadDriver/aot-gpu.cpp @@ -2,9 +2,7 @@ // Test with `--offload-new-driver` // // RUN: %clangxx -fsycl -fsycl-device-code-split=per_source \ -// RUN: -fsycl-targets=spir64_gen \ -// RUN: -Xsycl-target-backend=spir64_gen \ -// RUN: "-device dg2" -I %S/Inputs -o %t.out \ +// RUN: -fsycl-targets=intel_gpu_dg2 -I %S/Inputs -o %t.out \ // RUN: %S/split-per-source-main.cpp \ // RUN: %S/Inputs/split-per-source-second-file.cpp \ // RUN: -fsycl-dead-args-optimization --offload-new-driver diff --git a/sycl/test-e2e/NewOffloadDriver/aot-multiple-devices.cpp b/sycl/test-e2e/NewOffloadDriver/aot-multiple-devices.cpp index cbee51ad49877..5639ff8711d31 100644 --- a/sycl/test-e2e/NewOffloadDriver/aot-multiple-devices.cpp +++ b/sycl/test-e2e/NewOffloadDriver/aot-multiple-devices.cpp @@ -17,16 +17,16 @@ // Produce a fat object for all targets (generic SPIR-V, CPU, GPU). // It is required to specify the device name when compiling AOT for GPU. -// RUN: %{run-aux} %clangxx --offload-new-driver -fsycl -fsycl-targets=spir64,spir64_x86_64,spir64_gen -Xsycl-target-backend=spir64_gen %gpu_aot_target_opts %S/Inputs/aot.cpp -c -o %t.o +// RUN: %{run-aux} %clangxx --offload-new-driver -fsycl -fsycl-targets=spir64,spir64_x86_64,%{gpu_aot_target} %S/Inputs/aot.cpp -c -o %t.o // Check that all targets compiled in the fat binary are selected during // linking. -// RUN: %{run-aux} %clangxx --offload-new-driver -fsycl -fsycl-targets=spir64_x86_64,spir64_gen -Xsycl-target-backend=spir64_gen %gpu_aot_target_opts -v %t.o -o %t_cpu_gpu.out 2>&1 | FileCheck %s --check-prefix=CHECK-DEVICE +// RUN: %{run-aux} %clangxx --offload-new-driver -fsycl -fsycl-targets=spir64_x86_64,%{gpu_aot_target} -v %t.o -o %t_cpu_gpu.out 2>&1 | FileCheck %s --check-prefix=CHECK-DEVICE // RUN: %{run} %t_cpu_gpu.out // Check that all targets compiled in the fat binary are selected during // linking (including CPU even though it is not specified during linking). -// RUN: %{run-aux} %clangxx --offload-new-driver -fsycl -fsycl-targets=spir64,spir64_gen -Xsycl-target-backend=spir64_gen %gpu_aot_target_opts -v %t.o -o %t_spv_gpu.out 2>&1 | FileCheck %s --check-prefix=CHECK-DEVICE +// RUN: %{run-aux} %clangxx --offload-new-driver -fsycl -fsycl-targets=spir64,%{gpu_aot_target} -v %t.o -o %t_spv_gpu.out 2>&1 | FileCheck %s --check-prefix=CHECK-DEVICE // Check that execution on AOT-compatible devices is unaffected // RUN: env ONEAPI_DEVICE_SELECTOR="*:gpu" %{run-unfiltered-devices} %t_spv_gpu.out diff --git a/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_several_targets.cpp b/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_several_targets.cpp index 51cfe332d1234..ee4bd416f4e65 100644 --- a/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_several_targets.cpp +++ b/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_several_targets.cpp @@ -1,5 +1,5 @@ // REQUIRES: opencl-aot, ocloc, opencl-cpu-rt -// RUN: %{run-aux} %clangxx -fsycl -fsycl-targets=spir64_x86_64,spir64_gen -Xsycl-target-backend=spir64_gen %gpu_aot_target_opts %S/Inputs/is_compatible_with_env.cpp -o %t.out +// RUN: %{run-aux} %clangxx -fsycl -fsycl-targets=spir64_x86_64,%{gpu_aot_target} %S/Inputs/is_compatible_with_env.cpp -o %t.out // RUN: %if !(level_zero || opencl) %{ not %} %{run} %t.out diff --git a/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_spir64_gen.cpp b/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_spir64_gen.cpp index 4e836a3bb6d4f..0f942089c8a80 100644 --- a/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_spir64_gen.cpp +++ b/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_spir64_gen.cpp @@ -1,5 +1,5 @@ // REQUIRES: ocloc, target-spir -// RUN: %clangxx -fsycl -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen "-device *" %S/Inputs/is_compatible_with_env.cpp -o %t.out +// RUN: %clangxx -fsycl -fsycl-targets=%{gpu_aot_target} %S/Inputs/is_compatible_with_env.cpp -o %t.out // RUN: %if !(level_zero || opencl && gpu) %{ not %} %{run} %t.out diff --git a/sycl/test-e2e/ProgramManager/uneven_kernel_split.cpp b/sycl/test-e2e/ProgramManager/uneven_kernel_split.cpp index 414fd33565d47..45f627e0939f6 100644 --- a/sycl/test-e2e/ProgramManager/uneven_kernel_split.cpp +++ b/sycl/test-e2e/ProgramManager/uneven_kernel_split.cpp @@ -5,9 +5,9 @@ // DEFINE: %{gpu_compile_flags} = %if !new-offload-model %{-Wno-unused-command-line-argument%} // RUN: %clangxx -fsycl -fsycl-targets=spir64_x86_64 -I %S/Inputs/ %S/uneven_kernel_split.cpp -c -o %t.o -// RUN: %clangxx -fsycl -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen %gpu_aot_target_opts -I %S/Inputs/ %S/Inputs/gpu_kernel1.cpp %{gpu_compile_flags} -c -o %t1.o -// RUN: %clangxx -fsycl -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen %gpu_aot_target_opts -I %S/Inputs/ %S/Inputs/gpu_kernel2.cpp %{gpu_compile_flags} -c -o %t2.o -// RUN: %clangxx -fsycl -fsycl-targets=spir64_x86_64,spir64_gen -Xsycl-target-backend=spir64_gen %gpu_aot_target_opts -Wno-unused-command-line-argument %t.o %t1.o %t2.o -o %t.out +// RUN: %clangxx -fsycl -fsycl-targets=%{gpu_aot_target} -I %S/Inputs/ %S/Inputs/gpu_kernel1.cpp %{gpu_compile_flags} -c -o %t1.o +// RUN: %clangxx -fsycl -fsycl-targets=%{gpu_aot_target} -I %S/Inputs/ %S/Inputs/gpu_kernel2.cpp %{gpu_compile_flags} -c -o %t2.o +// RUN: %clangxx -fsycl -fsycl-targets=spir64_x86_64,%{gpu_aot_target} -Wno-unused-command-line-argument %t.o %t1.o %t2.o -o %t.out // RUN: %{run} %t.out // Test require the following device image structure: cpu target device image diff --git a/sycl/test-e2e/README.md b/sycl/test-e2e/README.md index 61c6370719c51..0493850694ac3 100644 --- a/sycl/test-e2e/README.md +++ b/sycl/test-e2e/README.md @@ -207,9 +207,10 @@ may be set to "gfx906". Otherwise must be provided via the ***amd_arch*** LIT parameter (e.g., ***--param amd_arch=gfx906***) at runtime via the command line or via the ***LIT_OPTS*** environment variable. -***GPU_AOT_TARGET_OPTS*** - defines additional options which are passed to AOT -compilation command line for GPU device. If not specified "-device *" value is -used. +***GPU_AOT_TARGET*** - defines the Intel GPU target(s) used for AOT compilation, +as a `-fsycl-targets` value (e.g. "intel_gpu_pvc" or +"intel_gpu_pvc,intel_gpu_bmg_g21"). If not specified, the architectures of the +detected Intel GPU devices are used, or "intel_gpu_pvc" if none are detected. ***OFFLOAD_BUILD_TARGET*** - when testing the Offload backend, this must be set to specify the correct build target type for the available Offload device. @@ -292,9 +293,12 @@ configure specific single test execution in the command line: * **compatibility_testing** - forces LIT infra to skip the tests compilation to support compatibility testing (a SYCL application is built with one version of SYCL compiler and then run with different SYCL RT version); -* **gpu_aot_target_opts** - defines additional options which are passed to AOT - compilation command line for GPU device. It can be also set by CMake variable - GPU_AOT_TARGET_OPTS. If not specified "-device *" value is used. +* **gpu_aot_target** - defines the Intel GPU target(s) used for AOT + compilation, as a `-fsycl-targets` value. Tests use it via the + `%{gpu_aot_target}` substitution. It can be also set by CMake variable + GPU_AOT_TARGET. If not specified, the architectures of the detected Intel GPU + devices are used, or "intel_gpu_pvc" if none are detected (e.g. in + `build-only` mode). * **gpu-intel-dg2** - tells LIT infra that Intel GPU DG2 is present in the system. It is developer / CI infra responsibility to make sure that the device is available in the system. @@ -474,12 +478,11 @@ to only run the compilation stage, or the execution stage respectively. The triples to compile in this mode are set via the `sycl_build_targets` lit parameter. Valid build targets are: `spir`, `nvidia`, `amd`, `native_cpu`, `spir_gen`, `spir_x86_64`. These correspond to `spir64`, - `nvptx64-nvidia-cuda`, `amdgcn-amd-amdhsa`, `native_cpu`, `spir64_gen`, and - `spir64_x86_64` triples respectively. `spir_gen` and `spir_x86_64` select the - AOT-only GPU and CPU triples respectively, rather than the JIT `spir64` - triple; `spir_gen` compilation also passes the `GPU_AOT_TARGET_OPTS`/ - `gpu_aot_target_opts` options (see above) to - `-Xsycl-target-backend=spir64_gen`. Each build target should be separated + `nvptx64-nvidia-cuda`, `amdgcn-amd-amdhsa`, `native_cpu`, Intel GPU AOT and + `spir64_x86_64` targets respectively. `spir_gen` and `spir_x86_64` select the + AOT-only GPU and CPU targets respectively, rather than the JIT `spir64` + triple; `spir_gen` compiles for the `GPU_AOT_TARGET`/`gpu_aot_target` Intel + GPU target(s) (see above). Each build target should be separated with a semicolon. This parameter is set to `all` by default, which enables autodetection for the available build targets. A test can be marked as requiring, or not supporting a particular triple via the `target-*` features. diff --git a/sycl/test-e2e/SYCLBIN/aot_object_load.cpp b/sycl/test-e2e/SYCLBIN/aot_object_load.cpp index 494e4b1de20ab..1b61dae4378f0 100644 --- a/sycl/test-e2e/SYCLBIN/aot_object_load.cpp +++ b/sycl/test-e2e/SYCLBIN/aot_object_load.cpp @@ -16,9 +16,8 @@ // -- is in object state and was previously dropped by // -- SYCLBINBinaries::getBestCompatibleImages. // -// -- The test targets spir64_gen via ocloc with -device * -// (%gpu_aot_target_opts) -// -- so it is not locked to a single Intel GPU architecture. spir64_x86_64 +// -- The test targets the configured Intel GPU architecture(s) +// -- (%{gpu_aot_target}) so it is not locked to a single one. spir64_x86_64 // -- is intentionally not exercised here: the Intel OpenCL CPU runtime's // -- clBuildProgram does not accept unresolved imported symbols (no // -- equivalent of IGC's -library-compilation), and the -cmd=compile path @@ -30,8 +29,7 @@ // RUN: %clangxx --offload-new-driver -fsyclbin=object \ // RUN: -fsycl-allow-device-image-dependencies \ -// RUN: -fsycl-targets=spir64_gen \ -// RUN: -Xsycl-target-backend=spir64_gen %gpu_aot_target_opts \ +// RUN: -fsycl-targets=%{gpu_aot_target} \ // RUN: %S/Inputs/aot_object_with_imports.cpp -o %t.syclbin // RUN: %{build} -o %t.out // RUN: %{run} %t.out %t.syclbin diff --git a/sycl/test-e2e/SYCLBIN/aot_object_no_symbols.cpp b/sycl/test-e2e/SYCLBIN/aot_object_no_symbols.cpp index 238cc3c20b47f..21a70c84843dd 100644 --- a/sycl/test-e2e/SYCLBIN/aot_object_no_symbols.cpp +++ b/sycl/test-e2e/SYCLBIN/aot_object_no_symbols.cpp @@ -30,14 +30,13 @@ // -- the native-AOT partition (build with ALLOW_UNRESOLVED_SYMBOLS + // -- dynamicLink) to reach executable state. // -// -- The -device * wildcard (%gpu_aot_target_opts) keeps this off a single -// -- Intel GPU architecture. Complements the hardware-independent unit test +// -- Targeting the configured Intel GPU architecture(s) (%{gpu_aot_target}) +// -- keeps this off a single hard-coded one. Complements the hardware-independent unit test // -- SYCLBINSelector.NativeWithoutSymbolsSurfacedForObject, which covers the // -- selector in isolation but neither the producer nor the link-and-run path. // RUN: %clangxx --offload-new-driver -fsyclbin=object \ -// RUN: -fsycl-targets=spir64_gen \ -// RUN: -Xsycl-target-backend=spir64_gen %gpu_aot_target_opts \ +// RUN: -fsycl-targets=%{gpu_aot_target} \ // RUN: %S/Inputs/aot_object_no_symbols.cpp -o %t.syclbin // RUN: %{build} -o %t.out // RUN: %{run} %t.out %t.syclbin diff --git a/sycl/test-e2e/SYCLBIN/basic_executable_bmg_aot.cpp b/sycl/test-e2e/SYCLBIN/basic_executable_bmg_aot.cpp index ffc7c392f0ee5..4ca4a807154f2 100644 --- a/sycl/test-e2e/SYCLBIN/basic_executable_bmg_aot.cpp +++ b/sycl/test-e2e/SYCLBIN/basic_executable_bmg_aot.cpp @@ -3,7 +3,7 @@ // -- Basic test for compiling and loading a SYCLBIN kernel_bundle in executable // -- state using only AOT target for BMG G21. -// RUN: %clangxx --offload-new-driver -fsyclbin=executable -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen "-device bmg-g21" %S/Inputs/basic_kernel.cpp -o %t.syclbin +// RUN: %clangxx --offload-new-driver -fsyclbin=executable -fsycl-targets=intel_gpu_bmg_g21 %S/Inputs/basic_kernel.cpp -o %t.syclbin // RUN: %{build} -o %t.out // RUN: %{run} %t.out %t.syclbin diff --git a/sycl/test-e2e/SYCLBIN/dae_executable_bmg_aot.cpp b/sycl/test-e2e/SYCLBIN/dae_executable_bmg_aot.cpp index 5d43b4db128ec..c56eae62c1bfb 100644 --- a/sycl/test-e2e/SYCLBIN/dae_executable_bmg_aot.cpp +++ b/sycl/test-e2e/SYCLBIN/dae_executable_bmg_aot.cpp @@ -3,7 +3,7 @@ // -- Test for using a kernel from a SYCLBIN with a dead argument using only AOT // target for BMG G21. -// RUN: %clangxx --offload-new-driver -fsyclbin=executable -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen "-device bmg-g21" %S/Inputs/dae_kernel.cpp -o %t.syclbin +// RUN: %clangxx --offload-new-driver -fsyclbin=executable -fsycl-targets=intel_gpu_bmg_g21 %S/Inputs/dae_kernel.cpp -o %t.syclbin // RUN: %{build} -o %t.out // RUN: %{run} %t.out %t.syclbin diff --git a/sycl/test-e2e/SYCLBIN/dg_executable_bmg_aot.cpp b/sycl/test-e2e/SYCLBIN/dg_executable_bmg_aot.cpp index d9c5c16fbe5b8..9cfcc453aa7a3 100644 --- a/sycl/test-e2e/SYCLBIN/dg_executable_bmg_aot.cpp +++ b/sycl/test-e2e/SYCLBIN/dg_executable_bmg_aot.cpp @@ -9,7 +9,7 @@ // UNSUPPORTED: cuda // UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/19533 -// RUN: %clangxx --offload-new-driver -fsyclbin=executable -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen "-device bmg-g21" %S/Inputs/dg_kernel.cpp -o %t.syclbin +// RUN: %clangxx --offload-new-driver -fsyclbin=executable -fsycl-targets=intel_gpu_bmg_g21 %S/Inputs/dg_kernel.cpp -o %t.syclbin // RUN: %{build} -o %t.out // RUN: %{run} %t.out %t.syclbin diff --git a/sycl/test-e2e/SYCLBIN/link_object_aot.cpp b/sycl/test-e2e/SYCLBIN/link_object_aot.cpp index 1242b8d1bd5a1..c33539b4f97f2 100644 --- a/sycl/test-e2e/SYCLBIN/link_object_aot.cpp +++ b/sycl/test-e2e/SYCLBIN/link_object_aot.cpp @@ -26,12 +26,10 @@ // -- exercises the full AOT object-state link-and-run path. // RUN: %clangxx --offload-new-driver -fsyclbin=object \ -// RUN: -fsycl-targets=spir64_gen \ -// RUN: -Xsycl-target-backend=spir64_gen %gpu_aot_target_opts \ +// RUN: -fsycl-targets=%{gpu_aot_target} \ // RUN: %S/Inputs/exporting_function.cpp -o %t.export.syclbin // RUN: %clangxx --offload-new-driver -fsyclbin=object \ -// RUN: -fsycl-targets=spir64_gen \ -// RUN: -Xsycl-target-backend=spir64_gen %gpu_aot_target_opts \ +// RUN: -fsycl-targets=%{gpu_aot_target} \ // RUN: %S/Inputs/importing_kernel.cpp -o %t.import.syclbin // RUN: %{build} -o %t.out // RUN: %{run} %t.out %t.export.syclbin %t.import.syclbin diff --git a/sycl/test-e2e/SYCLBIN/lit.local.cfg b/sycl/test-e2e/SYCLBIN/lit.local.cfg index 5467332c99af5..44ca4075e1312 100644 --- a/sycl/test-e2e/SYCLBIN/lit.local.cfg +++ b/sycl/test-e2e/SYCLBIN/lit.local.cfg @@ -2,10 +2,8 @@ syclbin_exec_opts = "" if 'target-spir' in config.available_features: syclbin_exec_opts = '-fsycl-targets=spir64,spir64_x86_64' if 'ocloc' in config.available_features: - # Due to a bug in clang-linker-wrapper we cannot use -device * on Windows, so - # we pick specific targets for now. Once that is fixed we can switch to the - # wildcard. See https://github.com/intel/llvm/issues/19373. - syclbin_exec_opts = syclbin_exec_opts + ',spir64_gen -Xsycl-target-backend=spir64_gen "-device pvc,dg2,bmg-g21"' + # Build for a fixed set of Intel GPU architectures. + syclbin_exec_opts = syclbin_exec_opts + ',intel_gpu_pvc,intel_gpu_dg2,intel_gpu_bmg_g21' config.substitutions.append(('%{syclbin_exec_opts}', format(syclbin_exec_opts))) # https://github.com/intel/llvm/issues/20142 diff --git a/sycl/test-e2e/SYCLBIN/optional_kernel_features_executable_bmg_aot.cpp b/sycl/test-e2e/SYCLBIN/optional_kernel_features_executable_bmg_aot.cpp index f9d377975d741..ef585d7de8bf5 100644 --- a/sycl/test-e2e/SYCLBIN/optional_kernel_features_executable_bmg_aot.cpp +++ b/sycl/test-e2e/SYCLBIN/optional_kernel_features_executable_bmg_aot.cpp @@ -3,7 +3,7 @@ // -- Test for compiling and loading a kernel bundle with a SYCLBIN containing // the use of optional kernel features using only AOT target for BMG G21. -// RUN: %clangxx --offload-new-driver -fsyclbin=executable -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen "-device bmg-g21" %S/Inputs/optional_kernel_features.cpp -o %t.syclbin +// RUN: %clangxx --offload-new-driver -fsyclbin=executable -fsycl-targets=intel_gpu_bmg_g21 %S/Inputs/optional_kernel_features.cpp -o %t.syclbin // RUN: %{build} -o %t.out // RUN: %{run} %t.out %t.syclbin diff --git a/sycl/test-e2e/SYCLBIN/simple_kernel_aot.cpp b/sycl/test-e2e/SYCLBIN/simple_kernel_aot.cpp index 9c2bdc2ffbc43..099aebd3a345f 100644 --- a/sycl/test-e2e/SYCLBIN/simple_kernel_aot.cpp +++ b/sycl/test-e2e/SYCLBIN/simple_kernel_aot.cpp @@ -3,11 +3,11 @@ // REQUIRES: ocloc -// RUN: %clangxx --offload-new-driver -fsyclbin=input -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen "-device bmg-g21" -o %t.input_bmg_g21_gpu_device_arch.syclbin %s -// RUN: %clangxx --offload-new-driver -fsyclbin=input -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen "-device *" -o %t.input_all_gpu_device_archs.syclbin %s -// RUN: %clangxx --offload-new-driver -fsyclbin=object -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen "-device bmg-g21" -o %t.object_bmg_g21_gpu_device_arch.syclbin %s -// RUN: %clangxx --offload-new-driver -fsyclbin=executable -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen "-device bmg-g21" -o %t.executable_bmg_g21_gpu_device_arch.syclbin %s -// RUN: %clangxx --offload-new-driver -fsyclbin -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen "-device bmg-g21" -o %t.default_bmg_g21_gpu_device_arch.syclbin %s +// RUN: %clangxx --offload-new-driver -fsyclbin=input -fsycl-targets=intel_gpu_bmg_g21 -o %t.input_bmg_g21_gpu_device_arch.syclbin %s +// RUN: %clangxx --offload-new-driver -fsyclbin=input -fsycl-targets=%{gpu_aot_target} -o %t.input_all_gpu_device_archs.syclbin %s +// RUN: %clangxx --offload-new-driver -fsyclbin=object -fsycl-targets=intel_gpu_bmg_g21 -o %t.object_bmg_g21_gpu_device_arch.syclbin %s +// RUN: %clangxx --offload-new-driver -fsyclbin=executable -fsycl-targets=intel_gpu_bmg_g21 -o %t.executable_bmg_g21_gpu_device_arch.syclbin %s +// RUN: %clangxx --offload-new-driver -fsyclbin -fsycl-targets=intel_gpu_bmg_g21 -o %t.default_bmg_g21_gpu_device_arch.syclbin %s // RUN: syclbin-dump %t.input_bmg_g21_gpu_device_arch.syclbin | FileCheck %s --check-prefix CHECK-INPUT-BMG-G21-GPU-DEVICE-ARCH // RUN: syclbin-dump %t.input_all_gpu_device_archs.syclbin | FileCheck %s --check-prefix CHECK-INPUT-ALL-GPU-DEVICE-ARCHS diff --git a/sycl/test-e2e/SpecConstants/2020/image_selection.cpp b/sycl/test-e2e/SpecConstants/2020/image_selection.cpp index 89c32efe4a214..0f5a4634c6481 100644 --- a/sycl/test-e2e/SpecConstants/2020/image_selection.cpp +++ b/sycl/test-e2e/SpecConstants/2020/image_selection.cpp @@ -5,21 +5,21 @@ // are replaced with defaults and original images. // clang-format off -// RUN: %clangxx -fsycl-add-default-spec-consts-image -fsycl -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen %gpu_aot_target_opts %s -o %t1.out +// RUN: %clangxx -fsycl-add-default-spec-consts-image -fsycl -fsycl-targets=%{gpu_aot_target} %s -o %t1.out // RUN: env SYCL_UR_TRACE=2 %{run} %t1.out | FileCheck --match-full-lines --check-prefix=CHECK-ENABLED %s // clang-format on // Check the behaviour when -fsycl-add-default-spec-consts-image option is not // used. -// RUN: %clangxx -fsycl -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen %gpu_aot_target_opts %s -o %t2.out +// RUN: %clangxx -fsycl -fsycl-targets=%{gpu_aot_target} %s -o %t2.out // RUN: env SYCL_UR_TRACE=2 %{run} %t2.out | FileCheck --match-full-lines --check-prefix=CHECK-DEFAULT %s // Check the behaviour when -fsycl-add-default-spec-consts-image option is used // and we have spirv image in addition to AOT. // clang-format off -// RUN: %clangxx -fsycl -fsycl-targets=spir64,spir64_gen -Xsycl-target-backend=spir64_gen %gpu_aot_target_opts %s -o %t3.out +// RUN: %clangxx -fsycl -fsycl-targets=spir64,%{gpu_aot_target} %s -o %t3.out // RUN: env SYCL_UR_TRACE=2 %{run} %t3.out | FileCheck --match-full-lines --check-prefix=CHECK-MIX %s // clang-format on @@ -28,7 +28,7 @@ // choose images with inlined values in this case. // clang-format off -// RUN: %clangxx -fsycl-add-default-spec-consts-image -fsycl -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen %gpu_aot_target_opts %s -o %t3.out +// RUN: %clangxx -fsycl-add-default-spec-consts-image -fsycl -fsycl-targets=%{gpu_aot_target} %s -o %t3.out // RUN: env SYCL_UR_TRACE=2 %{run} %t3.out | FileCheck --match-full-lines --check-prefix=CHECK-DEFAULT-EXPLICIT-SET %s // clang-format on @@ -38,7 +38,7 @@ // this case. // clang-format off -// RUN: %clangxx -fsycl-add-default-spec-consts-image -fsycl -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen %gpu_aot_target_opts %s -o %t3.out +// RUN: %clangxx -fsycl-add-default-spec-consts-image -fsycl -fsycl-targets=%{gpu_aot_target} %s -o %t3.out // RUN: env SYCL_UR_TRACE=2 %{run} %t3.out | FileCheck --match-full-lines --check-prefix=CHECK-DEFAULT-BACK-TO-DEFAULT %s // clang-format on #include diff --git a/sycl/test-e2e/SpecConstants/2020/native_specialization_constant.cpp b/sycl/test-e2e/SpecConstants/2020/native_specialization_constant.cpp index d6bf10c393419..bb7f0085ee221 100644 --- a/sycl/test-e2e/SpecConstants/2020/native_specialization_constant.cpp +++ b/sycl/test-e2e/SpecConstants/2020/native_specialization_constant.cpp @@ -7,7 +7,7 @@ // RUN: %{build} -DJIT -o %t1.out // RUN: %{run} %t1.out -// RUN: %if any-device-is-gpu %{ %{run-aux} %clangxx -fsycl -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen %gpu_aot_target_opts %s -o %t2.out %} +// RUN: %if any-device-is-gpu %{ %{run-aux} %clangxx -fsycl -fsycl-targets=%{gpu_aot_target} %s -o %t2.out %} // RUN: %if gpu %{ %{run} %t2.out %} // RUN: %if any-device-is-cpu %{ %{run-aux} %clangxx -fsycl -fsycl-targets=spir64_x86_64 %s -o %t3.out %} diff --git a/sycl/test-e2e/SpecConstants/2020/non_native/gpu.cpp b/sycl/test-e2e/SpecConstants/2020/non_native/gpu.cpp index c7c8cfd8a9aa0..ac09141dbdcdf 100644 --- a/sycl/test-e2e/SpecConstants/2020/non_native/gpu.cpp +++ b/sycl/test-e2e/SpecConstants/2020/non_native/gpu.cpp @@ -1,6 +1,6 @@ // REQUIRES: ocloc, gpu, target-spir -// RUN: %clangxx -fsycl -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen %gpu_aot_target_opts %S/Inputs/common.cpp -o %t.out -fsycl-dead-args-optimization +// RUN: %clangxx -fsycl -fsycl-targets=%{gpu_aot_target} %S/Inputs/common.cpp -o %t.out -fsycl-dead-args-optimization // RUN: %{run} %t.out // This test checks correctness of SYCL2020 non-native specialization constants diff --git a/sycl/test-e2e/SpecConstants/2020/non_native/spec_const_alignment_test.cpp b/sycl/test-e2e/SpecConstants/2020/non_native/spec_const_alignment_test.cpp index 6d007586498f7..cd23123eb90eb 100644 --- a/sycl/test-e2e/SpecConstants/2020/non_native/spec_const_alignment_test.cpp +++ b/sycl/test-e2e/SpecConstants/2020/non_native/spec_const_alignment_test.cpp @@ -1,6 +1,6 @@ // REQUIRES: opencl-aot, ocloc, target-spir -// RUN: %if any-device-is-gpu %{ %{run-aux} %clangxx -fsycl -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen %gpu_aot_target_opts %s -o %t1.out %} +// RUN: %if any-device-is-gpu %{ %{run-aux} %clangxx -fsycl -fsycl-targets=%{gpu_aot_target} %s -o %t1.out %} // RUN: %if gpu %{ %{run} %t1.out %} // RUN: %if any-device-is-cpu %{ %{run-aux} %clangxx -fsycl -fsycl-targets=spir64_x86_64 %s -o %t2.out %} diff --git a/sycl/test-e2e/ThreadSanitizer/lit.local.cfg b/sycl/test-e2e/ThreadSanitizer/lit.local.cfg index 1d75725c0e5c6..b8bef8e9b8474 100644 --- a/sycl/test-e2e/ThreadSanitizer/lit.local.cfg +++ b/sycl/test-e2e/ThreadSanitizer/lit.local.cfg @@ -30,8 +30,7 @@ config.substitutions.append( config.substitutions.append( ( "%device_tsan_aot_gpu_flags", - "-Xarch_device -fsanitize=thread -fsycl-targets=spir64_gen " - "-Xsycl-target-backend=spir64_gen '-device pvc'", + "-Xarch_device -fsanitize=thread -fsycl-targets=intel_gpu_pvc", ) ) diff --git a/sycl/test-e2e/format.py b/sycl/test-e2e/format.py index 587b71758f1e3..619d9bb0950a5 100644 --- a/sycl/test-e2e/format.py +++ b/sycl/test-e2e/format.py @@ -298,7 +298,15 @@ def execute(self, test, litConfig): (backend, _) = sycl_device.split(":") build_targets.add(test.config.backend_to_target[backend]) - triples = set(test.config.target_to_triple[t] for t in build_targets) + # spir_gen builds for the configured Intel GPU architecture(s). + triples = set( + ( + test.config.gpu_aot_target + if t == "target-spir_gen" + else test.config.target_to_triple[t] + ) + for t in build_targets + ) features_for_test = test.config.available_features.union(build_targets) substitutions = lit.TestRunner.getDefaultSubstitutions(test, tmpDir, tmpBase) @@ -306,10 +314,6 @@ def execute(self, test, litConfig): substitutions.append(("%{sycl_triple}", format(",".join(triples)))) sycl_target_opts = "-fsycl-targets=%{sycl_triple}" - if "target-spir_gen" in build_targets: - sycl_target_opts += " -Xsycl-target-backend=spir64_gen {}".format( - test.config.gpu_aot_target_opts - ) if "target-amd" in build_targets: hip_arch_opts = ( " -Xsycl-target-backend=amdgcn-amd-amdhsa --offload-arch={}".format( diff --git a/sycl/test-e2e/lit.cfg.py b/sycl/test-e2e/lit.cfg.py index 99977019f962f..69812319374f9 100644 --- a/sycl/test-e2e/lit.cfg.py +++ b/sycl/test-e2e/lit.cfg.py @@ -698,11 +698,6 @@ def open_check_file(file_name): directx_libs = ["/clang:" + l for l in directx_libs] config.substitutions.append(("%link-directx", " ".join(directx_libs))) -if not config.gpu_aot_target_opts: - config.gpu_aot_target_opts = '"-device *"' - -config.substitutions.append(("%gpu_aot_target_opts", config.gpu_aot_target_opts)) - if config.dump_ir_supported: config.available_features.add("dump_ir") @@ -1293,6 +1288,18 @@ def get_sycl_ls_verbose(sycl_device, env): ) config.substitutions.append(("%amd_arch_options", amd_arch_flags)) +# Intel GPU AOT target(s), as a -fsycl-targets value. Uses the detected Intel +# GPU architectures or defaults to intel_gpu_pvc in build-only mode. +if not config.gpu_aot_target: + intel_gpu_archs = sorted( + feature[len("arch-") :] + for feature in set().union(*config.sycl_dev_features.values()) + if feature.startswith("arch-intel_gpu_") + ) + config.gpu_aot_target = ",".join(intel_gpu_archs) or "intel_gpu_pvc" +lit_config.note("Intel GPU AOT target: {}".format(config.gpu_aot_target)) +config.substitutions.append(("%{gpu_aot_target}", config.gpu_aot_target)) + if lit_config.params.get("compatibility_testing", "False") != "False": config.substitutions.append(("%clangxx", " true ")) config.substitutions.append(("%clang", " true ")) diff --git a/sycl/test-e2e/lit.site.cfg.py.in b/sycl/test-e2e/lit.site.cfg.py.in index 5ae02752917b1..682defae5e813 100644 --- a/sycl/test-e2e/lit.site.cfg.py.in +++ b/sycl/test-e2e/lit.site.cfg.py.in @@ -54,7 +54,7 @@ config.cxx_flags = lit_config.params.get("cxx_flags", "@SYCL_E2E_CLANG_CXX_FLAGS config.c_flags = "@CMAKE_C_FLAGS@" config.external_tests = "@SYCL_EXTERNAL_TESTS@" config.extra_include = "@CMAKE_CURRENT_SOURCE_DIR@/include" -config.gpu_aot_target_opts = lit_config.params.get("gpu_aot_target_opts", "@GPU_AOT_TARGET_OPTS@") +config.gpu_aot_target = lit_config.params.get("gpu_aot_target", "@GPU_AOT_TARGET@") config.vulkan_include_dir = "@Vulkan_INCLUDE_DIRS@" config.vulkan_lib = "@Vulkan_LIBRARY@" From 7569eea476af0a806235faf989c423863a74afb0 Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Fri, 9 Oct 2026 13:47:38 -0700 Subject: [PATCH 02/22] [SYCL][Driver] Deprecate -fsycl-targets=spir64_gen spir64_gen is deprecated in favor of -fsycl-targets=intel_gpu_, both as a -fsycl-targets value and as the triple of -Xsycl-target-backend/frontend/linker: - With the old offloading model it is still accepted, with a -Wdeprecated warning. Support is removed together with the old model. - With the new offloading model (--offload-new-driver) it is an error. Resolve intel_gpu_* aliases when selecting the device sanitizer library, so that e.g. intel_gpu_12_60_7 selects the same PVC library as '-device 12.60.7' did. New offloading model driver tests that used spir64_gen are converted to the equivalent intel_gpu_ form; cases that only exercised parsing of '-device' remain covered by the old offloading model tests. Update the user documentation accordingly. --- .../clang/Basic/DiagnosticDriverKinds.td | 2 + clang/lib/Driver/Driver.cpp | 36 ++++++++++ clang/lib/Driver/ToolChains/SYCL.cpp | 8 ++- clang/test/Driver/sycl-device-lib.cpp | 37 +++------- .../test/Driver/sycl-device-post-link-opt.cpp | 2 +- .../test/Driver/sycl-device-traits-macros.cpp | 2 +- .../Driver/sycl-foffload-fp32-prec-div.cpp | 4 +- .../Driver/sycl-foffload-fp32-prec-sqrt.cpp | 4 +- clang/test/Driver/sycl-fsyclbin.cpp | 4 +- .../test/Driver/sycl-ftarget-compile-fast.cpp | 4 +- .../sycl-ftarget-register-alloc-mode.cpp | 46 ++++-------- clang/test/Driver/sycl-no-rdc-new-driver.cpp | 2 +- .../sycl-nordc-old-model-include-binary.cpp | 2 +- clang/test/Driver/sycl-ocloc-path.cpp | 5 +- clang/test/Driver/sycl-offload-new-driver.cpp | 58 ++++----------- clang/test/Driver/sycl-offload-save-temps.cpp | 2 - clang/test/Driver/sycl-offload-with-split.cpp | 2 +- clang/test/Driver/sycl-offload.cpp | 11 ++- clang/test/Driver/sycl-pch-create.cpp | 4 +- clang/test/Driver/sycl-pch-use.cpp | 4 +- .../Driver/sycl-spir64-gen-deprecated.cpp | 70 +++++++++++++++++++ clang/test/Driver/sycl-spirv-opt.cpp | 2 +- clang/test/Driver/sycl-spirv-to-ir-opt.cpp | 2 +- .../sycl-suppress-undefined-func-warnings.cpp | 2 +- clang/test/Driver/sycl-unique-prefix.cpp | 2 +- sycl/doc/GetStartedGuide.md | 4 +- sycl/doc/UsersManual.md | 12 +++- .../fp64-conv-emu-convert.cpp | 2 +- .../fp64-conv-emu-test-call-convert.cpp | 2 +- .../fp64-conv-emu-test-call.cpp | 2 +- ...p64-conv-emu-with-recursive-type-decls.cpp | 2 +- .../fp64-conv-emu.cpp | 2 +- 32 files changed, 192 insertions(+), 151 deletions(-) create mode 100644 clang/test/Driver/sycl-spir64-gen-deprecated.cpp diff --git a/clang/include/clang/Basic/DiagnosticDriverKinds.td b/clang/include/clang/Basic/DiagnosticDriverKinds.td index 7eaecee9bb917..947ac096f1975 100644 --- a/clang/include/clang/Basic/DiagnosticDriverKinds.td +++ b/clang/include/clang/Basic/DiagnosticDriverKinds.td @@ -480,6 +480,8 @@ def err_drv_sycl_thinlto_split_off: Error< "'%0' is not supported when '%1' is set with '-fsycl'">; def err_drv_sycl_offload_arch_new_driver: Error< "'--offload-arch' is supported when '-fsycl' is set with '--offload-new-driver'">; +def err_drv_sycl_spir64_gen_new_driver : Error< + "'%0' is not supported with the new offloading model; use '%1' instead">; def err_drv_sycl_offload_arch_missing_value : Error< "must pass in an explicit cpu or gpu architecture to '--offload-arch'">; def warn_drv_offload_target_duplicate diff --git a/clang/lib/Driver/Driver.cpp b/clang/lib/Driver/Driver.cpp index 24a3817aa63f3..52f385a49b24b 100644 --- a/clang/lib/Driver/Driver.cpp +++ b/clang/lib/Driver/Driver.cpp @@ -1286,6 +1286,42 @@ static void diagnoseSYCLOptions(Compilation &C, bool IsSYCL) { << SYCLForceTarget->getAsString(C.getInputArgs()); } + // spir64_gen is deprecated in favor of intel_gpu_, both as a + // -fsycl-targets value and as the triple of -Xsycl-target-*. It is not + // supported with the new offloading model, and is removed together with the + // old offloading model. + if (IsSYCL) { + const Driver &D = C.getDriver(); + auto IsSPIRGenTarget = [&](StringRef Val) { + llvm::Triple TT = D.getSYCLDeviceTriple(Val); + return TT.isSPIR() && TT.getSubArch() == llvm::Triple::SPIRSubArch_gen; + }; + auto DiagnoseSPIRGenTarget = [&](StringRef Spelling, StringRef Val) { + std::string Opt = (Spelling + Val).str(); + std::string Replacement = (Spelling + "intel_gpu_").str(); + if (D.getUseNewOffloadingDriver()) + D.Diag(diag::err_drv_sycl_spir64_gen_new_driver) << Opt << Replacement; + else + D.Diag(diag::warn_drv_deprecated_option_release) + << Opt << /*HasReplacement=*/true << Replacement; + }; + + for (const Arg *A : + C.getInputArgs().filtered(options::OPT_offload_targets_EQ)) { + // Report the spelling the user wrote, e.g. -fsycl-targets=. + StringRef Spelling = + A->getAlias() ? A->getAlias()->getSpelling() : A->getSpelling(); + for (StringRef Val : A->getValues()) + if (IsSPIRGenTarget(Val)) + DiagnoseSPIRGenTarget(Spelling, Val); + } + for (const Arg *A : C.getInputArgs().filtered( + options::OPT_Xsycl_backend_EQ, options::OPT_Xsycl_frontend_EQ, + options::OPT_Xsycl_linker_EQ)) + if (IsSPIRGenTarget(A->getValue())) + DiagnoseSPIRGenTarget(A->getSpelling(), A->getValue()); + } + // Check if -fsycl-host-compiler is used in conjunction with -fsycl. Arg *SYCLHostCompiler = getArgRequiringSYCLRuntime(options::OPT_fsycl_host_compiler_EQ); diff --git a/clang/lib/Driver/ToolChains/SYCL.cpp b/clang/lib/Driver/ToolChains/SYCL.cpp index c439f17af5493..58a3b34ecd7a0 100644 --- a/clang/lib/Driver/ToolChains/SYCL.cpp +++ b/clang/lib/Driver/ToolChains/SYCL.cpp @@ -553,10 +553,14 @@ static void addSYCLDeviceSanitizerLibs( if (SYCLTargetStr.starts_with("spir64_x86_64")) return AOT_CPU; - if (SYCLTargetStr == "intel_gpu_pvc") + // Resolve intel_gpu_* so that every alias of a device (e.g. + // intel_gpu_12_60_7 for pvc) selects the same library. + StringRef GenDevice = SYCL::gen::resolveGenDevice(SYCLTargetStr); + std::string Temp; + if (!GenDevice.empty() && checkPVCDevice(GenDevice.str(), Temp)) return AOT_PVC; - if (SYCLTargetStr.starts_with("intel_gpu_dg2")) + if (GenDevice == "dg2" || GenDevice.starts_with("acm_g")) return AOT_DG2; if (SYCLTargetStr.starts_with("spir64_gen")) { diff --git a/clang/test/Driver/sycl-device-lib.cpp b/clang/test/Driver/sycl-device-lib.cpp index bdeeb8e868323..dd1f3ce3ebe48 100644 --- a/clang/test/Driver/sycl-device-lib.cpp +++ b/clang/test/Driver/sycl-device-lib.cpp @@ -49,21 +49,8 @@ /// test behavior of linking libsycl-asan-pvc for PVC target AOT compilation when asan flag is applied. // RUN: %clangxx -fsycl -fsycl-targets=intel_gpu_pvc --offload-new-driver %s --sysroot=%S/Inputs/SYCL \ // RUN: -Xarch_device -fsanitize=address -### 2>&1 | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_ASAN_PVC -// RUN: %clangxx -fsycl -fsycl-targets=spir64_gen --offload-new-driver %s --sysroot=%S/Inputs/SYCL \ -// RUN: -Xarch_device -fsanitize=address -Xsycl-target-backend "-device pvc" -### 2>&1 \ -// RUN: | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_ASAN_PVC -// RUN: %clangxx -fsycl -fsycl-targets=spir64_gen --offload-new-driver %s --sysroot=%S/Inputs/SYCL \ -// RUN: -Xarch_device -fsanitize=address -Xsycl-target-backend=spir64_gen "-device pvc" -### 2>&1 \ -// RUN: | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_ASAN_PVC -// RUN: %clangxx -fsycl -fsycl-targets=spir64_gen --offload-new-driver %s --sysroot=%S/Inputs/SYCL \ -// RUN: -Xarch_device -fsanitize=address -Xsycl-target-backend "-device 12.60.7" -### 2>&1 \ -// RUN: | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_ASAN_PVC -// RUN: %clangxx -fsycl -fsycl-targets=spir64_gen --offload-new-driver %s --sysroot=%S/Inputs/SYCL \ -// RUN: -Xarch_device -fsanitize=address -Xs "-device pvc" -### 2>&1 \ -// RUN: | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_ASAN_PVC -// RUN: %clangxx -fsycl -fsycl-targets=spir64_gen --offload-new-driver %s --sysroot=%S/Inputs/SYCL \ -// RUN: -Xarch_device -fsanitize=address -Xs "-device 12.60.7" -### 2>&1 \ -// RUN: | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_ASAN_PVC +// RUN: %clangxx -fsycl -fsycl-targets=intel_gpu_12_60_7 --offload-new-driver %s --sysroot=%S/Inputs/SYCL \ +// RUN: -Xarch_device -fsanitize=address -### 2>&1 | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_ASAN_PVC // SYCL_DEVICE_LIB_ASAN_PVC: clang{{.*}} "-mlink-builtin-bitcode" "{{.*}}libsycl-crt.bc" // SYCL_DEVICE_LIB_ASAN_PVC-SAME: {{.*}}libsycl-cmath.bc // SYCL_DEVICE_LIB_ASAN_PVC-SAME: {{.*}}libsycl-imf.bc @@ -98,15 +85,10 @@ /// test behavior of linking libsycl-asan-dg2 for DG2 target AOT compilation when asan flag is applied. // RUN: %clangxx -fsycl -fsycl-targets=intel_gpu_dg2_g10 --offload-new-driver %s --sysroot=%S/Inputs/SYCL \ // RUN: -Xarch_device -fsanitize=address -### 2>&1 | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_ASAN_DG2 -// RUN: %clangxx -fsycl -fsycl-targets=spir64_gen --offload-new-driver %s --sysroot=%S/Inputs/SYCL \ -// RUN: -Xarch_device -fsanitize=address -Xsycl-target-backend "-device dg2" -### 2>&1 \ -// RUN: | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_ASAN_DG2 -// RUN: %clangxx -fsycl -fsycl-targets=spir64_gen --offload-new-driver %s --sysroot=%S/Inputs/SYCL \ -// RUN: -Xarch_device -fsanitize=address -Xsycl-target-backend=spir64_gen "-device dg2" -### 2>&1 \ -// RUN: | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_ASAN_DG2 -// RUN: %clangxx -fsycl -fsycl-targets=spir64_gen --offload-new-driver %s --sysroot=%S/Inputs/SYCL \ -// RUN: -Xarch_device -fsanitize=address -Xs "-device dg2" -### 2>&1 \ -// RUN: | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_ASAN_DG2 +// RUN: %clangxx -fsycl -fsycl-targets=intel_gpu_dg2 --offload-new-driver %s --sysroot=%S/Inputs/SYCL \ +// RUN: -Xarch_device -fsanitize=address -### 2>&1 | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_ASAN_DG2 +// RUN: %clangxx -fsycl -fsycl-targets=intel_gpu_12_55_8 --offload-new-driver %s --sysroot=%S/Inputs/SYCL \ +// RUN: -Xarch_device -fsanitize=address -### 2>&1 | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_ASAN_DG2 // SYCL_DEVICE_LIB_ASAN_DG2: clang{{.*}} "-mlink-builtin-bitcode" "{{.*}}libsycl-crt.bc" // SYCL_DEVICE_LIB_ASAN_DG2-SAME: {{.*}}libsycl-cmath.bc // SYCL_DEVICE_LIB_ASAN_DG2-SAME: {{.*}}libsycl-imf.bc @@ -115,11 +97,8 @@ /// ########################################################################### /// test behavior of linking libsycl-asan for multiple targets AOT compilation /// when asan flag is applied. -// RUN: %clangxx -fsycl -fsycl-targets=spir64_gen --offload-new-driver %s --sysroot=%S/Inputs/SYCL \ -// RUN: -Xarch_device -fsanitize=address -Xsycl-target-backend "-device pvc,dg2" -### 2>&1 \ -// RUN: | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_ASAN_MUL -// RUN: %clangxx -fsycl -fsycl-targets=spir64_gen --offload-new-driver %s --sysroot=%S/Inputs/SYCL \ -// RUN: -Xarch_device -fsanitize=address -Xsycl-target-backend=spir64_gen "-device pvc,dg2" -### 2>&1 \ +// RUN: %clangxx -fsycl -fsycl-targets=intel_gpu_pvc,intel_gpu_dg2 --offload-new-driver %s \ +// RUN: --sysroot=%S/Inputs/SYCL -Xarch_device -fsanitize=address -### 2>&1 \ // RUN: | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_ASAN_MUL // SYCL_DEVICE_LIB_ASAN_MUL: clang{{.*}} "-mlink-builtin-bitcode" "{{.*}}libsycl-crt.bc" // SYCL_DEVICE_LIB_ASAN_MUL-SAME: {{.*}}libsycl-cmath.bc diff --git a/clang/test/Driver/sycl-device-post-link-opt.cpp b/clang/test/Driver/sycl-device-post-link-opt.cpp index cd42a8787b18b..1a54737f104bc 100644 --- a/clang/test/Driver/sycl-device-post-link-opt.cpp +++ b/clang/test/Driver/sycl-device-post-link-opt.cpp @@ -8,7 +8,7 @@ // RUN: %clangxx -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL -Xdevice-post-link=spir64_gen "foo" -### %s 2>&1 | \ // RUN: FileCheck %s -check-prefix CHECK-SINGLE-TARGET-UNUSED --implicit-check-not 'sycl-post-link{{.*}} "foo"' -// RUN: %clangxx -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL -fsycl-targets=spir64,spir64_gen -Xdevice-post-link=spir64_gen "foo" -Xdevice-post-link=spir64 "bar" -### %s 2>&1 | \ +// RUN: %clangxx -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL -fsycl-targets=spir64,intel_gpu_pvc -Xdevice-post-link=intel_gpu_pvc "foo" -Xdevice-post-link=spir64 "bar" -### %s 2>&1 | \ // RUN: FileCheck %s -check-prefix CHECK-MULTIPLE-TARGET --implicit-check-not 'sycl-post-link{{.*}} "foo" "bar"' // Each token is forwarded as its own --sycl-post-link-options occurrence. diff --git a/clang/test/Driver/sycl-device-traits-macros.cpp b/clang/test/Driver/sycl-device-traits-macros.cpp index 1dda50b94c383..377ded08e70fb 100644 --- a/clang/test/Driver/sycl-device-traits-macros.cpp +++ b/clang/test/Driver/sycl-device-traits-macros.cpp @@ -20,7 +20,7 @@ /// definition, one for host and one for each of the two devices. // RUN: %clang -fsycl --no-offload-new-driver -fsycl-targets=spir64,spir64_gen -### %s 2>&1 \ // RUN: | FileCheck -check-prefix=CHECK-SYCL-TARGETS %s -// RUN: %clang -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL -fsycl-targets=spir64,spir64_gen -### %s 2>&1 \ +// RUN: %clang -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL -fsycl-targets=spir64,intel_gpu_pvc -### %s 2>&1 \ // RUN: | FileCheck -check-prefix=CHECK-SYCL-TARGETS %s // CHECK-SYCL-TARGETS-COUNT-3: "-D__SYCL_ANY_DEVICE_HAS_ANY_ASPECT__=1" diff --git a/clang/test/Driver/sycl-foffload-fp32-prec-div.cpp b/clang/test/Driver/sycl-foffload-fp32-prec-div.cpp index 18710fbbd2921..e91c6d31dfffa 100644 --- a/clang/test/Driver/sycl-foffload-fp32-prec-div.cpp +++ b/clang/test/Driver/sycl-foffload-fp32-prec-div.cpp @@ -1,13 +1,13 @@ // Test SYCL -foffload-fp32-prec-div // RUN: %clang -### -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL \ -// RUN: -fsycl-targets=spir64_gen -foffload-fp32-prec-div %s 2>&1 \ +// RUN: -fsycl-targets=intel_gpu_pvc -foffload-fp32-prec-div %s 2>&1 \ // RUN: | FileCheck -check-prefix=AOT %s // RUN: %clang -### -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL \ // RUN: -foffload-fp32-prec-div %s 2>&1 \ // RUN: | FileCheck -check-prefix=JIT %s -// AOT: llvm-offload-binary{{.*}} "--image=file={{.*}}.bc,triple=spir64_gen-unknown-unknown,arch={{.*}},kind=sycl,compile-opts=-options -ze-fp32-correctly-rounded-divide-sqrt{{.*}}" +// AOT: llvm-offload-binary{{.*}} "--image=file={{.*}}.bc,triple=spir64_gen-unknown-unknown,arch={{.*}},kind=sycl,compile-opts={{.*}}-options -ze-fp32-correctly-rounded-divide-sqrt{{.*}}" // JIT: llvm-offload-binary{{.*}} "--image=file={{.*}}.bc,triple=spir64-unknown-unknown,arch={{.*}}compile-opts={{.*}}-foffload-fp32-prec-div" diff --git a/clang/test/Driver/sycl-foffload-fp32-prec-sqrt.cpp b/clang/test/Driver/sycl-foffload-fp32-prec-sqrt.cpp index e65476e049424..2d750ea875e2f 100644 --- a/clang/test/Driver/sycl-foffload-fp32-prec-sqrt.cpp +++ b/clang/test/Driver/sycl-foffload-fp32-prec-sqrt.cpp @@ -1,13 +1,13 @@ // Test SYCL -foffload-fp32-prec-sqrt // RUN: %clang -### -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL \ -// RUN: -fsycl-targets=spir64_gen -foffload-fp32-prec-sqrt %s 2>&1 \ +// RUN: -fsycl-targets=intel_gpu_pvc -foffload-fp32-prec-sqrt %s 2>&1 \ // RUN: | FileCheck -check-prefix=AOT %s // RUN: %clang -### -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL \ // RUN: -foffload-fp32-prec-sqrt %s 2>&1 \ // RUN: | FileCheck -check-prefix=JIT %s -// AOT: llvm-offload-binary{{.*}} "--image=file={{.*}}.bc,triple=spir64_gen-unknown-unknown,arch={{.*}},kind=sycl,compile-opts=-options -ze-fp32-correctly-rounded-divide-sqrt{{.*}}" +// AOT: llvm-offload-binary{{.*}} "--image=file={{.*}}.bc,triple=spir64_gen-unknown-unknown,arch={{.*}},kind=sycl,compile-opts={{.*}}-options -ze-fp32-correctly-rounded-divide-sqrt{{.*}}" // JIT: llvm-offload-binary{{.*}} "--image=file={{.*}}.bc,triple=spir64-unknown-unknown,arch={{.*}}compile-opts={{.*}}-foffload-fp32-prec-sqrt" diff --git a/clang/test/Driver/sycl-fsyclbin.cpp b/clang/test/Driver/sycl-fsyclbin.cpp index bec637662005e..2c88accf237f3 100644 --- a/clang/test/Driver/sycl-fsyclbin.cpp +++ b/clang/test/Driver/sycl-fsyclbin.cpp @@ -119,12 +119,12 @@ /// For AOT spir64_gen, -fsyclbin=input/object additionally implies /// -library-compilation in the backend options so symbols stay /// externally visible for runtime cross-image resolution. -// RUN: %clangxx -fsyclbin=object -fsycl-targets=spir64_gen \ +// RUN: %clangxx -fsyclbin=object -fsycl-targets=intel_gpu_pvc \ // RUN: --offload-new-driver --sysroot=%S/Inputs/SYCL %s -### 2>&1 \ // RUN: | FileCheck %s --check-prefix=CHECK_LIBCOMP_IMPLIED // CHECK_LIBCOMP_IMPLIED: -library-compilation -// RUN: %clangxx -fsyclbin=executable -fsycl-targets=spir64_gen \ +// RUN: %clangxx -fsyclbin=executable -fsycl-targets=intel_gpu_pvc \ // RUN: --offload-new-driver --sysroot=%S/Inputs/SYCL %s -### 2>&1 \ // RUN: | FileCheck %s --check-prefix=CHECK_LIBCOMP_NOT_IMPLIED // CHECK_LIBCOMP_NOT_IMPLIED-NOT: -library-compilation diff --git a/clang/test/Driver/sycl-ftarget-compile-fast.cpp b/clang/test/Driver/sycl-ftarget-compile-fast.cpp index 538831de60db1..f1558cd78d990 100644 --- a/clang/test/Driver/sycl-ftarget-compile-fast.cpp +++ b/clang/test/Driver/sycl-ftarget-compile-fast.cpp @@ -1,10 +1,10 @@ // Test -ftarget-compile-fast behaviors // RUN: %clang -### -target x86_64-unknown-linux-gnu -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL \ -// RUN: -fsycl-targets=spir64_gen -ftarget-compile-fast %s 2>&1 \ +// RUN: -fsycl-targets=intel_gpu_pvc -ftarget-compile-fast %s 2>&1 \ // RUN: | FileCheck -check-prefix=TARGET_COMPILE_FAST_GEN %s // RUN: %clang_cl -### --target=x86_64-pc-windows-msvc -fsycl --offload-new-driver /clang:--sysroot=%S/Inputs/SYCL \ -// RUN: -fsycl-targets=spir64_gen -ftarget-compile-fast %s 2>&1 \ +// RUN: -fsycl-targets=intel_gpu_pvc -ftarget-compile-fast %s 2>&1 \ // RUN: | FileCheck -check-prefix=TARGET_COMPILE_FAST_GEN %s // TARGET_COMPILE_FAST_GEN: llvm-offload-binary diff --git a/clang/test/Driver/sycl-ftarget-register-alloc-mode.cpp b/clang/test/Driver/sycl-ftarget-register-alloc-mode.cpp index 6dffb72de4443..ea3541b67fb33 100644 --- a/clang/test/Driver/sycl-ftarget-register-alloc-mode.cpp +++ b/clang/test/Driver/sycl-ftarget-register-alloc-mode.cpp @@ -1,47 +1,35 @@ // Test SYCL -ftarget-register-alloc-mode // RUN: %clang -### -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL \ -// RUN: -fsycl-targets=spir64_gen -ftarget-register-alloc-mode=pvc:auto %s 2>&1 \ +// RUN: -fsycl-targets=intel_gpu_pvc -ftarget-register-alloc-mode=pvc:auto %s 2>&1 \ // RUN: | FileCheck -check-prefix=AUTO_AOT %s -DDEVICE=pvc // RUN: %clang -### -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL \ -// RUN: -fsycl-targets=spir64_gen -ftarget-register-alloc-mode=pvc:large %s 2>&1 \ +// RUN: -fsycl-targets=intel_gpu_pvc -ftarget-register-alloc-mode=pvc:large %s 2>&1 \ // RUN: | FileCheck -check-prefix=LARGE_AOT %s // RUN: %clang -### -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL \ -// RUN: -fsycl-targets=spir64_gen -ftarget-register-alloc-mode=pvc:small %s 2>&1 \ +// RUN: -fsycl-targets=intel_gpu_pvc -ftarget-register-alloc-mode=pvc:small %s 2>&1 \ // RUN: | FileCheck -check-prefix=SMALL_AOT %s // RUN: %clang -### -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL \ -// RUN: -fsycl-targets=spir64_gen -ftarget-register-alloc-mode=pvc:default %s 2>&1 \ +// RUN: -fsycl-targets=intel_gpu_pvc -ftarget-register-alloc-mode=pvc:default %s 2>&1 \ // RUN: | FileCheck -check-prefix=DEFAULT_AOT %s // RUN: %clang -### -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL \ -// RUN: -fsycl-targets=spir64_gen -Xs "-device pvc" %s 2>&1 \ +// RUN: -fsycl-targets=intel_gpu_pvc %s 2>&1 \ // RUN: | FileCheck %if system-windows %{ -check-prefix=DEFAULT_AOT %} %else %{ -check-prefix=AUTO_AOT %} %s -DDEVICE=pvc // RUN: %clang -### -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL \ -// RUN: -fsycl-targets=spir64_gen -Xsycl-target-backend "-device pvc" %s 2>&1 \ +// RUN: -fsycl-targets=intel_gpu_12_60_7 %s 2>&1 \ // RUN: | FileCheck %if system-windows %{ -check-prefix=DEFAULT_AOT %} %else %{ -check-prefix=AUTO_AOT %} %s -DDEVICE=pvc // RUN: %clang -### -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL \ -// RUN: -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen "-device pvc" %s 2>&1 \ +// RUN: -fsycl-targets=intel_gpu_pvc,intel_gpu_mtl_s %s 2>&1 \ // RUN: | FileCheck %if system-windows %{ -check-prefix=DEFAULT_AOT %} %else %{ -check-prefix=AUTO_AOT %} %s -DDEVICE=pvc // RUN: %clang -### -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL \ -// RUN: -fsycl-targets=spir64_gen -Xs "-device 0x0BD5" %s 2>&1 \ -// RUN: | FileCheck %if system-windows %{ -check-prefix=DEFAULT_AOT %} %else %{ -check-prefix=AUTO_AOT %} %s -DDEVICE=pvc - -// RUN: %clang -### -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL \ -// RUN: -fsycl-targets=spir64_gen -Xs "-device 12.60.7" %s 2>&1 \ -// RUN: | FileCheck %if system-windows %{ -check-prefix=DEFAULT_AOT %} %else %{ -check-prefix=AUTO_AOT %} %s -DDEVICE=12.60.7 - -// RUN: %clang -### -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL \ -// RUN: -fsycl-targets=spir64_gen -Xs "-device pvc,mtl-s" %s 2>&1 \ -// RUN: | FileCheck %if system-windows %{ -check-prefix=DEFAULT_AOT %} %else %{ -check-prefix=AUTO_AOT %} %s -DDEVICE=pvc - -// RUN: %clang -### -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL \ -// RUN: -fsycl-targets=spir64_gen -ftarget-register-alloc-mode=pvc:small,pvc:large %s 2>&1 \ +// RUN: -fsycl-targets=intel_gpu_pvc -ftarget-register-alloc-mode=pvc:small,pvc:large %s 2>&1 \ // RUN: | FileCheck -check-prefix=MULTIPLE_ARGS_AOT %s // RUN: %clang -### -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL \ @@ -68,18 +56,18 @@ // RUN: | FileCheck -check-prefix=MULTIPLE_ARGS_JIT %s // RUN: not %clang -### -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL \ -// RUN: -fsycl-targets=spir64_gen -ftarget-register-alloc-mode=dg2:auto %s 2>&1 \ +// RUN: -fsycl-targets=intel_gpu_pvc -ftarget-register-alloc-mode=dg2:auto %s 2>&1 \ // RUN: | FileCheck -check-prefix=BAD_DEVICE %s // RUN: not %clang -### -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL \ -// RUN: -fsycl-targets=spir64_gen -ftarget-register-alloc-mode=pvc:superlarge %s 2>&1 \ +// RUN: -fsycl-targets=intel_gpu_pvc -ftarget-register-alloc-mode=pvc:superlarge %s 2>&1 \ // RUN: | FileCheck -check-prefix=BAD_MODE %s // RUN: not %clang -### -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL \ -// RUN: -fsycl-targets=spir64_gen -ftarget-register-alloc-mode=dg2:superlarge %s 2>&1 \ +// RUN: -fsycl-targets=intel_gpu_pvc -ftarget-register-alloc-mode=dg2:superlarge %s 2>&1 \ // RUN: | FileCheck -check-prefix=BAD_BOTH %s -// RUN: %clangxx -### -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL -fsycl-targets=spir64_gen -Xs "-device bdw" \ +// RUN: %clangxx -### -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL -fsycl-targets=intel_gpu_bdw \ // RUN: %s 2>&1 \ // RUN: | FileCheck -check-prefix=NO_PVC %s @@ -89,22 +77,18 @@ // RUNx: %s 2>&1 \ // RUNx: | FileCheck -check-prefix=NO_PVC %s -// RUN: %clangxx -### -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL -fsycl-targets=spir64_gen -Xs "-device pvc:mtl-s" \ -// RUN: %s 2>&1 \ -// RUN: | FileCheck -check-prefix=NO_PVC %s - // NO_PVC-NOT: -device_options // NO_PVC-NOT: -ze-opt-large-register-file // AUTO_AOT: llvm-offload-binary{{.*}} "--image=file={{.*}}.bc,triple=spir64_gen-unknown-unknown,arch={{.*}},kind=sycl,compile-opts=-device_options [[DEVICE]] -ze-intel-enable-auto-large-GRF-mode{{.*}}" -// LARGE_AOT: llvm-offload-binary{{.*}} "--image=file={{.*}}.bc,triple=spir64_gen-unknown-unknown,arch=generic,kind=sycl,compile-opts=-device_options pvc -ze-opt-large-register-file" +// LARGE_AOT: llvm-offload-binary{{.*}} "--image=file={{.*}}.bc,triple=spir64_gen-unknown-unknown,arch=pvc,kind=sycl,compile-opts=-device_options pvc -ze-opt-large-register-file" -// SMALL_AOT: llvm-offload-binary{{.*}} "--image=file={{.*}}.bc,triple=spir64_gen-unknown-unknown,arch=generic,kind=sycl,compile-opts=-device_options pvc -ze-intel-128-GRF-per-thread" +// SMALL_AOT: llvm-offload-binary{{.*}} "--image=file={{.*}}.bc,triple=spir64_gen-unknown-unknown,arch=pvc,kind=sycl,compile-opts=-device_options pvc -ze-intel-128-GRF-per-thread" // DEFAULT_AOT-NOT: -device_options -// MULTIPLE_ARGS_AOT: llvm-offload-binary{{.*}} "--image=file={{.*}}.bc,triple=spir64_gen-unknown-unknown,arch=generic,kind=sycl,compile-opts=-device_options pvc -ze-intel-128-GRF-per-thread -device_options pvc -ze-opt-large-register-file" +// MULTIPLE_ARGS_AOT: llvm-offload-binary{{.*}} "--image=file={{.*}}.bc,triple=spir64_gen-unknown-unknown,arch=pvc,kind=sycl,compile-opts=-device_options pvc -ze-intel-128-GRF-per-thread -device_options pvc -ze-opt-large-register-file" // AUTO_JIT: llvm-offload-binary{{.*}} "--image=file={{.*}}.bc,triple=spir64-unknown-unknown,arch={{.*}},kind=sycl,compile-opts=-ftarget-register-alloc-mode=pvc:-ze-intel-enable-auto-large-GRF-mode" diff --git a/clang/test/Driver/sycl-no-rdc-new-driver.cpp b/clang/test/Driver/sycl-no-rdc-new-driver.cpp index de576711b8ecd..ac8a825501cf3 100644 --- a/clang/test/Driver/sycl-no-rdc-new-driver.cpp +++ b/clang/test/Driver/sycl-no-rdc-new-driver.cpp @@ -28,7 +28,7 @@ // Test compilation step: with -c, -fno-sycl-rdc finalizes the device code of // the translation unit into a fat binary at compile time. -// RUN: %clang -### --offload-new-driver -Werror --target=x86_64-unknown-linux-gnu -fsycl -fsycl-targets=spir64_gen -fno-sycl-rdc --no-offloadlib -fno-sycl-instrument-device-code %t.cpp -c -o %t.o 2>&1 \ +// RUN: %clang -### --offload-new-driver -Werror --target=x86_64-unknown-linux-gnu -fsycl -fsycl-targets=intel_gpu_pvc -fno-sycl-rdc --no-offloadlib -fno-sycl-instrument-device-code %t.cpp -c -o %t.o 2>&1 \ // RUN: | FileCheck -check-prefix=CHK-COMPILE-STEP %s // CHK-COMPILE-STEP-NOT: argument unused during compilation diff --git a/clang/test/Driver/sycl-nordc-old-model-include-binary.cpp b/clang/test/Driver/sycl-nordc-old-model-include-binary.cpp index 879cd7e67069a..f2ee5b4000da0 100644 --- a/clang/test/Driver/sycl-nordc-old-model-include-binary.cpp +++ b/clang/test/Driver/sycl-nordc-old-model-include-binary.cpp @@ -10,6 +10,6 @@ /// The new offloading model embeds the finalized device image. // RUN: %clangxx -### --target=x86_64-unknown-linux-gnu -fsycl \ -// RUN: --offload-new-driver -fsycl-targets=spir64_gen -fno-sycl-rdc -c %s 2>&1 \ +// RUN: --offload-new-driver -fsycl-targets=intel_gpu_pvc -fno-sycl-rdc -c %s 2>&1 \ // RUN: | FileCheck -check-prefix=CHK-NEW %s // CHK-NEW: "-fsycl-is-host"{{.*}} "-foffload-include-binary" diff --git a/clang/test/Driver/sycl-ocloc-path.cpp b/clang/test/Driver/sycl-ocloc-path.cpp index c122fd0ee3612..1bc41ec9f739e 100644 --- a/clang/test/Driver/sycl-ocloc-path.cpp +++ b/clang/test/Driver/sycl-ocloc-path.cpp @@ -37,9 +37,6 @@ /// Check that --ocloc-path= is forwarded to the clang-linker-wrapper for the /// new offloading model. -// RUN: %clang -### -fsycl --offload-new-driver -fsycl-targets=spir64_gen \ -// RUN: --sysroot=%S/Inputs/SYCL --ocloc-path=/my/ocloc/dir %s 2>&1 \ -// RUN: | FileCheck -check-prefix=CHK-OCLOC-PATH-NEW %s // RUN: %clang -### -fsycl --offload-new-driver \ // RUN: -fsycl-targets=intel_gpu_pvc --sysroot=%S/Inputs/SYCL \ // RUN: --ocloc-path=/my/ocloc/dir %s 2>&1 \ @@ -97,7 +94,7 @@ // RUN: --ocloc-path="/my/ocloc dir/with spaces" %s 2>&1 \ // RUN: | FileCheck -check-prefix=CHK-OCLOC-PATH-SPACES-OLD %s // CHK-OCLOC-PATH-SPACES-OLD: "/my/ocloc dir/with spaces{{[/\\]+}}ocloc{{(\.exe)?}}" "-output" -// RUN: %clang -### -fsycl --offload-new-driver -fsycl-targets=spir64_gen \ +// RUN: %clang -### -fsycl --offload-new-driver -fsycl-targets=intel_gpu_pvc \ // RUN: --sysroot=%S/Inputs/SYCL --ocloc-path="/my/ocloc dir/with spaces" %s 2>&1 \ // RUN: | FileCheck -check-prefix=CHK-OCLOC-PATH-SPACES-NEW %s // CHK-OCLOC-PATH-SPACES-NEW: clang-linker-wrapper{{.*}} "--ocloc-path=/my/ocloc dir/with spaces" diff --git a/clang/test/Driver/sycl-offload-new-driver.cpp b/clang/test/Driver/sycl-offload-new-driver.cpp index 93fc6f65d7a63..5623d97506258 100644 --- a/clang/test/Driver/sycl-offload-new-driver.cpp +++ b/clang/test/Driver/sycl-offload-new-driver.cpp @@ -194,8 +194,8 @@ /// Test option passing behavior for clang-offload-wrapper options for AOT. // RUN: %clangxx --target=x86_64-unknown-linux-gnu -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL \ -// RUN: -fsycl-targets=spir64_gen,spir64_x86_64 \ -// RUN: -Xsycl-target-backend=spir64_gen -backend-gen-opt \ +// RUN: -fsycl-targets=intel_gpu_pvc,spir64_x86_64 \ +// RUN: -Xsycl-target-backend=intel_gpu_pvc -backend-gen-opt \ // RUN: -Xsycl-target-backend=spir64_x86_64 -backend-cpu-opt \ // RUN: -### %s 2>&1 \ // RUN: | FileCheck -check-prefix WRAPPER_OPTIONS_BACKEND_AOT %s @@ -216,8 +216,8 @@ // Clang's generic LTO mode reaches the wrapper, which drops it before // invoking ocloc or opencl-aot. // RUN: %clangxx --target=x86_64-unknown-linux-gnu -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL \ -// RUN: -fsycl-targets=spir64_gen,spir64_x86_64 -foffload-lto=full \ -// RUN: -Xsycl-target-linker=spir64_gen -gpu-link \ +// RUN: -fsycl-targets=intel_gpu_pvc,spir64_x86_64 -foffload-lto=full \ +// RUN: -Xsycl-target-linker=intel_gpu_pvc -gpu-link \ // RUN: -Xsycl-target-linker=spir64_x86_64 -cpu-link -### %s 2>&1 \ // RUN: | FileCheck -check-prefix WRAPPER_AOT_LTO %s // WRAPPER_AOT_LTO: clang-linker-wrapper{{.*}} "--device-compiler=spir64_gen-unknown-unknown=-flto=full" @@ -259,71 +259,37 @@ // RUN: | FileCheck -check-prefix CHECK_EMBED_IR %s // CHECK_EMBED_IR: clang-linker-wrapper{{.*}} "-sycl-embed-ir" -/// Verify the filename being passed to the packager does not contain commas -/// that are used in -device settings. -// RUN: %clangxx -fsycl -### -fsycl-targets=spir64_gen --offload-new-driver --sysroot=%S/Inputs/SYCL \ -// RUN: -Xsycl-target-backend=spir64_gen "-device pvc,bdw" %s 2>&1 \ -// RUN: | FileCheck -check-prefix COMMA_FILE %s -// COMMA_FILE: llvm-offload-binary{{.*}} "--image=file={{.*}}pvc@bdw{{.*}},triple=spir64_gen-unknown-unknown,arch=pvc,arch=bdw,kind=sycl,compile-opts=-device_options pvc -ze-intel-enable-auto-large-GRF-mode -device pvc,compile-opts=bdw" - -/// Verify the arch value for the packager is populated with different -/// scenarios for spir64_gen -// RUN: %clangxx -fsycl -### -fsycl-targets=spir64_gen --offload-new-driver --sysroot=%S/Inputs/SYCL \ -// RUN: -Xsycl-target-backend=spir64_gen "-device bdw" %s 2>&1 \ -// RUN: | FileCheck -check-prefix ARCH_CHECK %s -// RUN: %clangxx -fsycl -### -fsycl-targets=spir64_gen --offload-new-driver --sysroot=%S/Inputs/SYCL \ -// RUN: -Xsycl-target-backend "-device bdw" %s 2>&1 \ -// RUN: | FileCheck -check-prefix ARCH_CHECK %s +/// Verify the arch value for the packager is populated for intel_gpu_* // RUN: %clangxx -fsycl -### -fsycl-targets=intel_gpu_bdw \ // RUN: --offload-new-driver --sysroot=%S/Inputs/SYCL %s 2>&1 \ // RUN: | FileCheck -check-prefix ARCH_CHECK %s -// RUN: %clang_cl -fsycl -### -fsycl-targets=spir64_gen --offload-new-driver /clang:--sysroot=%S/Inputs/SYCL \ -// RUN: -Xsycl-target-backend=spir64_gen "-device bdw" %s 2>&1 \ -// RUN: | FileCheck -check-prefix ARCH_CHECK %s -// RUN: %clang_cl -fsycl -### -fsycl-targets=spir64_gen --offload-new-driver /clang:--sysroot=%S/Inputs/SYCL \ -// RUN: -Xsycl-target-backend "-device bdw" %s 2>&1 \ -// RUN: | FileCheck -check-prefix ARCH_CHECK %s // RUN: %clang_cl -fsycl -### -fsycl-targets=intel_gpu_bdw \ // RUN: --offload-new-driver /clang:--sysroot=%S/Inputs/SYCL %s 2>&1 \ // RUN: | FileCheck -check-prefix ARCH_CHECK %s // ARCH_CHECK: llvm-offload-binary{{.*}} "--image=file={{.*}}triple=spir64_gen-unknown-unknown,arch=bdw,kind=sycl{{.*}}" -// Verify when a comma-separated list of architectures is provided in -device, they are -// passed to llvm-offload-binary correctly. -// RUN: %clangxx -fsycl -### -fsycl-targets=spir64_gen --offload-new-driver --sysroot=%S/Inputs/SYCL \ -// RUN: -Xsycl-target-backend "-device pvc,bdw" %s 2>&1 \ -// RUN: | FileCheck -check-prefix MULTI_ARCH %s -// RUN: %clangxx -fsycl -### -fsycl-targets=spir64_gen --offload-new-driver --sysroot=%S/Inputs/SYCL \ -// RUN: -Xsycl-target-backend=spir64_gen "-device pvc,bdw" %s 2>&1 \ -// RUN: | FileCheck -check-prefix MULTI_ARCH %s -// RUN: %clangxx -fsycl -### -fsycl-targets=spir64_gen --offload-new-driver --sysroot=%S/Inputs/SYCL \ -// RUN: -Xs "-device pvc,bdw" %s 2>&1 \ -// RUN: | FileCheck -check-prefix MULTI_ARCH %s -// MULTI_ARCH: llvm-offload-binary{{.*}} "--image=file={{.*}}triple=spir64_gen-unknown-unknown,arch=pvc,arch=bdw,kind=sycl -// MULTI_ARCH-SAME: compile-opts=-device_options pvc -ze-intel-enable-auto-large-GRF-mode -device pvc,compile-opts=bdw" - // Verify for multiple targets with -Xsycl-target-backend= with commas in the values // are passed correctly to llvm-offload-binary. // RUN: %clangxx -fsycl -### --offload-new-driver --sysroot=%S/Inputs/SYCL -fno-sycl-libspirv \ -// RUN: -fsycl-targets=nvptx64-nvidia-cuda,amdgcn-amd-amdhsa,spir64_gen \ +// RUN: -fsycl-targets=nvptx64-nvidia-cuda,amdgcn-amd-amdhsa,intel_gpu_pvc \ // RUN: -Xsycl-target-backend=amdgcn-amd-amdhsa --offload-arch=gfx908,gfx1010 \ // RUN: -Xsycl-target-backend=nvptx64-nvidia-cuda --offload-arch=sm_86,sm_87,sm_89 \ -// RUN: -Xsycl-target-backend=spir64_gen "-device pvc,bdw" \ -// RUN: -Xsycl-target-linker=spir64_gen "-DFOO,BAR" \ +// RUN: -Xsycl-target-backend=intel_gpu_pvc "-FOO a,b" \ +// RUN: -Xsycl-target-linker=intel_gpu_pvc "-DFOO,BAR" \ // RUN: -nogpulib %s 2>&1 | FileCheck -check-prefix=MULTI_ARCH2 %s // MULTI_ARCH2: llvm-offload-binary{{.*}} "--image=file={{.*}}triple=amdgpu10.10-amd-amdhsa,arch=gfx1010,kind=sycl,compile-opts=--offload-arch=gfx908,compile-opts=gfx1010" // MULTI_ARCH2-SAME: "--image=file={{.*}}triple=amdgpu9.08-amd-amdhsa,arch=gfx908,kind=sycl,compile-opts=--offload-arch=gfx908,compile-opts=gfx1010" // MULTI_ARCH2-SAME: "--image=file={{.*}}triple=nvptx64-nvidia-cuda,arch=sm_86,kind=sycl,compile-opts=--offload-arch=sm_86,compile-opts=sm_87,compile-opts=sm_89" // MULTI_ARCH2-SAME: "--image=file={{.*}}triple=nvptx64-nvidia-cuda,arch=sm_87,kind=sycl,compile-opts=--offload-arch=sm_86,compile-opts=sm_87,compile-opts=sm_89" // MULTI_ARCH2-SAME: "--image=file={{.*}}triple=nvptx64-nvidia-cuda,arch=sm_89,kind=sycl,compile-opts=--offload-arch=sm_86,compile-opts=sm_87,compile-opts=sm_89" -// MULTI_ARCH2-SAME: "--image=file={{.*}}triple=spir64_gen-unknown-unknown,arch=pvc,arch=bdw,kind=sycl,compile-opts=-device_options pvc -ze-intel-enable-auto-large-GRF-mode -device pvc,compile-opts=bdw,link-opts=-DFOO,link-opts=BAR" +// MULTI_ARCH2-SAME: "--image=file={{.*}}triple=spir64_gen-unknown-unknown,arch=pvc,kind=sycl,compile-opts=-device_options pvc -ze-intel-enable-auto-large-GRF-mode -FOO a,compile-opts=b,link-opts=-DFOO,link-opts=BAR" // Verify that the driver correctly handles link-opt and compile-opt values with commas -// RUN: %clangxx -fsycl -### -fsycl-targets=spir64_gen --offload-new-driver --sysroot=%S/Inputs/SYCL \ -// RUN: -Xsycl-target-backend "-device bdw -FOO a,b" \ +// RUN: %clangxx -fsycl -### -fsycl-targets=intel_gpu_bdw --offload-new-driver --sysroot=%S/Inputs/SYCL \ +// RUN: -Xsycl-target-backend "-FOO a,b" \ // RUN: -Xsycl-target-linker "-BAR x,y" %s 2>&1 \ // RUN: | FileCheck -check-prefix COMMA_OPTS %s -// COMMA_OPTS: llvm-offload-binary{{.*}} "--image=file={{.*}}triple=spir64_gen-unknown-unknown,arch=bdw,kind=sycl,compile-opts=-device bdw -FOO a,compile-opts=b,link-opts=-BAR x,link-opts=y" +// COMMA_OPTS: llvm-offload-binary{{.*}} "--image=file={{.*}}triple=spir64_gen-unknown-unknown,arch=bdw,kind=sycl,compile-opts=-FOO a,compile-opts=b,link-opts=-BAR x,link-opts=y" /// Verify that --cuda-path is passed to clang-linker-wrapper for SYCL offload // RUN: %clangxx -fsycl -### -fsycl-targets=nvptx64-nvidia-cuda -fno-sycl-libspirv \ diff --git a/clang/test/Driver/sycl-offload-save-temps.cpp b/clang/test/Driver/sycl-offload-save-temps.cpp index e5b2ad0280067..e6d3d2c447e62 100644 --- a/clang/test/Driver/sycl-offload-save-temps.cpp +++ b/clang/test/Driver/sycl-offload-save-temps.cpp @@ -36,8 +36,6 @@ /// consumed by the host compile and passing it (an alias for --offload-targets=) /// triggers spurious OpenMP target triple validation on SYCL-specific names /// like 'intel_gpu_pvc', causing: error: OpenMP target is invalid: 'intel_gpu_pvc' -// RUN: %clangxx -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL -fsycl-targets=spir64_gen-unknown-unknown -target x86_64-unknown-linux-gnu -save-temps %s -### 2>&1 \ -// RUN: | FileCheck %s --check-prefix=CHK-NO-FSYCL-TARGETS-IN-HOST-BC // RUN: %clangxx -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL -fsycl-targets=intel_gpu_pvc -target x86_64-unknown-linux-gnu -save-temps %s -### 2>&1 \ // RUN: | FileCheck %s --check-prefix=CHK-NO-FSYCL-TARGETS-IN-HOST-BC /// Verify the host compile step emitting BC is present but without -fsycl-targets=. diff --git a/clang/test/Driver/sycl-offload-with-split.cpp b/clang/test/Driver/sycl-offload-with-split.cpp index 3a1f5440e1498..6fbad916af9b5 100644 --- a/clang/test/Driver/sycl-offload-with-split.cpp +++ b/clang/test/Driver/sycl-offload-with-split.cpp @@ -7,7 +7,7 @@ /// ########################################################################### /// Ahead of Time compilation for gen, cpu - tool invocation -// RUN: %clang -target x86_64-unknown-linux-gnu -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL -fno-sycl-instrument-device-code --no-offloadlib -fsycl-device-code-split -fsycl-targets=spir64_gen-unknown-unknown %s -### 2>&1 \ +// RUN: %clang -target x86_64-unknown-linux-gnu -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL -fno-sycl-instrument-device-code --no-offloadlib -fsycl-device-code-split -fsycl-targets=intel_gpu_pvc %s -### 2>&1 \ // RUN: | FileCheck %s -check-prefix=CHK-TOOLS-AOT // RUN: %clang -target x86_64-unknown-linux-gnu -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL -fno-sycl-instrument-device-code --no-offloadlib -fsycl-device-code-split -fsycl-targets=spir64_x86_64-unknown-unknown %s -### 2>&1 \ // RUN: | FileCheck %s -check-prefix=CHK-TOOLS-AOT diff --git a/clang/test/Driver/sycl-offload.cpp b/clang/test/Driver/sycl-offload.cpp index 30d401677f388..345a2deaf9a0f 100644 --- a/clang/test/Driver/sycl-offload.cpp +++ b/clang/test/Driver/sycl-offload.cpp @@ -419,8 +419,7 @@ // RUN: %clang -target x86_64-unknown-linux-gnu -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL \ // RUN: -fno-sycl-instrument-device-code --no-offloadlib \ -// RUN: -fsycl-targets=nvptx64-nvidia-cuda,spir64_gen \ -// RUN: -Xsycl-target-backend=spir64_gen "-device skl" \ +// RUN: -fsycl-targets=nvptx64-nvidia-cuda,intel_gpu_skl \ // RUN: -ccc-print-phases %s 2>&1 \ // RUN: | FileCheck -check-prefix=CHK-PHASE-MULTI-TARG-BOUND-ARCH2 %s // CHK-PHASE-MULTI-TARG-BOUND-ARCH2: 0: input, "[[INPUT:.+\.cpp]]", c++, (host-sycl) @@ -650,7 +649,7 @@ // RUN: touch %t_empty.o // RUN: %clangxx -### -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL -target x86_64-unknown-linux-gnu -fsycl-targets=spir64_x86_64 %t_empty.o %s 2>&1 \ // RUN: | FileCheck -check-prefix DEFAULT_LINK %s -// RUN: %clangxx -### -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL -target x86_64-unknown-linux-gnu -fsycl-targets=spir64_gen %t_empty.o %s 2>&1 \ +// RUN: %clangxx -### -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL -target x86_64-unknown-linux-gnu -fsycl-targets=intel_gpu_pvc %t_empty.o %s 2>&1 \ // RUN: | FileCheck -check-prefix DEFAULT_LINK %s // DEFAULT_LINK: clang-linker-wrapper{{.*}} @@ -668,7 +667,7 @@ // RUN: %clangxx -### -target x86_64-unknown-linux-gnu -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL -fsycl-targets=spir64_x86_64 -Xsycl-target-backend -DFOO %S/Inputs/SYCL/objlin64.o 2>&1 \ // RUN: | FileCheck -check-prefixes=SYCL_TARGET_OPT_AOT,SYCL_TARGET_OPT_CPU %s -// RUN: %clangxx -### -target x86_64-unknown-linux-gnu -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL -fsycl-targets=spir64_gen -Xsycl-target-backend -DFOO %S/Inputs/SYCL/objlin64.o 2>&1 \ +// RUN: %clangxx -### -target x86_64-unknown-linux-gnu -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL -fsycl-targets=intel_gpu_pvc -Xsycl-target-backend -DFOO %S/Inputs/SYCL/objlin64.o 2>&1 \ // RUN: | FileCheck -check-prefixes=SYCL_TARGET_OPT_AOT,SYCL_TARGET_OPT_GPU %s // SYCL_TARGET_OPT_AOT-NOT: error: cannot deduce implicit triple value for '-Xsycl-target-backend' // SYCL_TARGET_OPT_CPU: clang-linker-wrapper{{.*}} "--device-linker=sycl:spir64_x86_64-unknown-unknown=--opencl-aot-options=-DFOO" @@ -711,11 +710,11 @@ // CHECK_FSYCL_FP64_CONV_EMU-NOT: clang{{.*}} "-cc1" "-triple x86_64-unknown-linux-gnu" {{.*}} "-fsycl-fp64-conv-emu" // CHECK_FSYCL_FP64_CONV_EMU-DAG: clang{{.*}} "-cc1" "-triple" "spir64_gen{{.*}}" "-fsycl-fp64-conv-emu" // CHECK_FSYCL_FP64_CONV_EMU-DAG: llvm-offload-binary{{.*}} "--image=file={{.*}}.bc,triple=spir64_gen-unknown-unknown,arch=pvc,kind=sycl,compile-opts={{.*}}-options -ze-fp64-gen-conv-emu{{.*}}" -// RUN: %clang_cl -fsycl --offload-new-driver /clang:--sysroot=%S/Inputs/SYCL -fsycl-targets=spir64_gen-unknown-unknown %s -fsycl-fp64-conv-emu -### 2>&1 \ +// RUN: %clang_cl -fsycl --offload-new-driver /clang:--sysroot=%S/Inputs/SYCL -fsycl-targets=intel_gpu_pvc %s -fsycl-fp64-conv-emu -### 2>&1 \ // RUN: | FileCheck %s -check-prefixes=CHECK_FSYCL_FP64_CONV_EMU_WIN // CHECK_FSYCL_FP64_CONV_EMU_WIN-NOT: clang{{.*}} "-cc1" "-triple x86_64-unknown-linux-gnu" {{.*}} "-fsycl-fp64-conv-emu" // CHECK_FSYCL_FP64_CONV_EMU_WIN-DAG: clang{{.*}} "-cc1" "-triple" "spir64_gen{{.*}}" "-fsycl-fp64-conv-emu" -// CHECK_FSYCL_FP64_CONV_EMU_WIN-DAG: llvm-offload-binary{{.*}} "--image=file={{.*}}.bc,triple=spir64_gen-unknown-unknown,arch=generic,kind=sycl,compile-opts={{.*}}-options -ze-fp64-gen-conv-emu{{.*}}" +// CHECK_FSYCL_FP64_CONV_EMU_WIN-DAG: llvm-offload-binary{{.*}} "--image=file={{.*}}.bc,triple=spir64_gen-unknown-unknown,arch=pvc,kind=sycl,compile-opts={{.*}}-options -ze-fp64-gen-conv-emu{{.*}}" /// ########################################################################### diff --git a/clang/test/Driver/sycl-pch-create.cpp b/clang/test/Driver/sycl-pch-create.cpp index 36c4f199b3957..549675a18520f 100644 --- a/clang/test/Driver/sycl-pch-create.cpp +++ b/clang/test/Driver/sycl-pch-create.cpp @@ -41,7 +41,7 @@ // LX_CREATE_TARGETS: clang-offload-bundler{{.*}} "-type=pch" // LX_CREATE_TARGETS: "-targets=sycl-spir64{{.*}},sycl-spir64_gen{{.*}},host-x86_64{{.*}}" "-output={{.*}}.pch" "-input=[[PCHFILE1]]" "-input=[[PCHFILE2]]" "-input=[[PCHFILE3]]" -// RUN: %clang --offload-new-driver --no-offloadlib -fno-sycl-instrument-device-code -fsycl -x c++-header -fsycl-targets=spir64,spir64_gen -c %t.h -### 2>&1 | FileCheck -check-prefix=LX_CREATE_TARGETS_OND %s +// RUN: %clang --offload-new-driver --no-offloadlib -fno-sycl-instrument-device-code -fsycl -x c++-header -fsycl-targets=spir64,intel_gpu_pvc -c %t.h -### 2>&1 | FileCheck -check-prefix=LX_CREATE_TARGETS_OND %s // LX_CREATE_TARGETS_OND: clang{{.*}} "-triple" "spir64{{.*}}"{{.*}} "-fsycl-is-device" // LX_CREATE_TARGETS_OND-SAME: "-fsycl-int-header=[[HEADER1:.+\.h]]"{{.*}} "-fsycl-int-footer=[[FOOTER1:.+\.h]]"{{.*}} "-sycl-std=2020" // LX_CREATE_TARGETS_OND-SAME: "-o" "[[PCHFILE1:.+pch]]" @@ -114,7 +114,7 @@ // WS_CREATE_TARGETS: clang-offload-bundler{{.*}} "-type=pch" // WS_CREATE_TARGETS: "-targets=sycl-spir64{{.*}},sycl-spir64_gen{{.*}},host-x86_64{{.*}}" "-output={{.*}}.pch" "-input=[[PCHFILE1]]" "-input=[[PCHFILE2]]" "-input=[[PCHFILE3]]" -// RUN: %clang_cl --offload-new-driver --no-offloadlib -fno-sycl-instrument-device-code -fsycl -x c++-header -fsycl-targets=spir64,spir64_gen -c %t.h -### -- %s 2>&1 | FileCheck -check-prefix=WS_CREATE_TARGETS_OND %s +// RUN: %clang_cl --offload-new-driver --no-offloadlib -fno-sycl-instrument-device-code -fsycl -x c++-header -fsycl-targets=spir64,intel_gpu_pvc -c %t.h -### -- %s 2>&1 | FileCheck -check-prefix=WS_CREATE_TARGETS_OND %s // WS_CREATE_TARGETS_OND: clang{{.*}} "-triple" "spir64{{.*}}"{{.*}} "-fsycl-is-device" // WS_CREATE_TARGETS_OND-SAME: "-fsycl-int-header=[[HEADER1:.+\.h]]"{{.*}} "-fsycl-int-footer=[[FOOTER1:.+\.h]]"{{.*}} "-sycl-std=2020" // WS_CREATE_TARGETS_OND-SAME: "-o" "[[PCHFILE1:.+pch]]" diff --git a/clang/test/Driver/sycl-pch-use.cpp b/clang/test/Driver/sycl-pch-use.cpp index 05a570b83c72a..6ed078a1e20fb 100644 --- a/clang/test/Driver/sycl-pch-use.cpp +++ b/clang/test/Driver/sycl-pch-use.cpp @@ -33,7 +33,7 @@ // LX_USE_TARGETS: clang{{.*}} "-triple" "x86_64{{.*}}"{{.*}} "-fsycl-is-host" // LX_USE_TARGETS-SAME: "-include-pch" "[[PCHFILE3]]"{{.*}} -// RUN: %clang --offload-new-driver --no-offloadlib -fno-sycl-instrument-device-code -fsycl -c -include-pch %t.h.pch -fsycl-targets=spir64,spir64_gen %s -### 2>&1 | FileCheck -check-prefix=LX_USE_TARGETS_OND %s +// RUN: %clang --offload-new-driver --no-offloadlib -fno-sycl-instrument-device-code -fsycl -c -include-pch %t.h.pch -fsycl-targets=spir64,intel_gpu_pvc %s -### 2>&1 | FileCheck -check-prefix=LX_USE_TARGETS_OND %s // LX_USE_TARGETS_OND: llvm-offload-binary{{.*}} "[[MAINPCHFILE:.+\.pch]]" // LX_USE_TARGETS_OND-SAME: "--image=file=[[PCHFILE1:.+\.pch]],triple=x86_64{{.*}},arch=generic,kind=host" // LX_USE_TARGETS_OND: llvm-offload-binary{{.*}} "[[MAINPCHFILE:.+\.pch]]" @@ -76,7 +76,7 @@ // WS_USE_TARGETS: clang{{.*}} "-triple" "x86_64{{.*}}"{{.*}} "-fsycl-is-host" // WS_USE_TARGETS-SAME: "-include-pch" "[[PCHFILE3]]"{{.*}} -// RUN: %clang_cl --offload-new-driver --no-offloadlib -fno-sycl-instrument-device-code -fsycl -fsycl-targets=spir64,spir64_gen -c /Yu%t.h -### -- %s 2>&1 | FileCheck -check-prefix=WS_USE_TARGETS_OND %s +// RUN: %clang_cl --offload-new-driver --no-offloadlib -fno-sycl-instrument-device-code -fsycl -fsycl-targets=spir64,intel_gpu_pvc -c /Yu%t.h -### -- %s 2>&1 | FileCheck -check-prefix=WS_USE_TARGETS_OND %s // WS_USE_TARGETS_OND: llvm-offload-binary{{.*}} "[[MAINPCHFILE:.+\.pch]]" // WS_USE_TARGETS_OND-SAME: "--image=file=[[PCHFILE1:.+\.pch]],triple=x86_64{{.*}},arch=generic,kind=host" // WS_USE_TARGETS_OND: llvm-offload-binary{{.*}} "[[MAINPCHFILE]]" diff --git a/clang/test/Driver/sycl-spir64-gen-deprecated.cpp b/clang/test/Driver/sycl-spir64-gen-deprecated.cpp new file mode 100644 index 0000000000000..33fb9ecf5da9d --- /dev/null +++ b/clang/test/Driver/sycl-spir64-gen-deprecated.cpp @@ -0,0 +1,70 @@ +/// Test the diagnostics for spir64_gen used as a SYCL target, which is +/// deprecated in favor of intel_gpu_. + +/// Old offloading model: deprecation warning. +// RUN: %clangxx -### -fsycl --no-offload-new-driver -fsycl-targets=spir64_gen \ +// RUN: -Xsycl-target-backend "-device pvc" %s 2>&1 \ +// RUN: | FileCheck %s --check-prefix=WARN -DOPT=-fsycl-targets= +// RUN: %clangxx -### -fsycl --no-offload-new-driver \ +// RUN: -fsycl-targets=spir64_gen-unknown-unknown \ +// RUN: -Xsycl-target-backend "-device pvc" %s 2>&1 \ +// RUN: | FileCheck %s --check-prefix=WARN-TRIPLE +// RUN: %clangxx -### -fsycl --no-offload-new-driver \ +// RUN: --offload-targets=spir64_gen -Xsycl-target-backend "-device pvc" %s 2>&1 \ +// RUN: | FileCheck %s --check-prefix=WARN -DOPT=--offload-targets= +// RUN: %clangxx -### -fsycl --no-offload-new-driver \ +// RUN: -fsycl-targets=intel_gpu_pvc,spir64_gen \ +// RUN: -Xsycl-target-backend=spir64_gen "-device dg2" %s 2>&1 \ +// RUN: | FileCheck %s --check-prefixes=WARN,WARN-XS -DOPT=-fsycl-targets= +// WARN: warning: option '[[OPT]]spir64_gen' is deprecated and will be removed in a future release, use '[[OPT]]intel_gpu_' instead [-Wdeprecated] +// WARN-TRIPLE: warning: option '-fsycl-targets=spir64_gen-unknown-unknown' is deprecated and will be removed in a future release, use '-fsycl-targets=intel_gpu_' instead [-Wdeprecated] +// WARN-XS: warning: option '-Xsycl-target-backend=spir64_gen' is deprecated and will be removed in a future release, use '-Xsycl-target-backend=intel_gpu_' instead [-Wdeprecated] + +/// -Xsycl-target-* with a spir64_gen triple is diagnosed even when the target +/// is spelled intel_gpu_. +// RUN: %clangxx -### -fsycl --no-offload-new-driver -fsycl-targets=intel_gpu_pvc \ +// RUN: -Xsycl-target-backend=spir64_gen "-options -extra" \ +// RUN: -Xsycl-target-frontend=spir64_gen -DFOO \ +// RUN: -Xsycl-target-linker=spir64_gen "-foo" %s 2>&1 \ +// RUN: | FileCheck %s --check-prefix=WARN-X +// WARN-X-DAG: warning: option '-Xsycl-target-backend=spir64_gen' is deprecated and will be removed in a future release, use '-Xsycl-target-backend=intel_gpu_' instead [-Wdeprecated] +// WARN-X-DAG: warning: option '-Xsycl-target-frontend=spir64_gen' is deprecated and will be removed in a future release, use '-Xsycl-target-frontend=intel_gpu_' instead [-Wdeprecated] +// WARN-X-DAG: warning: option '-Xsycl-target-linker=spir64_gen' is deprecated and will be removed in a future release, use '-Xsycl-target-linker=intel_gpu_' instead [-Wdeprecated] + +/// The deprecation warning is controlled by -Wdeprecated. +// RUN: %clangxx -### -fsycl --no-offload-new-driver -fsycl-targets=spir64_gen \ +// RUN: -Xsycl-target-backend "-device pvc" -Wno-deprecated %s 2>&1 \ +// RUN: | FileCheck %s --check-prefix=NO-DIAG +// RUN: not %clangxx -### -fsycl --no-offload-new-driver -fsycl-targets=spir64_gen \ +// RUN: -Xsycl-target-backend "-device pvc" -Werror=deprecated %s 2>&1 \ +// RUN: | FileCheck %s --check-prefix=WERROR +// WERROR: error: option '-fsycl-targets=spir64_gen' is deprecated + +/// New offloading model: error. +// RUN: not %clangxx -### -fsycl --offload-new-driver -fsycl-targets=spir64_gen \ +// RUN: -Xsycl-target-backend "-device pvc" %s 2>&1 \ +// RUN: | FileCheck %s --check-prefix=ERR +// RUN: not %clangxx -### -fsycl --offload-new-driver \ +// RUN: -fsycl-targets=spir64_gen-unknown-unknown \ +// RUN: -Xsycl-target-backend "-device pvc" %s 2>&1 \ +// RUN: | FileCheck %s --check-prefix=ERR-TRIPLE +// RUN: not %clangxx -### -fsycl --offload-new-driver -fsycl-targets=intel_gpu_pvc \ +// RUN: -Xsycl-target-backend=spir64_gen "-options -extra" %s 2>&1 \ +// RUN: | FileCheck %s --check-prefix=ERR-XS +// ERR: error: '-fsycl-targets=spir64_gen' is not supported with the new offloading model; use '-fsycl-targets=intel_gpu_' instead +// ERR-TRIPLE: error: '-fsycl-targets=spir64_gen-unknown-unknown' is not supported with the new offloading model; use '-fsycl-targets=intel_gpu_' instead +// ERR-XS: error: '-Xsycl-target-backend=spir64_gen' is not supported with the new offloading model; use '-Xsycl-target-backend=intel_gpu_' instead + +/// The error cannot be downgraded with -Wno-deprecated. +// RUN: not %clangxx -### -fsycl --offload-new-driver -fsycl-targets=spir64_gen \ +// RUN: -Xsycl-target-backend "-device pvc" -Wno-deprecated %s 2>&1 \ +// RUN: | FileCheck %s --check-prefix=ERR + +/// intel_gpu_ is not diagnosed in either model. +// RUN: %clangxx -### -fsycl --no-offload-new-driver -fsycl-targets=intel_gpu_pvc \ +// RUN: -Xsycl-target-backend=intel_gpu_pvc "-options -extra" %s 2>&1 \ +// RUN: | FileCheck %s --check-prefix=NO-DIAG +// RUN: %clangxx -### -fsycl --offload-new-driver -fsycl-targets=intel_gpu_pvc \ +// RUN: -Xsycl-target-backend=intel_gpu_pvc "-options -extra" %s 2>&1 \ +// RUN: | FileCheck %s --check-prefix=NO-DIAG +// NO-DIAG-NOT: spir64_gen' is diff --git a/clang/test/Driver/sycl-spirv-opt.cpp b/clang/test/Driver/sycl-spirv-opt.cpp index 7d89f9c9fa829..6fb4741df4c53 100644 --- a/clang/test/Driver/sycl-spirv-opt.cpp +++ b/clang/test/Driver/sycl-spirv-opt.cpp @@ -8,7 +8,7 @@ // RUN: %clangxx -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL -Xspirv-translator=spir64_gen "foo" -### %s 2>&1 | \ // RUN: FileCheck %s -check-prefix CHECK-SINGLE-TARGET-UNUSED --implicit-check-not 'llvm-spirv{{.*}} "foo"' -// RUN: %clangxx -fsycl -fsycl-targets=spir64,spir64_gen --offload-new-driver --sysroot=%S/Inputs/SYCL -Xspirv-translator=spir64_gen "foo" -Xspirv-translator=spir64 "bar" -### %s 2>&1 | \ +// RUN: %clangxx -fsycl -fsycl-targets=spir64,intel_gpu_pvc --offload-new-driver --sysroot=%S/Inputs/SYCL -Xspirv-translator=intel_gpu_pvc "foo" -Xspirv-translator=spir64 "bar" -### %s 2>&1 | \ // RUN: FileCheck %s -check-prefix CHECK-MULTIPLE-TARGET // Each token is forwarded as its own --llvm-spirv-options occurrence. diff --git a/clang/test/Driver/sycl-spirv-to-ir-opt.cpp b/clang/test/Driver/sycl-spirv-to-ir-opt.cpp index f476d14f9ad22..5af098b824452 100644 --- a/clang/test/Driver/sycl-spirv-to-ir-opt.cpp +++ b/clang/test/Driver/sycl-spirv-to-ir-opt.cpp @@ -8,7 +8,7 @@ // RUN: %clangxx -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL -Xspirv-to-ir-wrapper=spir64_gen "foo" -### %s 2>&1 | \ // RUN: FileCheck %s -check-prefix CHECK-SINGLE-TARGET-UNUSED --implicit-check-not 'spirv-to-ir-wrapper-options{{.*}}=foo' -// RUN: %clangxx -fsycl -fsycl-targets=spir64,spir64_gen --offload-new-driver --sysroot=%S/Inputs/SYCL -Xspirv-to-ir-wrapper=spir64_gen "foo" -Xspirv-to-ir-wrapper=spir64 "bar" -### %s 2>&1 | \ +// RUN: %clangxx -fsycl -fsycl-targets=spir64,intel_gpu_pvc --offload-new-driver --sysroot=%S/Inputs/SYCL -Xspirv-to-ir-wrapper=intel_gpu_pvc "foo" -Xspirv-to-ir-wrapper=spir64 "bar" -### %s 2>&1 | \ // RUN: FileCheck %s -check-prefix CHECK-MULTIPLE-TARGET // Each token is forwarded as its own --spirv-to-ir-wrapper-options occurrence. diff --git a/clang/test/Driver/sycl-suppress-undefined-func-warnings.cpp b/clang/test/Driver/sycl-suppress-undefined-func-warnings.cpp index 2a85d3b4f95da..f59e3eb772e6f 100644 --- a/clang/test/Driver/sycl-suppress-undefined-func-warnings.cpp +++ b/clang/test/Driver/sycl-suppress-undefined-func-warnings.cpp @@ -64,7 +64,7 @@ // AOT target (spir64_gen): forwarding is target-independent. // RUN: %clangxx -### -fsycl -fsycl-targets=spir64_gen -Wno-sycl-undefined-func-in-image %s 2>&1 \ // RUN: | FileCheck --check-prefix=WNO %s -// RUN: %clangxx -### -fsycl --offload-new-driver -fsycl-targets=spir64_gen \ +// RUN: %clangxx -### -fsycl --offload-new-driver -fsycl-targets=intel_gpu_pvc \ // RUN: -fno-sycl-instrument-device-code --no-offloadlib \ // RUN: -Wno-sycl-undefined-func-in-image %s 2>&1 \ // RUN: | FileCheck --check-prefix=WNO-NEW %s diff --git a/clang/test/Driver/sycl-unique-prefix.cpp b/clang/test/Driver/sycl-unique-prefix.cpp index 8ec1e0c19429f..57416f5ed7b02 100644 --- a/clang/test/Driver/sycl-unique-prefix.cpp +++ b/clang/test/Driver/sycl-unique-prefix.cpp @@ -1,7 +1,7 @@ /// Test for Unique prefix setting for SYCL compilations // RUN: touch %t_file1.cpp // RUN: touch %t_file2.cpp -// RUN: %clangxx -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL -fsycl-targets=spir64-unknown-unknown,spir64_gen-unknown-unknown -c %t_file1.cpp %t_file2.cpp -### 2>&1 \ +// RUN: %clangxx -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL -fsycl-targets=spir64-unknown-unknown,intel_gpu_pvc -c %t_file1.cpp %t_file2.cpp -### 2>&1 \ // RUN: | FileCheck -check-prefix=CHECK_PREFIX %s // CHECK_PREFIX: clang{{.*}} "-triple" "spir64-unknown-unknown"{{.*}} "-fsycl-is-device"{{.*}} "-fsycl-unique-prefix=[[PREFIX1:uid([A-z0-9]){16}]]"{{.*}} "{{.*}}_file1.cpp" // CHECK_PREFIX: clang{{.*}} "-triple" "spir64_gen-unknown-unknown"{{.*}} "-fsycl-is-device"{{.*}} "-fsycl-unique-prefix=[[PREFIX1]]"{{.*}} "{{.*}}_file1.cpp" diff --git a/sycl/doc/GetStartedGuide.md b/sycl/doc/GetStartedGuide.md index 9c66d7981f5a6..baed5a54b4c79 100644 --- a/sycl/doc/GetStartedGuide.md +++ b/sycl/doc/GetStartedGuide.md @@ -730,7 +730,7 @@ specify the target architecture. The examples provided use a supported alias for the target, representing a full triple. Additional details can be found in the [Users Manual](UsersManual.md#generic-options). -```-fsycl-targets=spir64_gen``` for GPU, +```-fsycl-targets=intel_gpu_``` (e.g. ```intel_gpu_pvc```) for GPU, ```-fsycl-targets=spir64_x86_64``` for CPU. Multiple target architectures are supported. @@ -739,7 +739,7 @@ E.g., this command builds simple-sycl-app for GPU and CPU devices in ahead of time mode: ```bash -clang++ -fsycl -fsycl-targets=spir64_gen,spir64_x86_64 simple-sycl-app.cpp -o simple-sycl-app-aot.exe +clang++ -fsycl -fsycl-targets=intel_gpu_pvc,spir64_x86_64 simple-sycl-app.cpp -o simple-sycl-app-aot.exe ``` Additionally, user can pass specific options of AOT compiler to diff --git a/sycl/doc/UsersManual.md b/sycl/doc/UsersManual.md index f6035f35b67d6..22e5b26a73893 100644 --- a/sycl/doc/UsersManual.md +++ b/sycl/doc/UsersManual.md @@ -33,7 +33,10 @@ and not recommended to use in production environment. The following triples are supported by default: * spir64 - this is the default generic SPIR-V target; * spir64_x86_64 - generate code ahead of time for x86_64 CPUs; - * spir64_gen - generate code ahead of time for Intel Processor Graphics; + * spir64_gen - generate code ahead of time for Intel Processor Graphics. + Deprecated: use the intel_gpu_* values listed below instead. It emits a + deprecation warning with the old offloading model and is not supported + with the new offloading model (--offload-new-driver); Full target triples can also be used: * spir64-unknown-unknown, spir64_x86_64-unknown-unknown, spir64_gen-unknown-unknown @@ -45,7 +48,9 @@ and not recommended to use in production environment. currently overrides all the other specified SYCL targets when enabled.) Special target values specific to Intel, NVIDIA and AMD Processor Graphics - support are accepted, providing a streamlined interface for AOT. + support are accepted, providing a streamlined interface for AOT. The + intel_gpu_* values are the supported way to target Intel GPUs ahead of + time. A comma-separated list of valid Intel, NVIDIA and AMD Processor Graphics values is supported. * intel_gpu_cri, intel_gpu_35_11_0 - Crescent Island Intel graphics architecture * intel_gpu_nvl_p, intel_gpu_35_10_0 - Nova Lake P Intel graphics architecture @@ -478,7 +483,8 @@ and not recommended to use in production environment. Exposes exported symbols in a generated target library to allow for visibility to other modules. - NOTE: This flag is only supported for spir64_gen AOT targets. + NOTE: This flag is only supported for Intel GPU AOT targets + (intel_gpu_*). **`-ftarget-register-alloc-mode=`** diff --git a/sycl/test/check_device_code/optional_kernel_features/fp64-conv-emu-convert.cpp b/sycl/test/check_device_code/optional_kernel_features/fp64-conv-emu-convert.cpp index 5aeb711a8445f..bd28b6b155daa 100644 --- a/sycl/test/check_device_code/optional_kernel_features/fp64-conv-emu-convert.cpp +++ b/sycl/test/check_device_code/optional_kernel_features/fp64-conv-emu-convert.cpp @@ -1,6 +1,6 @@ // UNSUPPORTED: cuda, hip -// RUN: %clangxx -O0 -fsycl -fsycl-device-only -fsycl-targets=spir64_gen -fsycl-fp64-conv-emu %s -c -S -emit-llvm -o- | FileCheck %s +// RUN: %clangxx -O0 -fsycl -fsycl-device-only -fsycl-targets=intel_gpu_dg2 -fsycl-fp64-conv-emu %s -c -S -emit-llvm -o- | FileCheck %s // CHECK-NOT: define {{.*}} spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlvE_(){{.*}} !sycl_used_aspects diff --git a/sycl/test/check_device_code/optional_kernel_features/fp64-conv-emu-test-call-convert.cpp b/sycl/test/check_device_code/optional_kernel_features/fp64-conv-emu-test-call-convert.cpp index 8043dce78328e..fb17af53d343a 100644 --- a/sycl/test/check_device_code/optional_kernel_features/fp64-conv-emu-test-call-convert.cpp +++ b/sycl/test/check_device_code/optional_kernel_features/fp64-conv-emu-test-call-convert.cpp @@ -1,6 +1,6 @@ // UNSUPPORTED: cuda, hip -// RUN: %clangxx -O0 -fsycl -fsycl-device-only -fsycl-targets=spir64_gen -fsycl-fp64-conv-emu %s -c -S -emit-llvm -o- | FileCheck %s +// RUN: %clangxx -O0 -fsycl -fsycl-device-only -fsycl-targets=intel_gpu_dg2 -fsycl-fp64-conv-emu %s -c -S -emit-llvm -o- | FileCheck %s // CHECK-NOT: define {{.*}} spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlvE_(){{.*}} !sycl_used_aspects diff --git a/sycl/test/check_device_code/optional_kernel_features/fp64-conv-emu-test-call.cpp b/sycl/test/check_device_code/optional_kernel_features/fp64-conv-emu-test-call.cpp index ecdb761c17ec2..4f15b705897c4 100644 --- a/sycl/test/check_device_code/optional_kernel_features/fp64-conv-emu-test-call.cpp +++ b/sycl/test/check_device_code/optional_kernel_features/fp64-conv-emu-test-call.cpp @@ -1,6 +1,6 @@ // UNSUPPORTED: cuda, hip -// RUN: %clangxx -O0 -fsycl -fsycl-device-only -fsycl-targets=spir64_gen -fsycl-fp64-conv-emu %s -c -S -emit-llvm -o- | FileCheck %s +// RUN: %clangxx -O0 -fsycl -fsycl-device-only -fsycl-targets=intel_gpu_dg2 -fsycl-fp64-conv-emu %s -c -S -emit-llvm -o- | FileCheck %s // CHECK: define {{.*}} spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlvE_(){{.*}} !sycl_used_aspects ![[ASPECTFP64:[0-9]+]] // CHECK-DAG: define {{.*}} spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE0_clES2_EUlvE_(){{.*}} !sycl_used_aspects ![[ASPECTFP64]] diff --git a/sycl/test/check_device_code/optional_kernel_features/fp64-conv-emu-with-recursive-type-decls.cpp b/sycl/test/check_device_code/optional_kernel_features/fp64-conv-emu-with-recursive-type-decls.cpp index faa695d460e93..3f836b2911523 100644 --- a/sycl/test/check_device_code/optional_kernel_features/fp64-conv-emu-with-recursive-type-decls.cpp +++ b/sycl/test/check_device_code/optional_kernel_features/fp64-conv-emu-with-recursive-type-decls.cpp @@ -1,6 +1,6 @@ // UNSUPPORTED: cuda, hip -// RUN: %clangxx -O0 -fsycl -fsycl-device-only -fsycl-targets=spir64_gen -fsycl-fp64-conv-emu %s -c -S -emit-llvm -o- | FileCheck %s +// RUN: %clangxx -O0 -fsycl -fsycl-device-only -fsycl-targets=intel_gpu_dg2 -fsycl-fp64-conv-emu %s -c -S -emit-llvm -o- | FileCheck %s // CHECK-NOT: define {{.*}} spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlvE_(){{.*}} !sycl_used_aspects diff --git a/sycl/test/check_device_code/optional_kernel_features/fp64-conv-emu.cpp b/sycl/test/check_device_code/optional_kernel_features/fp64-conv-emu.cpp index 82a65559f4172..d681cba122319 100644 --- a/sycl/test/check_device_code/optional_kernel_features/fp64-conv-emu.cpp +++ b/sycl/test/check_device_code/optional_kernel_features/fp64-conv-emu.cpp @@ -1,6 +1,6 @@ // UNSUPPORTED: cuda, hip -// RUN: %clangxx -O0 -fsycl -fsycl-device-only -fsycl-targets=spir64_gen -fsycl-fp64-conv-emu %s -c -S -emit-llvm -o- | FileCheck %s +// RUN: %clangxx -O0 -fsycl -fsycl-device-only -fsycl-targets=intel_gpu_dg2 -fsycl-fp64-conv-emu %s -c -S -emit-llvm -o- | FileCheck %s // CHECK: define {{.*}} spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlvE_(){{.*}} !sycl_used_aspects ![[ASPECTFP64:[0-9]+]] // CHECK-DAG: ![[ASPECTFP64]] = !{![[TMP:[0-9]+]]} // CHECK-DAG: ![[TMP]] = !{!"fp64", i32 6} From 2a6c3a38a0c59f8bc7f8cccd343972c3b1053c0d Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Fri, 9 Oct 2026 14:34:36 -0700 Subject: [PATCH 03/22] [SYCL] Address review comments for spir64_gen deprecation - E2E: in build-only mode, default gpu_aot_target to the architectures of the CI run-only runners instead of only intel_gpu_pvc, since binaries built once are run on all of them. - Driver: also diagnose spir64_gen as the triple of -Xspirv-translator=, -Xdevice-post-link= and -Xspirv-to-ir-wrapper=. - Docs: describe how to migrate -device values that have no single intel_gpu_ equivalent. - Convert the remaining new offloading model spir64_gen RUN lines, including the Windows-only sycl-windows-device-filename.cpp. --- clang/lib/Driver/Driver.cpp | 6 ++++-- clang/test/Driver/sycl-device-post-link-opt.cpp | 4 ++-- clang/test/Driver/sycl-spir64-gen-deprecated.cpp | 16 ++++++++++++++++ clang/test/Driver/sycl-spirv-opt.cpp | 4 ++-- clang/test/Driver/sycl-spirv-to-ir-opt.cpp | 4 ++-- .../test/Driver/sycl-windows-device-filename.cpp | 4 ++-- sycl/doc/UsersManual.md | 11 ++++++++++- sycl/test-e2e/README.md | 8 +++++--- sycl/test-e2e/lit.cfg.py | 15 +++++++++++++-- 9 files changed, 56 insertions(+), 16 deletions(-) diff --git a/clang/lib/Driver/Driver.cpp b/clang/lib/Driver/Driver.cpp index c2dacfb09fa79..f463320f84fd8 100644 --- a/clang/lib/Driver/Driver.cpp +++ b/clang/lib/Driver/Driver.cpp @@ -1287,7 +1287,7 @@ static void diagnoseSYCLOptions(Compilation &C, bool IsSYCL) { } // spir64_gen is deprecated in favor of intel_gpu_, both as a - // -fsycl-targets value and as the triple of -Xsycl-target-*. It is not + // -fsycl-targets value and as the triple of the -X= options. It is not // supported with the new offloading model, and is removed together with the // old offloading model. if (IsSYCL) { @@ -1317,7 +1317,9 @@ static void diagnoseSYCLOptions(Compilation &C, bool IsSYCL) { } for (const Arg *A : C.getInputArgs().filtered( options::OPT_Xsycl_backend_EQ, options::OPT_Xsycl_frontend_EQ, - options::OPT_Xsycl_linker_EQ)) + options::OPT_Xsycl_linker_EQ, options::OPT_Xspirv_translator_EQ, + options::OPT_Xdevice_post_link_EQ, + options::OPT_Xspirv_to_ir_wrapper_EQ)) if (IsSPIRGenTarget(A->getValue())) DiagnoseSPIRGenTarget(A->getSpelling(), A->getValue()); } diff --git a/clang/test/Driver/sycl-device-post-link-opt.cpp b/clang/test/Driver/sycl-device-post-link-opt.cpp index 1a54737f104bc..c043f44488fa8 100644 --- a/clang/test/Driver/sycl-device-post-link-opt.cpp +++ b/clang/test/Driver/sycl-device-post-link-opt.cpp @@ -5,7 +5,7 @@ // RUN: %clangxx -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL -Xdevice-post-link "foo" -### %s 2>&1 | \ // RUN: FileCheck %s -check-prefix CHECK-SINGLE-TARGET -// RUN: %clangxx -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL -Xdevice-post-link=spir64_gen "foo" -### %s 2>&1 | \ +// RUN: %clangxx -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL -Xdevice-post-link=intel_gpu_pvc "foo" -### %s 2>&1 | \ // RUN: FileCheck %s -check-prefix CHECK-SINGLE-TARGET-UNUSED --implicit-check-not 'sycl-post-link{{.*}} "foo"' // RUN: %clangxx -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL -fsycl-targets=spir64,intel_gpu_pvc -Xdevice-post-link=intel_gpu_pvc "foo" -Xdevice-post-link=spir64 "bar" -### %s 2>&1 | \ @@ -20,7 +20,7 @@ // CHECK-SINGLE-TARGET: clang-linker-wrapper{{.*}} {{.*}}--sycl-post-link-options={{.*}}foo{{.*}} -// CHECK-SINGLE-TARGET-UNUSED: argument unused during compilation: '-Xdevice-post-link=spir64_gen foo' +// CHECK-SINGLE-TARGET-UNUSED: argument unused during compilation: '-Xdevice-post-link=intel_gpu_pvc foo' // CHECK-MULTIPLE-TARGET: clang-linker-wrapper{{.*}} {{.*}}"--sycl-post-link-options=sycl:spir64-unknown-unknown=bar"{{.*}}"--sycl-post-link-options=sycl:spir64_gen-unknown-unknown=foo" diff --git a/clang/test/Driver/sycl-spir64-gen-deprecated.cpp b/clang/test/Driver/sycl-spir64-gen-deprecated.cpp index 33fb9ecf5da9d..4d1152874e2e1 100644 --- a/clang/test/Driver/sycl-spir64-gen-deprecated.cpp +++ b/clang/test/Driver/sycl-spir64-gen-deprecated.cpp @@ -55,6 +55,22 @@ // ERR-TRIPLE: error: '-fsycl-targets=spir64_gen-unknown-unknown' is not supported with the new offloading model; use '-fsycl-targets=intel_gpu_' instead // ERR-XS: error: '-Xsycl-target-backend=spir64_gen' is not supported with the new offloading model; use '-Xsycl-target-backend=intel_gpu_' instead +/// The other triple-scoped tool options are diagnosed the same way. +// RUN: not %clangxx -### -fsycl --offload-new-driver -fsycl-targets=intel_gpu_pvc \ +// RUN: -Xspirv-translator=spir64_gen "-foo" -Xdevice-post-link=spir64_gen "-foo" \ +// RUN: -Xspirv-to-ir-wrapper=spir64_gen "-foo" %s 2>&1 \ +// RUN: | FileCheck %s --check-prefix=ERR-XTOOL +// RUN: %clangxx -### -fsycl --no-offload-new-driver -fsycl-targets=intel_gpu_pvc \ +// RUN: -Xspirv-translator=spir64_gen "-foo" -Xdevice-post-link=spir64_gen "-foo" \ +// RUN: -Xspirv-to-ir-wrapper=spir64_gen "-foo" %s 2>&1 \ +// RUN: | FileCheck %s --check-prefix=WARN-XTOOL +// ERR-XTOOL-DAG: error: '-Xspirv-translator=spir64_gen' is not supported with the new offloading model; use '-Xspirv-translator=intel_gpu_' instead +// ERR-XTOOL-DAG: error: '-Xdevice-post-link=spir64_gen' is not supported with the new offloading model; use '-Xdevice-post-link=intel_gpu_' instead +// ERR-XTOOL-DAG: error: '-Xspirv-to-ir-wrapper=spir64_gen' is not supported with the new offloading model; use '-Xspirv-to-ir-wrapper=intel_gpu_' instead +// WARN-XTOOL-DAG: warning: option '-Xspirv-translator=spir64_gen' is deprecated and will be removed in a future release, use '-Xspirv-translator=intel_gpu_' instead [-Wdeprecated] +// WARN-XTOOL-DAG: warning: option '-Xdevice-post-link=spir64_gen' is deprecated and will be removed in a future release, use '-Xdevice-post-link=intel_gpu_' instead [-Wdeprecated] +// WARN-XTOOL-DAG: warning: option '-Xspirv-to-ir-wrapper=spir64_gen' is deprecated and will be removed in a future release, use '-Xspirv-to-ir-wrapper=intel_gpu_' instead [-Wdeprecated] + /// The error cannot be downgraded with -Wno-deprecated. // RUN: not %clangxx -### -fsycl --offload-new-driver -fsycl-targets=spir64_gen \ // RUN: -Xsycl-target-backend "-device pvc" -Wno-deprecated %s 2>&1 \ diff --git a/clang/test/Driver/sycl-spirv-opt.cpp b/clang/test/Driver/sycl-spirv-opt.cpp index 6fb4741df4c53..7ab9869acdde6 100644 --- a/clang/test/Driver/sycl-spirv-opt.cpp +++ b/clang/test/Driver/sycl-spirv-opt.cpp @@ -5,7 +5,7 @@ // RUN: %clangxx -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL -Xspirv-translator "foo" -### %s 2>&1 | \ // RUN: FileCheck %s -check-prefix CHECK-SINGLE-TARGET -// RUN: %clangxx -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL -Xspirv-translator=spir64_gen "foo" -### %s 2>&1 | \ +// RUN: %clangxx -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL -Xspirv-translator=intel_gpu_pvc "foo" -### %s 2>&1 | \ // RUN: FileCheck %s -check-prefix CHECK-SINGLE-TARGET-UNUSED --implicit-check-not 'llvm-spirv{{.*}} "foo"' // RUN: %clangxx -fsycl -fsycl-targets=spir64,intel_gpu_pvc --offload-new-driver --sysroot=%S/Inputs/SYCL -Xspirv-translator=intel_gpu_pvc "foo" -Xspirv-translator=spir64 "bar" -### %s 2>&1 | \ @@ -19,7 +19,7 @@ // RUN: FileCheck %s -check-prefix CHECK-SPACE // CHECK-SINGLE-TARGET: clang-linker-wrapper{{.*}} "--llvm-spirv-options=sycl:spir64-unknown-unknown=foo{{.*}} -// CHECK-SINGLE-TARGET-UNUSED: argument unused during compilation: '-Xspirv-translator=spir64_gen foo' +// CHECK-SINGLE-TARGET-UNUSED: argument unused during compilation: '-Xspirv-translator=intel_gpu_pvc foo' // CHECK-MULTIPLE-TARGET: clang-linker-wrapper{{.*}} "--llvm-spirv-options=sycl:spir64-unknown-unknown=bar"{{.*}}"--llvm-spirv-options=sycl:spir64_gen-unknown-unknown=foo" // CHECK-MULTIPLE-TOKENS: clang-linker-wrapper{{.*}} "--llvm-spirv-options=sycl:spir64-unknown-unknown=foo"{{.*}}"--llvm-spirv-options=sycl:spir64-unknown-unknown=bar" // CHECK-SPACE: clang-linker-wrapper{{.*}} "--llvm-spirv-options=sycl:spir64-unknown-unknown=foo bar" diff --git a/clang/test/Driver/sycl-spirv-to-ir-opt.cpp b/clang/test/Driver/sycl-spirv-to-ir-opt.cpp index 5af098b824452..a5a6b4bfea2ab 100644 --- a/clang/test/Driver/sycl-spirv-to-ir-opt.cpp +++ b/clang/test/Driver/sycl-spirv-to-ir-opt.cpp @@ -5,7 +5,7 @@ // RUN: %clangxx -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL -Xspirv-to-ir-wrapper "foo" -### %s 2>&1 | \ // RUN: FileCheck %s -check-prefix CHECK-SINGLE-TARGET -// RUN: %clangxx -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL -Xspirv-to-ir-wrapper=spir64_gen "foo" -### %s 2>&1 | \ +// RUN: %clangxx -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL -Xspirv-to-ir-wrapper=intel_gpu_pvc "foo" -### %s 2>&1 | \ // RUN: FileCheck %s -check-prefix CHECK-SINGLE-TARGET-UNUSED --implicit-check-not 'spirv-to-ir-wrapper-options{{.*}}=foo' // RUN: %clangxx -fsycl -fsycl-targets=spir64,intel_gpu_pvc --offload-new-driver --sysroot=%S/Inputs/SYCL -Xspirv-to-ir-wrapper=intel_gpu_pvc "foo" -Xspirv-to-ir-wrapper=spir64 "bar" -### %s 2>&1 | \ @@ -19,7 +19,7 @@ // RUN: FileCheck %s -check-prefix CHECK-SPACE // CHECK-SINGLE-TARGET: clang-linker-wrapper{{.*}} "--spirv-to-ir-wrapper-options=sycl:spir64-unknown-unknown=foo{{.*}} -// CHECK-SINGLE-TARGET-UNUSED: argument unused during compilation: '-Xspirv-to-ir-wrapper=spir64_gen foo' +// CHECK-SINGLE-TARGET-UNUSED: argument unused during compilation: '-Xspirv-to-ir-wrapper=intel_gpu_pvc foo' // CHECK-MULTIPLE-TARGET: clang-linker-wrapper{{.*}} "--spirv-to-ir-wrapper-options=sycl:spir64-unknown-unknown=bar"{{.*}}"--spirv-to-ir-wrapper-options=sycl:spir64_gen-unknown-unknown=foo" // CHECK-MULTIPLE-TOKENS: clang-linker-wrapper{{.*}} "--spirv-to-ir-wrapper-options=sycl:spir64-unknown-unknown=foo"{{.*}}"--spirv-to-ir-wrapper-options=sycl:spir64-unknown-unknown=bar" // CHECK-SPACE: clang-linker-wrapper{{.*}} "--spirv-to-ir-wrapper-options=sycl:spir64-unknown-unknown=foo bar" diff --git a/clang/test/Driver/sycl-windows-device-filename.cpp b/clang/test/Driver/sycl-windows-device-filename.cpp index af6550bf1a547..f54c043b88463 100644 --- a/clang/test/Driver/sycl-windows-device-filename.cpp +++ b/clang/test/Driver/sycl-windows-device-filename.cpp @@ -2,7 +2,7 @@ // REQUIRES: system-windows // RUN: %clang -### --target=x86_64-pc-windows-msvc -fsycl \ -// RUN: -fsycl-targets=spir64_gen --offload-new-driver \ +// RUN: -fsycl-targets=intel_gpu_pvc --offload-new-driver \ // RUN: --sysroot=%S/Inputs/SYCL \ // RUN: -Xsycl-target-backend "-device arch1:arch2" %s 2>&1 \ // RUN: | FileCheck %s -check-prefix=CHECK_COLON @@ -12,7 +12,7 @@ // CHECK_COLON-NOT: sycl-windows-device-filename-arch1:arch2 // RUN: %clang -### --target=x86_64-pc-windows-msvc -fsycl \ -// RUN: -fsycl-targets=spir64_gen --offload-new-driver \ +// RUN: -fsycl-targets=intel_gpu_pvc --offload-new-driver \ // RUN: --sysroot=%S/Inputs/SYCL \ // RUN: -Xsycl-target-backend "-device *" %s 2>&1 \ // RUN: | FileCheck %s -check-prefix=CHECK_STAR diff --git a/sycl/doc/UsersManual.md b/sycl/doc/UsersManual.md index 22e5b26a73893..61e66834174ca 100644 --- a/sycl/doc/UsersManual.md +++ b/sycl/doc/UsersManual.md @@ -36,7 +36,16 @@ and not recommended to use in production environment. * spir64_gen - generate code ahead of time for Intel Processor Graphics. Deprecated: use the intel_gpu_* values listed below instead. It emits a deprecation warning with the old offloading model and is not supported - with the new offloading model (--offload-new-driver); + with the new offloading model (--offload-new-driver). To migrate + '-fsycl-targets=spir64_gen -Xsycl-target-backend "-device "': + * a device name, e.g. "-device pvc": use -fsycl-targets=intel_gpu_pvc; + * a full IP version, e.g. "-device 12.60.7": use + -fsycl-targets=intel_gpu_12_60_7; + * a list, e.g. "-device pvc,dg2": use + -fsycl-targets=intel_gpu_pvc,intel_gpu_dg2; + * "-device *", hex device IDs (e.g. 0x0BD5), partial versions (e.g. + 12.60) and devices without an intel_gpu_* value (e.g. pvc-sdv) have + no equivalent; list the intel_gpu_* values of the devices to target; Full target triples can also be used: * spir64-unknown-unknown, spir64_x86_64-unknown-unknown, spir64_gen-unknown-unknown diff --git a/sycl/test-e2e/README.md b/sycl/test-e2e/README.md index 0493850694ac3..42d0f2cb67e74 100644 --- a/sycl/test-e2e/README.md +++ b/sycl/test-e2e/README.md @@ -210,7 +210,9 @@ or via the ***LIT_OPTS*** environment variable. ***GPU_AOT_TARGET*** - defines the Intel GPU target(s) used for AOT compilation, as a `-fsycl-targets` value (e.g. "intel_gpu_pvc" or "intel_gpu_pvc,intel_gpu_bmg_g21"). If not specified, the architectures of the -detected Intel GPU devices are used, or "intel_gpu_pvc" if none are detected. +detected Intel GPU devices are used. If none are detected, the architectures of +the CI run-only runners are used, so that binaries built in `build-only` mode +run on all of them. ***OFFLOAD_BUILD_TARGET*** - when testing the Offload backend, this must be set to specify the correct build target type for the available Offload device. @@ -297,8 +299,8 @@ configure specific single test execution in the command line: compilation, as a `-fsycl-targets` value. Tests use it via the `%{gpu_aot_target}` substitution. It can be also set by CMake variable GPU_AOT_TARGET. If not specified, the architectures of the detected Intel GPU - devices are used, or "intel_gpu_pvc" if none are detected (e.g. in - `build-only` mode). + devices are used. If none are detected (e.g. in `build-only` mode), the + architectures of the CI run-only runners are used. * **gpu-intel-dg2** - tells LIT infra that Intel GPU DG2 is present in the system. It is developer / CI infra responsibility to make sure that the device is available in the system. diff --git a/sycl/test-e2e/lit.cfg.py b/sycl/test-e2e/lit.cfg.py index 5a36a3e81de7b..3dd96be5e35e5 100644 --- a/sycl/test-e2e/lit.cfg.py +++ b/sycl/test-e2e/lit.cfg.py @@ -1303,14 +1303,25 @@ def get_sycl_ls_verbose(sycl_device, env): config.substitutions.append(("%amd_arch_options", amd_arch_flags)) # Intel GPU AOT target(s), as a -fsycl-targets value. Uses the detected Intel -# GPU architectures or defaults to intel_gpu_pvc in build-only mode. +# GPU architectures. Without any (e.g. build-only mode), it uses the +# architectures of the CI run-only runners, since binaries built once are run +# on all of them. +CI_RUNNER_GPU_AOT_TARGETS = ",".join( + [ + "intel_gpu_tgllp", + "intel_gpu_acm_g10", + "intel_gpu_pvc", + "intel_gpu_bmg_g21", + "intel_gpu_mtl_u", + ] +) if not config.gpu_aot_target: intel_gpu_archs = sorted( feature[len("arch-") :] for feature in set().union(*config.sycl_dev_features.values()) if feature.startswith("arch-intel_gpu_") ) - config.gpu_aot_target = ",".join(intel_gpu_archs) or "intel_gpu_pvc" + config.gpu_aot_target = ",".join(intel_gpu_archs) or CI_RUNNER_GPU_AOT_TARGETS lit_config.note("Intel GPU AOT target: {}".format(config.gpu_aot_target)) config.substitutions.append(("%{gpu_aot_target}", config.gpu_aot_target)) From 8c7322317ba6a3473f53b3cf80f29b42b055266a Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Fri, 9 Oct 2026 14:53:30 -0700 Subject: [PATCH 04/22] [SYCL][E2E] Fix comment formatting in aot_object_no_symbols.cpp --- sycl/test-e2e/SYCLBIN/aot_object_no_symbols.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/sycl/test-e2e/SYCLBIN/aot_object_no_symbols.cpp b/sycl/test-e2e/SYCLBIN/aot_object_no_symbols.cpp index 21a70c84843dd..224df2e222ed9 100644 --- a/sycl/test-e2e/SYCLBIN/aot_object_no_symbols.cpp +++ b/sycl/test-e2e/SYCLBIN/aot_object_no_symbols.cpp @@ -31,7 +31,8 @@ // -- dynamicLink) to reach executable state. // // -- Targeting the configured Intel GPU architecture(s) (%{gpu_aot_target}) -// -- keeps this off a single hard-coded one. Complements the hardware-independent unit test +// -- keeps this off a single hard-coded one. Complements the +// -- hardware-independent unit test // -- SYCLBINSelector.NativeWithoutSymbolsSurfacedForObject, which covers the // -- selector in isolation but neither the producer nor the link-and-run path. From 905a197300d22333cec123ed3fce997bd99f6a87 Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Mon, 12 Oct 2026 11:41:22 -0700 Subject: [PATCH 05/22] [SYCL][E2E] Exclude Gen12 from the Windows build-only gpu_aot_target default The Windows ocloc used in CI no longer supports Gen12 and older GPUs, so building for intel_gpu_tgllp fails there. Windows run-only CI only uses DG2 and BMG runners, so build for those on Windows. Linux still supports Gen12 and keeps intel_gpu_tgllp for its gen12 runner. --- sycl/test-e2e/lit.cfg.py | 22 +++++++++++++--------- 1 file changed, 13 insertions(+), 9 deletions(-) diff --git a/sycl/test-e2e/lit.cfg.py b/sycl/test-e2e/lit.cfg.py index 3dd96be5e35e5..2877945e87bd2 100644 --- a/sycl/test-e2e/lit.cfg.py +++ b/sycl/test-e2e/lit.cfg.py @@ -1306,15 +1306,19 @@ def get_sycl_ls_verbose(sycl_device, env): # GPU architectures. Without any (e.g. build-only mode), it uses the # architectures of the CI run-only runners, since binaries built once are run # on all of them. -CI_RUNNER_GPU_AOT_TARGETS = ",".join( - [ - "intel_gpu_tgllp", - "intel_gpu_acm_g10", - "intel_gpu_pvc", - "intel_gpu_bmg_g21", - "intel_gpu_mtl_u", - ] -) +if platform.system() == "Windows": + # Windows ocloc no longer supports Gen12 and older GPUs, such as tgllp. + CI_RUNNER_GPU_AOT_TARGETS = "intel_gpu_acm_g10,intel_gpu_bmg_g21" +else: + CI_RUNNER_GPU_AOT_TARGETS = ",".join( + [ + "intel_gpu_tgllp", + "intel_gpu_acm_g10", + "intel_gpu_pvc", + "intel_gpu_bmg_g21", + "intel_gpu_mtl_u", + ] + ) if not config.gpu_aot_target: intel_gpu_archs = sorted( feature[len("arch-") :] From 10bb403e014d451459d422dea1c8c47a70c083ed Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Mon, 12 Oct 2026 12:20:59 -0700 Subject: [PATCH 06/22] [Clang][OffloadBundler] Match the target ID of SYCL bundle entry IDs SYCL bundle entry IDs use a 3-component triple followed by the target ID, e.g. sycl-spir64_gen-unknown-unknown-pvc. The bundler parsed the target ID as the triple environment and left the target ID empty, so all entries with the same triple matched each other. With several -fsycl-targets=intel_gpu_* (or nvidia_gpu_*) targets and separate compile and link in the old offloading model, every target was then unbundled to the first entry's device code. Take the last component as the target ID when it is not a known environment but is a known offload architecture. A SYCL code object without a target ID stays compatible with any target ID of the same triple, as before. --- clang/lib/Driver/OffloadBundler.cpp | 26 ++++++++ .../clang-offload-bundler/sycl-target-id.c | 60 +++++++++++++++++++ 2 files changed, 86 insertions(+) create mode 100644 clang/test/OffloadTools/clang-offload-bundler/sycl-target-id.c diff --git a/clang/lib/Driver/OffloadBundler.cpp b/clang/lib/Driver/OffloadBundler.cpp index eca7840113674..644af33ab03df 100644 --- a/clang/lib/Driver/OffloadBundler.cpp +++ b/clang/lib/Driver/OffloadBundler.cpp @@ -134,6 +134,20 @@ OffloadTargetInfo::OffloadTargetInfo(const StringRef Target, this->OffloadKind = Components.front(); ArrayRef TripleSlice{&Components[1], /*length=*/4}; llvm::Triple T = llvm::Triple(llvm::join(TripleSlice, "-")); + + // SYCL bundle entry IDs use a 3-component triple followed by the target ID, + // e.g. sycl-spir64_gen-unknown-unknown-pvc, so the last component is the + // target ID rather than the environment. + if (Components.size() == 5 && + T.getEnvironment() == llvm::Triple::UnknownEnvironment && + !clang::StringToOffloadArch(Components.back().split(':').first) + .isUnknown()) { + this->TargetID = Components.back(); + this->Triple = + llvm::Triple(T.getArchName(), T.getVendorName(), T.getOSName()); + return; + } + this->Triple = llvm::Triple(T.getArchName(), T.getVendorName(), T.getOSName(), T.getEnvironmentName()); } @@ -1714,6 +1728,18 @@ bool isCodeObjectCompatible(const OffloadTargetInfo &CodeObjectInfo, return false; } + // SYCL device code without a target ID (e.g. -fsycl-targets=spir64_gen) + // stays compatible with any target ID of the same triple. + if (CodeObjectInfo.OffloadKind == "sycl" && + (CodeObjectInfo.TargetID.empty() || TargetInfo.TargetID.empty())) { + DEBUG_WITH_TYPE("CodeObjectCompatibility", + dbgs() << "Compatible: SYCL code object without target ID " + "\t[CodeObject: " + << CodeObjectInfo.str() + << "]\t:\t[Target: " << TargetInfo.str() << "]\n"); + return true; + } + // Incompatible if Processors mismatch. llvm::StringMap CodeObjectFeatureMap, TargetFeatureMap; std::optional CodeObjectProc = clang::parseTargetID( diff --git a/clang/test/OffloadTools/clang-offload-bundler/sycl-target-id.c b/clang/test/OffloadTools/clang-offload-bundler/sycl-target-id.c new file mode 100644 index 0000000000000..0e96a6ba3f025 --- /dev/null +++ b/clang/test/OffloadTools/clang-offload-bundler/sycl-target-id.c @@ -0,0 +1,60 @@ +// REQUIRES: x86-registered-target + +// SYCL bundle entry IDs use a 3-component triple followed by the target ID, +// e.g. sycl-spir64_gen-unknown-unknown-pvc. Check that each target ID is +// unbundled to its own code object. + +// RUN: %clang -O0 -target x86_64-unknown-linux-gnu %s -c -o %t.o +// RUN: echo 'host' > %t.host +// RUN: echo 'tgt1' > %t.tgt1 +// RUN: echo 'tgt2' > %t.tgt2 + +// Text bundles. +// RUN: clang-offload-bundler -type=ii -input=%t.host -input=%t.tgt1 -input=%t.tgt2 -output=%t.bundle.ii \ +// RUN: -targets=host-x86_64-unknown-linux-gnu,sycl-spir64_gen-unknown-unknown-tgllp,sycl-spir64_gen-unknown-unknown-pvc +// RUN: clang-offload-bundler -type=ii -input=%t.bundle.ii -output=%t.res1 -output=%t.res2 -unbundle \ +// RUN: -targets=sycl-spir64_gen-unknown-unknown-tgllp,sycl-spir64_gen-unknown-unknown-pvc +// RUN: diff %t.tgt1 %t.res1 +// RUN: diff %t.tgt2 %t.res2 +// RUN: clang-offload-bundler -type=ii -input=%t.bundle.ii -output=%t.res -unbundle \ +// RUN: -targets=sycl-spir64_gen-unknown-unknown-pvc +// RUN: diff %t.tgt2 %t.res + +// RUN: clang-offload-bundler -type=ii -input=%t.host -input=%t.tgt1 -input=%t.tgt2 -output=%t.bundle.nv.ii \ +// RUN: -targets=host-x86_64-unknown-linux-gnu,sycl-nvptx64-nvidia-cuda-sm_80,sycl-nvptx64-nvidia-cuda-sm_90 +// RUN: clang-offload-bundler -type=ii -input=%t.bundle.nv.ii -output=%t.res -unbundle \ +// RUN: -targets=sycl-nvptx64-nvidia-cuda-sm_90 +// RUN: diff %t.tgt2 %t.res + +// Object bundles, as produced for -fsycl-targets=intel_gpu_* with and without +// -fno-sycl-rdc. +// RUN: clang-offload-bundler -type=o -input=%t.o -input=%t.tgt1 -input=%t.tgt2 -output=%t.bundle.o \ +// RUN: -targets=host-x86_64-unknown-linux-gnu,sycl-spir64_gen-unknown-unknown-tgllp,sycl-spir64_gen-unknown-unknown-pvc +// RUN: clang-offload-bundler -type=o -input=%t.bundle.o -output=%t.res -unbundle \ +// RUN: -targets=sycl-spir64_gen-unknown-unknown-pvc +// RUN: diff %t.tgt2 %t.res +// RUN: clang-offload-bundler -type=o -input=%t.o -input=%t.tgt1 -input=%t.tgt2 -output=%t.bundle.image.o \ +// RUN: -targets=host-x86_64-unknown-linux-gnu,sycl-spir64_gen_image-unknown-unknown-tgllp,sycl-spir64_gen_image-unknown-unknown-pvc +// RUN: clang-offload-bundler -type=o -input=%t.bundle.image.o -output=%t.res -unbundle \ +// RUN: -targets=sycl-spir64_gen_image-unknown-unknown-pvc +// RUN: diff %t.tgt2 %t.res + +// A SYCL code object without a target ID is still compatible with any target +// ID of the same triple, and vice versa. +// RUN: clang-offload-bundler -type=ii -input=%t.host -input=%t.tgt1 -output=%t.bundle.noid.ii \ +// RUN: -targets=host-x86_64-unknown-linux-gnu,sycl-spir64_gen-unknown-unknown +// RUN: clang-offload-bundler -type=ii -input=%t.bundle.noid.ii -output=%t.res -unbundle \ +// RUN: -targets=sycl-spir64_gen-unknown-unknown-pvc +// RUN: diff %t.tgt1 %t.res +// RUN: clang-offload-bundler -type=ii -input=%t.bundle.ii -output=%t.res -unbundle \ +// RUN: -targets=sycl-spir64_gen-unknown-unknown +// RUN: diff %t.tgt1 %t.res + +// 4-component triples are unaffected. +// RUN: clang-offload-bundler -type=ii -input=%t.host -input=%t.tgt1 -input=%t.tgt2 -output=%t.bundle.hip.ii \ +// RUN: -targets=host-x86_64-unknown-linux-gnu,hipv4-amdgcn-amd-amdhsa--gfx908,hipv4-amdgcn-amd-amdhsa--gfx90a +// RUN: clang-offload-bundler -type=ii -input=%t.bundle.hip.ii -output=%t.res -unbundle \ +// RUN: -targets=hipv4-amdgcn-amd-amdhsa--gfx90a +// RUN: diff %t.tgt2 %t.res + +int A = 0; From 29293ce0ba997c9ce40ad1d92456c3ca33587502 Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Mon, 12 Oct 2026 12:28:27 -0700 Subject: [PATCH 07/22] [SYCL][E2E] Fix AOT tests that relied on spir64_gen images running on other devices Images built for intel_gpu_ carry a compile_target and are only selected for that device, unlike spir64_gen -device images: - Basic/build_log.cpp needs an image built for another device, so it keeps spir64_gen with -Wno-deprecated and is not supported with the new offloading model. - DeviceCodeSplit/aot-gpu.cpp, NewOffloadDriver/aot-gpu.cpp and KernelAndProgram/test_cache_jit_aot.cpp don't depend on a specific device, so build them for %{gpu_aot_target}. Drop the second AOT block of test_cache_jit_aot.cpp, which only differed by building one image for two devices. --- sycl/test-e2e/Basic/build_log.cpp | 4 +++- sycl/test-e2e/DeviceCodeSplit/aot-gpu.cpp | 2 +- .../KernelAndProgram/test_cache_jit_aot.cpp | 21 +------------------ sycl/test-e2e/NewOffloadDriver/aot-gpu.cpp | 2 +- 4 files changed, 6 insertions(+), 23 deletions(-) diff --git a/sycl/test-e2e/Basic/build_log.cpp b/sycl/test-e2e/Basic/build_log.cpp index 359584f1e777a..dc589ba40c8bc 100644 --- a/sycl/test-e2e/Basic/build_log.cpp +++ b/sycl/test-e2e/Basic/build_log.cpp @@ -1,8 +1,10 @@ // REQUIRES: opencl || level_zero, gpu, ocloc, target-spir // UNSUPPORTED: arch-intel_gpu_dg2 // UNSUPPORTED-INTENDED: see https://github.com/intel/llvm/pull/20643 +// UNSUPPORTED: new-offload-model +// UNSUPPORTED-INTENDED: needs a spir64_gen image built for another device. -// RUN: %clangxx -fsycl -fsycl-targets=intel_gpu_dg2 %s -o %t.out +// RUN: %clangxx -fsycl -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen "-device dg2" -Wno-deprecated %s -o %t.out // RUN: env SYCL_RT_WARNING_LEVEL=2 %{run} %t.out 2>&1 | FileCheck %s #include diff --git a/sycl/test-e2e/DeviceCodeSplit/aot-gpu.cpp b/sycl/test-e2e/DeviceCodeSplit/aot-gpu.cpp index 06801b0f60bc3..0ad21f5b1b046 100644 --- a/sycl/test-e2e/DeviceCodeSplit/aot-gpu.cpp +++ b/sycl/test-e2e/DeviceCodeSplit/aot-gpu.cpp @@ -1,7 +1,7 @@ // REQUIRES: ocloc, gpu, spir-family, !gpu-intel-gen12 // // RUN: %clangxx -fsycl -fsycl-device-code-split=per_source \ -// RUN: -fsycl-targets=intel_gpu_dg2 -I %S/Inputs -o %t.out \ +// RUN: -fsycl-targets=%{gpu_aot_target} -I %S/Inputs -o %t.out \ // RUN: %S/split-per-source-main.cpp \ // RUN: %S/Inputs/split-per-source-second-file.cpp \ // RUN: -fsycl-dead-args-optimization diff --git a/sycl/test-e2e/KernelAndProgram/test_cache_jit_aot.cpp b/sycl/test-e2e/KernelAndProgram/test_cache_jit_aot.cpp index 24a6b92cb12cc..213b610769046 100644 --- a/sycl/test-e2e/KernelAndProgram/test_cache_jit_aot.cpp +++ b/sycl/test-e2e/KernelAndProgram/test_cache_jit_aot.cpp @@ -27,7 +27,7 @@ // RUN: %{cache_vars} %{run-unfiltered-devices} %t.out 2>&1 | FileCheck %s --check-prefixes RESULT1 // ****************************** // -// REDEFINE: %{build_cmd} = %clangxx -fsycl -fsycl-targets=intel_gpu_acm_g10 %s +// REDEFINE: %{build_cmd} = %clangxx -fsycl -fsycl-targets=%{gpu_aot_target} %s // ****************************** // Check the logs first. // RUN: %{run-aux} %{build_cmd} -DVALUE=1 -o %t.out @@ -36,25 +36,6 @@ // RUN: %{cache_vars} %{run-unfiltered-devices} %t.out 2>&1 | FileCheck %s --check-prefixes=CHECK-CACHE-READ // RUN: %{cache_vars} %{run-unfiltered-devices} %t.out 2>&1 | FileCheck %s --check-prefixes=CHECK-CACHE-READ -// -// Now try to substitute the cached image and verify it is actually taken and -// the code/binary there is executed. -// RUN: %{run-aux} mv %t/cache_dir/*/*/*/*/*.bin %t.value1.bin -// RUN: %{run-aux} rm -rf %t/cache_dir/* -// RUN: %{run-aux} %{build_cmd} -DVALUE=2 -o %t.out -// RUN: %{cache_vars} %{run-unfiltered-devices} %t.out 2>&1 | FileCheck %s --check-prefixes RESULT2 -// RUN: %{run-aux} mv %t.value1.bin %t/cache_dir/*/*/*/*/*.bin -// RUN: %{cache_vars} %{run-unfiltered-devices} %t.out 2>&1 | FileCheck %s --check-prefixes RESULT1 -// ****************************** -// -// REDEFINE: %{build_cmd} = %clangxx -fsycl -fsycl-targets=intel_gpu_acm_g10,intel_gpu_acm_g11 %s -// ****************************** -// Check the logs first. -// RUN: %{run-aux} %{build_cmd} -DVALUE=1 -o %t.out -// RUN: %{run-aux} rm -rf %t/cache_dir/* -// RUN: %{cache_vars} %{run-unfiltered-devices} %t.out 2>&1 | FileCheck %s --check-prefixes=CHECK-CACHE-WRITE -// RUN: %{cache_vars} %{run-unfiltered-devices} %t.out 2>&1 | FileCheck %s --check-prefixes=CHECK-CACHE-READ -// RUN: %{cache_vars} %{run-unfiltered-devices} %t.out 2>&1 | FileCheck %s --check-prefixes=CHECK-CACHE-READ // // Now try to substitute the cached image and verify it is actually taken and // the code/binary there is executed. diff --git a/sycl/test-e2e/NewOffloadDriver/aot-gpu.cpp b/sycl/test-e2e/NewOffloadDriver/aot-gpu.cpp index 626bb3bb59c70..bf9deca2a28a2 100644 --- a/sycl/test-e2e/NewOffloadDriver/aot-gpu.cpp +++ b/sycl/test-e2e/NewOffloadDriver/aot-gpu.cpp @@ -2,7 +2,7 @@ // Test with `--offload-new-driver` // // RUN: %clangxx -fsycl -fsycl-device-code-split=per_source \ -// RUN: -fsycl-targets=intel_gpu_dg2 -I %S/Inputs -o %t.out \ +// RUN: -fsycl-targets=%{gpu_aot_target} -I %S/Inputs -o %t.out \ // RUN: %S/split-per-source-main.cpp \ // RUN: %S/Inputs/split-per-source-second-file.cpp \ // RUN: -fsycl-dead-args-optimization --offload-new-driver From bb2a27add4bd80397dda65a796849206044eb373 Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Mon, 12 Oct 2026 12:29:15 -0700 Subject: [PATCH 08/22] [SYCL][E2E] Accept one abstract module per architecture in simple_kernel_aot.cpp spir64_gen with -device * produced a single abstract module. A list of intel_gpu_ targets produces one per architecture. --- sycl/test-e2e/SYCLBIN/simple_kernel_aot.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test-e2e/SYCLBIN/simple_kernel_aot.cpp b/sycl/test-e2e/SYCLBIN/simple_kernel_aot.cpp index 099aebd3a345f..10a2b632dd7cc 100644 --- a/sycl/test-e2e/SYCLBIN/simple_kernel_aot.cpp +++ b/sycl/test-e2e/SYCLBIN/simple_kernel_aot.cpp @@ -48,7 +48,7 @@ void TestKernel(int *Ptr, int Size) { // CHECK-INPUT-ALL-GPU-DEVICE-ARCHS-NEXT: Global metadata: // CHECK-INPUT-ALL-GPU-DEVICE-ARCHS-NEXT: SYCLBIN/global metadata: // CHECK-INPUT-ALL-GPU-DEVICE-ARCHS-NEXT: state: 0 -// CHECK-INPUT-ALL-GPU-DEVICE-ARCHS-NEXT: Number of Abstract Modules: 1 +// CHECK-INPUT-ALL-GPU-DEVICE-ARCHS-NEXT: Number of Abstract Modules: {{[1-9][0-9]*}} // CHECK-INPUT-ALL-GPU-DEVICE-ARCHS-NEXT: Abstract Module 0: // CHECK-INPUT-ALL-GPU-DEVICE-ARCHS-NEXT: Metadata: // CHECK-INPUT-ALL-GPU-DEVICE-ARCHS: Number of IR Modules: 0 From 3560b03b2374becdab20ec29442abb53ef7e8e9c Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Mon, 12 Oct 2026 14:31:30 -0700 Subject: [PATCH 09/22] [SYCL] Only use AOT images built for the device's architecture When an application is compiled for several -fsycl-targets=intel_gpu_* architectures, each device image is present once per architecture, tagged with a compile_target. Device image dependencies and kernel bundle images were only checked against the image triple, so an image built for another architecture could be picked, e.g. a PVC kernel image linked with the TGLLP image of its dependency, failing with "Unable to merge incompatible images". Skip images whose compile_target does not match the device, as is already done when picking the image for a kernel. Images without a compile_target are unaffected. Make the mock device in SYCLBINSerialize a PVC GPU, since its executable bundle is built from an intel_gpu_pvc image. --- .../program_manager/program_manager.cpp | 21 ++- sycl/unittests/SYCL2020/SYCLBINSerialize.cpp | 14 ++ .../DynamicLinking/CMakeLists.txt | 1 + .../DynamicLinkingCompileTarget.cpp | 140 ++++++++++++++++++ 4 files changed, 173 insertions(+), 3 deletions(-) create mode 100644 sycl/unittests/program_manager/DynamicLinking/DynamicLinkingCompileTarget.cpp diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index fca8164e90970..c7ef89c5644c6 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -573,6 +573,18 @@ static bool compatibleWithDevice(const RTDeviceBinaryImage *BinImage, return (0 == SuitableImageID); } +// An image with a compile_target was built for that device architecture only. +// Images without one are matched by compatibleWithDevice. +static bool doesCompileTargetMatchDevice(const RTDeviceBinaryImage &Img, + const device_impl &Dev) { + const RTDeviceBinaryImage::PropertyRange &Range = Img.getDeviceRequirements(); + const bool HasCompileTarget = + std::any_of(Range.begin(), Range.end(), [](const auto &Prop) { + return Prop->Name == std::string_view("compile_target"); + }); + return !HasCompileTarget || doesImageTargetMatchDevice(Img, Dev); +} + // Check if the device image is a BF16 devicelib image. bool ProgramManager::isBfloat16DeviceImage( const RTDeviceBinaryImage *BinImage) { @@ -693,7 +705,8 @@ ProgramManager::collectDeviceImageDepsForImportedSymbols( const RTDeviceBinaryImage *Img = It->second; if (!doesDevSupportDeviceRequirements(Dev, *Img) || - !compatibleWithDevice(Img, Dev)) + !compatibleWithDevice(Img, Dev) || + !doesCompileTargetMatchDevice(*Img, Dev)) continue; // If the image is a BF16 device image, we need to check if it @@ -791,7 +804,8 @@ ProgramManager::collectDependentDeviceImagesForVirtualFunctions( // virtual function set and it is incompatible, then we should link its // "dummy" version to avoid link errors about unresolved external // symbols. - if (doesDevSupportDeviceRequirements(Dev, *BinImage)) + if (doesDevSupportDeviceRequirements(Dev, *BinImage) && + doesCompileTargetMatchDevice(*BinImage, Dev)) DeviceImagesToLink.insert(BinImage); } } @@ -2316,7 +2330,8 @@ ProgramManager::getSYCLDeviceImagesWithCompatibleState( for (const RTDeviceBinaryImage *BinImage : BinImages) { if (!compatibleWithDevice(BinImage, Dev) || - !doesDevSupportDeviceRequirements(Dev, *BinImage)) + !doesDevSupportDeviceRequirements(Dev, *BinImage) || + !doesCompileTargetMatchDevice(*BinImage, Dev)) continue; auto InsertRes = ImageInfoMap.insert({BinImage, {}}); diff --git a/sycl/unittests/SYCL2020/SYCLBINSerialize.cpp b/sycl/unittests/SYCL2020/SYCLBINSerialize.cpp index 43bfea488c44b..7732a2087af0d 100644 --- a/sycl/unittests/SYCL2020/SYCLBINSerialize.cpp +++ b/sycl/unittests/SYCL2020/SYCLBINSerialize.cpp @@ -62,6 +62,16 @@ sycl::unittest::MockDeviceImage Imgs[] = { makeImageWithCompileTarget("SYCLBINSerializeKernel", "intel_gpu_pvc")}; sycl::unittest::MockDeviceImageArray ImgArray{Imgs}; +// Makes the mock device a PVC GPU, so that the intel_gpu_pvc image is used. +ur_result_t redefinedDeviceGetInfoPVC(void *pParams) { + auto Params = *static_cast(pParams); + if (*Params.ppropName == UR_DEVICE_INFO_IP_VERSION && *Params.ppPropValue) + *static_cast(*Params.ppPropValue) = 0x030f0000; + if (*Params.ppropName == UR_DEVICE_INFO_TYPE && *Params.ppPropValue) + *static_cast(*Params.ppPropValue) = UR_DEVICE_TYPE_GPU; + return UR_RESULT_SUCCESS; +} + // Strips the OffloadBinary wrapper that ext_oneapi_get_content emits and // returns a string_view over the contained raw SYCLBIN bytes. Returns // std::nullopt if the wrapper does not parse cleanly; callers should ASSERT @@ -115,6 +125,8 @@ stripOffloadWrapper(const std::vector &Bytes) { // confirm the property survived. TEST(SYCLBINSerialize, CompileTargetRoundTrip) { sycl::unittest::UrMock<> Mock; + mock::getCallbacks().set_after_callback("urDeviceGetInfo", + &redefinedDeviceGetInfoPVC); const sycl::device Dev = sycl::platform().get_devices()[0]; sycl::context Ctx{Dev}; @@ -222,6 +234,8 @@ TEST(SYCLBINSerialize, BundleStateRoundTripObject) { // inspection alone (e.g. CompileTargetRoundTrip) does not catch. TEST(SYCLBINSerialize, ImageBytesRoundTrip) { sycl::unittest::UrMock<> Mock; + mock::getCallbacks().set_after_callback("urDeviceGetInfo", + &redefinedDeviceGetInfoPVC); const sycl::device Dev = sycl::platform().get_devices()[0]; sycl::context Ctx{Dev}; diff --git a/sycl/unittests/program_manager/DynamicLinking/CMakeLists.txt b/sycl/unittests/program_manager/DynamicLinking/CMakeLists.txt index d57334faf6dd5..8f19b170eb37d 100644 --- a/sycl/unittests/program_manager/DynamicLinking/CMakeLists.txt +++ b/sycl/unittests/program_manager/DynamicLinking/CMakeLists.txt @@ -4,5 +4,6 @@ if (NOT LLVM_LIBCXX_USED) add_sycl_unittest(DynamicLinkingTests OBJECT DynamicLinking.cpp + DynamicLinkingCompileTarget.cpp ) endif() diff --git a/sycl/unittests/program_manager/DynamicLinking/DynamicLinkingCompileTarget.cpp b/sycl/unittests/program_manager/DynamicLinking/DynamicLinkingCompileTarget.cpp new file mode 100644 index 0000000000000..afe7bd5a94b46 --- /dev/null +++ b/sycl/unittests/program_manager/DynamicLinking/DynamicLinkingCompileTarget.cpp @@ -0,0 +1,140 @@ +#include + +#include +#include +#include +#include + +#include + +// Check that the dependencies of an AOT image are resolved with images built +// for the same device when there are images for several devices. + +namespace DynamicLinkingCompileTargetTest { +class MultiArchKernel; +} // namespace DynamicLinkingCompileTargetTest + +namespace sycl { +inline namespace _V1 { +namespace detail { +template <> +struct KernelInfo + : public unittest::MockKernelInfoBase { + static constexpr const char *getName() { return "MultiArchKernel"; } +}; +} // namespace detail +} // namespace _V1 +} // namespace sycl + +namespace { +static constexpr unsigned TGLLP_PRG = 31; +static constexpr unsigned TGLLP_PRG_DEP = 37; +static constexpr unsigned PVC_PRG = 41; +static constexpr unsigned PVC_PRG_DEP = 43; + +// IP version of a PVC device. +static constexpr uint32_t PvcIp = 0x030f0000; + +sycl::unittest::MockDeviceImage +generateImage(std::initializer_list KernelNames, + const std::string &ExportedSymbol, + const std::string &ImportedSymbol, unsigned char Magic, + const std::string &CompileTarget) { + sycl::unittest::MockPropertySet PropSet; + + std::vector Data(8 + CompileTarget.size()); + std::copy(CompileTarget.begin(), CompileTarget.end(), Data.data() + 8); + PropSet.insert(__SYCL_PROPERTY_SET_SYCL_DEVICE_REQUIREMENTS, + sycl::unittest::MockProperty("compile_target", Data, + SYCL_PROPERTY_TYPE_BYTE_ARRAY)); + + auto makeSymbolProps = [](const std::string &Symbol) { + std::vector Props; + Props.emplace_back(Symbol, std::vector(sizeof(uint32_t)), + SYCL_PROPERTY_TYPE_UINT32); + return Props; + }; + if (!ExportedSymbol.empty()) + PropSet.insert(__SYCL_PROPERTY_SET_SYCL_EXPORTED_SYMBOLS, + makeSymbolProps(ExportedSymbol)); + if (!ImportedSymbol.empty()) + PropSet.insert(__SYCL_PROPERTY_SET_SYCL_IMPORTED_SYMBOLS, + makeSymbolProps(ImportedSymbol)); + + return sycl::unittest::MockDeviceImage{ + SYCL_DEVICE_BINARY_TYPE_NATIVE, + __SYCL_DEVICE_BINARY_TARGET_SPIRV64_GEN, + "", // Compile options + "", // Link options + std::vector{Magic}, + sycl::unittest::makeEmptyKernels(KernelNames), + std::move(PropSet)}; +} + +// The TGLLP images come first, so that picking the first image exporting the +// symbol would choose the wrong dependency for a PVC device. +static sycl::unittest::MockDeviceImage Imgs[] = { + generateImage({"MultiArchKernel"}, "", "MultiArchKernelDep", TGLLP_PRG, + "intel_gpu_tgllp"), + generateImage({"MultiArchKernelDep"}, "MultiArchKernelDep", "", + TGLLP_PRG_DEP, "intel_gpu_tgllp"), + generateImage({"MultiArchKernel"}, "", "MultiArchKernelDep", PVC_PRG, + "intel_gpu_pvc"), + generateImage({"MultiArchKernelDep"}, "MultiArchKernelDep", "", PVC_PRG_DEP, + "intel_gpu_pvc")}; +static sycl::unittest::MockDeviceImageArray<4> ImgArray{Imgs}; + +ur_result_t redefinedDeviceGetInfo(void *pParams) { + auto Params = *static_cast(pParams); + if (*Params.ppropName == UR_DEVICE_INFO_IP_VERSION && *Params.ppPropValue) + *static_cast(*Params.ppPropValue) = PvcIp; + if (*Params.ppropName == UR_DEVICE_INFO_TYPE && *Params.ppPropValue) + *static_cast(*Params.ppPropValue) = UR_DEVICE_TYPE_GPU; + return UR_RESULT_SUCCESS; +} +} // namespace + +TEST(DynamicLinking, AheadOfTimeDependencyForDevice) { + sycl::unittest::UrMock Mock; + setupRuntimeLinkingMock(); + mock::getCallbacks().set_after_callback("urDeviceGetInfo", + &redefinedDeviceGetInfo); + + sycl::platform Plt = sycl::platform(); + sycl::queue Q(Plt.get_devices()[0]); + + CapturedLinkingData.clear(); + + Q.single_task([=]() {}); + ASSERT_EQ(CapturedLinkingData.NumOfUrProgramLinkCalls, 1u); + EXPECT_TRUE( + CapturedLinkingData.LinkedProgramsContains({PVC_PRG, PVC_PRG_DEP})); + EXPECT_EQ(CapturedLinkingData.ProgramUsedToCreateKernel, + PVC_PRG * PVC_PRG_DEP); +} + +TEST(DynamicLinking, AheadOfTimeKernelBundleForDevice) { + sycl::unittest::UrMock Mock; + setupRuntimeLinkingMock(); + mock::getCallbacks().set_after_callback("urDeviceGetInfo", + &redefinedDeviceGetInfo); + + sycl::platform Plt = sycl::platform(); + sycl::queue Q(Plt.get_devices()[0]); + + CapturedLinkingData.clear(); + + auto Bundle = sycl::get_kernel_bundle( + Q.get_context(), {Q.get_device()}, + {sycl::get_kernel_id< + DynamicLinkingCompileTargetTest::MultiArchKernel>()}); + Q.submit([&](sycl::handler &CGH) { + CGH.use_kernel_bundle(Bundle); + CGH.single_task([=]() {}); + }); + ASSERT_EQ(CapturedLinkingData.NumOfUrProgramLinkCalls, 1u); + EXPECT_TRUE( + CapturedLinkingData.LinkedProgramsContains({PVC_PRG, PVC_PRG_DEP})); + EXPECT_EQ(CapturedLinkingData.ProgramUsedToCreateKernel, + PVC_PRG * PVC_PRG_DEP); +} From 08d46ad973d62638a82bfe822549fa665c9df140 Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Mon, 12 Oct 2026 15:31:58 -0700 Subject: [PATCH 10/22] [SYCL][E2E] Add a SPIR-V image to the AOT block of test_cache_jit_aot.cpp The test runs with %{run-unfiltered-devices}, so on a machine with several GPUs it can run on a device other than the one %{gpu_aot_target} was detected for, e.g. the discrete GPU of the ARL runner. Add spir64 so that such a device falls back to the SPIR-V image, while the AOT image is still preferred on the device it was built for. --- sycl/test-e2e/KernelAndProgram/test_cache_jit_aot.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test-e2e/KernelAndProgram/test_cache_jit_aot.cpp b/sycl/test-e2e/KernelAndProgram/test_cache_jit_aot.cpp index 213b610769046..2ac63b10c467c 100644 --- a/sycl/test-e2e/KernelAndProgram/test_cache_jit_aot.cpp +++ b/sycl/test-e2e/KernelAndProgram/test_cache_jit_aot.cpp @@ -27,7 +27,7 @@ // RUN: %{cache_vars} %{run-unfiltered-devices} %t.out 2>&1 | FileCheck %s --check-prefixes RESULT1 // ****************************** // -// REDEFINE: %{build_cmd} = %clangxx -fsycl -fsycl-targets=%{gpu_aot_target} %s +// REDEFINE: %{build_cmd} = %clangxx -fsycl -fsycl-targets=%{gpu_aot_target},spir64 %s // ****************************** // Check the logs first. // RUN: %{run-aux} %{build_cmd} -DVALUE=1 -o %t.out From f4a274db4bd85efa424275bb8ce4480e8967b647 Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Tue, 13 Oct 2026 09:53:09 -0700 Subject: [PATCH 11/22] [SYCL][E2E] Set the build-only gpu_aot_target from the CI workflows Rather than listing the architectures of the CI runners in lit.cfg.py, have the CI build workflows pass them via --param gpu_aot_target, next to the runner lists they need to match. Without a detected Intel GPU or the parameter, lit now defaults to intel_gpu_dg2 and says so. --- .github/workflows/sycl-linux-build.yml | 11 +++++---- .github/workflows/sycl-linux-precommit.yml | 2 ++ .github/workflows/sycl-windows-build.yml | 9 +++++-- .github/workflows/sycl-windows-precommit.yml | 2 ++ sycl/test-e2e/README.md | 9 ++++--- sycl/test-e2e/lit.cfg.py | 25 +++++++------------- 6 files changed, 30 insertions(+), 28 deletions(-) diff --git a/.github/workflows/sycl-linux-build.yml b/.github/workflows/sycl-linux-build.yml index 54c5bcbaf9cea..19a80f14d5ecd 100644 --- a/.github/workflows/sycl-linux-build.yml +++ b/.github/workflows/sycl-linux-build.yml @@ -152,6 +152,9 @@ jobs: env: CCACHE_DIR: ${{ inputs.build_cache_root }}/build_cache_${{ inputs.build_cache_suffix }} CCACHE_MAXSIZE: 8G + # Intel GPU architectures of the run-only E2E runners in + # sycl-linux-precommit.yml. Keep in sync when adding or removing one. + GPU_AOT_TARGET: intel_gpu_tgllp,intel_gpu_acm_g10,intel_gpu_pvc,intel_gpu_bmg_g21,intel_gpu_mtl_u steps: - name: Deduce artifact archive params # To reduce number of inputs parameters that is limited for manual triggers. @@ -330,7 +333,7 @@ jobs: target_devices: all binaries_artifact: ${{ inputs.e2e_binaries_artifact }} sycl_compiler: $GITHUB_WORKSPACE/toolchain/bin/dpclang++ - extra_lit_opts: --param sycl_build_targets="spir;nvidia" + extra_lit_opts: --param sycl_build_targets="spir;nvidia" --param gpu_aot_target=${{ env.GPU_AOT_TARGET }} - name: Build E2E tests with NewOffloadModel if: | @@ -345,7 +348,7 @@ jobs: target_devices: all binaries_artifact: ${{ inputs.e2e_binaries_new_offload_model_artifact }} sycl_compiler: $GITHUB_WORKSPACE/toolchain/bin/dpclang++ - extra_lit_opts: --param sycl_build_targets="spir;nvidia" --param enable_new_offload_model=True + extra_lit_opts: --param sycl_build_targets="spir;nvidia" --param enable_new_offload_model=True --param gpu_aot_target=${{ env.GPU_AOT_TARGET }} - name: Build E2E tests with SPIR-V Backend if: ${{ inputs.e2e_binaries_spirv_backend_artifact && !cancelled() && steps.build.conclusion == 'success' }} @@ -356,7 +359,7 @@ jobs: target_devices: all binaries_artifact: ${{ inputs.e2e_binaries_spirv_backend_artifact }} sycl_compiler: $GITHUB_WORKSPACE/toolchain/bin/dpclang++ - extra_lit_opts: --param spirv-backend=True + extra_lit_opts: --param spirv-backend=True --param gpu_aot_target=${{ env.GPU_AOT_TARGET }} - name: Build E2E tests in Preview Mode if: ${{ inputs.e2e_binaries_preview_artifact && !cancelled() && steps.build.conclusion == 'success' && !startsWith(github.base_ref, 'sycl-rel-') }} @@ -367,4 +370,4 @@ jobs: target_devices: all binaries_artifact: ${{ inputs.e2e_binaries_preview_artifact }} sycl_compiler: $GITHUB_WORKSPACE/toolchain/bin/dpclang++ - extra_lit_opts: --param test-preview-mode=True + extra_lit_opts: --param test-preview-mode=True --param gpu_aot_target=${{ env.GPU_AOT_TARGET }} diff --git a/.github/workflows/sycl-linux-precommit.yml b/.github/workflows/sycl-linux-precommit.yml index 68a6ce9f939ee..c473b3830140d 100644 --- a/.github/workflows/sycl-linux-precommit.yml +++ b/.github/workflows/sycl-linux-precommit.yml @@ -142,6 +142,8 @@ jobs: echo FILTER_6_2="$(grep -v '^#\|^\W*$' devops/compat_ci_exclude.sycl-rel-6_2 | paste -sd '|')" >> $GITHUB_OUTPUT echo FILTER_6_3="$(grep -v '^#\|^\W*$' devops/compat_ci_exclude.sycl-rel-6_3 | paste -sd '|')" >> $GITHUB_OUTPUT + # The E2E binaries run here are built for GPU_AOT_TARGET in + # sycl-linux-build.yml; update it when adding or removing an Intel GPU runner. E2E: needs: [build, detect_changes, compat_read_exclude] if: ${{ !cancelled() && needs.build.outputs.build_conclusion == 'success' }} diff --git a/.github/workflows/sycl-windows-build.yml b/.github/workflows/sycl-windows-build.yml index 5d8edfde7bf45..da370d946fe8a 100644 --- a/.github/workflows/sycl-windows-build.yml +++ b/.github/workflows/sycl-windows-build.yml @@ -104,6 +104,11 @@ jobs: runs-on: [Windows, build] outputs: build_conclusion: ${{ steps.build.conclusion }} + env: + # Intel GPU architectures of the run-only E2E runners in + # sycl-windows-precommit.yml. Keep in sync when adding or removing one. + # Gen12 and older (gen12 runner) are not supported by Windows ocloc. + GPU_AOT_TARGET: intel_gpu_acm_g10,intel_gpu_bmg_g21 steps: - name: Detect hung tests if: always() @@ -257,7 +262,7 @@ jobs: testing_mode: build-only target_devices: all binaries_artifact: ${{ inputs.e2e_binaries_artifact }} - extra_lit_opts: --param sycl_build_targets="spir" + extra_lit_opts: --param sycl_build_targets="spir" --param gpu_aot_target=${{ env.GPU_AOT_TARGET }} cxx: ${{ inputs.cxx }} - name: Build E2E tests with New Offload Model @@ -272,7 +277,7 @@ jobs: testing_mode: build-only target_devices: all binaries_artifact: ${{ inputs.e2e_binaries_new_offload_model_artifact }} - extra_lit_opts: --param sycl_build_targets="spir" --param enable_new_offload_model=True + extra_lit_opts: --param sycl_build_targets="spir" --param enable_new_offload_model=True --param gpu_aot_target=${{ env.GPU_AOT_TARGET }} cxx: ${{ inputs.cxx }} - name: Detect hung tests diff --git a/.github/workflows/sycl-windows-precommit.yml b/.github/workflows/sycl-windows-precommit.yml index 4c747545df5bf..2290861ce13eb 100644 --- a/.github/workflows/sycl-windows-precommit.yml +++ b/.github/workflows/sycl-windows-precommit.yml @@ -60,6 +60,8 @@ jobs: e2e_binaries_artifact: sycl_windows_e2ebin e2e_binaries_new_offload_model_artifact: sycl_windows_e2ebin_with_new_offload_model + # The E2E binaries run here are built for GPU_AOT_TARGET in + # sycl-windows-build.yml; update it when adding or removing an Intel GPU runner. run_prebuilt_e2e_tests: needs: build # Continue if build was successful. diff --git a/sycl/test-e2e/README.md b/sycl/test-e2e/README.md index 42d0f2cb67e74..c063496c55962 100644 --- a/sycl/test-e2e/README.md +++ b/sycl/test-e2e/README.md @@ -210,9 +210,7 @@ or via the ***LIT_OPTS*** environment variable. ***GPU_AOT_TARGET*** - defines the Intel GPU target(s) used for AOT compilation, as a `-fsycl-targets` value (e.g. "intel_gpu_pvc" or "intel_gpu_pvc,intel_gpu_bmg_g21"). If not specified, the architectures of the -detected Intel GPU devices are used. If none are detected, the architectures of -the CI run-only runners are used, so that binaries built in `build-only` mode -run on all of them. +detected Intel GPU devices are used, or "intel_gpu_dg2" if none are detected. ***OFFLOAD_BUILD_TARGET*** - when testing the Offload backend, this must be set to specify the correct build target type for the available Offload device. @@ -299,8 +297,9 @@ configure specific single test execution in the command line: compilation, as a `-fsycl-targets` value. Tests use it via the `%{gpu_aot_target}` substitution. It can be also set by CMake variable GPU_AOT_TARGET. If not specified, the architectures of the detected Intel GPU - devices are used. If none are detected (e.g. in `build-only` mode), the - architectures of the CI run-only runners are used. + devices are used, or "intel_gpu_dg2" if none are detected (e.g. in + `build-only` mode). Binaries built in `build-only` mode only run on the + devices they are built for, so set it to cover the devices they run on. * **gpu-intel-dg2** - tells LIT infra that Intel GPU DG2 is present in the system. It is developer / CI infra responsibility to make sure that the device is available in the system. diff --git a/sycl/test-e2e/lit.cfg.py b/sycl/test-e2e/lit.cfg.py index 2877945e87bd2..89070d6eefb9a 100644 --- a/sycl/test-e2e/lit.cfg.py +++ b/sycl/test-e2e/lit.cfg.py @@ -1303,29 +1303,20 @@ def get_sycl_ls_verbose(sycl_device, env): config.substitutions.append(("%amd_arch_options", amd_arch_flags)) # Intel GPU AOT target(s), as a -fsycl-targets value. Uses the detected Intel -# GPU architectures. Without any (e.g. build-only mode), it uses the -# architectures of the CI run-only runners, since binaries built once are run -# on all of them. -if platform.system() == "Windows": - # Windows ocloc no longer supports Gen12 and older GPUs, such as tgllp. - CI_RUNNER_GPU_AOT_TARGETS = "intel_gpu_acm_g10,intel_gpu_bmg_g21" -else: - CI_RUNNER_GPU_AOT_TARGETS = ",".join( - [ - "intel_gpu_tgllp", - "intel_gpu_acm_g10", - "intel_gpu_pvc", - "intel_gpu_bmg_g21", - "intel_gpu_mtl_u", - ] - ) +# GPU architectures, or intel_gpu_dg2 if none are detected (e.g. build-only +# mode). CI passes the architectures of its run-only runners. if not config.gpu_aot_target: intel_gpu_archs = sorted( feature[len("arch-") :] for feature in set().union(*config.sycl_dev_features.values()) if feature.startswith("arch-intel_gpu_") ) - config.gpu_aot_target = ",".join(intel_gpu_archs) or CI_RUNNER_GPU_AOT_TARGETS + config.gpu_aot_target = ",".join(intel_gpu_archs) + if not config.gpu_aot_target: + config.gpu_aot_target = "intel_gpu_dg2" + lit_config.note( + "No Intel GPU detected, defaulting the AOT target to intel_gpu_dg2" + ) lit_config.note("Intel GPU AOT target: {}".format(config.gpu_aot_target)) config.substitutions.append(("%{gpu_aot_target}", config.gpu_aot_target)) From b3ef50090d37d006656912769c259efba1ec48de Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Tue, 13 Oct 2026 11:41:33 -0700 Subject: [PATCH 12/22] [SYCL][Doc] Describe spir64_gen deprecation in terms of --offload-new-driver --- sycl/doc/UsersManual.md | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/doc/UsersManual.md b/sycl/doc/UsersManual.md index 61e66834174ca..ea8fafa991174 100644 --- a/sycl/doc/UsersManual.md +++ b/sycl/doc/UsersManual.md @@ -35,8 +35,8 @@ and not recommended to use in production environment. * spir64_x86_64 - generate code ahead of time for x86_64 CPUs; * spir64_gen - generate code ahead of time for Intel Processor Graphics. Deprecated: use the intel_gpu_* values listed below instead. It emits a - deprecation warning with the old offloading model and is not supported - with the new offloading model (--offload-new-driver). To migrate + deprecation warning by default, and is an error with + --offload-new-driver. To migrate '-fsycl-targets=spir64_gen -Xsycl-target-backend "-device "': * a device name, e.g. "-device pvc": use -fsycl-targets=intel_gpu_pvc; * a full IP version, e.g. "-device 12.60.7": use From 0dff57cc0c0b55284cdfadc8010825c2dc4d751f Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Tue, 13 Oct 2026 11:41:34 -0700 Subject: [PATCH 13/22] [SYCL][E2E] Clarify why build_log.cpp is unsupported with the new offloading model --- sycl/test-e2e/Basic/build_log.cpp | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/sycl/test-e2e/Basic/build_log.cpp b/sycl/test-e2e/Basic/build_log.cpp index dc589ba40c8bc..060f0ac125323 100644 --- a/sycl/test-e2e/Basic/build_log.cpp +++ b/sycl/test-e2e/Basic/build_log.cpp @@ -2,7 +2,9 @@ // UNSUPPORTED: arch-intel_gpu_dg2 // UNSUPPORTED-INTENDED: see https://github.com/intel/llvm/pull/20643 // UNSUPPORTED: new-offload-model -// UNSUPPORTED-INTENDED: needs a spir64_gen image built for another device. +// UNSUPPORTED-INTENDED: runs a kernel built for DG2 on another GPU to get a +// build log. intel_gpu_dg2 images are only used on DG2, and spir64_gen is an +// error with --offload-new-driver. // RUN: %clangxx -fsycl -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen "-device dg2" -Wno-deprecated %s -o %t.out // RUN: env SYCL_RT_WARNING_LEVEL=2 %{run} %t.out 2>&1 | FileCheck %s From 424490c5dc10c3050004610dfe2a5f2e97a03873 Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Tue, 13 Oct 2026 11:46:07 -0700 Subject: [PATCH 14/22] [SYCL][E2E] Keep gpu_aot_target_opts for extra AOT options gpu_aot_target_opts (GPU_AOT_TARGET_OPTS) passes additional backend options to the AOT compilation of each target in gpu_aot_target, e.g. -cl-opt-disable, for the spir_gen build target and %aot_options. The device is now selected by gpu_aot_target, so passing -device is an error. --- sycl/test-e2e/README.md | 8 ++++++++ sycl/test-e2e/format.py | 9 ++++++++- sycl/test-e2e/lit.cfg.py | 14 ++++++++++++++ sycl/test-e2e/lit.site.cfg.py.in | 1 + 4 files changed, 31 insertions(+), 1 deletion(-) diff --git a/sycl/test-e2e/README.md b/sycl/test-e2e/README.md index c063496c55962..4a9b532bf322b 100644 --- a/sycl/test-e2e/README.md +++ b/sycl/test-e2e/README.md @@ -207,6 +207,10 @@ may be set to "gfx906". Otherwise must be provided via the ***amd_arch*** LIT parameter (e.g., ***--param amd_arch=gfx906***) at runtime via the command line or via the ***LIT_OPTS*** environment variable. +***GPU_AOT_TARGET_OPTS*** - defines additional options passed to the AOT +compilation of each Intel GPU target in GPU_AOT_TARGET (e.g. "-cl-opt-disable"). +It must not contain "-device"; use GPU_AOT_TARGET to select the targets. + ***GPU_AOT_TARGET*** - defines the Intel GPU target(s) used for AOT compilation, as a `-fsycl-targets` value (e.g. "intel_gpu_pvc" or "intel_gpu_pvc,intel_gpu_bmg_g21"). If not specified, the architectures of the @@ -293,6 +297,10 @@ configure specific single test execution in the command line: * **compatibility_testing** - forces LIT infra to skip the tests compilation to support compatibility testing (a SYCL application is built with one version of SYCL compiler and then run with different SYCL RT version); +* **gpu_aot_target_opts** - defines additional options passed to the AOT + compilation of each target in `gpu_aot_target`, for tests built with the + `spir_gen` build target or `%aot_options`. It can be also set by CMake + variable GPU_AOT_TARGET_OPTS. It must not contain "-device". * **gpu_aot_target** - defines the Intel GPU target(s) used for AOT compilation, as a `-fsycl-targets` value. Tests use it via the `%{gpu_aot_target}` substitution. It can be also set by CMake variable diff --git a/sycl/test-e2e/format.py b/sycl/test-e2e/format.py index fd161f1f62203..1f02e7d64f079 100644 --- a/sycl/test-e2e/format.py +++ b/sycl/test-e2e/format.py @@ -345,12 +345,19 @@ def execute(self, test, litConfig): if requires_cpu: aot_options = "-fsycl-targets=spir64_x86_64" else: - aot_options = "-fsycl-targets=" + test.config.gpu_aot_target + aot_options = ( + "-fsycl-targets=" + + test.config.gpu_aot_target + + " " + + test.config.gpu_aot_backend_opts + ) else: aot_options = "" substitutions.append(("%aot_options", aot_options)) sycl_target_opts = "-fsycl-targets=%{sycl_triple}" + if "target-spir_gen" in build_targets: + sycl_target_opts += " " + test.config.gpu_aot_backend_opts if "target-amd" in build_targets: hip_arch_opts = ( " -Xsycl-target-backend=amdgcn-amd-amdhsa --offload-arch={}".format( diff --git a/sycl/test-e2e/lit.cfg.py b/sycl/test-e2e/lit.cfg.py index 89070d6eefb9a..3fc031d281dea 100644 --- a/sycl/test-e2e/lit.cfg.py +++ b/sycl/test-e2e/lit.cfg.py @@ -1320,6 +1320,20 @@ def get_sycl_ls_verbose(sycl_device, env): lit_config.note("Intel GPU AOT target: {}".format(config.gpu_aot_target)) config.substitutions.append(("%{gpu_aot_target}", config.gpu_aot_target)) +# Extra backend options for Intel GPU AOT compilation, passed for each target in +# gpu_aot_target. The device is selected by gpu_aot_target instead. +if "-device" in config.gpu_aot_target_opts.split(): + lit_config.error( + "gpu_aot_target_opts must not contain -device; use gpu_aot_target to " + "select the Intel GPU AOT target(s)" + ) +config.gpu_aot_backend_opts = "" +if config.gpu_aot_target_opts: + config.gpu_aot_backend_opts = " ".join( + "-Xsycl-target-backend={} '{}'".format(target, config.gpu_aot_target_opts) + for target in config.gpu_aot_target.split(",") + ) + if lit_config.params.get("compatibility_testing", "False") != "False": config.substitutions.append(("%clangxx", " true ")) config.substitutions.append(("%clang", " true ")) diff --git a/sycl/test-e2e/lit.site.cfg.py.in b/sycl/test-e2e/lit.site.cfg.py.in index 682defae5e813..f4636bb323f30 100644 --- a/sycl/test-e2e/lit.site.cfg.py.in +++ b/sycl/test-e2e/lit.site.cfg.py.in @@ -55,6 +55,7 @@ config.c_flags = "@CMAKE_C_FLAGS@" config.external_tests = "@SYCL_EXTERNAL_TESTS@" config.extra_include = "@CMAKE_CURRENT_SOURCE_DIR@/include" config.gpu_aot_target = lit_config.params.get("gpu_aot_target", "@GPU_AOT_TARGET@") +config.gpu_aot_target_opts = lit_config.params.get("gpu_aot_target_opts", "@GPU_AOT_TARGET_OPTS@") config.vulkan_include_dir = "@Vulkan_INCLUDE_DIRS@" config.vulkan_lib = "@Vulkan_LIBRARY@" From ebc840f5a4251aab0dc03a7fa460fd19de9e54c5 Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Tue, 13 Oct 2026 15:10:18 -0700 Subject: [PATCH 15/22] [SYCL][clang-linker-wrapper] Wrap all architectures in one module for -fno-sycl-rdc The -fno-sycl-rdc compile step (--emit-fatbin-only) wrapped the images of each target group separately and then expected a single output, so a compile for several architectures of one triple, e.g. -fsycl-targets=intel_gpu_pvc,intel_gpu_dg2, failed with "Expect single output from the SYCL device linker". Collect the images of all groups and wrap them into a single wrapper module, as is done for the other offload kinds. The per-architecture information, such as the compile_target, is carried by each image. Inputs with different target triples, which the driver already rejects with -fno-sycl-rdc, are diagnosed rather than labelled with one triple. --- .../linker-wrapper-sycl.cpp | 11 +++++ .../ClangLinkerWrapper.cpp | 46 ++++++++++++++----- 2 files changed, 46 insertions(+), 11 deletions(-) diff --git a/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-sycl.cpp b/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-sycl.cpp index 1b380896381cf..4c9642072cd3e 100644 --- a/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-sycl.cpp +++ b/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-sycl.cpp @@ -6,6 +6,7 @@ // // RUN: %clang %s -fsycl -fsycl-targets=spir64-unknown-unknown -c --offload-new-driver --no-offloadlib -fno-sycl-instrument-device-code -o %t.o // RUN: %clang %s -fsycl -fsycl-targets=intel_gpu_pvc -c --offload-new-driver --no-offloadlib -fno-sycl-instrument-device-code -o %t_aot_gpu.o +// RUN: %clang %s -fsycl -fsycl-targets=intel_gpu_pvc,intel_gpu_dg2 -c --offload-new-driver --no-offloadlib -fno-sycl-instrument-device-code -o %t_aot_gpu_multi.o // RUN: %clang %s -fsycl -fsycl-targets=spir64_x86_64 -c --offload-new-driver --no-offloadlib -fno-sycl-instrument-device-code -o %t_aot_cpu.o // RUN: %clang %s -fsycl -fsycl-targets=nvptx64-nvidia-cuda --cuda-gpu-arch=sm_50 -nocudalib -fno-sycl-libspirv -c --offload-new-driver -o %t_nvptx.o // RUN: %clang %s -fsycl -fsycl-targets=amdgcn-amd-amdhsa -Xsycl-target-backend=amdgcn-amd-amdhsa --offload-arch=gfx803 -fgpu-rdc -nogpulib -fno-sycl-libspirv -c --offload-new-driver -o %t_amdgcn.o @@ -348,6 +349,16 @@ // CHK-FATBIN-ONLY-NEXT: "{{.*cp|copy}}" [[WRAPPEROUT]].bc {{.*}}.syclfb // CHK-FATBIN-ONLY-NOT: "{{.*}}/ld" +// With several GPU architectures, all their images are wrapped into a single +// wrapper module. +// RUN: clang-linker-wrapper --emit-fatbin-only --no-sycl-rdc --host-triple=x86_64-unknown-linux-gnu --linker-path=/usr/bin/ld -o %t.multi.syclfb %t_aot_gpu_multi.o --dry-run 2>&1 | FileCheck -check-prefix=CHK-FATBIN-ONLY-MULTI %s +// CHK-FATBIN-ONLY-MULTI: offload-wrapper: output: [[WRAPPEROUT:.*]].bc, input: {{.*}}, input: {{.*}} +// CHK-FATBIN-ONLY-MULTI-NOT: offload-wrapper: +// CHK-FATBIN-ONLY-MULTI: "{{.*cp|copy}}" [[WRAPPEROUT]].bc {{.*}}.multi.syclfb +// CHK-FATBIN-ONLY-MULTI-NOT: "{{.*}}/ld" +// RUN: not clang-linker-wrapper --emit-fatbin-only --no-sycl-rdc --host-triple=x86_64-unknown-linux-gnu --linker-path=/usr/bin/ld -o %t.mixed.syclfb %t_aot_gpu.o %t_aot_cpu.o --dry-run 2>&1 | FileCheck -check-prefix=CHK-FATBIN-ONLY-MIXED %s +// CHK-FATBIN-ONLY-MIXED: error: --no-sycl-rdc does not support multiple SYCL target triples + // Verify list of commands when syclbin is used // RUN: clang-linker-wrapper --bitcode-library=spir64-unknown-unknown=%t.devicelib.bc -sycl-post-link-options=SYCL_POST_LINK_OPTIONS --host-triple=x86_64-unknown-linux-gnu --linker-path=/usr/bin/ld -o /dev/null %t.o --dry-run -syclbin=executable 2>&1 | FileCheck --check-prefix CHK-SYCLBIN-CMDS %s // CHK-SYCLBIN-CMDS: spirv-to-ir-wrapper{{.*}} -o [[FIRSTLLVMLINKIN:.*]].bc --llvm-spirv-opts --spirv-preserve-auxdata --spirv-target-env=SPV-IR --spirv-builtin-format=global diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index d2e66ccf44dbf..9948bb5abcecb 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -1216,12 +1216,14 @@ static Error compressImages(SmallVectorImpl &Images, } /// Reads device images from the given \p InputFile and wraps them -/// in one LLVM IR Module as a constant data. +/// in one LLVM IR Module as a constant data. \p Target overrides the target +/// triple given by \p Args. /// /// \returns A path to the LLVM Module that contains wrapped images. Expected wrapSYCLBinariesFromFile(ArrayRef SplitModules, - const ArgList &Args, bool IsEmbeddedIR) { + const ArgList &Args, bool IsEmbeddedIR, + StringRef Target = "") { auto OutputFileOrErr = createOutputFile( sys::path::filename(ExecutableName) + ".sycl.image.wrapper", "bc"); if (!OutputFileOrErr) @@ -1246,7 +1248,8 @@ wrapSYCLBinariesFromFile(ArrayRef SplitModules, return OutputFilePath; } - StringRef Target = Args.getLastArgValue(OPT_triple_EQ); + if (Target.empty()) + Target = Args.getLastArgValue(OPT_triple_EQ); if (Target.empty()) return createStringError( inconvertibleErrorCode(), @@ -2955,6 +2958,12 @@ linkAndWrapDeviceFiles(ArrayRef> LinkerInputFiles, std::mutex SYCLBINModulesMtx; SmallVector SYCLBINModules; + // The -fno-sycl-rdc compile step wraps the modules of all architectures into + // a single wrapper module. Keep them per input group for a stable order. + std::vector> NoRDCModules( + LinkerInputFiles.size()); + std::string NoRDCTriple; + // Initialize the images with any overriding inputs. if (Args.hasArg(OPT_override_image)) if (Error Err = handleOverrideImages(Args, Images)) @@ -3060,14 +3069,15 @@ linkAndWrapDeviceFiles(ArrayRef> LinkerInputFiles, std::scoped_lock Guard(SYCLBINModulesMtx); SYCLBINModules.emplace_back(std::move(MD)); } else if (!NeedsWrapping) { - // -fno-sycl-rdc compile step: emit the wrapper module as bitcode; the - // host compilation links it in via -foffload-include-binary. - Expected OutputFile = - sycl::wrapSYCLBinariesFromFile(Modules, LinkerArgs, - /*IsEmbeddedIR=*/false); - if (!OutputFile) - return OutputFile.takeError(); - AppendImageToWrapperOutput(*OutputFile); + // -fno-sycl-rdc compile step: the modules are wrapped once all inputs + // are processed. + StringRef Triple = LinkerArgs.getLastArgValue(OPT_triple_EQ); + std::scoped_lock Guard(ImageMtx); + if (!NoRDCTriple.empty() && NoRDCTriple != Triple) + return createStringError( + "--no-sycl-rdc does not support multiple SYCL target triples"); + NoRDCModules[&Input - LinkerInputFiles.data()] = std::move(Modules); + NoRDCTriple = Triple.str(); } else { // TODO(NOM7): Remove this call and use community flow for bundle/wrap Expected OutputFile = @@ -3140,6 +3150,20 @@ linkAndWrapDeviceFiles(ArrayRef> LinkerInputFiles, if (Err) return std::move(Err); + // -fno-sycl-rdc compile step: emit the wrapper module as bitcode; the host + // compilation links it in via -foffload-include-binary. + if (!NoRDCTriple.empty()) { + std::vector Modules; + for (std::vector &GroupModules : NoRDCModules) + std::move(GroupModules.begin(), GroupModules.end(), + std::back_inserter(Modules)); + Expected OutputFile = sycl::wrapSYCLBinariesFromFile( + Modules, Args, /*IsEmbeddedIR=*/false, NoRDCTriple); + if (!OutputFile) + return OutputFile.takeError(); + WrappedOutput.push_back(*OutputFile); + } + if (OutputSYCLBIN) { auto OutputOrErr = sycl::packageSYCLBIN(SYCLBINState, SYCLBINModules); if (!OutputOrErr) From bc79849db3b7f6b3bfa402cec4bcd4fa22a680a2 Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Wed, 14 Oct 2026 09:20:30 -0700 Subject: [PATCH 16/22] [SYCL] Check the compile_target in compatibleWithDevice Move the compile_target check into compatibleWithDevice, so that every user of it only considers AOT images built for the device's architecture: device image dependencies, kernel bundle images, and hasCompatibleImage. hasCompatibleImage is used by the default device selector to prefer devices with an image. An image built only for intel_gpu_pvc counted as compatible with any Intel GPU, so on a machine with a TGLLP and a PVC GPU the TGLLP one could be selected, and launching the kernel then failed. Apply the same check when picking an image from a set, as done for device_global accesses, which only asked the backend and so could pick an image built for another architecture. --- .../program_manager/program_manager.cpp | 57 ++++++----- sycl/unittests/program_manager/CMakeLists.txt | 1 + .../CompileTargetSelection/CMakeLists.txt | 5 + .../CompileTargetSelection.cpp | 97 +++++++++++++++++++ .../DynamicLinkingCompileTarget.cpp | 41 ++++++++ 5 files changed, 176 insertions(+), 25 deletions(-) create mode 100644 sycl/unittests/program_manager/CompileTargetSelection/CMakeLists.txt create mode 100644 sycl/unittests/program_manager/CompileTargetSelection/CompileTargetSelection.cpp diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index c7ef89c5644c6..580ad93463ac8 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -547,8 +547,23 @@ static const char *getUrDeviceTarget(const char *URDeviceTarget) { return UR_DEVICE_BINARY_TARGET_UNKNOWN; } +// An image with a compile_target was built for that device architecture only. +// Images without one are matched by their target. +static bool doesCompileTargetMatchDevice(const RTDeviceBinaryImage &Img, + const device_impl &Dev) { + const RTDeviceBinaryImage::PropertyRange &Range = Img.getDeviceRequirements(); + const bool HasCompileTarget = + std::any_of(Range.begin(), Range.end(), [](const auto &Prop) { + return Prop->Name == std::string_view("compile_target"); + }); + return !HasCompileTarget || doesImageTargetMatchDevice(Img, Dev); +} + static bool compatibleWithDevice(const RTDeviceBinaryImage *BinImage, const device_impl &DeviceImpl) { + if (!doesCompileTargetMatchDevice(*BinImage, DeviceImpl)) + return false; + adapter_impl &Adapter = DeviceImpl.getAdapter(); const ur_device_handle_t &URDeviceHandle = DeviceImpl.getHandleRef(); @@ -573,18 +588,6 @@ static bool compatibleWithDevice(const RTDeviceBinaryImage *BinImage, return (0 == SuitableImageID); } -// An image with a compile_target was built for that device architecture only. -// Images without one are matched by compatibleWithDevice. -static bool doesCompileTargetMatchDevice(const RTDeviceBinaryImage &Img, - const device_impl &Dev) { - const RTDeviceBinaryImage::PropertyRange &Range = Img.getDeviceRequirements(); - const bool HasCompileTarget = - std::any_of(Range.begin(), Range.end(), [](const auto &Prop) { - return Prop->Name == std::string_view("compile_target"); - }); - return !HasCompileTarget || doesImageTargetMatchDevice(Img, Dev); -} - // Check if the device image is a BF16 devicelib image. bool ProgramManager::isBfloat16DeviceImage( const RTDeviceBinaryImage *BinImage) { @@ -705,8 +708,7 @@ ProgramManager::collectDeviceImageDepsForImportedSymbols( const RTDeviceBinaryImage *Img = It->second; if (!doesDevSupportDeviceRequirements(Dev, *Img) || - !compatibleWithDevice(Img, Dev) || - !doesCompileTargetMatchDevice(*Img, Dev)) + !compatibleWithDevice(Img, Dev)) continue; // If the image is a BF16 device image, we need to check if it @@ -1456,11 +1458,18 @@ const RTDeviceBinaryImage &ProgramManager::getDeviceImage( } std::lock_guard ImgMapsGuard(m_ImgMapsMutex); - std::vector RawImgs(ImageSet.size()); - auto ImageIterator = ImageSet.begin(); - for (size_t i = 0; i < ImageSet.size(); i++, ImageIterator++) + // Only consider images built for the device's architecture, if any. + std::vector Imgs; + for (const RTDeviceBinaryImage *Img : ImageSet) + if (doesCompileTargetMatchDevice(*Img, DeviceImpl)) + Imgs.push_back(Img); + if (Imgs.empty()) + Imgs.assign(ImageSet.begin(), ImageSet.end()); + + std::vector RawImgs(Imgs.size()); + for (size_t i = 0; i < Imgs.size(); i++) RawImgs[i] = reinterpret_cast( - const_cast(&(*ImageIterator)->getRawData())); + const_cast(&Imgs[i]->getRawData())); uint32_t ImgInd = 0; // Ask the native runtime under the given context to choose the device image // it prefers. @@ -1474,15 +1483,14 @@ const RTDeviceBinaryImage &ProgramManager::getDeviceImage( ContextImpl.getAdapter().call( DeviceImpl.getHandleRef(), UrBinaries.data(), UrBinaries.size(), &ImgInd); - ImageIterator = ImageSet.begin(); - std::advance(ImageIterator, ImgInd); + const RTDeviceBinaryImage *SelectedImg = Imgs[ImgInd]; if constexpr (DbgProgMgr > 0) { - std::cerr << "selected device image: " << &(*ImageIterator)->getRawData() + std::cerr << "selected device image: " << &SelectedImg->getRawData() << "\n"; - (*ImageIterator)->print(); + SelectedImg->print(); } - return **ImageIterator; + return *SelectedImg; } // Check if device image is compressed. @@ -2330,8 +2338,7 @@ ProgramManager::getSYCLDeviceImagesWithCompatibleState( for (const RTDeviceBinaryImage *BinImage : BinImages) { if (!compatibleWithDevice(BinImage, Dev) || - !doesDevSupportDeviceRequirements(Dev, *BinImage) || - !doesCompileTargetMatchDevice(*BinImage, Dev)) + !doesDevSupportDeviceRequirements(Dev, *BinImage)) continue; auto InsertRes = ImageInfoMap.insert({BinImage, {}}); diff --git a/sycl/unittests/program_manager/CMakeLists.txt b/sycl/unittests/program_manager/CMakeLists.txt index a88f3495768c0..d88c60bceea66 100644 --- a/sycl/unittests/program_manager/CMakeLists.txt +++ b/sycl/unittests/program_manager/CMakeLists.txt @@ -21,3 +21,4 @@ endif() add_subdirectory(arg_mask) add_subdirectory(DynamicLinking) +add_subdirectory(CompileTargetSelection) diff --git a/sycl/unittests/program_manager/CompileTargetSelection/CMakeLists.txt b/sycl/unittests/program_manager/CompileTargetSelection/CMakeLists.txt new file mode 100644 index 0000000000000..51cad31e87fe3 --- /dev/null +++ b/sycl/unittests/program_manager/CompileTargetSelection/CMakeLists.txt @@ -0,0 +1,5 @@ +# Uses its own test binary: the registered images decide which devices have a +# compatible image, so other tests' images must not be visible. +add_sycl_unittest(CompileTargetSelectionTests OBJECT + CompileTargetSelection.cpp +) diff --git a/sycl/unittests/program_manager/CompileTargetSelection/CompileTargetSelection.cpp b/sycl/unittests/program_manager/CompileTargetSelection/CompileTargetSelection.cpp new file mode 100644 index 0000000000000..5eb6315459f49 --- /dev/null +++ b/sycl/unittests/program_manager/CompileTargetSelection/CompileTargetSelection.cpp @@ -0,0 +1,97 @@ +#include + +#include +#include +#include + +#include + +// Check that an AOT image built for one GPU architecture only makes devices of +// that architecture preferred by the default device selector. + +class CompileTargetSelectionKernel; + +namespace sycl { +inline namespace _V1 { +namespace detail { +template <> +struct KernelInfo + : public unittest::MockKernelInfoBase { + static constexpr const char *getName() { + return "CompileTargetSelectionKernel"; + } +}; +} // namespace detail +} // namespace _V1 +} // namespace sycl + +namespace { +namespace exp = sycl::ext::oneapi::experimental; + +sycl::unittest::MockDeviceImage generatePVCImage() { + std::string CompileTarget = "intel_gpu_pvc"; + std::vector Data(8 + CompileTarget.size()); + std::copy(CompileTarget.begin(), CompileTarget.end(), Data.data() + 8); + sycl::unittest::MockPropertySet PropSet; + PropSet.insert(__SYCL_PROPERTY_SET_SYCL_DEVICE_REQUIREMENTS, + sycl::unittest::MockProperty("compile_target", Data, + SYCL_PROPERTY_TYPE_BYTE_ARRAY)); + return sycl::unittest::MockDeviceImage{ + SYCL_DEVICE_BINARY_TYPE_NATIVE, + __SYCL_DEVICE_BINARY_TARGET_SPIRV64_GEN, + "", // Compile options + "", // Link options + std::vector{1}, + sycl::unittest::makeEmptyKernels({"CompileTargetSelectionKernel"}), + std::move(PropSet)}; +} +sycl::unittest::MockDeviceImage Imgs[] = {generatePVCImage()}; +sycl::unittest::MockDeviceImageArray<1> ImgArray{Imgs}; + +// IP versions of a TGLLP and a PVC GPU. The TGLLP one comes first so that it +// would be chosen if both devices were preferred equally. +uint32_t MockDeviceIPs[] = {0x03000000, 0x030f0000}; + +ur_result_t redefinedDeviceGet(void *pParams) { + auto Params = *static_cast(pParams); + if (*Params.ppNumDevices) { + **Params.ppNumDevices = std::size(MockDeviceIPs); + return UR_RESULT_SUCCESS; + } + if (*Params.pphDevices) + for (uint32_t I = 0; I < *Params.pNumEntries; ++I) + (*Params.pphDevices)[I] = + reinterpret_cast(&MockDeviceIPs[I]); + return UR_RESULT_SUCCESS; +} + +ur_result_t redefinedDeviceGetInfo(void *pParams) { + auto Params = *static_cast(pParams); + if (*Params.ppropName == UR_DEVICE_INFO_IP_VERSION && *Params.ppPropValue) + *static_cast(*Params.ppPropValue) = + *reinterpret_cast(*Params.phDevice); + if (*Params.ppropName == UR_DEVICE_INFO_TYPE && *Params.ppPropValue) + *static_cast(*Params.ppPropValue) = UR_DEVICE_TYPE_GPU; + return UR_RESULT_SUCCESS; +} +} // namespace + +TEST(CompileTargetSelection, DefaultSelectorPrefersDeviceWithImage) { + sycl::unittest::UrMock Mock; + mock::getCallbacks().set_replace_callback("urDeviceGet", &redefinedDeviceGet); + mock::getCallbacks().set_after_callback("urDeviceGetInfo", + &redefinedDeviceGetInfo); + + std::vector Devs = sycl::platform().get_devices(); + ASSERT_EQ(Devs.size(), 2u); + ASSERT_EQ(Devs[0].get_info(), + exp::architecture::intel_gpu_tgllp); + ASSERT_EQ(Devs[1].get_info(), + exp::architecture::intel_gpu_pvc); + EXPECT_LT(sycl::default_selector_v(Devs[0]), + sycl::default_selector_v(Devs[1])); + + sycl::device Dev{sycl::default_selector_v}; + EXPECT_EQ(Dev.get_info(), + exp::architecture::intel_gpu_pvc); +} diff --git a/sycl/unittests/program_manager/DynamicLinking/DynamicLinkingCompileTarget.cpp b/sycl/unittests/program_manager/DynamicLinking/DynamicLinkingCompileTarget.cpp index afe7bd5a94b46..228a7b28a7c01 100644 --- a/sycl/unittests/program_manager/DynamicLinking/DynamicLinkingCompileTarget.cpp +++ b/sycl/unittests/program_manager/DynamicLinking/DynamicLinkingCompileTarget.cpp @@ -1,5 +1,7 @@ #include +#include + #include #include #include @@ -138,3 +140,42 @@ TEST(DynamicLinking, AheadOfTimeKernelBundleForDevice) { EXPECT_EQ(CapturedLinkingData.ProgramUsedToCreateKernel, PVC_PRG * PVC_PRG_DEP); } + +// Images are also picked from a given set, e.g. for device_global accesses. +// The selection by the backend only sees the image targets, which are the same +// for all images here, so make it pick the TGLLP image whenever it is offered. +std::vector OfferedImgs; +ur_result_t redefinedDeviceSelectBinary(void *pParams) { + auto Params = *static_cast(pParams); + **Params.ppSelectedBinary = 0; + for (uint32_t I = 0; I < *Params.pNumBinaries; ++I) + if (OfferedImgs.size() == *Params.pNumBinaries && + *OfferedImgs[I]->getRawData().BinaryStart == TGLLP_PRG) + **Params.ppSelectedBinary = I; + return UR_RESULT_SUCCESS; +} + +TEST(DynamicLinking, AheadOfTimeImageFromSetForDevice) { + sycl::unittest::UrMock Mock; + mock::getCallbacks().set_after_callback("urDeviceGetInfo", + &redefinedDeviceGetInfo); + + sycl::queue Q(sycl::platform().get_devices()[0]); + auto &PM = sycl::detail::ProgramManager::getInstance(); + std::set KernelImgs = + PM.getRawDeviceImages({sycl::get_kernel_id< + DynamicLinkingCompileTargetTest::MultiArchKernel>()}); + ASSERT_EQ(KernelImgs.size(), 2u); + + std::unordered_set ImgSet( + KernelImgs.begin(), KernelImgs.end()); + // The images offered to the backend if none are filtered out. + OfferedImgs.assign(ImgSet.begin(), ImgSet.end()); + mock::getCallbacks().set_replace_callback("urDeviceSelectBinary", + &redefinedDeviceSelectBinary); + + const sycl::detail::RTDeviceBinaryImage &Img = + PM.getDeviceImage(ImgSet, *sycl::detail::getSyclObjImpl(Q.get_context()), + *sycl::detail::getSyclObjImpl(Q.get_device())); + EXPECT_EQ(*Img.getRawData().BinaryStart, PVC_PRG); +} From 5c2cc9698e9a4e280a5247c48fe86a3ad422cf5d Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Wed, 14 Oct 2026 11:31:07 -0700 Subject: [PATCH 17/22] [SYCL][E2E] Remove the stale XFAIL from the AddressSanitizer AOT GPU tests They were expected to fail on PVC with the new offloading model because clang-linker-wrapper crashed for -fsycl-targets=spir64_gen with -Xsycl-target-backend=spir64_gen '-device pvc,dg2' (#23265). The tests now build with -fsycl-targets=intel_gpu_pvc,intel_gpu_dg2 and pass there. --- sycl/test-e2e/AddressSanitizer/aot/gpu.cpp | 3 --- .../test-e2e/AddressSanitizer/common/kernel-filter-aot-gpu.cpp | 3 --- 2 files changed, 6 deletions(-) diff --git a/sycl/test-e2e/AddressSanitizer/aot/gpu.cpp b/sycl/test-e2e/AddressSanitizer/aot/gpu.cpp index 4f302070e9113..f2dfab2a5a8ef 100644 --- a/sycl/test-e2e/AddressSanitizer/aot/gpu.cpp +++ b/sycl/test-e2e/AddressSanitizer/aot/gpu.cpp @@ -1,9 +1,6 @@ // REQUIRES: linux, gpu && level_zero // REQUIRES: (arch-intel_gpu_pvc || gpu-intel-dg2) -// XFAIL: arch-intel_gpu_pvc && new-offload-model -// XFAIL-TRACKER: https://github.com/intel/llvm/issues/23265 - // RUN: %{run-aux} %{build} %device_asan_aot_gpu_flags -O0 -g %S/Inputs/host-usm-oob.cpp -o %t.out // RUN: %{run} not --crash %t.out 2>&1 | FileCheck %S/Inputs/host-usm-oob.cpp diff --git a/sycl/test-e2e/AddressSanitizer/common/kernel-filter-aot-gpu.cpp b/sycl/test-e2e/AddressSanitizer/common/kernel-filter-aot-gpu.cpp index 9e9cb59c46581..c3bb6d6f8b9bb 100644 --- a/sycl/test-e2e/AddressSanitizer/common/kernel-filter-aot-gpu.cpp +++ b/sycl/test-e2e/AddressSanitizer/common/kernel-filter-aot-gpu.cpp @@ -1,9 +1,6 @@ // REQUIRES: linux, gpu && level_zero // REQUIRES: (arch-intel_gpu_pvc || gpu-intel-dg2) -// XFAIL: arch-intel_gpu_pvc && new-offload-model -// XFAIL-TRACKER: https://github.com/intel/llvm/issues/23265 - // RUN: %{run-aux} %{build} %device_asan_aot_gpu_flags %S/kernel-filter.cpp -g -O2 -fsanitize-ignorelist=%p/ignorelist.txt -o %t1 // RUN: %{run} %t1 2>&1 | FileCheck %S/kernel-filter.cpp --check-prefixes CHECK-IGNORE // RUN: %{run-aux} %{build} %device_asan_aot_gpu_flags %S/kernel-filter.cpp -g -O2 -o %t2 From 4fdb976633d2c0a52fa76069c6cef8e0a51e685c Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Wed, 14 Oct 2026 11:44:54 -0700 Subject: [PATCH 18/22] [SYCL][E2E] Build bfloat16_example_aot.cpp for gpu_aot_target The test isn't specific to DG2 but was converted from spir64_gen -device dg2 to intel_gpu_dg2. Unlike the spir64_gen image, an intel_gpu_dg2 image is only used on DG2, so the test failed on other GPUs, e.g. on BMG with "has no image for the selected device". --- sycl/test-e2e/BFloat16/bfloat16_example_aot.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/test-e2e/BFloat16/bfloat16_example_aot.cpp b/sycl/test-e2e/BFloat16/bfloat16_example_aot.cpp index 2d3de11d6f75e..c63841b4a9a72 100644 --- a/sycl/test-e2e/BFloat16/bfloat16_example_aot.cpp +++ b/sycl/test-e2e/BFloat16/bfloat16_example_aot.cpp @@ -8,11 +8,11 @@ // RUN: %clangxx -fsycl -fsycl-targets=spir64 %s -o %t.out // RUN: %{run} %t.out -// RUN: %clangxx -fsycl -fsycl-targets=spir64,intel_gpu_dg2 %s -o %t.out +// RUN: %clangxx -fsycl -fsycl-targets=spir64,%{gpu_aot_target} %s -o %t.out // RUN: %{run} %t.out // CPU AOT targets host isa, so we compile on the run system instead. -// RUN: %{run-aux} %clangxx -fsycl -fsycl-targets=spir64_x86_64,intel_gpu_dg2 %s -o %t.out +// RUN: %{run-aux} %clangxx -fsycl -fsycl-targets=spir64_x86_64,%{gpu_aot_target} %s -o %t.out // RUN: %{run} %t.out #include "bfloat16_example.hpp" From 0c428da71de627bfa132f7c7f3259600bdbced6c Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Wed, 14 Oct 2026 12:05:22 -0700 Subject: [PATCH 19/22] [SYCL][Driver] Rename err_drv_sycl_spir64_gen_new_driver to err_drv_not_supported_new_model The diagnostic text is not specific to spir64_gen. --- clang/include/clang/Basic/DiagnosticDriverKinds.td | 2 +- clang/lib/Driver/Driver.cpp | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/clang/include/clang/Basic/DiagnosticDriverKinds.td b/clang/include/clang/Basic/DiagnosticDriverKinds.td index 58a95a955efb2..6aa759e5cc49e 100644 --- a/clang/include/clang/Basic/DiagnosticDriverKinds.td +++ b/clang/include/clang/Basic/DiagnosticDriverKinds.td @@ -483,7 +483,7 @@ def err_drv_sycl_thinlto_split_off: Error< "'%0' is not supported when '%1' is set with '-fsycl'">; def err_drv_sycl_offload_arch_new_driver: Error< "'--offload-arch' is supported when '-fsycl' is set with '--offload-new-driver'">; -def err_drv_sycl_spir64_gen_new_driver : Error< +def err_drv_not_supported_new_model : Error< "'%0' is not supported with the new offloading model; use '%1' instead">; def err_drv_sycl_offload_arch_missing_value : Error< "must pass in an explicit cpu or gpu architecture to '--offload-arch'">; diff --git a/clang/lib/Driver/Driver.cpp b/clang/lib/Driver/Driver.cpp index a96810be505b0..01cb7352ed635 100644 --- a/clang/lib/Driver/Driver.cpp +++ b/clang/lib/Driver/Driver.cpp @@ -1300,7 +1300,7 @@ static void diagnoseSYCLOptions(Compilation &C, bool IsSYCL) { std::string Opt = (Spelling + Val).str(); std::string Replacement = (Spelling + "intel_gpu_").str(); if (D.getUseNewOffloadingDriver()) - D.Diag(diag::err_drv_sycl_spir64_gen_new_driver) << Opt << Replacement; + D.Diag(diag::err_drv_not_supported_new_model) << Opt << Replacement; else D.Diag(diag::warn_drv_deprecated_option_release) << Opt << /*HasReplacement=*/true << Replacement; From 4e9ca412bde1a0a41907ab4b4d6cdf1332805990 Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Tue, 20 Oct 2026 10:41:28 -0700 Subject: [PATCH 20/22] [SYCL][Doc] Document intel_gpu_ instead of spir64_gen in the users manual Mention spir64_gen only in a deprecation note at the end of the -fsycl-targets description. --- sycl/doc/UsersManual.md | 35 +++++++++++++++++------------------ 1 file changed, 17 insertions(+), 18 deletions(-) diff --git a/sycl/doc/UsersManual.md b/sycl/doc/UsersManual.md index ea8fafa991174..965fb2f1e4590 100644 --- a/sycl/doc/UsersManual.md +++ b/sycl/doc/UsersManual.md @@ -33,22 +33,10 @@ and not recommended to use in production environment. The following triples are supported by default: * spir64 - this is the default generic SPIR-V target; * spir64_x86_64 - generate code ahead of time for x86_64 CPUs; - * spir64_gen - generate code ahead of time for Intel Processor Graphics. - Deprecated: use the intel_gpu_* values listed below instead. It emits a - deprecation warning by default, and is an error with - --offload-new-driver. To migrate - '-fsycl-targets=spir64_gen -Xsycl-target-backend "-device "': - * a device name, e.g. "-device pvc": use -fsycl-targets=intel_gpu_pvc; - * a full IP version, e.g. "-device 12.60.7": use - -fsycl-targets=intel_gpu_12_60_7; - * a list, e.g. "-device pvc,dg2": use - -fsycl-targets=intel_gpu_pvc,intel_gpu_dg2; - * "-device *", hex device IDs (e.g. 0x0BD5), partial versions (e.g. - 12.60) and devices without an intel_gpu_* value (e.g. pvc-sdv) have - no equivalent; list the intel_gpu_* values of the devices to target; + * intel_gpu_ - generate code ahead of time for Intel Processor + Graphics; see the list of intel_gpu_* values below; Full target triples can also be used: - * spir64-unknown-unknown, spir64_x86_64-unknown-unknown, - spir64_gen-unknown-unknown + * spir64-unknown-unknown, spir64_x86_64-unknown-unknown Available in special build configuration: * nvptx64-nvidia-cuda - generate code ahead of time for CUDA target; * amdgcn-amd-amdhsa - generate code ahead of time for HIP target; @@ -57,9 +45,7 @@ and not recommended to use in production environment. currently overrides all the other specified SYCL targets when enabled.) Special target values specific to Intel, NVIDIA and AMD Processor Graphics - support are accepted, providing a streamlined interface for AOT. The - intel_gpu_* values are the supported way to target Intel GPUs ahead of - time. + support are accepted, providing a streamlined interface for AOT. A comma-separated list of valid Intel, NVIDIA and AMD Processor Graphics values is supported. * intel_gpu_cri, intel_gpu_35_11_0 - Crescent Island Intel graphics architecture * intel_gpu_nvl_p, intel_gpu_35_10_0 - Nova Lake P Intel graphics architecture @@ -160,6 +146,19 @@ and not recommended to use in production environment. * amd_gpu_gfx1200 - GCN GFX12 (RDNA 4) architecture * amd_gpu_gfx1201 - GCN GFX12 (RDNA 4) architecture + Deprecated: spir64_gen (spir64_gen-unknown-unknown) is deprecated in favor + of the intel_gpu_* values above. It emits a deprecation warning by default, + and is an error with --offload-new-driver. To migrate + '-fsycl-targets=spir64_gen -Xsycl-target-backend "-device "': + * a device name, e.g. "-device pvc": use -fsycl-targets=intel_gpu_pvc; + * a full IP version, e.g. "-device 12.60.7": use + -fsycl-targets=intel_gpu_12_60_7; + * a list, e.g. "-device pvc,dg2": use + -fsycl-targets=intel_gpu_pvc,intel_gpu_dg2; + * "-device *", hex device IDs (e.g. 0x0BD5), partial versions (e.g. + 12.60) and devices without an intel_gpu_* value (e.g. pvc-sdv) have + no equivalent; list the intel_gpu_* values of the devices to target. + ## Language options **`-sycl-std=`** [EXPERIMENTAL] From 0af2c3248d4378fb43bae39112c819a15270aee8 Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Tue, 20 Oct 2026 10:48:03 -0700 Subject: [PATCH 21/22] [SYCL][Driver] Name the SYCL target in the spir64_gen check and test it without -device --- clang/lib/Driver/Driver.cpp | 16 +++++++++------- clang/test/Driver/sycl-spir64-gen-deprecated.cpp | 5 +++++ 2 files changed, 14 insertions(+), 7 deletions(-) diff --git a/clang/lib/Driver/Driver.cpp b/clang/lib/Driver/Driver.cpp index 01cb7352ed635..7693c36c2f9bf 100644 --- a/clang/lib/Driver/Driver.cpp +++ b/clang/lib/Driver/Driver.cpp @@ -1292,12 +1292,14 @@ static void diagnoseSYCLOptions(Compilation &C, bool IsSYCL) { // old offloading model. if (IsSYCL) { const Driver &D = C.getDriver(); - auto IsSPIRGenTarget = [&](StringRef Val) { - llvm::Triple TT = D.getSYCLDeviceTriple(Val); + // Target is a SYCL target as given on the command line, e.g. spir64_gen, + // spir64_gen-unknown-unknown or intel_gpu_pvc. + auto IsSPIRGenTarget = [&](StringRef Target) { + llvm::Triple TT = D.getSYCLDeviceTriple(Target); return TT.isSPIR() && TT.getSubArch() == llvm::Triple::SPIRSubArch_gen; }; - auto DiagnoseSPIRGenTarget = [&](StringRef Spelling, StringRef Val) { - std::string Opt = (Spelling + Val).str(); + auto DiagnoseSPIRGenTarget = [&](StringRef Spelling, StringRef Target) { + std::string Opt = (Spelling + Target).str(); std::string Replacement = (Spelling + "intel_gpu_").str(); if (D.getUseNewOffloadingDriver()) D.Diag(diag::err_drv_not_supported_new_model) << Opt << Replacement; @@ -1311,9 +1313,9 @@ static void diagnoseSYCLOptions(Compilation &C, bool IsSYCL) { // Report the spelling the user wrote, e.g. -fsycl-targets=. StringRef Spelling = A->getAlias() ? A->getAlias()->getSpelling() : A->getSpelling(); - for (StringRef Val : A->getValues()) - if (IsSPIRGenTarget(Val)) - DiagnoseSPIRGenTarget(Spelling, Val); + for (StringRef Target : A->getValues()) + if (IsSPIRGenTarget(Target)) + DiagnoseSPIRGenTarget(Spelling, Target); } for (const Arg *A : C.getInputArgs().filtered( options::OPT_Xsycl_backend_EQ, options::OPT_Xsycl_frontend_EQ, diff --git a/clang/test/Driver/sycl-spir64-gen-deprecated.cpp b/clang/test/Driver/sycl-spir64-gen-deprecated.cpp index 4d1152874e2e1..0f6da18c43e10 100644 --- a/clang/test/Driver/sycl-spir64-gen-deprecated.cpp +++ b/clang/test/Driver/sycl-spir64-gen-deprecated.cpp @@ -16,6 +16,9 @@ // RUN: -fsycl-targets=intel_gpu_pvc,spir64_gen \ // RUN: -Xsycl-target-backend=spir64_gen "-device dg2" %s 2>&1 \ // RUN: | FileCheck %s --check-prefixes=WARN,WARN-XS -DOPT=-fsycl-targets= +// Diagnosed also without a -device option, e.g. at a compile-only (-c) step. +// RUN: %clangxx -### -fsycl --no-offload-new-driver -fsycl-targets=spir64_gen \ +// RUN: -c %s 2>&1 | FileCheck %s --check-prefix=WARN -DOPT=-fsycl-targets= // WARN: warning: option '[[OPT]]spir64_gen' is deprecated and will be removed in a future release, use '[[OPT]]intel_gpu_' instead [-Wdeprecated] // WARN-TRIPLE: warning: option '-fsycl-targets=spir64_gen-unknown-unknown' is deprecated and will be removed in a future release, use '-fsycl-targets=intel_gpu_' instead [-Wdeprecated] // WARN-XS: warning: option '-Xsycl-target-backend=spir64_gen' is deprecated and will be removed in a future release, use '-Xsycl-target-backend=intel_gpu_' instead [-Wdeprecated] @@ -51,6 +54,8 @@ // RUN: not %clangxx -### -fsycl --offload-new-driver -fsycl-targets=intel_gpu_pvc \ // RUN: -Xsycl-target-backend=spir64_gen "-options -extra" %s 2>&1 \ // RUN: | FileCheck %s --check-prefix=ERR-XS +// RUN: not %clangxx -### -fsycl --offload-new-driver -fsycl-targets=spir64_gen \ +// RUN: -c %s 2>&1 | FileCheck %s --check-prefix=ERR // ERR: error: '-fsycl-targets=spir64_gen' is not supported with the new offloading model; use '-fsycl-targets=intel_gpu_' instead // ERR-TRIPLE: error: '-fsycl-targets=spir64_gen-unknown-unknown' is not supported with the new offloading model; use '-fsycl-targets=intel_gpu_' instead // ERR-XS: error: '-Xsycl-target-backend=spir64_gen' is not supported with the new offloading model; use '-Xsycl-target-backend=intel_gpu_' instead From fadb0a3d3eac7701bf9b19cbd9afa7eef74fcbc5 Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Tue, 20 Oct 2026 13:58:23 -0700 Subject: [PATCH 22/22] [SYCL][Driver] Say spir64_gen will be removed with the old offloading model --- clang/lib/Driver/Driver.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/clang/lib/Driver/Driver.cpp b/clang/lib/Driver/Driver.cpp index 7693c36c2f9bf..792b0f5f21388 100644 --- a/clang/lib/Driver/Driver.cpp +++ b/clang/lib/Driver/Driver.cpp @@ -1288,8 +1288,8 @@ static void diagnoseSYCLOptions(Compilation &C, bool IsSYCL) { // spir64_gen is deprecated in favor of intel_gpu_, both as a // -fsycl-targets value and as the triple of the -X= options. It is not - // supported with the new offloading model, and is removed together with the - // old offloading model. + // supported with the new offloading model, and will be removed together with + // the old offloading model. if (IsSYCL) { const Driver &D = C.getDriver(); // Target is a SYCL target as given on the command line, e.g. spir64_gen,