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/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 2a9950d2154be..f931b4a275571 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -160,6 +160,12 @@ static bool CanonicalPrefixes = true; using OffloadingImage = OffloadBinary::OffloadingImage; +// 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. +static bool UseClangSYCLLinker = false; + namespace llvm { // Provide DenseMapInfo so that OffloadKind can be used in a DenseMap. template <> struct DenseMapInfo { @@ -1676,11 +1682,85 @@ Expected bundleDeviceModule(StringRef ClangOutput, llvm_unreachable("bundleDeviceModule called for non-NVPTX/AMDGCN triple"); } +/// 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). +/// \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, + 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, + }; + for (OptSpecifier Opt : DirectOpts) { + 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, + }; + 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()); + } +} + } // namespace sycl namespace generic { Expected clang(ArrayRef InputFiles, const ArgList &Args, - bool IsSYCLKind = false) { + uint16_t ActiveOffloadKindMask) { llvm::TimeTraceScope TimeScope("Clang"); // Use `clang` to invoke the appropriate device tools. Expected ClangPath = @@ -1734,16 +1814,28 @@ 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 && (ActiveOffloadKindMask & OFK_SYCL) && + Triple.isSPIROrSPIRV()) { + CmdArgs.push_back("-fsycl"); + CmdArgs.push_back("--sycl-link"); + sycl::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); @@ -1846,7 +1938,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()) { @@ -1860,7 +1953,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: @@ -1872,7 +1965,16 @@ 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 + // 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, ActiveOffloadKindMask); + auto SPVFile = sycl::runLLVMToSPIRVTranslation(InputFiles[0], Args); if (!SPVFile) return SPVFile.takeError(); @@ -1891,10 +1993,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: @@ -2370,9 +2472,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()) { @@ -2432,7 +2533,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(); @@ -2757,6 +2859,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">;