From 20f67e14beeae239ea6262323da8a5f43c5656e4 Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Fri, 8 May 2026 09:57:25 -0700 Subject: [PATCH 1/5] Only have option forwarding and test changes. --- clang/test/Driver/clang-sycl-linker-test.cpp | 49 +++++++++++++- .../ClangLinkerWrapper.cpp | 67 +++++++++++++++++++ 2 files changed, 115 insertions(+), 1 deletion(-) diff --git a/clang/test/Driver/clang-sycl-linker-test.cpp b/clang/test/Driver/clang-sycl-linker-test.cpp index 74543885e32af..5e0afa1bddc7b 100644 --- a/clang/test/Driver/clang-sycl-linker-test.cpp +++ b/clang/test/Driver/clang-sycl-linker-test.cpp @@ -20,7 +20,7 @@ // RUN: touch %t.dir/lib2.bc // RUN: clang-sycl-linker --dry-run -v -triple=spirv64 %t_1.bc %t_2.bc --library-path=%t.dir --device-libs=lib1.bc,lib2.bc -o a.spv 2>&1 \ // RUN: | FileCheck %s --check-prefix=DEVLIBS -// DEVLIBS: sycl-device-link: inputs: {{.*}}.bc libfiles: {{.*}}lib1.bc, {{.*}}lib2.bc output: [[LLVMLINKOUT:.*]].bc +// DEVLIBS: sycl-device-link: inputs: {{.*}}.bc, {{.*}}.bc libfiles: {{.*}}lib1.bc, {{.*}}lib2.bc output: [[LLVMLINKOUT:.*]].bc // DEVLIBS-NEXT: LLVM backend: input: [[LLVMLINKOUT]].bc, output: a_0.spv // // Test a simple case with a random file (not bitcode) as input. @@ -69,3 +69,50 @@ // RUN: not clang-sycl-linker --dry-run %t_1.bc %t_2.bc -o a.out 2>&1 \ // RUN: | FileCheck %s --check-prefix=NOTARGET // NOTARGET: Target triple must be specified +// +// Test that driver options are correctly passed through clang-linker-wrapper. +// RUN: %clangxx -### -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown \ +// RUN: -Xsycl-target-backend=spir64-unknown-unknown "-backend-opt" \ +// RUN: %t_1.bc %t_2.bc 2>&1 | FileCheck %s --check-prefix=CHECK-BACKEND-OPT +// CHECK-BACKEND-OPT: clang-linker-wrapper{{.*}} "--device-compiler=sycl:spir64-unknown-unknown=-backend-opt" +// +// Test multiple backend options. +// RUN: %clangxx -### -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown \ +// RUN: -Xsycl-target-backend=spir64-unknown-unknown "-opt1 -opt2" \ +// RUN: %t_1.bc %t_2.bc 2>&1 | FileCheck %s --check-prefix=CHECK-MULTI-OPTS +// CHECK-MULTI-OPTS: clang-linker-wrapper{{.*}} "--device-compiler=sycl:spir64-unknown-unknown=-opt1 -opt2" +// +// Test linker options are forwarded. +// RUN: %clangxx -### -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown \ +// RUN: -Xsycl-target-linker=spir64-unknown-unknown "-linker-opt" \ +// RUN: %t_1.bc %t_2.bc 2>&1 | FileCheck %s --check-prefix=CHECK-LINKER-OPT +// CHECK-LINKER-OPT: clang-linker-wrapper{{.*}} "--device-linker=sycl:spir64-unknown-unknown=-linker-opt" +// +// Test both backend and linker options together. +// RUN: %clangxx -### -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown \ +// RUN: -Xsycl-target-backend=spir64-unknown-unknown "-backend-opt" \ +// RUN: -Xsycl-target-linker=spir64-unknown-unknown "-linker-opt" \ +// RUN: %t_1.bc %t_2.bc 2>&1 | FileCheck %s --check-prefix=CHECK-BOTH-OPTS +// CHECK-BOTH-OPTS: clang-linker-wrapper{{.*}} "--device-compiler=sycl:spir64-unknown-unknown=-backend-opt" +// CHECK-BOTH-OPTS-SAME: "--device-linker=sycl:spir64-unknown-unknown=-linker-opt" +// +// Test AOT GPU with backend options. +// RUN: %clangxx -### -fsycl --offload-new-driver -fsycl-targets=spir64_gen-unknown-unknown \ +// RUN: -Xsycl-target-backend=spir64_gen "-device pvc" \ +// RUN: %t_1.bc %t_2.bc 2>&1 | FileCheck %s --check-prefix=CHECK-AOT-GPU +// CHECK-AOT-GPU: clang-linker-wrapper{{.*}} "--device-compiler=sycl:spir64_gen-unknown-unknown=-device pvc" +// +// Test AOT CPU with backend options. +// RUN: %clangxx -### -fsycl --offload-new-driver -fsycl-targets=spir64_x86_64-unknown-unknown \ +// RUN: -Xsycl-target-backend=spir64_x86_64 "-march=skylake" \ +// RUN: %t_1.bc %t_2.bc 2>&1 | FileCheck %s --check-prefix=CHECK-AOT-CPU +// CHECK-AOT-CPU: clang-linker-wrapper{{.*}} "--device-compiler=sycl:spir64_x86_64-unknown-unknown=-march=skylake" +// +// Test multi-target with target-specific options (both in same invocation). +// RUN: %clangxx -### -fsycl --offload-new-driver \ +// RUN: -fsycl-targets=spir64_gen-unknown-unknown,spir64_x86_64-unknown-unknown \ +// RUN: -Xsycl-target-backend=spir64_gen "-device pvc" \ +// RUN: -Xsycl-target-backend=spir64_x86_64 "-march=skylake" \ +// RUN: %t_1.bc %t_2.bc 2>&1 | FileCheck %s --check-prefix=CHECK-MULTI-TARGET +// CHECK-MULTI-TARGET: clang-linker-wrapper{{.*}} "--device-compiler=sycl:spir64_gen-unknown-unknown=-device pvc" +// CHECK-MULTI-TARGET-SAME: "--device-compiler=sycl:spir64_x86_64-unknown-unknown=-march=skylake" diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index 2a9950d2154be..07ea58d0292eb 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -160,6 +160,10 @@ static bool CanonicalPrefixes = true; using OffloadingImage = OffloadBinary::OffloadingImage; +// TODO: Remove this once the linking and backend compilation have been fully +// migrated to clang-sycl-linker. +bool UseClangSYCLLinker = false; + namespace llvm { // Provide DenseMapInfo so that OffloadKind can be used in a DenseMap. template <> struct DenseMapInfo { @@ -1679,6 +1683,58 @@ Expected bundleDeviceModule(StringRef ClangOutput, } // namespace sycl namespace generic { +// Forwards all required flags to the clang-sycl-linker via -Xlinker. +static void appendClangSYCLLinkerArgs(const ArgList &Args, + SmallVectorImpl &CmdArgs, + const llvm::Triple &Triple, + StringRef Arch) { + auto XLinker = [&](const Twine &Arg) { + CmdArgs.append({"-Xlinker", Args.MakeArgString(Arg)}); + }; + + XLinker("--triple=" + Triple.getTriple()); + if (!Arch.empty()) + XLinker("--arch=" + Arch); + + static const OptSpecifier DirectOpts[] = { + OPT_save_temps, + OPT_dry_run, + OPT_print_wrapped_module, + OPT_sycl_thin_lto, + OPT_sycl_allow_device_image_dependencies, + OPT_sycl_remove_unused_external_funcs, + OPT_no_sycl_remove_unused_external_funcs, + OPT_sycl_device_code_split_esimd, + OPT_no_sycl_device_code_split_esimd, + OPT_sycl_add_default_spec_consts_image, + OPT_no_sycl_add_default_spec_consts_image, + }; + for (OptSpecifier Opt : DirectOpts) { + if (Arg *A = Args.getLastArg(Opt)) { + XLinker("--" + A->getOption().getName().str()); + } + } + + static const OptSpecifier ValueOpts[] = { + OPT_sycl_dump_device_code_EQ, OPT_llvm_spirv_options_EQ, + OPT_sycl_post_link_options_EQ, OPT_syclbin_EQ, + OPT_bitcode_library_EQ, + }; + for (OptSpecifier Opt : ValueOpts) { + if (const opt::Arg *A = Args.getLastArg(Opt)) { + XLinker("--" + A->getOption().getName().str() + A->getValue()); + } + } + + // Forward all device compiler and linker options (can have multiple) + for (const opt::Arg *A : Args.filtered(OPT_device_compiler_args_EQ)) { + XLinker("--" + A->getOption().getName().str() + A->getValue()); + } + for (const opt::Arg *A : Args.filtered(OPT_device_linker_args_EQ)) { + XLinker("--" + A->getOption().getName().str() + A->getValue()); + } +} + Expected clang(ArrayRef InputFiles, const ArgList &Args, bool IsSYCLKind = false) { llvm::TimeTraceScope TimeScope("Clang"); @@ -1744,6 +1800,17 @@ Expected clang(ArrayRef InputFiles, const ArgList &Args, CmdArgs.push_back("-c"); } + if (UseClangSYCLLinker) { + CmdArgs.push_back("-fsycl"); + CmdArgs.push_back("--sycl-link"); + appendClangSYCLLinkerArgs(Args, CmdArgs, Triple, Arch); + for (StringRef InputFile : InputFiles) + CmdArgs.append({"-Xlinker", Args.MakeArgString(InputFile)}); + if (Error Err = executeCommands(*ClangPath, CmdArgs)) + return std::move(Err); + return *TempFileOrErr; + } + for (StringRef InputFile : InputFiles) CmdArgs.push_back(InputFile); From 0b40efc53963cd327663de8081047d8c384c6484 Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Tue, 12 May 2026 17:57:33 -0700 Subject: [PATCH 2/5] Check option forwarding to CSL. --- .../Driver/clang-linker-wrapper-sycl-link.cpp | 42 +++++++ clang/test/Driver/clang-sycl-linker-test.cpp | 103 +++++++++++++++--- .../Driver/sycl-clang-linker-wrapper-e2e.cpp | 64 +++++++++++ .../ClangLinkerWrapper.cpp | 26 ++++- .../clang-linker-wrapper/LinkerWrapperOpts.td | 3 + 5 files changed, 219 insertions(+), 19 deletions(-) create mode 100644 clang/test/Driver/clang-linker-wrapper-sycl-link.cpp create mode 100644 clang/test/Driver/sycl-clang-linker-wrapper-e2e.cpp diff --git a/clang/test/Driver/clang-linker-wrapper-sycl-link.cpp b/clang/test/Driver/clang-linker-wrapper-sycl-link.cpp new file mode 100644 index 0000000000000..f195ca2856a84 --- /dev/null +++ b/clang/test/Driver/clang-linker-wrapper-sycl-link.cpp @@ -0,0 +1,42 @@ +// Test that clang-linker-wrapper correctly invokes clang-sycl-linker when --use-clang-sycl-linker is enabled +// +// REQUIRES: spirv-registered-target +// +// Create SYCL offload object file +// RUN: %clangxx -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown -c %s -o %t.o +// +// Test that clang-sycl-linker is invoked with basic options +// RUN: clang-linker-wrapper --dry-run --use-clang-sycl-linker \ +// RUN: --linker-path=/usr/bin/ld --host-triple=x86_64-unknown-linux-gnu \ +// RUN: %t.o -o %t.out 2>&1 \ +// RUN: | FileCheck %s --check-prefix=CHECK-BASIC +// CHECK-BASIC: clang{{.*}}-fsycl --sycl-link +// CHECK-BASIC-SAME: -Xlinker --triple=spir64-unknown-unknown +// +// Test that --save-temps is forwarded +// RUN: clang-linker-wrapper --dry-run --use-clang-sycl-linker --save-temps \ +// RUN: --linker-path=/usr/bin/ld --host-triple=x86_64-unknown-linux-gnu \ +// RUN: %t.o -o %t.out 2>&1 \ +// RUN: | FileCheck %s --check-prefix=CHECK-SAVE-TEMPS +// CHECK-SAVE-TEMPS: clang{{.*}}-fsycl --sycl-link +// CHECK-SAVE-TEMPS-SAME: -Xlinker --save-temps +// +// Test that device compiler options are forwarded +// RUN: clang-linker-wrapper --dry-run --use-clang-sycl-linker \ +// RUN: --linker-path=/usr/bin/ld --host-triple=x86_64-unknown-linux-gnu \ +// RUN: --device-compiler=sycl:spir64-unknown-unknown=-my-backend-opt \ +// RUN: %t.o -o %t.out 2>&1 \ +// RUN: | FileCheck %s --check-prefix=CHECK-DEVICE-COMPILER +// CHECK-DEVICE-COMPILER: clang{{.*}}-fsycl --sycl-link +// CHECK-DEVICE-COMPILER-SAME: -Xlinker --device-compiler=sycl:spir64-unknown-unknown=-my-backend-opt +// +// Test that device linker options are forwarded +// RUN: clang-linker-wrapper --dry-run --use-clang-sycl-linker \ +// RUN: --linker-path=/usr/bin/ld --host-triple=x86_64-unknown-linux-gnu \ +// RUN: --device-linker=sycl:spir64-unknown-unknown=-my-linker-opt \ +// RUN: %t.o -o %t.out 2>&1 \ +// RUN: | FileCheck %s --check-prefix=CHECK-DEVICE-LINKER +// CHECK-DEVICE-LINKER: clang{{.*}}-fsycl --sycl-link +// CHECK-DEVICE-LINKER-SAME: -Xlinker --device-linker=sycl:spir64-unknown-unknown=-my-linker-opt + +int foo() { return 42; } diff --git a/clang/test/Driver/clang-sycl-linker-test.cpp b/clang/test/Driver/clang-sycl-linker-test.cpp index 5e0afa1bddc7b..a84d27eaf1efe 100644 --- a/clang/test/Driver/clang-sycl-linker-test.cpp +++ b/clang/test/Driver/clang-sycl-linker-test.cpp @@ -2,13 +2,16 @@ // // REQUIRES: spirv-registered-target // -// Test the dry run of a simple case to link two input files. +// Create bitcode files for clang-sycl-linker direct tests // RUN: %clangxx -emit-llvm -c -target spirv64 %s -o %t_1.bc // RUN: %clangxx -emit-llvm -c -target spirv64 %s -o %t_2.bc +// Create SYCL offload object files for driver tests +// RUN: %clangxx -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown -c %s -o %t_1.o +// RUN: %clangxx -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown -c %s -o %t_2.o // RUN: clang-sycl-linker --dry-run -v -triple=spirv64 %t_1.bc %t_2.bc -o %t-spirv.out 2>&1 \ // RUN: | FileCheck %s --check-prefix=SIMPLE-FO -// SIMPLE-FO: sycl-device-link: inputs: {{.*}}.bc, {{.*}}.bc libfiles: output: [[LLVMLINKOUT:.*]].bc -// SIMPLE-FO-NEXT: LLVM backend: input: [[LLVMLINKOUT]].bc, output: {{.*}}_0.spv +// SIMPLE-FO: sycl-device-link: inputs: {{.*}}.bc, {{.*}}.bc{{.*}}libfiles:{{.*}}output: [[LLVMLINKOUT:.*]].bc +// SIMPLE-FO: LLVM backend: input: {{.*}}.bc, output: {{.*}}_0.spv // // Test that IMG_SPIRV image kind is set for non-AOT compilation. // RUN: llvm-objdump --offloading %t-spirv.out | FileCheck %s --check-prefix=IMAGE-KIND-SPIRV @@ -20,8 +23,8 @@ // RUN: touch %t.dir/lib2.bc // RUN: clang-sycl-linker --dry-run -v -triple=spirv64 %t_1.bc %t_2.bc --library-path=%t.dir --device-libs=lib1.bc,lib2.bc -o a.spv 2>&1 \ // RUN: | FileCheck %s --check-prefix=DEVLIBS -// DEVLIBS: sycl-device-link: inputs: {{.*}}.bc, {{.*}}.bc libfiles: {{.*}}lib1.bc, {{.*}}lib2.bc output: [[LLVMLINKOUT:.*]].bc -// DEVLIBS-NEXT: LLVM backend: input: [[LLVMLINKOUT]].bc, output: a_0.spv +// DEVLIBS: sycl-device-link: inputs: {{.*}}.bc, {{.*}}.bc{{.*}}libfiles:{{.*}}lib1.bc, {{.*}}lib2.bc{{.*}}output: [[LLVMLINKOUT:.*]].bc +// DEVLIBS: LLVM backend: input: {{.*}}.bc, output: a_0.spv // // Test a simple case with a random file (not bitcode) as input. // RUN: touch %t.o @@ -41,8 +44,8 @@ // RUN: clang-sycl-linker --dry-run -v -triple=spirv64 -arch=bmg_g21 %t_1.bc %t_2.bc -o %t-aot-gpu.out 2>&1 \ // RUN: --ocloc-options="-a -b" \ // RUN: | FileCheck %s --check-prefix=AOT-INTEL-GPU -// AOT-INTEL-GPU: sycl-device-link: inputs: {{.*}}.bc, {{.*}}.bc libfiles: output: [[LLVMLINKOUT:.*]].bc -// AOT-INTEL-GPU-NEXT: LLVM backend: input: [[LLVMLINKOUT]].bc, output: [[SPIRVTRANSLATIONOUT:.*]]_0.spv +// AOT-INTEL-GPU: sycl-device-link: inputs: {{.*}}.bc, {{.*}}.bc{{.*}}libfiles:{{.*}}output: [[LLVMLINKOUT:.*]].bc +// AOT-INTEL-GPU: LLVM backend: input: {{.*}}.bc, output: [[SPIRVTRANSLATIONOUT:.*]]_0.spv // AOT-INTEL-GPU-NEXT: "{{.*}}ocloc{{.*}}" {{.*}}-device bmg_g21 -a -b {{.*}}-output [[SPIRVTRANSLATIONOUT]]_0.out -file [[SPIRVTRANSLATIONOUT]]_0.spv // // Test that IMG_Object image kind is set for AOT compilation (Intel GPU). @@ -53,8 +56,8 @@ // RUN: clang-sycl-linker --dry-run -v -triple=spirv64 -arch=graniterapids %t_1.bc %t_2.bc -o %t-aot-cpu.out 2>&1 \ // RUN: --opencl-aot-options="-a -b" \ // RUN: | FileCheck %s --check-prefix=AOT-INTEL-CPU -// AOT-INTEL-CPU: sycl-device-link: inputs: {{.*}}.bc, {{.*}}.bc libfiles: output: [[LLVMLINKOUT:.*]].bc -// AOT-INTEL-CPU-NEXT: LLVM backend: input: [[LLVMLINKOUT]].bc, output: [[SPIRVTRANSLATIONOUT:.*]]_0.spv +// AOT-INTEL-CPU: sycl-device-link: inputs: {{.*}}.bc, {{.*}}.bc{{.*}}libfiles:{{.*}}output: [[LLVMLINKOUT:.*]].bc +// AOT-INTEL-CPU: LLVM backend: input: {{.*}}.bc, output: [[SPIRVTRANSLATIONOUT:.*]]_0.spv // AOT-INTEL-CPU-NEXT: "{{.*}}opencl-aot{{.*}}" {{.*}}--device=cpu -a -b {{.*}}-o [[SPIRVTRANSLATIONOUT]]_0.out [[SPIRVTRANSLATIONOUT]]_0.spv // // Test that IMG_Object image kind is set for AOT compilation (Intel CPU). @@ -73,39 +76,39 @@ // Test that driver options are correctly passed through clang-linker-wrapper. // RUN: %clangxx -### -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown \ // RUN: -Xsycl-target-backend=spir64-unknown-unknown "-backend-opt" \ -// RUN: %t_1.bc %t_2.bc 2>&1 | FileCheck %s --check-prefix=CHECK-BACKEND-OPT +// RUN: %t_1.o %t_2.o 2>&1 | FileCheck %s --check-prefix=CHECK-BACKEND-OPT // CHECK-BACKEND-OPT: clang-linker-wrapper{{.*}} "--device-compiler=sycl:spir64-unknown-unknown=-backend-opt" // // Test multiple backend options. // RUN: %clangxx -### -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown \ // RUN: -Xsycl-target-backend=spir64-unknown-unknown "-opt1 -opt2" \ -// RUN: %t_1.bc %t_2.bc 2>&1 | FileCheck %s --check-prefix=CHECK-MULTI-OPTS +// RUN: %t_1.o %t_2.o 2>&1 | FileCheck %s --check-prefix=CHECK-MULTI-OPTS // CHECK-MULTI-OPTS: clang-linker-wrapper{{.*}} "--device-compiler=sycl:spir64-unknown-unknown=-opt1 -opt2" // // Test linker options are forwarded. // RUN: %clangxx -### -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown \ // RUN: -Xsycl-target-linker=spir64-unknown-unknown "-linker-opt" \ -// RUN: %t_1.bc %t_2.bc 2>&1 | FileCheck %s --check-prefix=CHECK-LINKER-OPT +// RUN: %t_1.o %t_2.o 2>&1 | FileCheck %s --check-prefix=CHECK-LINKER-OPT // CHECK-LINKER-OPT: clang-linker-wrapper{{.*}} "--device-linker=sycl:spir64-unknown-unknown=-linker-opt" // // Test both backend and linker options together. // RUN: %clangxx -### -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown \ // RUN: -Xsycl-target-backend=spir64-unknown-unknown "-backend-opt" \ // RUN: -Xsycl-target-linker=spir64-unknown-unknown "-linker-opt" \ -// RUN: %t_1.bc %t_2.bc 2>&1 | FileCheck %s --check-prefix=CHECK-BOTH-OPTS +// RUN: %t_1.o %t_2.o 2>&1 | FileCheck %s --check-prefix=CHECK-BOTH-OPTS // CHECK-BOTH-OPTS: clang-linker-wrapper{{.*}} "--device-compiler=sycl:spir64-unknown-unknown=-backend-opt" // CHECK-BOTH-OPTS-SAME: "--device-linker=sycl:spir64-unknown-unknown=-linker-opt" // // Test AOT GPU with backend options. // RUN: %clangxx -### -fsycl --offload-new-driver -fsycl-targets=spir64_gen-unknown-unknown \ // RUN: -Xsycl-target-backend=spir64_gen "-device pvc" \ -// RUN: %t_1.bc %t_2.bc 2>&1 | FileCheck %s --check-prefix=CHECK-AOT-GPU +// RUN: %t_1.o %t_2.o 2>&1 | FileCheck %s --check-prefix=CHECK-AOT-GPU // CHECK-AOT-GPU: clang-linker-wrapper{{.*}} "--device-compiler=sycl:spir64_gen-unknown-unknown=-device pvc" // // Test AOT CPU with backend options. // RUN: %clangxx -### -fsycl --offload-new-driver -fsycl-targets=spir64_x86_64-unknown-unknown \ // RUN: -Xsycl-target-backend=spir64_x86_64 "-march=skylake" \ -// RUN: %t_1.bc %t_2.bc 2>&1 | FileCheck %s --check-prefix=CHECK-AOT-CPU +// RUN: %t_1.o %t_2.o 2>&1 | FileCheck %s --check-prefix=CHECK-AOT-CPU // CHECK-AOT-CPU: clang-linker-wrapper{{.*}} "--device-compiler=sycl:spir64_x86_64-unknown-unknown=-march=skylake" // // Test multi-target with target-specific options (both in same invocation). @@ -113,6 +116,74 @@ // RUN: -fsycl-targets=spir64_gen-unknown-unknown,spir64_x86_64-unknown-unknown \ // RUN: -Xsycl-target-backend=spir64_gen "-device pvc" \ // RUN: -Xsycl-target-backend=spir64_x86_64 "-march=skylake" \ -// RUN: %t_1.bc %t_2.bc 2>&1 | FileCheck %s --check-prefix=CHECK-MULTI-TARGET +// RUN: %t_1.o %t_2.o 2>&1 | FileCheck %s --check-prefix=CHECK-MULTI-TARGET // CHECK-MULTI-TARGET: clang-linker-wrapper{{.*}} "--device-compiler=sycl:spir64_gen-unknown-unknown=-device pvc" // CHECK-MULTI-TARGET-SAME: "--device-compiler=sycl:spir64_x86_64-unknown-unknown=-march=skylake" +// +// Test that --use-clang-sycl-linker flag is properly forwarded. +// RUN: %clangxx -### -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown \ +// RUN: -Xlinker --use-clang-sycl-linker \ +// RUN: %t_1.o %t_2.o 2>&1 | FileCheck %s --check-prefix=CHECK-USE-CSL +// CHECK-USE-CSL: clang-linker-wrapper{{.*}} "--use-clang-sycl-linker" +// +// Tests for appendClangSYCLLinkerArgs functionality when UseClangSYCLLinker=true +// +// Test that backend options are forwarded when using clang-sycl-linker. +// RUN: %clangxx -### -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown \ +// RUN: -Xlinker --use-clang-sycl-linker \ +// RUN: -Xsycl-target-backend=spir64-unknown-unknown "-backend-opt" \ +// RUN: %t_1.o %t_2.o 2>&1 | FileCheck %s --check-prefix=CHECK-CSL-BACKEND +// CHECK-CSL-BACKEND: clang-linker-wrapper{{.*}} "--device-compiler=sycl:spir64-unknown-unknown=-backend-opt"{{.*}} "--use-clang-sycl-linker" +// +// Test that linker options are forwarded when using clang-sycl-linker. +// RUN: %clangxx -### -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown \ +// RUN: -Xlinker --use-clang-sycl-linker \ +// RUN: -Xsycl-target-linker=spir64-unknown-unknown "-linker-opt" \ +// RUN: %t_1.o %t_2.o 2>&1 | FileCheck %s --check-prefix=CHECK-CSL-LINKER +// CHECK-CSL-LINKER: clang-linker-wrapper{{.*}} "--device-linker=sycl:spir64-unknown-unknown=-linker-opt"{{.*}} "--use-clang-sycl-linker" +// +// Test that both backend and linker options work together with clang-sycl-linker. +// RUN: %clangxx -### -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown \ +// RUN: -Xlinker --use-clang-sycl-linker \ +// RUN: -Xsycl-target-backend=spir64-unknown-unknown "-backend-opt" \ +// RUN: -Xsycl-target-linker=spir64-unknown-unknown "-linker-opt" \ +// RUN: %t_1.o %t_2.o 2>&1 | FileCheck %s --check-prefix=CHECK-CSL-BOTH +// CHECK-CSL-BOTH: clang-linker-wrapper{{.*}} "--device-compiler=sycl:spir64-unknown-unknown=-backend-opt"{{.*}} "--device-linker=sycl:spir64-unknown-unknown=-linker-opt"{{.*}} "--use-clang-sycl-linker" +// +// Test AOT GPU with clang-sycl-linker. +// RUN: %clangxx -### -fsycl --offload-new-driver -fsycl-targets=spir64_gen-unknown-unknown \ +// RUN: -Xlinker --use-clang-sycl-linker \ +// RUN: -Xsycl-target-backend=spir64_gen "-device pvc" \ +// RUN: %t_1.o %t_2.o 2>&1 | FileCheck %s --check-prefix=CHECK-CSL-AOT-GPU +// CHECK-CSL-AOT-GPU: clang-linker-wrapper{{.*}} "--device-compiler=sycl:spir64_gen-unknown-unknown=-device pvc"{{.*}} "--use-clang-sycl-linker" +// +// Test AOT CPU with clang-sycl-linker. +// RUN: %clangxx -### -fsycl --offload-new-driver -fsycl-targets=spir64_x86_64-unknown-unknown \ +// RUN: -Xlinker --use-clang-sycl-linker \ +// RUN: -Xsycl-target-backend=spir64_x86_64 "-march=skylake" \ +// RUN: %t_1.o %t_2.o 2>&1 | FileCheck %s --check-prefix=CHECK-CSL-AOT-CPU +// CHECK-CSL-AOT-CPU: clang-linker-wrapper{{.*}} "--device-compiler=sycl:spir64_x86_64-unknown-unknown=-march=skylake"{{.*}} "--use-clang-sycl-linker" +// +// Test multi-target with clang-sycl-linker. +// RUN: %clangxx -### -fsycl --offload-new-driver \ +// RUN: -fsycl-targets=spir64_gen-unknown-unknown,spir64_x86_64-unknown-unknown \ +// RUN: -Xlinker --use-clang-sycl-linker \ +// RUN: -Xsycl-target-backend=spir64_gen "-device pvc" \ +// RUN: -Xsycl-target-backend=spir64_x86_64 "-march=skylake" \ +// RUN: %t_1.o %t_2.o 2>&1 | FileCheck %s --check-prefix=CHECK-CSL-MULTI +// CHECK-CSL-MULTI: clang-linker-wrapper{{.*}} "--device-compiler=sycl:spir64_gen-unknown-unknown=-device pvc"{{.*}} "--device-compiler=sycl:spir64_x86_64-unknown-unknown=-march=skylake"{{.*}} "--use-clang-sycl-linker" +// +// Test that --save-temps is forwarded when using clang-sycl-linker. +// RUN: %clangxx -### -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown \ +// RUN: -Xlinker --use-clang-sycl-linker -Xlinker --save-temps \ +// RUN: %t_1.o %t_2.o 2>&1 | FileCheck %s --check-prefix=CHECK-CSL-SAVE-TEMPS +// CHECK-CSL-SAVE-TEMPS: clang-linker-wrapper{{.*}} "--use-clang-sycl-linker"{{.*}} "--save-temps" +// +// Test that --dry-run is forwarded when using clang-sycl-linker. +// RUN: %clangxx -### -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown \ +// RUN: -Xlinker --use-clang-sycl-linker -Xlinker --dry-run \ +// RUN: %t_1.o %t_2.o 2>&1 | FileCheck %s --check-prefix=CHECK-CSL-DRY +// CHECK-CSL-DRY: clang-linker-wrapper{{.*}} "--use-clang-sycl-linker"{{.*}} "--dry-run" + + + diff --git a/clang/test/Driver/sycl-clang-linker-wrapper-e2e.cpp b/clang/test/Driver/sycl-clang-linker-wrapper-e2e.cpp new file mode 100644 index 0000000000000..9b11675a9d206 --- /dev/null +++ b/clang/test/Driver/sycl-clang-linker-wrapper-e2e.cpp @@ -0,0 +1,64 @@ +// End-to-end test for --use-clang-sycl-linker flag with actual SYCL device code +// This test verifies that options flow correctly from the driver through +// clang-linker-wrapper to clang-sycl-linker. +// +// REQUIRES: x86-registered-target +// REQUIRES: spirv-registered-target +// +// RUN: %clangxx -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown -c %s -o %t.o +// RUN: %clangxx -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown %t.o -o %t.out \ +// RUN: -### 2>&1 | FileCheck %s --check-prefix=CHECK-DEFAULT +// CHECK-DEFAULT: clang-linker-wrapper +// CHECK-DEFAULT-NOT: "--use-clang-sycl-linker" +// +// Test that --use-clang-sycl-linker flag is passed through +// RUN: %clangxx -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown %t.o -o %t.out \ +// RUN: -Xlinker --use-clang-sycl-linker -### 2>&1 \ +// RUN: | FileCheck %s --check-prefix=CHECK-FLAG +// CHECK-FLAG: clang-linker-wrapper{{.*}} "--use-clang-sycl-linker" +// +// Test with backend options +// RUN: %clangxx -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown %t.o -o %t.out \ +// RUN: -Xlinker --use-clang-sycl-linker \ +// RUN: -Xsycl-target-backend=spir64-unknown-unknown "-test-backend-opt" \ +// RUN: -### 2>&1 | FileCheck %s --check-prefix=CHECK-BACKEND +// CHECK-BACKEND: clang-linker-wrapper{{.*}} "--device-compiler=sycl:spir64-unknown-unknown=-test-backend-opt"{{.*}} "--use-clang-sycl-linker" +// +// Test with linker options +// RUN: %clangxx -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown %t.o -o %t.out \ +// RUN: -Xlinker --use-clang-sycl-linker \ +// RUN: -Xsycl-target-linker=spir64-unknown-unknown "-test-linker-opt" \ +// RUN: -### 2>&1 | FileCheck %s --check-prefix=CHECK-LINKER +// CHECK-LINKER: clang-linker-wrapper{{.*}} "--device-linker=sycl:spir64-unknown-unknown=-test-linker-opt"{{.*}} "--use-clang-sycl-linker" +// +// Test with both backend and linker options +// RUN: %clangxx -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown %t.o -o %t.out \ +// RUN: -Xlinker --use-clang-sycl-linker \ +// RUN: -Xsycl-target-backend=spir64-unknown-unknown "-test-backend" \ +// RUN: -Xsycl-target-linker=spir64-unknown-unknown "-test-linker" \ +// RUN: -### 2>&1 | FileCheck %s --check-prefix=CHECK-BOTH +// CHECK-BOTH: clang-linker-wrapper{{.*}} "--device-compiler=sycl:spir64-unknown-unknown=-test-backend"{{.*}} "--device-linker=sycl:spir64-unknown-unknown=-test-linker"{{.*}} "--use-clang-sycl-linker" +// +// Test with --save-temps +// RUN: %clangxx -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown %t.o -o %t.out \ +// RUN: -Xlinker --use-clang-sycl-linker -Xlinker --save-temps \ +// RUN: -### 2>&1 | FileCheck %s --check-prefix=CHECK-SAVE-TEMPS +// CHECK-SAVE-TEMPS: clang-linker-wrapper{{.*}} "--use-clang-sycl-linker"{{.*}} "--save-temps" + +#include + +int main() { + sycl::queue q; + + int *data = sycl::malloc_shared(1, q); + *data = 0; + + q.parallel_for(sycl::range<1>(1), [=](sycl::id<1> idx) { + data[idx] = 42; + }).wait(); + + int result = *data; + sycl::free(data, q); + + return !(result == 42); +} diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index 07ea58d0292eb..dfe3c28aa9c87 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -160,8 +160,10 @@ static bool CanonicalPrefixes = true; using OffloadingImage = OffloadBinary::OffloadingImage; -// TODO: Remove this once the linking and backend compilation have been fully -// migrated to clang-sycl-linker. +// Controls whether to use clang-sycl-linker for SYCL device linking. +// Set via --use-clang-sycl-linker flag. +// TODO: Remove this flag once the linking and backend compilation have been +// fully migrated to clang-sycl-linker and make it the default behavior. bool UseClangSYCLLinker = false; namespace llvm { @@ -1683,7 +1685,15 @@ Expected bundleDeviceModule(StringRef ClangOutput, } // namespace sycl namespace generic { -// Forwards all required flags to the clang-sycl-linker via -Xlinker. +/// Append arguments for clang-sycl-linker invocation via -Xlinker. +/// This function forwards all necessary options from clang-linker-wrapper to +/// clang-sycl-linker by wrapping them with -Xlinker when invoking +/// "clang -fsycl --sycl-link". +/// +/// \param Args The parsed command-line arguments from clang-linker-wrapper. +/// \param CmdArgs The command arguments list to append to (for clang invocation). +/// \param Triple The target triple for the device. +/// \param Arch The target architecture (e.g., "pvc", "graniterapids"). static void appendClangSYCLLinkerArgs(const ArgList &Args, SmallVectorImpl &CmdArgs, const llvm::Triple &Triple, @@ -1940,6 +1950,15 @@ Expected linkDevice(ArrayRef InputFiles, "For SPIR targets, Linking is supported only for JIT compilations " "and AOT compilations for Intel CPUs/GPUs"); if (IsSYCLKind) { + // When UseClangSYCLLinker is enabled, route spirv64 through + // generic::clang() which invokes "clang -fsycl --sycl-link" and forwards + // options via -Xlinker. Without this check, spirv64 uses the old + // sycl::runLLVMToSPIRVTranslation() path below. Other architectures + // (nvptx, amdgcn, etc.) already route through generic::clang() and check + // UseClangSYCLLinker there. + if (UseClangSYCLLinker) + return generic::clang(InputFiles, Args, IsSYCLKind); + auto SPVFile = sycl::runLLVMToSPIRVTranslation(InputFiles[0], Args); if (!SPVFile) return SPVFile.takeError(); @@ -2824,6 +2843,7 @@ int main(int Argc, char **Argv) { SaveTemps = Args.hasArg(OPT_save_temps); CudaBinaryPath = Args.getLastArgValue(OPT_cuda_path_EQ).str(); CanonicalPrefixes = !Args.hasArg(OPT_no_canonical_prefixes); + UseClangSYCLLinker = Args.hasArg(OPT_use_clang_sycl_linker); llvm::Triple Triple( Args.getLastArgValue(OPT_host_triple_EQ, sys::getDefaultTargetTriple())); diff --git a/clang/tools/clang-linker-wrapper/LinkerWrapperOpts.td b/clang/tools/clang-linker-wrapper/LinkerWrapperOpts.td index c85ee0483cea7..3327d02edee61 100644 --- a/clang/tools/clang-linker-wrapper/LinkerWrapperOpts.td +++ b/clang/tools/clang-linker-wrapper/LinkerWrapperOpts.td @@ -36,6 +36,9 @@ def device_compiler_args_EQ : Joined<["--"], "device-compiler=">, def clang_backend : Flag<["--"], "clang-backend">, Flags<[WrapperOnlyOption]>, HelpText<"Run the backend using clang rather than the LTO backend">; +def use_clang_sycl_linker : Flag<["--"], "use-clang-sycl-linker">, + Flags<[WrapperOnlyOption]>, + HelpText<"Use clang-sycl-linker for SYCL device linking">; def dry_run : Flag<["--"], "dry-run">, Flags<[WrapperOnlyOption]>, HelpText<"Print program arguments without running">; From f5327ea1a33b6fd23638e1e82b1bb0f96109db6f Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Tue, 12 May 2026 18:11:04 -0700 Subject: [PATCH 3/5] Fix clang format. --- clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index dfe3c28aa9c87..9544d06bc44eb 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -1691,7 +1691,8 @@ namespace generic { /// "clang -fsycl --sycl-link". /// /// \param Args The parsed command-line arguments from clang-linker-wrapper. -/// \param CmdArgs The command arguments list to append to (for clang invocation). +/// \param CmdArgs The command arguments list to append to (for clang +/// invocation). /// \param Triple The target triple for the device. /// \param Arch The target architecture (e.g., "pvc", "graniterapids"). static void appendClangSYCLLinkerArgs(const ArgList &Args, From 50fddb164ceddaa7d20bdc965bf5390a7fb46525 Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Wed, 13 May 2026 10:41:25 -0700 Subject: [PATCH 4/5] Migrate to uint16_t ActiveOffloadKindMask and update test. --- clang/test/Driver/clang-sycl-linker-test.cpp | 8 ++--- .../ClangLinkerWrapper.cpp | 31 ++++++++++--------- 2 files changed, 20 insertions(+), 19 deletions(-) diff --git a/clang/test/Driver/clang-sycl-linker-test.cpp b/clang/test/Driver/clang-sycl-linker-test.cpp index a84d27eaf1efe..3d2aeaf3b2b1d 100644 --- a/clang/test/Driver/clang-sycl-linker-test.cpp +++ b/clang/test/Driver/clang-sycl-linker-test.cpp @@ -11,7 +11,7 @@ // RUN: clang-sycl-linker --dry-run -v -triple=spirv64 %t_1.bc %t_2.bc -o %t-spirv.out 2>&1 \ // RUN: | FileCheck %s --check-prefix=SIMPLE-FO // SIMPLE-FO: sycl-device-link: inputs: {{.*}}.bc, {{.*}}.bc{{.*}}libfiles:{{.*}}output: [[LLVMLINKOUT:.*]].bc -// SIMPLE-FO: LLVM backend: input: {{.*}}.bc, output: {{.*}}_0.spv +// SIMPLE-FO-NEXT: LLVM backend: input: [[LLVMLINKOUT]].bc, output: {{.*}}_0.spv // // Test that IMG_SPIRV image kind is set for non-AOT compilation. // RUN: llvm-objdump --offloading %t-spirv.out | FileCheck %s --check-prefix=IMAGE-KIND-SPIRV @@ -24,7 +24,7 @@ // RUN: clang-sycl-linker --dry-run -v -triple=spirv64 %t_1.bc %t_2.bc --library-path=%t.dir --device-libs=lib1.bc,lib2.bc -o a.spv 2>&1 \ // RUN: | FileCheck %s --check-prefix=DEVLIBS // DEVLIBS: sycl-device-link: inputs: {{.*}}.bc, {{.*}}.bc{{.*}}libfiles:{{.*}}lib1.bc, {{.*}}lib2.bc{{.*}}output: [[LLVMLINKOUT:.*]].bc -// DEVLIBS: LLVM backend: input: {{.*}}.bc, output: a_0.spv +// DEVLIBS-NEXT: LLVM backend: input: [[LLVMLINKOUT]].bc, output: a_0.spv // // Test a simple case with a random file (not bitcode) as input. // RUN: touch %t.o @@ -45,7 +45,7 @@ // RUN: --ocloc-options="-a -b" \ // RUN: | FileCheck %s --check-prefix=AOT-INTEL-GPU // AOT-INTEL-GPU: sycl-device-link: inputs: {{.*}}.bc, {{.*}}.bc{{.*}}libfiles:{{.*}}output: [[LLVMLINKOUT:.*]].bc -// AOT-INTEL-GPU: LLVM backend: input: {{.*}}.bc, output: [[SPIRVTRANSLATIONOUT:.*]]_0.spv +// AOT-INTEL-GPU-NEXT: LLVM backend: input: [[LLVMLINKOUT]].bc, output: [[SPIRVTRANSLATIONOUT:.*]]_0.spv // AOT-INTEL-GPU-NEXT: "{{.*}}ocloc{{.*}}" {{.*}}-device bmg_g21 -a -b {{.*}}-output [[SPIRVTRANSLATIONOUT]]_0.out -file [[SPIRVTRANSLATIONOUT]]_0.spv // // Test that IMG_Object image kind is set for AOT compilation (Intel GPU). @@ -57,7 +57,7 @@ // RUN: --opencl-aot-options="-a -b" \ // RUN: | FileCheck %s --check-prefix=AOT-INTEL-CPU // AOT-INTEL-CPU: sycl-device-link: inputs: {{.*}}.bc, {{.*}}.bc{{.*}}libfiles:{{.*}}output: [[LLVMLINKOUT:.*]].bc -// AOT-INTEL-CPU: LLVM backend: input: {{.*}}.bc, output: [[SPIRVTRANSLATIONOUT:.*]]_0.spv +// AOT-INTEL-CPU-NEXT: LLVM backend: input: [[LLVMLINKOUT]].bc, output: [[SPIRVTRANSLATIONOUT:.*]]_0.spv // AOT-INTEL-CPU-NEXT: "{{.*}}opencl-aot{{.*}}" {{.*}}--device=cpu -a -b {{.*}}-o [[SPIRVTRANSLATIONOUT]]_0.out [[SPIRVTRANSLATIONOUT]]_0.spv // // Test that IMG_Object image kind is set for AOT compilation (Intel CPU). diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index 9544d06bc44eb..f1e50eed39abb 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -1747,7 +1747,7 @@ static void appendClangSYCLLinkerArgs(const ArgList &Args, } Expected clang(ArrayRef InputFiles, const ArgList &Args, - bool IsSYCLKind = false) { + uint16_t ActiveOffloadKindMask = 0) { llvm::TimeTraceScope TimeScope("Clang"); // Use `clang` to invoke the appropriate device tools. Expected ClangPath = @@ -1801,17 +1801,17 @@ Expected clang(ArrayRef InputFiles, const ArgList &Args, if (!Triple.isNVPTX() && !Triple.isSPIRV() && !Triple.isNativeCPU()) CmdArgs.push_back("-Wl,--no-undefined"); - if (IsSYCLKind && Triple.isNVPTX()) + if ((ActiveOffloadKindMask & OFK_SYCL) && Triple.isNVPTX()) CmdArgs.push_back("-S"); - if (IsSYCLKind && Triple.isNativeCPU()) { + if ((ActiveOffloadKindMask & OFK_SYCL) && Triple.isNativeCPU()) { CmdArgs.push_back("-Wno-override-module"); CmdArgs.push_back("-mllvm"); CmdArgs.push_back("-sycl-native-cpu-backend"); CmdArgs.push_back("-c"); } - if (UseClangSYCLLinker) { + if (UseClangSYCLLinker && (ActiveOffloadKindMask & OFK_SYCL)) { CmdArgs.push_back("-fsycl"); CmdArgs.push_back("--sycl-link"); appendClangSYCLLinkerArgs(Args, CmdArgs, Triple, Arch); @@ -1924,7 +1924,8 @@ Expected clang(ArrayRef InputFiles, const ArgList &Args, } // namespace generic Expected linkDevice(ArrayRef InputFiles, - const ArgList &Args, bool IsSYCLKind = false, + const ArgList &Args, + uint16_t ActiveOffloadKindMask = 0, StringRef SYCLBackendOptions = StringRef()) { const llvm::Triple Triple(Args.getLastArgValue(OPT_triple_EQ)); switch (Triple.getArch()) { @@ -1938,7 +1939,7 @@ Expected linkDevice(ArrayRef InputFiles, case Triple::ppc64: case Triple::ppc64le: case Triple::systemz: - return generic::clang(InputFiles, Args, IsSYCLKind); + return generic::clang(InputFiles, Args, ActiveOffloadKindMask); case Triple::spirv32: case Triple::spirv64: case Triple::spir: @@ -1950,7 +1951,7 @@ Expected linkDevice(ArrayRef InputFiles, inconvertibleErrorCode(), "For SPIR targets, Linking is supported only for JIT compilations " "and AOT compilations for Intel CPUs/GPUs"); - if (IsSYCLKind) { + if (ActiveOffloadKindMask & OFK_SYCL) { // When UseClangSYCLLinker is enabled, route spirv64 through // generic::clang() which invokes "clang -fsycl --sycl-link" and forwards // options via -Xlinker. Without this check, spirv64 uses the old @@ -1958,7 +1959,7 @@ Expected linkDevice(ArrayRef InputFiles, // (nvptx, amdgcn, etc.) already route through generic::clang() and check // UseClangSYCLLinker there. if (UseClangSYCLLinker) - return generic::clang(InputFiles, Args, IsSYCLKind); + return generic::clang(InputFiles, Args, ActiveOffloadKindMask); auto SPVFile = sycl::runLLVMToSPIRVTranslation(InputFiles[0], Args); if (!SPVFile) @@ -1978,10 +1979,10 @@ Expected linkDevice(ArrayRef InputFiles, return StringRef(""); } case Triple::loongarch64: - return generic::clang(InputFiles, Args, IsSYCLKind); + return generic::clang(InputFiles, Args, ActiveOffloadKindMask); case Triple::native_cpu: - if (IsSYCLKind) - return generic::clang(InputFiles, Args, IsSYCLKind); + if (ActiveOffloadKindMask & OFK_SYCL) + return generic::clang(InputFiles, Args, ActiveOffloadKindMask); return createStringError(Triple.getArchName() + " linking is not supported other than for SYCL"); default: @@ -2457,9 +2458,8 @@ linkAndWrapDeviceFiles(ArrayRef> LinkerInputFiles, } for (size_t I = 0, E = SplitModules.size(); I != E; ++I) { SmallVector Files = {SplitModules[I].ModuleFilePath}; - auto ClangOutputOrErr = - linkDevice(Files, LinkerArgs, true /* IsSYCLKind */, - CompileLinkOptionsOrErr->first); + auto ClangOutputOrErr = linkDevice(Files, LinkerArgs, OFK_SYCL, + CompileLinkOptionsOrErr->first); if (!ClangOutputOrErr) return ClangOutputOrErr.takeError(); if (Triple.isNVPTX() || Triple.isAMDGCN()) { @@ -2519,7 +2519,8 @@ linkAndWrapDeviceFiles(ArrayRef> LinkerInputFiles, } // Link the remaining device files using the device linker. - auto OutputOrErr = linkDevice(InputFiles, LinkerArgs); + auto OutputOrErr = + linkDevice(InputFiles, LinkerArgs, ActiveOffloadKindMask); if (!OutputOrErr) return OutputOrErr.takeError(); From 92df82d4d9eebf89a92db5cf5ca905f31aba92ba Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Fri, 15 May 2026 15:06:36 -0700 Subject: [PATCH 5/5] Address review comments. --- clang/test/Driver/clang-sycl-linker-test.cpp | 128 +----------------- .../Driver/sycl-clang-linker-wrapper-e2e.cpp | 64 --------- .../Driver/sycl-linker-wrapper-options.cpp | 125 +++++++++++++++++ .../ClangLinkerWrapper.cpp | 56 +++++--- 4 files changed, 165 insertions(+), 208 deletions(-) delete mode 100644 clang/test/Driver/sycl-clang-linker-wrapper-e2e.cpp create mode 100644 clang/test/Driver/sycl-linker-wrapper-options.cpp diff --git a/clang/test/Driver/clang-sycl-linker-test.cpp b/clang/test/Driver/clang-sycl-linker-test.cpp index 3d2aeaf3b2b1d..74543885e32af 100644 --- a/clang/test/Driver/clang-sycl-linker-test.cpp +++ b/clang/test/Driver/clang-sycl-linker-test.cpp @@ -2,15 +2,12 @@ // // REQUIRES: spirv-registered-target // -// Create bitcode files for clang-sycl-linker direct tests +// Test the dry run of a simple case to link two input files. // RUN: %clangxx -emit-llvm -c -target spirv64 %s -o %t_1.bc // RUN: %clangxx -emit-llvm -c -target spirv64 %s -o %t_2.bc -// Create SYCL offload object files for driver tests -// RUN: %clangxx -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown -c %s -o %t_1.o -// RUN: %clangxx -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown -c %s -o %t_2.o // RUN: clang-sycl-linker --dry-run -v -triple=spirv64 %t_1.bc %t_2.bc -o %t-spirv.out 2>&1 \ // RUN: | FileCheck %s --check-prefix=SIMPLE-FO -// SIMPLE-FO: sycl-device-link: inputs: {{.*}}.bc, {{.*}}.bc{{.*}}libfiles:{{.*}}output: [[LLVMLINKOUT:.*]].bc +// SIMPLE-FO: sycl-device-link: inputs: {{.*}}.bc, {{.*}}.bc libfiles: output: [[LLVMLINKOUT:.*]].bc // SIMPLE-FO-NEXT: LLVM backend: input: [[LLVMLINKOUT]].bc, output: {{.*}}_0.spv // // Test that IMG_SPIRV image kind is set for non-AOT compilation. @@ -23,7 +20,7 @@ // RUN: touch %t.dir/lib2.bc // RUN: clang-sycl-linker --dry-run -v -triple=spirv64 %t_1.bc %t_2.bc --library-path=%t.dir --device-libs=lib1.bc,lib2.bc -o a.spv 2>&1 \ // RUN: | FileCheck %s --check-prefix=DEVLIBS -// DEVLIBS: sycl-device-link: inputs: {{.*}}.bc, {{.*}}.bc{{.*}}libfiles:{{.*}}lib1.bc, {{.*}}lib2.bc{{.*}}output: [[LLVMLINKOUT:.*]].bc +// DEVLIBS: sycl-device-link: inputs: {{.*}}.bc libfiles: {{.*}}lib1.bc, {{.*}}lib2.bc output: [[LLVMLINKOUT:.*]].bc // DEVLIBS-NEXT: LLVM backend: input: [[LLVMLINKOUT]].bc, output: a_0.spv // // Test a simple case with a random file (not bitcode) as input. @@ -44,7 +41,7 @@ // RUN: clang-sycl-linker --dry-run -v -triple=spirv64 -arch=bmg_g21 %t_1.bc %t_2.bc -o %t-aot-gpu.out 2>&1 \ // RUN: --ocloc-options="-a -b" \ // RUN: | FileCheck %s --check-prefix=AOT-INTEL-GPU -// AOT-INTEL-GPU: sycl-device-link: inputs: {{.*}}.bc, {{.*}}.bc{{.*}}libfiles:{{.*}}output: [[LLVMLINKOUT:.*]].bc +// AOT-INTEL-GPU: sycl-device-link: inputs: {{.*}}.bc, {{.*}}.bc libfiles: output: [[LLVMLINKOUT:.*]].bc // AOT-INTEL-GPU-NEXT: LLVM backend: input: [[LLVMLINKOUT]].bc, output: [[SPIRVTRANSLATIONOUT:.*]]_0.spv // AOT-INTEL-GPU-NEXT: "{{.*}}ocloc{{.*}}" {{.*}}-device bmg_g21 -a -b {{.*}}-output [[SPIRVTRANSLATIONOUT]]_0.out -file [[SPIRVTRANSLATIONOUT]]_0.spv // @@ -56,7 +53,7 @@ // RUN: clang-sycl-linker --dry-run -v -triple=spirv64 -arch=graniterapids %t_1.bc %t_2.bc -o %t-aot-cpu.out 2>&1 \ // RUN: --opencl-aot-options="-a -b" \ // RUN: | FileCheck %s --check-prefix=AOT-INTEL-CPU -// AOT-INTEL-CPU: sycl-device-link: inputs: {{.*}}.bc, {{.*}}.bc{{.*}}libfiles:{{.*}}output: [[LLVMLINKOUT:.*]].bc +// AOT-INTEL-CPU: sycl-device-link: inputs: {{.*}}.bc, {{.*}}.bc libfiles: output: [[LLVMLINKOUT:.*]].bc // AOT-INTEL-CPU-NEXT: LLVM backend: input: [[LLVMLINKOUT]].bc, output: [[SPIRVTRANSLATIONOUT:.*]]_0.spv // AOT-INTEL-CPU-NEXT: "{{.*}}opencl-aot{{.*}}" {{.*}}--device=cpu -a -b {{.*}}-o [[SPIRVTRANSLATIONOUT]]_0.out [[SPIRVTRANSLATIONOUT]]_0.spv // @@ -72,118 +69,3 @@ // RUN: not clang-sycl-linker --dry-run %t_1.bc %t_2.bc -o a.out 2>&1 \ // RUN: | FileCheck %s --check-prefix=NOTARGET // NOTARGET: Target triple must be specified -// -// Test that driver options are correctly passed through clang-linker-wrapper. -// RUN: %clangxx -### -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown \ -// RUN: -Xsycl-target-backend=spir64-unknown-unknown "-backend-opt" \ -// RUN: %t_1.o %t_2.o 2>&1 | FileCheck %s --check-prefix=CHECK-BACKEND-OPT -// CHECK-BACKEND-OPT: clang-linker-wrapper{{.*}} "--device-compiler=sycl:spir64-unknown-unknown=-backend-opt" -// -// Test multiple backend options. -// RUN: %clangxx -### -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown \ -// RUN: -Xsycl-target-backend=spir64-unknown-unknown "-opt1 -opt2" \ -// RUN: %t_1.o %t_2.o 2>&1 | FileCheck %s --check-prefix=CHECK-MULTI-OPTS -// CHECK-MULTI-OPTS: clang-linker-wrapper{{.*}} "--device-compiler=sycl:spir64-unknown-unknown=-opt1 -opt2" -// -// Test linker options are forwarded. -// RUN: %clangxx -### -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown \ -// RUN: -Xsycl-target-linker=spir64-unknown-unknown "-linker-opt" \ -// RUN: %t_1.o %t_2.o 2>&1 | FileCheck %s --check-prefix=CHECK-LINKER-OPT -// CHECK-LINKER-OPT: clang-linker-wrapper{{.*}} "--device-linker=sycl:spir64-unknown-unknown=-linker-opt" -// -// Test both backend and linker options together. -// RUN: %clangxx -### -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown \ -// RUN: -Xsycl-target-backend=spir64-unknown-unknown "-backend-opt" \ -// RUN: -Xsycl-target-linker=spir64-unknown-unknown "-linker-opt" \ -// RUN: %t_1.o %t_2.o 2>&1 | FileCheck %s --check-prefix=CHECK-BOTH-OPTS -// CHECK-BOTH-OPTS: clang-linker-wrapper{{.*}} "--device-compiler=sycl:spir64-unknown-unknown=-backend-opt" -// CHECK-BOTH-OPTS-SAME: "--device-linker=sycl:spir64-unknown-unknown=-linker-opt" -// -// Test AOT GPU with backend options. -// RUN: %clangxx -### -fsycl --offload-new-driver -fsycl-targets=spir64_gen-unknown-unknown \ -// RUN: -Xsycl-target-backend=spir64_gen "-device pvc" \ -// RUN: %t_1.o %t_2.o 2>&1 | FileCheck %s --check-prefix=CHECK-AOT-GPU -// CHECK-AOT-GPU: clang-linker-wrapper{{.*}} "--device-compiler=sycl:spir64_gen-unknown-unknown=-device pvc" -// -// Test AOT CPU with backend options. -// RUN: %clangxx -### -fsycl --offload-new-driver -fsycl-targets=spir64_x86_64-unknown-unknown \ -// RUN: -Xsycl-target-backend=spir64_x86_64 "-march=skylake" \ -// RUN: %t_1.o %t_2.o 2>&1 | FileCheck %s --check-prefix=CHECK-AOT-CPU -// CHECK-AOT-CPU: clang-linker-wrapper{{.*}} "--device-compiler=sycl:spir64_x86_64-unknown-unknown=-march=skylake" -// -// Test multi-target with target-specific options (both in same invocation). -// RUN: %clangxx -### -fsycl --offload-new-driver \ -// RUN: -fsycl-targets=spir64_gen-unknown-unknown,spir64_x86_64-unknown-unknown \ -// RUN: -Xsycl-target-backend=spir64_gen "-device pvc" \ -// RUN: -Xsycl-target-backend=spir64_x86_64 "-march=skylake" \ -// RUN: %t_1.o %t_2.o 2>&1 | FileCheck %s --check-prefix=CHECK-MULTI-TARGET -// CHECK-MULTI-TARGET: clang-linker-wrapper{{.*}} "--device-compiler=sycl:spir64_gen-unknown-unknown=-device pvc" -// CHECK-MULTI-TARGET-SAME: "--device-compiler=sycl:spir64_x86_64-unknown-unknown=-march=skylake" -// -// Test that --use-clang-sycl-linker flag is properly forwarded. -// RUN: %clangxx -### -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown \ -// RUN: -Xlinker --use-clang-sycl-linker \ -// RUN: %t_1.o %t_2.o 2>&1 | FileCheck %s --check-prefix=CHECK-USE-CSL -// CHECK-USE-CSL: clang-linker-wrapper{{.*}} "--use-clang-sycl-linker" -// -// Tests for appendClangSYCLLinkerArgs functionality when UseClangSYCLLinker=true -// -// Test that backend options are forwarded when using clang-sycl-linker. -// RUN: %clangxx -### -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown \ -// RUN: -Xlinker --use-clang-sycl-linker \ -// RUN: -Xsycl-target-backend=spir64-unknown-unknown "-backend-opt" \ -// RUN: %t_1.o %t_2.o 2>&1 | FileCheck %s --check-prefix=CHECK-CSL-BACKEND -// CHECK-CSL-BACKEND: clang-linker-wrapper{{.*}} "--device-compiler=sycl:spir64-unknown-unknown=-backend-opt"{{.*}} "--use-clang-sycl-linker" -// -// Test that linker options are forwarded when using clang-sycl-linker. -// RUN: %clangxx -### -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown \ -// RUN: -Xlinker --use-clang-sycl-linker \ -// RUN: -Xsycl-target-linker=spir64-unknown-unknown "-linker-opt" \ -// RUN: %t_1.o %t_2.o 2>&1 | FileCheck %s --check-prefix=CHECK-CSL-LINKER -// CHECK-CSL-LINKER: clang-linker-wrapper{{.*}} "--device-linker=sycl:spir64-unknown-unknown=-linker-opt"{{.*}} "--use-clang-sycl-linker" -// -// Test that both backend and linker options work together with clang-sycl-linker. -// RUN: %clangxx -### -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown \ -// RUN: -Xlinker --use-clang-sycl-linker \ -// RUN: -Xsycl-target-backend=spir64-unknown-unknown "-backend-opt" \ -// RUN: -Xsycl-target-linker=spir64-unknown-unknown "-linker-opt" \ -// RUN: %t_1.o %t_2.o 2>&1 | FileCheck %s --check-prefix=CHECK-CSL-BOTH -// CHECK-CSL-BOTH: clang-linker-wrapper{{.*}} "--device-compiler=sycl:spir64-unknown-unknown=-backend-opt"{{.*}} "--device-linker=sycl:spir64-unknown-unknown=-linker-opt"{{.*}} "--use-clang-sycl-linker" -// -// Test AOT GPU with clang-sycl-linker. -// RUN: %clangxx -### -fsycl --offload-new-driver -fsycl-targets=spir64_gen-unknown-unknown \ -// RUN: -Xlinker --use-clang-sycl-linker \ -// RUN: -Xsycl-target-backend=spir64_gen "-device pvc" \ -// RUN: %t_1.o %t_2.o 2>&1 | FileCheck %s --check-prefix=CHECK-CSL-AOT-GPU -// CHECK-CSL-AOT-GPU: clang-linker-wrapper{{.*}} "--device-compiler=sycl:spir64_gen-unknown-unknown=-device pvc"{{.*}} "--use-clang-sycl-linker" -// -// Test AOT CPU with clang-sycl-linker. -// RUN: %clangxx -### -fsycl --offload-new-driver -fsycl-targets=spir64_x86_64-unknown-unknown \ -// RUN: -Xlinker --use-clang-sycl-linker \ -// RUN: -Xsycl-target-backend=spir64_x86_64 "-march=skylake" \ -// RUN: %t_1.o %t_2.o 2>&1 | FileCheck %s --check-prefix=CHECK-CSL-AOT-CPU -// CHECK-CSL-AOT-CPU: clang-linker-wrapper{{.*}} "--device-compiler=sycl:spir64_x86_64-unknown-unknown=-march=skylake"{{.*}} "--use-clang-sycl-linker" -// -// Test multi-target with clang-sycl-linker. -// RUN: %clangxx -### -fsycl --offload-new-driver \ -// RUN: -fsycl-targets=spir64_gen-unknown-unknown,spir64_x86_64-unknown-unknown \ -// RUN: -Xlinker --use-clang-sycl-linker \ -// RUN: -Xsycl-target-backend=spir64_gen "-device pvc" \ -// RUN: -Xsycl-target-backend=spir64_x86_64 "-march=skylake" \ -// RUN: %t_1.o %t_2.o 2>&1 | FileCheck %s --check-prefix=CHECK-CSL-MULTI -// CHECK-CSL-MULTI: clang-linker-wrapper{{.*}} "--device-compiler=sycl:spir64_gen-unknown-unknown=-device pvc"{{.*}} "--device-compiler=sycl:spir64_x86_64-unknown-unknown=-march=skylake"{{.*}} "--use-clang-sycl-linker" -// -// Test that --save-temps is forwarded when using clang-sycl-linker. -// RUN: %clangxx -### -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown \ -// RUN: -Xlinker --use-clang-sycl-linker -Xlinker --save-temps \ -// RUN: %t_1.o %t_2.o 2>&1 | FileCheck %s --check-prefix=CHECK-CSL-SAVE-TEMPS -// CHECK-CSL-SAVE-TEMPS: clang-linker-wrapper{{.*}} "--use-clang-sycl-linker"{{.*}} "--save-temps" -// -// Test that --dry-run is forwarded when using clang-sycl-linker. -// RUN: %clangxx -### -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown \ -// RUN: -Xlinker --use-clang-sycl-linker -Xlinker --dry-run \ -// RUN: %t_1.o %t_2.o 2>&1 | FileCheck %s --check-prefix=CHECK-CSL-DRY -// CHECK-CSL-DRY: clang-linker-wrapper{{.*}} "--use-clang-sycl-linker"{{.*}} "--dry-run" - - - diff --git a/clang/test/Driver/sycl-clang-linker-wrapper-e2e.cpp b/clang/test/Driver/sycl-clang-linker-wrapper-e2e.cpp deleted file mode 100644 index 9b11675a9d206..0000000000000 --- a/clang/test/Driver/sycl-clang-linker-wrapper-e2e.cpp +++ /dev/null @@ -1,64 +0,0 @@ -// End-to-end test for --use-clang-sycl-linker flag with actual SYCL device code -// This test verifies that options flow correctly from the driver through -// clang-linker-wrapper to clang-sycl-linker. -// -// REQUIRES: x86-registered-target -// REQUIRES: spirv-registered-target -// -// RUN: %clangxx -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown -c %s -o %t.o -// RUN: %clangxx -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown %t.o -o %t.out \ -// RUN: -### 2>&1 | FileCheck %s --check-prefix=CHECK-DEFAULT -// CHECK-DEFAULT: clang-linker-wrapper -// CHECK-DEFAULT-NOT: "--use-clang-sycl-linker" -// -// Test that --use-clang-sycl-linker flag is passed through -// RUN: %clangxx -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown %t.o -o %t.out \ -// RUN: -Xlinker --use-clang-sycl-linker -### 2>&1 \ -// RUN: | FileCheck %s --check-prefix=CHECK-FLAG -// CHECK-FLAG: clang-linker-wrapper{{.*}} "--use-clang-sycl-linker" -// -// Test with backend options -// RUN: %clangxx -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown %t.o -o %t.out \ -// RUN: -Xlinker --use-clang-sycl-linker \ -// RUN: -Xsycl-target-backend=spir64-unknown-unknown "-test-backend-opt" \ -// RUN: -### 2>&1 | FileCheck %s --check-prefix=CHECK-BACKEND -// CHECK-BACKEND: clang-linker-wrapper{{.*}} "--device-compiler=sycl:spir64-unknown-unknown=-test-backend-opt"{{.*}} "--use-clang-sycl-linker" -// -// Test with linker options -// RUN: %clangxx -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown %t.o -o %t.out \ -// RUN: -Xlinker --use-clang-sycl-linker \ -// RUN: -Xsycl-target-linker=spir64-unknown-unknown "-test-linker-opt" \ -// RUN: -### 2>&1 | FileCheck %s --check-prefix=CHECK-LINKER -// CHECK-LINKER: clang-linker-wrapper{{.*}} "--device-linker=sycl:spir64-unknown-unknown=-test-linker-opt"{{.*}} "--use-clang-sycl-linker" -// -// Test with both backend and linker options -// RUN: %clangxx -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown %t.o -o %t.out \ -// RUN: -Xlinker --use-clang-sycl-linker \ -// RUN: -Xsycl-target-backend=spir64-unknown-unknown "-test-backend" \ -// RUN: -Xsycl-target-linker=spir64-unknown-unknown "-test-linker" \ -// RUN: -### 2>&1 | FileCheck %s --check-prefix=CHECK-BOTH -// CHECK-BOTH: clang-linker-wrapper{{.*}} "--device-compiler=sycl:spir64-unknown-unknown=-test-backend"{{.*}} "--device-linker=sycl:spir64-unknown-unknown=-test-linker"{{.*}} "--use-clang-sycl-linker" -// -// Test with --save-temps -// RUN: %clangxx -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown %t.o -o %t.out \ -// RUN: -Xlinker --use-clang-sycl-linker -Xlinker --save-temps \ -// RUN: -### 2>&1 | FileCheck %s --check-prefix=CHECK-SAVE-TEMPS -// CHECK-SAVE-TEMPS: clang-linker-wrapper{{.*}} "--use-clang-sycl-linker"{{.*}} "--save-temps" - -#include - -int main() { - sycl::queue q; - - int *data = sycl::malloc_shared(1, q); - *data = 0; - - q.parallel_for(sycl::range<1>(1), [=](sycl::id<1> idx) { - data[idx] = 42; - }).wait(); - - int result = *data; - sycl::free(data, q); - - return !(result == 42); -} diff --git a/clang/test/Driver/sycl-linker-wrapper-options.cpp b/clang/test/Driver/sycl-linker-wrapper-options.cpp new file mode 100644 index 0000000000000..e92d14a68357d --- /dev/null +++ b/clang/test/Driver/sycl-linker-wrapper-options.cpp @@ -0,0 +1,125 @@ +// Tests that clang driver correctly forwards options to clang-linker-wrapper +// for SYCL offload compilation. +// +// REQUIRES: spirv-registered-target +// +// RUN: %clangxx -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown -c %s -o %t_1.o +// RUN: %clangxx -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown -c %s -o %t_2.o +// +// Test that --use-clang-sycl-linker is not passed by default. +// RUN: %clangxx -### -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown \ +// RUN: %t_1.o 2>&1 | FileCheck %s --check-prefix=CHECK-DEFAULT +// CHECK-DEFAULT: clang-linker-wrapper +// CHECK-DEFAULT-NOT: "--use-clang-sycl-linker" +// +// Test that driver options are correctly passed through clang-linker-wrapper. +// RUN: %clangxx -### -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown \ +// RUN: -Xsycl-target-backend=spir64-unknown-unknown "-backend-opt" \ +// RUN: %t_1.o %t_2.o 2>&1 | FileCheck %s --check-prefix=CHECK-BACKEND-OPT +// CHECK-BACKEND-OPT: clang-linker-wrapper{{.*}} "--device-compiler=sycl:spir64-unknown-unknown=-backend-opt" +// +// Test multiple backend options. +// RUN: %clangxx -### -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown \ +// RUN: -Xsycl-target-backend=spir64-unknown-unknown "-opt1 -opt2" \ +// RUN: %t_1.o %t_2.o 2>&1 | FileCheck %s --check-prefix=CHECK-MULTI-OPTS +// CHECK-MULTI-OPTS: clang-linker-wrapper{{.*}} "--device-compiler=sycl:spir64-unknown-unknown=-opt1 -opt2" +// +// Test linker options are forwarded. +// RUN: %clangxx -### -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown \ +// RUN: -Xsycl-target-linker=spir64-unknown-unknown "-linker-opt" \ +// RUN: %t_1.o %t_2.o 2>&1 | FileCheck %s --check-prefix=CHECK-LINKER-OPT +// CHECK-LINKER-OPT: clang-linker-wrapper{{.*}} "--device-linker=sycl:spir64-unknown-unknown=-linker-opt" +// +// Test both backend and linker options together. +// RUN: %clangxx -### -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown \ +// RUN: -Xsycl-target-backend=spir64-unknown-unknown "-backend-opt" \ +// RUN: -Xsycl-target-linker=spir64-unknown-unknown "-linker-opt" \ +// RUN: %t_1.o %t_2.o 2>&1 | FileCheck %s --check-prefix=CHECK-BOTH-OPTS +// CHECK-BOTH-OPTS: clang-linker-wrapper{{.*}} "--device-compiler=sycl:spir64-unknown-unknown=-backend-opt" +// CHECK-BOTH-OPTS-SAME: "--device-linker=sycl:spir64-unknown-unknown=-linker-opt" +// +// Test AOT GPU with backend options. +// RUN: %clangxx -### -fsycl --offload-new-driver -fsycl-targets=spir64_gen-unknown-unknown \ +// RUN: -Xsycl-target-backend=spir64_gen "-device pvc" \ +// RUN: %t_1.o %t_2.o 2>&1 | FileCheck %s --check-prefix=CHECK-AOT-GPU +// CHECK-AOT-GPU: clang-linker-wrapper{{.*}} "--device-compiler=sycl:spir64_gen-unknown-unknown=-device pvc" +// +// Test AOT CPU with backend options. +// RUN: %clangxx -### -fsycl --offload-new-driver -fsycl-targets=spir64_x86_64-unknown-unknown \ +// RUN: -Xsycl-target-backend=spir64_x86_64 "-march=skylake" \ +// RUN: %t_1.o %t_2.o 2>&1 | FileCheck %s --check-prefix=CHECK-AOT-CPU +// CHECK-AOT-CPU: clang-linker-wrapper{{.*}} "--device-compiler=sycl:spir64_x86_64-unknown-unknown=-march=skylake" +// +// Test multi-target with target-specific options (both in same invocation). +// RUN: %clangxx -### -fsycl --offload-new-driver \ +// RUN: -fsycl-targets=spir64_gen-unknown-unknown,spir64_x86_64-unknown-unknown \ +// RUN: -Xsycl-target-backend=spir64_gen "-device pvc" \ +// RUN: -Xsycl-target-backend=spir64_x86_64 "-march=skylake" \ +// RUN: %t_1.o %t_2.o 2>&1 | FileCheck %s --check-prefix=CHECK-MULTI-TARGET +// CHECK-MULTI-TARGET: clang-linker-wrapper{{.*}} "--device-compiler=sycl:spir64_gen-unknown-unknown=-device pvc" +// CHECK-MULTI-TARGET-SAME: "--device-compiler=sycl:spir64_x86_64-unknown-unknown=-march=skylake" +// +// Test that --use-clang-sycl-linker flag is properly forwarded. +// RUN: %clangxx -### -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown \ +// RUN: -Xlinker --use-clang-sycl-linker \ +// RUN: %t_1.o %t_2.o 2>&1 | FileCheck %s --check-prefix=CHECK-USE-CSL +// CHECK-USE-CSL: clang-linker-wrapper{{.*}} "--use-clang-sycl-linker" +// +// Test that backend options are forwarded when using clang-sycl-linker. +// RUN: %clangxx -### -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown \ +// RUN: -Xlinker --use-clang-sycl-linker \ +// RUN: -Xsycl-target-backend=spir64-unknown-unknown "-backend-opt" \ +// RUN: %t_1.o %t_2.o 2>&1 | FileCheck %s --check-prefix=CHECK-CSL-BACKEND +// CHECK-CSL-BACKEND: clang-linker-wrapper{{.*}} "--device-compiler=sycl:spir64-unknown-unknown=-backend-opt"{{.*}} "--use-clang-sycl-linker" +// +// Test that linker options are forwarded when using clang-sycl-linker. +// RUN: %clangxx -### -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown \ +// RUN: -Xlinker --use-clang-sycl-linker \ +// RUN: -Xsycl-target-linker=spir64-unknown-unknown "-linker-opt" \ +// RUN: %t_1.o %t_2.o 2>&1 | FileCheck %s --check-prefix=CHECK-CSL-LINKER +// CHECK-CSL-LINKER: clang-linker-wrapper{{.*}} "--device-linker=sycl:spir64-unknown-unknown=-linker-opt"{{.*}} "--use-clang-sycl-linker" +// +// Test that both backend and linker options work together with clang-sycl-linker. +// RUN: %clangxx -### -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown \ +// RUN: -Xlinker --use-clang-sycl-linker \ +// RUN: -Xsycl-target-backend=spir64-unknown-unknown "-backend-opt" \ +// RUN: -Xsycl-target-linker=spir64-unknown-unknown "-linker-opt" \ +// RUN: %t_1.o %t_2.o 2>&1 | FileCheck %s --check-prefix=CHECK-CSL-BOTH +// CHECK-CSL-BOTH: clang-linker-wrapper{{.*}} "--device-compiler=sycl:spir64-unknown-unknown=-backend-opt"{{.*}} "--device-linker=sycl:spir64-unknown-unknown=-linker-opt"{{.*}} "--use-clang-sycl-linker" +// +// Test AOT GPU with clang-sycl-linker. +// RUN: %clangxx -### -fsycl --offload-new-driver -fsycl-targets=spir64_gen-unknown-unknown \ +// RUN: -Xlinker --use-clang-sycl-linker \ +// RUN: -Xsycl-target-backend=spir64_gen "-device pvc" \ +// RUN: %t_1.o %t_2.o 2>&1 | FileCheck %s --check-prefix=CHECK-CSL-AOT-GPU +// CHECK-CSL-AOT-GPU: clang-linker-wrapper{{.*}} "--device-compiler=sycl:spir64_gen-unknown-unknown=-device pvc"{{.*}} "--use-clang-sycl-linker" +// +// Test AOT CPU with clang-sycl-linker. +// RUN: %clangxx -### -fsycl --offload-new-driver -fsycl-targets=spir64_x86_64-unknown-unknown \ +// RUN: -Xlinker --use-clang-sycl-linker \ +// RUN: -Xsycl-target-backend=spir64_x86_64 "-march=skylake" \ +// RUN: %t_1.o %t_2.o 2>&1 | FileCheck %s --check-prefix=CHECK-CSL-AOT-CPU +// CHECK-CSL-AOT-CPU: clang-linker-wrapper{{.*}} "--device-compiler=sycl:spir64_x86_64-unknown-unknown=-march=skylake"{{.*}} "--use-clang-sycl-linker" +// +// Test multi-target with clang-sycl-linker. +// RUN: %clangxx -### -fsycl --offload-new-driver \ +// RUN: -fsycl-targets=spir64_gen-unknown-unknown,spir64_x86_64-unknown-unknown \ +// RUN: -Xlinker --use-clang-sycl-linker \ +// RUN: -Xsycl-target-backend=spir64_gen "-device pvc" \ +// RUN: -Xsycl-target-backend=spir64_x86_64 "-march=skylake" \ +// RUN: %t_1.o %t_2.o 2>&1 | FileCheck %s --check-prefix=CHECK-CSL-MULTI +// CHECK-CSL-MULTI: clang-linker-wrapper{{.*}} "--device-compiler=sycl:spir64_gen-unknown-unknown=-device pvc"{{.*}} "--device-compiler=sycl:spir64_x86_64-unknown-unknown=-march=skylake"{{.*}} "--use-clang-sycl-linker" +// +// Test that --save-temps is forwarded when using clang-sycl-linker. +// RUN: %clangxx -### -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown \ +// RUN: -Xlinker --use-clang-sycl-linker -Xlinker --save-temps \ +// RUN: %t_1.o %t_2.o 2>&1 | FileCheck %s --check-prefix=CHECK-CSL-SAVE-TEMPS +// CHECK-CSL-SAVE-TEMPS: clang-linker-wrapper{{.*}} "--use-clang-sycl-linker"{{.*}} "--save-temps" +// +// Test that --dry-run is forwarded when using clang-sycl-linker. +// RUN: %clangxx -### -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown \ +// RUN: -Xlinker --use-clang-sycl-linker -Xlinker --dry-run \ +// RUN: %t_1.o %t_2.o 2>&1 | FileCheck %s --check-prefix=CHECK-CSL-DRY +// CHECK-CSL-DRY: clang-linker-wrapper{{.*}} "--use-clang-sycl-linker"{{.*}} "--dry-run" + +int foo() { return 42; } diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index f1e50eed39abb..f931b4a275571 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -164,7 +164,7 @@ using OffloadingImage = OffloadBinary::OffloadingImage; // Set via --use-clang-sycl-linker flag. // TODO: Remove this flag once the linking and backend compilation have been // fully migrated to clang-sycl-linker and make it the default behavior. -bool UseClangSYCLLinker = false; +static bool UseClangSYCLLinker = false; namespace llvm { // Provide DenseMapInfo so that OffloadKind can be used in a DenseMap. @@ -1682,14 +1682,22 @@ Expected bundleDeviceModule(StringRef ClangOutput, llvm_unreachable("bundleDeviceModule called for non-NVPTX/AMDGCN triple"); } -} // namespace sycl - -namespace generic { /// Append arguments for clang-sycl-linker invocation via -Xlinker. /// This function forwards all necessary options from clang-linker-wrapper to /// clang-sycl-linker by wrapping them with -Xlinker when invoking /// "clang -fsycl --sycl-link". /// +/// TODO: The following options are forwarded but not yet defined in +/// SYCLLinkOpts.td, so clang-sycl-linker currently ignores them. Each should +/// be added to SYCLLinkOpts.td and wired up in ClangSYCLLinker.cpp as the +/// corresponding functionality is migrated into clang-sycl-linker: +/// --print-wrapped-module, --sycl-thin-lto, +/// --sycl-allow-device-image-dependencies, +/// --sycl-remove-unused-external-funcs, --sycl-device-code-split-esimd, +/// --sycl-add-default-spec-consts-image, --sycl-dump-device-code=, +/// --llvm-spirv-options=, --sycl-post-link-options=, +/// --device-compiler=, --device-linker= +/// /// \param Args The parsed command-line arguments from clang-linker-wrapper. /// \param CmdArgs The command arguments list to append to (for clang /// invocation). @@ -1713,28 +1721,30 @@ static void appendClangSYCLLinkerArgs(const ArgList &Args, OPT_print_wrapped_module, OPT_sycl_thin_lto, OPT_sycl_allow_device_image_dependencies, - OPT_sycl_remove_unused_external_funcs, - OPT_no_sycl_remove_unused_external_funcs, - OPT_sycl_device_code_split_esimd, - OPT_no_sycl_device_code_split_esimd, - OPT_sycl_add_default_spec_consts_image, - OPT_no_sycl_add_default_spec_consts_image, }; for (OptSpecifier Opt : DirectOpts) { - if (Arg *A = Args.getLastArg(Opt)) { + if (Arg *A = Args.getLastArg(Opt)) XLinker("--" + A->getOption().getName().str()); - } } - + if (Arg *A = Args.getLastArg(OPT_sycl_remove_unused_external_funcs, + OPT_no_sycl_remove_unused_external_funcs)) + XLinker("--" + A->getOption().getName().str()); + if (Arg *A = Args.getLastArg(OPT_sycl_device_code_split_esimd, + OPT_no_sycl_device_code_split_esimd)) + XLinker("--" + A->getOption().getName().str()); + if (Arg *A = Args.getLastArg(OPT_sycl_add_default_spec_consts_image, + OPT_no_sycl_add_default_spec_consts_image)) + XLinker("--" + A->getOption().getName().str()); + + // All options below are single-valued; only the last occurrence is forwarded. static const OptSpecifier ValueOpts[] = { - OPT_sycl_dump_device_code_EQ, OPT_llvm_spirv_options_EQ, - OPT_sycl_post_link_options_EQ, OPT_syclbin_EQ, - OPT_bitcode_library_EQ, + OPT_sycl_dump_device_code_EQ, + OPT_llvm_spirv_options_EQ, + OPT_sycl_post_link_options_EQ, }; for (OptSpecifier Opt : ValueOpts) { - if (const opt::Arg *A = Args.getLastArg(Opt)) { + if (const opt::Arg *A = Args.getLastArg(Opt)) XLinker("--" + A->getOption().getName().str() + A->getValue()); - } } // Forward all device compiler and linker options (can have multiple) @@ -1746,8 +1756,11 @@ static void appendClangSYCLLinkerArgs(const ArgList &Args, } } +} // namespace sycl + +namespace generic { Expected clang(ArrayRef InputFiles, const ArgList &Args, - uint16_t ActiveOffloadKindMask = 0) { + uint16_t ActiveOffloadKindMask) { llvm::TimeTraceScope TimeScope("Clang"); // Use `clang` to invoke the appropriate device tools. Expected ClangPath = @@ -1811,10 +1824,11 @@ Expected clang(ArrayRef InputFiles, const ArgList &Args, CmdArgs.push_back("-c"); } - if (UseClangSYCLLinker && (ActiveOffloadKindMask & OFK_SYCL)) { + if (UseClangSYCLLinker && (ActiveOffloadKindMask & OFK_SYCL) && + Triple.isSPIROrSPIRV()) { CmdArgs.push_back("-fsycl"); CmdArgs.push_back("--sycl-link"); - appendClangSYCLLinkerArgs(Args, CmdArgs, Triple, Arch); + sycl::appendClangSYCLLinkerArgs(Args, CmdArgs, Triple, Arch); for (StringRef InputFile : InputFiles) CmdArgs.append({"-Xlinker", Args.MakeArgString(InputFile)}); if (Error Err = executeCommands(*ClangPath, CmdArgs))