From 90fe7b683051ba8301fcb602448bbea7f244c5f5 Mon Sep 17 00:00:00 2001 From: "Sabianin, Maksim" Date: Fri, 8 May 2026 09:41:50 -0700 Subject: [PATCH 01/25] [SYCL][NewOffloadModel] Enable -no-sycl-rdc in Clang Driver for NewOffloadModel. --- clang/lib/Driver/ToolChains/Clang.cpp | 8 +++ clang/test/Driver/sycl-no-rdc-new-driver.cpp | 64 ++++++++++++++++++++ sycl/doc/design/NonRelocatableDeviceCode.md | 2 + sycl/doc/design/OffloadDesign.md | 20 ++++++ sycl/test-e2e/AOT/early_aot.cpp | 12 ++-- 5 files changed, 101 insertions(+), 5 deletions(-) create mode 100644 clang/test/Driver/sycl-no-rdc-new-driver.cpp diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index 15b81a8f94adc..d50eadbfee7de 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -12154,6 +12154,14 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA, if (Args.hasArg(options::OPT_fsycl_link_EQ)) CmdArgs.push_back(Args.MakeArgString("--sycl-device-link")); + // Propagate [no-]rdc mode to the linker wrapper for the SYCL case. + // The default behaviour is rdc mode ON, which requires no special flags. + // In order to enable non-rdc mode, we pass --no-sycl-rdc to the linker + // wrapper. Note: -f[no-]sycl-rdc is an alias of [no-]gpu_rdc. + if (!Args.hasFlag(options::OPT_fgpu_rdc, options::OPT_fno_gpu_rdc, + /*default=*/true)) + CmdArgs.push_back("--no-sycl-rdc"); + // -sycl-device-library-location= provides the location in which the // SYCL device libraries can be found. SmallString<128> DeviceLibDir(D.Dir); diff --git a/clang/test/Driver/sycl-no-rdc-new-driver.cpp b/clang/test/Driver/sycl-no-rdc-new-driver.cpp new file mode 100644 index 0000000000000..0ec31c67c5f2f --- /dev/null +++ b/clang/test/Driver/sycl-no-rdc-new-driver.cpp @@ -0,0 +1,64 @@ +/// Tests for -f[no-]sycl-rdc with --offload-new-driver. + +// Verifies that --no-sycl-rdc is propagated to clang-linker-wrapper when +// -fno-sycl-rdc is passed. RDC is ON by default; --no-sycl-rdc signals +// RDC is OFF. + +// RUN: touch %t.cpp + +// Default (no flag): RDC is ON by default for SYCL, so --no-sycl-rdc should NOT appear. +// RUN: %clang -### --offload-new-driver --target=x86_64-unknown-linux-gnu -fsycl %t.cpp 2>&1 \ +// RUN: | FileCheck -check-prefix=CHK-DEFAULT %s +// CHK-DEFAULT-NOT: --no-sycl-rdc + +// -fno-sycl-rdc: --no-sycl-rdc should appear. +// RUN: %clang -### --offload-new-driver -Werror --target=x86_64-unknown-linux-gnu -fsycl -fno-sycl-rdc %t.cpp 2>&1 \ +// RUN: | FileCheck -check-prefix=CHK-NO-RDC %s +// CHK-NO-RDC: clang-linker-wrapper{{.*}} "--no-sycl-rdc" + +// AOT Intel GPU target, default RDC: --no-sycl-rdc should NOT appear. +// RUN: %clang -### --offload-new-driver --target=x86_64-unknown-linux-gnu -fsycl -fsycl-targets=intel_gpu_pvc %t.cpp 2>&1 \ +// RUN: | FileCheck -check-prefix=CHK-AOT-RDC %s +// CHK-AOT-RDC-NOT: --no-sycl-rdc + +// AOT Intel GPU target + -fno-sycl-rdc: --no-sycl-rdc should appear. +// RUN: %clang -### --offload-new-driver -Werror --target=x86_64-unknown-linux-gnu -fsycl -fsycl-targets=intel_gpu_pvc -fno-sycl-rdc %t.cpp 2>&1 \ +// RUN: | FileCheck -check-prefix=CHK-AOT-NO-RDC %s +// CHK-AOT-NO-RDC: clang-linker-wrapper{{.*}} "--no-sycl-rdc" + +// Test compilation step. +// RUN: not %clang -### --offload-new-driver -Werror --target=x86_64-unknown-linux-gnu -fsycl -fsycl-targets=spir64_gen -fno-sycl-rdc %t.cpp -c -o %t.o 2>&1 \ +// RUN: | FileCheck -check-prefix=CHK-COMPILE-STEP-ERROR %s + +// CHK-COMPILE-STEP-ERROR: error: argument unused during compilation: '-fno-sycl-rdc' [-Werror,-Wunused-command-line-argument] + +// Verify pipeline with --offload-new-driver -fno-sycl-rdc. +// RUN: touch %t1.cpp +// RUN: touch %t2.cpp +// RUN: %clang -### --offload-new-driver -fsycl -fno-sycl-rdc %t1.cpp %t2.cpp 2>&1 -ccc-print-phases | FileCheck %s --check-prefix=CHECK-PIPELINE + +// CHECK-PIPELINE: 0: input, "{{.*}}1.cpp", c++, (host-sycl) +// CHECK-PIPELINE: 1: preprocessor, {0}, c++-cpp-output, (host-sycl) +// CHECK-PIPELINE: 2: compiler, {1}, ir, (host-sycl) +// CHECK-PIPELINE: 3: input, "{{.*}}1.cpp", c++, (device-sycl) +// CHECK-PIPELINE: 4: preprocessor, {3}, c++-cpp-output, (device-sycl) +// CHECK-PIPELINE: 5: compiler, {4}, ir, (device-sycl) +// CHECK-PIPELINE: 6: backend, {5}, ir, (device-sycl) +// CHECK-PIPELINE: 7: offload, "device-sycl (spir64-unknown-unknown)" {6}, ir +// CHECK-PIPELINE: 8: llvm-offload-binary, {7}, image, (device-sycl) +// CHECK-PIPELINE: 9: offload, "host-sycl (x86_64-unknown-linux-gnu)" {2}, "device-sycl (x86_64-unknown-linux-gnu)" {8}, ir +// CHECK-PIPELINE: 10: backend, {9}, assembler, (host-sycl) +// CHECK-PIPELINE: 11: assembler, {10}, object, (host-sycl) +// CHECK-PIPELINE: 12: input, "{{.*}}2.cpp", c++, (host-sycl) +// CHECK-PIPELINE: 13: preprocessor, {12}, c++-cpp-output, (host-sycl) +// CHECK-PIPELINE: 14: compiler, {13}, ir, (host-sycl) +// CHECK-PIPELINE: 15: input, "{{.*}}2.cpp", c++, (device-sycl) +// CHECK-PIPELINE: 16: preprocessor, {15}, c++-cpp-output, (device-sycl) +// CHECK-PIPELINE: 17: compiler, {16}, ir, (device-sycl) +// CHECK-PIPELINE: 18: backend, {17}, ir, (device-sycl) +// CHECK-PIPELINE: 19: offload, "device-sycl (spir64-unknown-unknown)" {18}, ir +// CHECK-PIPELINE: 20: llvm-offload-binary, {19}, image, (device-sycl) +// CHECK-PIPELINE: 21: offload, "host-sycl (x86_64-unknown-linux-gnu)" {14}, "device-sycl (x86_64-unknown-linux-gnu)" {20}, ir +// CHECK-PIPELINE: 22: backend, {21}, assembler, (host-sycl) +// CHECK-PIPELINE: 23: assembler, {22}, object, (host-sycl) +// CHECK-PIPELINE: 24: clang-linker-wrapper, {11, 23}, image, (host-sycl) diff --git a/sycl/doc/design/NonRelocatableDeviceCode.md b/sycl/doc/design/NonRelocatableDeviceCode.md index 8b74beca63be0..edb843185a98c 100644 --- a/sycl/doc/design/NonRelocatableDeviceCode.md +++ b/sycl/doc/design/NonRelocatableDeviceCode.md @@ -1,5 +1,7 @@ # Non-Relocatable Device Code +Note: This document reflects the design of NoRDC Mode support in the Old Offload Model. + ## Overview By default, SYCL allows device code to be relocatable, where function calls outside of the current translation unit are allowed using the `SYCL_EXTERNAL` attribute. diff --git a/sycl/doc/design/OffloadDesign.md b/sycl/doc/design/OffloadDesign.md index 35a86d2184ea2..9fee5f7cc4d94 100644 --- a/sycl/doc/design/OffloadDesign.md +++ b/sycl/doc/design/OffloadDesign.md @@ -243,6 +243,26 @@ are needed to pass along this information. *Table: Ahead of Time Info* +### NoRDC Mode support + +For the Old Offload Model support of NoRDC Mode see [NonRelocatableDeviceCode.md](NonRelocatableDeviceCode.md). + +The default compiler behavior is -fsycl-rdc, which incorporates linking of device code. If -fno-sycl-rdc is specified, the compiler skips linking of device code and performs offload processing on every module individually. + +The usage scenario for the old offload model is: +``` +clang++ --no-offload-old-driver -fsycl -fsycl-targets=T1,T2 input1.cpp -fno-sycl-rdc -c -o object1.o # -fno-sycl-rdc is specified +clang++ --no-offload-old-driver -fsycl -fsycl-targets=T1,T2 input2.cpp -fno-sycl-rdc -c -o object2.o # -fno-sycl-rdc is specified +clang++ --no-offload-old-driver -fsycl -fsycl-targets=T1,T2 object1.o object2.o -o a.out # -fno-sycl-rdc is NOT specified +``` + +Currently, SYCL offload processing resides in clang-linker-wrapper. That leads to the following usage scenario, which is inverted compared to the old offload model: +``` +clang++ --offload-new-driver -fsycl -fsycl-targets=T1,T2 input1.cpp -c -o object1.o # -fno-sycl-rdc is NOT specified +clang++ --offload-new-driver -fsycl -fsycl-targets=T1,T2 input2.cpp -c -o object2.o # -fno-sycl-rdc is NOT specified +clang++ --offload-new-driver -fsycl -fsycl-targets=T1,T2 -fno-sycl-rdc object1.o object2.o -o a.out # -fno-sycl-rdc is specified +``` + #### Format of the --device-compiler Option The `--device-compiler` option uses the format `--device-compiler=[:][=]` where: - `` : specifies the offloading kind (e.g., sycl, hip, openmp) and is optional. diff --git a/sycl/test-e2e/AOT/early_aot.cpp b/sycl/test-e2e/AOT/early_aot.cpp index 49385d7857fa5..c7eddb56832a6 100644 --- a/sycl/test-e2e/AOT/early_aot.cpp +++ b/sycl/test-e2e/AOT/early_aot.cpp @@ -1,13 +1,15 @@ // Test early-AOT behaviors with -fsycl -fno-sycl-rdc. This targets spir64_gen // REQUIRES: ocloc, gpu, target-spir -// XFAIL: new-offload-model -// XFAIL-TRACKER: CMPLRLLVM-51875 + +// Note: Temporary. New Offload Model requires -fno-sycl-rdc to be specified +// at linking step but Old Offload Model requires -fno-sycl-rdc to be +// specified at compilation step. // Build the early AOT device binaries -// RUN: %clangxx -fsycl -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen %gpu_aot_target_opts -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 -fno-sycl-rdc -c -DSUB_CPP %s -o %t_sub.o -// RUN: %clangxx -fsycl -DMAIN_CPP %s %t_add.o %t_sub.o -o %t.out +// 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 %if new-offload-model %{ -fno-sycl-rdc %} -DMAIN_CPP %s %t_add.o %t_sub.o -o %t.out // RUN: %{run} %t.out From 998d20f475cd090a0d5d556e458049165188ce8d Mon Sep 17 00:00:00 2001 From: "Sabianin, Maksim" Date: Thu, 2 Jul 2026 08:33:48 -0700 Subject: [PATCH 02/25] fix test for cross-platform execution --- clang/test/Driver/sycl-no-rdc-new-driver.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/test/Driver/sycl-no-rdc-new-driver.cpp b/clang/test/Driver/sycl-no-rdc-new-driver.cpp index 0ec31c67c5f2f..2f63c65ad6bcd 100644 --- a/clang/test/Driver/sycl-no-rdc-new-driver.cpp +++ b/clang/test/Driver/sycl-no-rdc-new-driver.cpp @@ -35,7 +35,7 @@ // Verify pipeline with --offload-new-driver -fno-sycl-rdc. // RUN: touch %t1.cpp // RUN: touch %t2.cpp -// RUN: %clang -### --offload-new-driver -fsycl -fno-sycl-rdc %t1.cpp %t2.cpp 2>&1 -ccc-print-phases | FileCheck %s --check-prefix=CHECK-PIPELINE +// RUN: %clang -### --offload-new-driver --target=x86_64-unknown-linux-gnu -fsycl -fno-sycl-rdc %t1.cpp %t2.cpp 2>&1 -ccc-print-phases | FileCheck %s --check-prefix=CHECK-PIPELINE // CHECK-PIPELINE: 0: input, "{{.*}}1.cpp", c++, (host-sycl) // CHECK-PIPELINE: 1: preprocessor, {0}, c++-cpp-output, (host-sycl) From 80d7d10986a4273452a1aa65377a749a767e774d Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Thu, 30 Jul 2026 13:46:27 -0700 Subject: [PATCH 03/25] [SYCL][NewOffloadModel] Address review comments on PR #21973 Fix --no-offload-old-driver typo to --no-offload-new-driver in the old offload model usage example. Add a note that compile-step support for -fno-sycl-rdc (matching the old model's -c usage pattern) will be implemented in a follow-up patch via clang-linker-wrapper --sycl-device-link --no-sycl-rdc per translation unit. --- sycl/doc/design/OffloadDesign.md | 12 +++++++++--- 1 file changed, 9 insertions(+), 3 deletions(-) diff --git a/sycl/doc/design/OffloadDesign.md b/sycl/doc/design/OffloadDesign.md index 9fee5f7cc4d94..887d995a49f94 100644 --- a/sycl/doc/design/OffloadDesign.md +++ b/sycl/doc/design/OffloadDesign.md @@ -251,9 +251,9 @@ The default compiler behavior is -fsycl-rdc, which incorporates linking of devic The usage scenario for the old offload model is: ``` -clang++ --no-offload-old-driver -fsycl -fsycl-targets=T1,T2 input1.cpp -fno-sycl-rdc -c -o object1.o # -fno-sycl-rdc is specified -clang++ --no-offload-old-driver -fsycl -fsycl-targets=T1,T2 input2.cpp -fno-sycl-rdc -c -o object2.o # -fno-sycl-rdc is specified -clang++ --no-offload-old-driver -fsycl -fsycl-targets=T1,T2 object1.o object2.o -o a.out # -fno-sycl-rdc is NOT specified +clang++ --no-offload-new-driver -fsycl -fsycl-targets=T1,T2 input1.cpp -fno-sycl-rdc -c -o object1.o # -fno-sycl-rdc is specified +clang++ --no-offload-new-driver -fsycl -fsycl-targets=T1,T2 input2.cpp -fno-sycl-rdc -c -o object2.o # -fno-sycl-rdc is specified +clang++ --no-offload-new-driver -fsycl -fsycl-targets=T1,T2 object1.o object2.o -o a.out # -fno-sycl-rdc is NOT specified ``` Currently, SYCL offload processing resides in clang-linker-wrapper. That leads to the following usage scenario, which is inverted compared to the old offload model: @@ -263,6 +263,12 @@ clang++ --offload-new-driver -fsycl -fsycl-targets=T1,T2 input2.cpp -c -o object clang++ --offload-new-driver -fsycl -fsycl-targets=T1,T2 -fno-sycl-rdc object1.o object2.o -o a.out # -fno-sycl-rdc is specified ``` +A follow-up patch will add support for specifying `-fno-sycl-rdc` at the compile step +(i.e. `clang++ --offload-new-driver -fsycl -fno-sycl-rdc -c`), matching the old offload +model's usage pattern. This will be implemented by invoking `clang-linker-wrapper +--sycl-device-link --no-sycl-rdc` per translation unit at compile time to finalize each +TU's device code independently, embedding the result directly into the host object. + #### Format of the --device-compiler Option The `--device-compiler` option uses the format `--device-compiler=[:][=]` where: - `` : specifies the offloading kind (e.g., sycl, hip, openmp) and is optional. From dbd5c446dcde460edc4b49e2927f9c602c0e7e3f Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Thu, 30 Jul 2026 14:17:35 -0700 Subject: [PATCH 04/25] [SYCL][NewOffloadModel] Add -fno-sycl-rdc support at compile step (-c) Adds support for specifying -fno-sycl-rdc at the compile step in the new offload model, matching the old offload model's usage pattern. Previously, -fno-sycl-rdc was only meaningful at the link step in the new model. Design (mirrors CUDA/HIP -fno-gpu-rdc per-TU finalize model): - BuildOffloadingActions: when -fno-sycl-rdc is passed at compile time, insert a per-TU LinkerWrapperJobAction (OffloadPackager + clang-linker-wrapper --sycl-device-link --no-sycl-rdc) to finalize the TU's device code immediately, producing a self-contained device image. - Clang::ConstructJob: route the finalized image to the new -fsycl-include-target-binary CC1 option instead of -fembed-offload-object, so the host CodeGen embeds and registers the image at compile time rather than deferring it to link time. - CodeGenModule::Release: when -fsycl-include-target-binary is set on the host cc1, read the device image file and call wrapSYCLBinaries to embed the binary and emit __sycl_register_lib/__sycl_unregister_lib ctors into the host module. The final clang-linker-wrapper invocation then finds no SYCL device input and performs a plain host link. - Downstream uses clang-linker-wrapper --sycl-device-link instead of the upstream's clang-sycl-linker, as all Intel SYCL pipeline logic resides in clang-linker-wrapper. New files/options: - -fsycl-include-target-binary : new CC1 option (analog of -fcuda-include-gpubinary) marshalled into CodeGenOpts.SYCLTargetBinaryFileName - clang/test/Driver/sycl-no-rdc-compile-step.cpp: driver tests verifying per-TU pipeline shape, RDC default regression, and link-step regression. Implements compile-step support tracked in CMPLRLLVM-51875. Co-Authored-By: Claude Sonnet 4.6 --- clang/include/clang/Basic/CodeGenOptions.h | 5 +++ clang/include/clang/Options/Options.td | 3 ++ clang/lib/CodeGen/CodeGenModule.cpp | 27 ++++++++++++++ clang/lib/Driver/Driver.cpp | 22 ++++++++++- clang/lib/Driver/ToolChains/Clang.cpp | 28 +++++++++++++- .../test/Driver/sycl-no-rdc-compile-step.cpp | 37 +++++++++++++++++++ 6 files changed, 118 insertions(+), 4 deletions(-) create mode 100644 clang/test/Driver/sycl-no-rdc-compile-step.cpp diff --git a/clang/include/clang/Basic/CodeGenOptions.h b/clang/include/clang/Basic/CodeGenOptions.h index 8b25279045441..d0b62f2adeccf 100644 --- a/clang/include/clang/Basic/CodeGenOptions.h +++ b/clang/include/clang/Basic/CodeGenOptions.h @@ -401,6 +401,11 @@ class CodeGenOptions : public CodeGenOptionsBase { /// CUDA runtime back-end for incorporating them into host-side object file. std::string CudaGpuBinaryFileName; + /// Name of file passed with -fsycl-include-target-binary option. Used in + /// -fno-sycl-rdc mode to embed and register a per-TU finalized SYCL device + /// image into the host object at compile time. + std::string SYCLTargetBinaryFileName; + /// List of filenames passed in using the -fembed-offload-object option. These /// are offloading binaries containing device images and metadata. std::vector OffloadObjects; diff --git a/clang/include/clang/Options/Options.td b/clang/include/clang/Options/Options.td index 837407f0ba574..153629c786c9e 100644 --- a/clang/include/clang/Options/Options.td +++ b/clang/include/clang/Options/Options.td @@ -9469,6 +9469,9 @@ def fcuda_is_device : Flag<["-"], "fcuda-is-device">, def fcuda_include_gpubinary : Separate<["-"], "fcuda-include-gpubinary">, HelpText<"Incorporate CUDA device-side binary into host object file.">, MarshallingInfoString>; +def fsycl_include_target_binary : Separate<["-"], "fsycl-include-target-binary">, + HelpText<"Incorporate SYCL device-side binary into host object file.">, + MarshallingInfoString>; def fcuda_allow_variadic_functions : Flag<["-"], "fcuda-allow-variadic-functions">, HelpText<"Deprecated; Allow variadic functions in CUDA device code.">; def fno_cuda_host_device_constexpr : Flag<["-"], "fno-cuda-host-device-constexpr">, diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 894781d6a1b77..3c728ed4f4d54 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -84,6 +84,7 @@ #include "llvm/Transforms/Utils/BuildLibCalls.h" #include "llvm/Transforms/Utils/KCFIHash.h" #include "llvm/Transforms/Utils/ModuleUtils.h" +#include "llvm/Frontend/Offloading/OffloadWrapper.h" #include #include @@ -1257,6 +1258,32 @@ void CodeGenModule::Release() { if (llvm::Function *CudaCtorFunction = CUDARuntime->finalizeModule()) AddGlobalCtor(CudaCtorFunction); } + // SYCL -fno-sycl-rdc: the per-TU device image was finalized at compile time + // by clang-linker-wrapper --sycl-device-link. Read it here and call + // wrapSYCLBinaries to embed the binary and emit __sycl_register_lib / + // __sycl_unregister_lib into the host module, mirroring what CUDA does via + // CGCUDANV::makeModuleCtorFunction for -fno-gpu-rdc. + // SYCLJITOptions is left default (empty) because device code is already + // fully finalized -- there are no JIT compile/link options to encode. + if (LangOpts.SYCLIsHost && + !getCodeGenOpts().SYCLTargetBinaryFileName.empty()) { + auto VFS = getFileSystem(); + auto BinaryOrErr = + VFS->getBufferForFile(getCodeGenOpts().SYCLTargetBinaryFileName, + -1, false); + if (std::error_code EC = BinaryOrErr.getError()) { + getDiags().Report(diag::err_cannot_open_file) + << getCodeGenOpts().SYCLTargetBinaryFileName << EC.message(); + } else { + llvm::ArrayRef Buffer((*BinaryOrErr)->getBufferStart(), + (*BinaryOrErr)->getBufferSize()); + if (llvm::Error E = llvm::offloading::wrapSYCLBinaries( + getModule(), Buffer, llvm::offloading::SYCLJITOptions{})) + getDiags().Report(diag::err_cannot_open_file) + << getCodeGenOpts().SYCLTargetBinaryFileName + << llvm::toString(std::move(E)); + } + } if (OpenMPRuntime) { OpenMPRuntime->createOffloadEntriesAndInfoMetadata(); OpenMPRuntime->clear(); diff --git a/clang/lib/Driver/Driver.cpp b/clang/lib/Driver/Driver.cpp index 7a0d81674601d..859492e8c2b22 100644 --- a/clang/lib/Driver/Driver.cpp +++ b/clang/lib/Driver/Driver.cpp @@ -8339,9 +8339,27 @@ Driver::BuildOffloadingActions(Compilation &C, llvm::opt::DerivedArgList &Args, DDep.add(*PackagerAction, *C.getOffloadToolChains().first->second, /*BA=*/{}, Action::OFK_HIP); + } else if (C.isOffloadingHostKind(Action::OFK_SYCL) && + tools::SYCL::shouldDoPerObjectFileLinking(C) && + !isa(HostAction)) { + // SYCL -fno-sycl-rdc at compile time (-c): finalize this TU's device code + // immediately via clang-linker-wrapper --sycl-device-link --no-sycl-rdc, + // producing a self-contained device image. The image is passed to the host + // cc1 via -fsycl-include-target-binary and embedded+registered at compile + // time. The final link step does no SYCL device work. This mirrors the + // CUDA/HIP -fno-gpu-rdc per-TU finalize model; downstream uses + // clang-linker-wrapper instead of clang-sycl-linker. + Action *PackagerAction = + C.MakeAction(OffloadActions, types::TY_Image); + ActionList AL{PackagerAction}; + Action *FinalizeAction = + C.MakeAction(AL, types::TY_Image); + DDep.add(*FinalizeAction, + *C.getSingleOffloadToolChain(), + /*BA=*/{}, Action::OFK_SYCL); } else { - // Package all the offloading actions into a single output that can be - // embedded in the host and linked. + // RDC (default): package raw device bitcode to be embedded in the host + // object and device-linked across all TUs at final link time. Action *PackagerAction = C.MakeAction(OffloadActions, types::TY_Image); DDep.add(*PackagerAction, *C.getSingleOffloadToolChain(), diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index d50eadbfee7de..0fc8040058043 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -9149,10 +9149,23 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA, CmdArgs.push_back(CudaDeviceInput->getFilename()); } else if (!HostOffloadingInputs.empty()) { if ((IsCuda || IsHIP) && !IsRDCMode) { + // CUDA/HIP no-RDC: device image finalized per-TU, embed at compile time. assert(HostOffloadingInputs.size() == 1 && "Only one input expected"); CmdArgs.push_back("-fcuda-include-gpubinary"); CmdArgs.push_back(HostOffloadingInputs.front().getFilename()); + } else if (IsSYCL && !IsRDCMode) { + // SYCL no-RDC (-fno-sycl-rdc): device image was finalized per-TU by + // clang-linker-wrapper --sycl-device-link. Pass it to host CodeGen via + // -fsycl-include-target-binary so wrapSYCLBinaries embeds and registers + // it at compile time. -fembed-offload-object is intentionally NOT used: + // that would defer processing to link time, defeating the purpose. + // Note: this block is only reached during the host cc1 compile step; + // at final link time HostOffloadingInputs is empty for SYCL no-RDC. + assert(HostOffloadingInputs.size() == 1 && "One finalized image per TU"); + CmdArgs.push_back("-fsycl-include-target-binary"); + CmdArgs.push_back(HostOffloadingInputs.front().getFilename()); } else { + // RDC: embed raw device bitcode for cross-TU device link at link time. for (const InputInfo Input : HostOffloadingInputs) CmdArgs.push_back(Args.MakeArgString("-fembed-offload-object=" + TC.getInputFilename(Input))); @@ -12158,9 +12171,20 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA, // The default behaviour is rdc mode ON, which requires no special flags. // In order to enable non-rdc mode, we pass --no-sycl-rdc to the linker // wrapper. Note: -f[no-]sycl-rdc is an alias of [no-]gpu_rdc. - if (!Args.hasFlag(options::OPT_fgpu_rdc, options::OPT_fno_gpu_rdc, - /*default=*/true)) + bool IsSYCLNoRDC = !Args.hasFlag(options::OPT_fgpu_rdc, + options::OPT_fno_gpu_rdc, + /*default=*/true); + if (IsSYCLNoRDC) { CmdArgs.push_back("--no-sycl-rdc"); + // When invoked as a per-TU device finalizer at compile time + // (-fno-sycl-rdc -c), the linker wrapper must run in device-link-only + // mode: execute the full SYCL pipeline but skip the host linker. + // TY_Image output identifies this as the per-TU finalize action + // (BuildOffloadingActions) rather than the final link-time invocation. + if (!Args.hasArg(options::OPT_fsycl_link_EQ) && + JA.getType() == types::TY_Image) + CmdArgs.push_back("--sycl-device-link"); + } // -sycl-device-library-location= provides the location in which the // SYCL device libraries can be found. diff --git a/clang/test/Driver/sycl-no-rdc-compile-step.cpp b/clang/test/Driver/sycl-no-rdc-compile-step.cpp new file mode 100644 index 0000000000000..0d7ef2853e42e --- /dev/null +++ b/clang/test/Driver/sycl-no-rdc-compile-step.cpp @@ -0,0 +1,37 @@ +/// Tests for -fno-sycl-rdc at the compile step (-c) with --offload-new-driver. +/// Verifies that the driver inserts a per-TU clang-linker-wrapper finalize +/// action and routes the result to -fsycl-include-target-binary on the host +/// cc1, instead of deferring device processing to link time. + +// RUN: touch %t.cpp + +// --- CHECK 1: -fno-sycl-rdc -c --- +// Per-TU finalize: clang-linker-wrapper must be invoked with --sycl-device-link +// and --no-sycl-rdc. Host cc1 must receive -fsycl-include-target-binary. +// -fembed-offload-object must NOT appear (device code is not deferred). +// RUN: %clang -### --offload-new-driver -Werror --target=x86_64-unknown-linux-gnu \ +// RUN: -fsycl -fno-sycl-rdc -c %t.cpp 2>&1 \ +// RUN: | FileCheck -check-prefix=CHK-COMPILE %s +// CHK-COMPILE: clang-linker-wrapper{{.*}} "--sycl-device-link"{{.*}} "--no-sycl-rdc" +// CHK-COMPILE: "-fsycl-include-target-binary" +// CHK-COMPILE-NOT: -fembed-offload-object + +// --- CHECK 2: default RDC -c (no flag) --- +// Default RDC path: -fembed-offload-object must appear (raw bitcode deferred to +// link time). -fsycl-include-target-binary and --no-sycl-rdc must NOT appear. +// RUN: %clang -### --offload-new-driver --target=x86_64-unknown-linux-gnu \ +// RUN: -fsycl -c %t.cpp 2>&1 \ +// RUN: | FileCheck -check-prefix=CHK-RDC %s +// CHK-RDC: -fembed-offload-object +// CHK-RDC-NOT: -fsycl-include-target-binary +// CHK-RDC-NOT: --no-sycl-rdc + +// --- CHECK 3: -fno-sycl-rdc at link step (regression for PR #22832) --- +// At link time --no-sycl-rdc is forwarded to clang-linker-wrapper. +// --sycl-device-link must NOT appear for the final link invocation. +// RUN: touch %t.o +// RUN: %clang -### --offload-new-driver -Werror --target=x86_64-unknown-linux-gnu \ +// RUN: -fsycl -fno-sycl-rdc %t.o 2>&1 \ +// RUN: | FileCheck -check-prefix=CHK-LINK %s +// CHK-LINK: clang-linker-wrapper{{.*}} "--no-sycl-rdc" +// CHK-LINK-NOT: clang-linker-wrapper{{.*}} "--sycl-device-link" From 07a16e7e803bb2756f0f302252299eaa65e9beaa Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Thu, 30 Jul 2026 14:32:42 -0700 Subject: [PATCH 05/25] [SYCL][Driver] Fix tests for -fno-sycl-rdc compile-step support Co-Authored-By: Claude Sonnet 4.6 --- .../test/Driver/sycl-no-rdc-compile-step.cpp | 2 +- clang/test/Driver/sycl-no-rdc-new-driver.cpp | 40 +++++++++---------- 2 files changed, 19 insertions(+), 23 deletions(-) diff --git a/clang/test/Driver/sycl-no-rdc-compile-step.cpp b/clang/test/Driver/sycl-no-rdc-compile-step.cpp index 0d7ef2853e42e..ee9fc6e874c86 100644 --- a/clang/test/Driver/sycl-no-rdc-compile-step.cpp +++ b/clang/test/Driver/sycl-no-rdc-compile-step.cpp @@ -12,7 +12,7 @@ // RUN: %clang -### --offload-new-driver -Werror --target=x86_64-unknown-linux-gnu \ // RUN: -fsycl -fno-sycl-rdc -c %t.cpp 2>&1 \ // RUN: | FileCheck -check-prefix=CHK-COMPILE %s -// CHK-COMPILE: clang-linker-wrapper{{.*}} "--sycl-device-link"{{.*}} "--no-sycl-rdc" +// CHK-COMPILE: clang-linker-wrapper{{.*}} "--no-sycl-rdc"{{.*}} "--sycl-device-link" // CHK-COMPILE: "-fsycl-include-target-binary" // CHK-COMPILE-NOT: -fembed-offload-object diff --git a/clang/test/Driver/sycl-no-rdc-new-driver.cpp b/clang/test/Driver/sycl-no-rdc-new-driver.cpp index 2f63c65ad6bcd..15a09f46f8a73 100644 --- a/clang/test/Driver/sycl-no-rdc-new-driver.cpp +++ b/clang/test/Driver/sycl-no-rdc-new-driver.cpp @@ -26,12 +26,6 @@ // RUN: | FileCheck -check-prefix=CHK-AOT-NO-RDC %s // CHK-AOT-NO-RDC: clang-linker-wrapper{{.*}} "--no-sycl-rdc" -// Test compilation step. -// RUN: not %clang -### --offload-new-driver -Werror --target=x86_64-unknown-linux-gnu -fsycl -fsycl-targets=spir64_gen -fno-sycl-rdc %t.cpp -c -o %t.o 2>&1 \ -// RUN: | FileCheck -check-prefix=CHK-COMPILE-STEP-ERROR %s - -// CHK-COMPILE-STEP-ERROR: error: argument unused during compilation: '-fno-sycl-rdc' [-Werror,-Wunused-command-line-argument] - // Verify pipeline with --offload-new-driver -fno-sycl-rdc. // RUN: touch %t1.cpp // RUN: touch %t2.cpp @@ -46,19 +40,21 @@ // CHECK-PIPELINE: 6: backend, {5}, ir, (device-sycl) // CHECK-PIPELINE: 7: offload, "device-sycl (spir64-unknown-unknown)" {6}, ir // CHECK-PIPELINE: 8: llvm-offload-binary, {7}, image, (device-sycl) -// CHECK-PIPELINE: 9: offload, "host-sycl (x86_64-unknown-linux-gnu)" {2}, "device-sycl (x86_64-unknown-linux-gnu)" {8}, ir -// CHECK-PIPELINE: 10: backend, {9}, assembler, (host-sycl) -// CHECK-PIPELINE: 11: assembler, {10}, object, (host-sycl) -// CHECK-PIPELINE: 12: input, "{{.*}}2.cpp", c++, (host-sycl) -// CHECK-PIPELINE: 13: preprocessor, {12}, c++-cpp-output, (host-sycl) -// CHECK-PIPELINE: 14: compiler, {13}, ir, (host-sycl) -// CHECK-PIPELINE: 15: input, "{{.*}}2.cpp", c++, (device-sycl) -// CHECK-PIPELINE: 16: preprocessor, {15}, c++-cpp-output, (device-sycl) -// CHECK-PIPELINE: 17: compiler, {16}, ir, (device-sycl) -// CHECK-PIPELINE: 18: backend, {17}, ir, (device-sycl) -// CHECK-PIPELINE: 19: offload, "device-sycl (spir64-unknown-unknown)" {18}, ir -// CHECK-PIPELINE: 20: llvm-offload-binary, {19}, image, (device-sycl) -// CHECK-PIPELINE: 21: offload, "host-sycl (x86_64-unknown-linux-gnu)" {14}, "device-sycl (x86_64-unknown-linux-gnu)" {20}, ir -// CHECK-PIPELINE: 22: backend, {21}, assembler, (host-sycl) -// CHECK-PIPELINE: 23: assembler, {22}, object, (host-sycl) -// CHECK-PIPELINE: 24: clang-linker-wrapper, {11, 23}, image, (host-sycl) +// CHECK-PIPELINE: 9: clang-linker-wrapper, {8}, image, (device-sycl) +// CHECK-PIPELINE: 10: offload, "host-sycl (x86_64-unknown-linux-gnu)" {2}, "device-sycl (x86_64-unknown-linux-gnu)" {9}, ir +// CHECK-PIPELINE: 11: backend, {10}, assembler, (host-sycl) +// CHECK-PIPELINE: 12: assembler, {11}, object, (host-sycl) +// CHECK-PIPELINE: 13: input, "{{.*}}2.cpp", c++, (host-sycl) +// CHECK-PIPELINE: 14: preprocessor, {13}, c++-cpp-output, (host-sycl) +// CHECK-PIPELINE: 15: compiler, {14}, ir, (host-sycl) +// CHECK-PIPELINE: 16: input, "{{.*}}2.cpp", c++, (device-sycl) +// CHECK-PIPELINE: 17: preprocessor, {16}, c++-cpp-output, (device-sycl) +// CHECK-PIPELINE: 18: compiler, {17}, ir, (device-sycl) +// CHECK-PIPELINE: 19: backend, {18}, ir, (device-sycl) +// CHECK-PIPELINE: 20: offload, "device-sycl (spir64-unknown-unknown)" {19}, ir +// CHECK-PIPELINE: 21: llvm-offload-binary, {20}, image, (device-sycl) +// CHECK-PIPELINE: 22: clang-linker-wrapper, {21}, image, (device-sycl) +// CHECK-PIPELINE: 23: offload, "host-sycl (x86_64-unknown-linux-gnu)" {15}, "device-sycl (x86_64-unknown-linux-gnu)" {22}, ir +// CHECK-PIPELINE: 24: backend, {23}, assembler, (host-sycl) +// CHECK-PIPELINE: 25: assembler, {24}, object, (host-sycl) +// CHECK-PIPELINE: 26: clang-linker-wrapper, {12, 25}, image, (host-sycl) From 2a41c6f19906747de724b90a9d2f21ad190f4f1e Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Thu, 30 Jul 2026 15:33:24 -0700 Subject: [PATCH 06/25] [SYCL] Fix clang-format issues --- clang/lib/CodeGen/CodeGenModule.cpp | 9 ++++----- clang/lib/Driver/Driver.cpp | 3 +-- clang/lib/Driver/ToolChains/Clang.cpp | 6 +++--- 3 files changed, 8 insertions(+), 10 deletions(-) diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 3c728ed4f4d54..74a242d4762f5 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -50,9 +50,9 @@ #include "clang/CodeGen/BackendUtil.h" #include "clang/CodeGen/ConstantInitBuilder.h" #include "clang/Frontend/FrontendDiagnostic.h" +#include "clang/Lex/Preprocessor.h" #include "clang/Sema/Sema.h" #include "clang/Sema/SemaSYCL.h" -#include "clang/Lex/Preprocessor.h" #include "llvm/ABI/IRTypeMapper.h" #include "llvm/ABI/TargetInfo.h" #include "llvm/ADT/STLExtras.h" @@ -60,6 +60,7 @@ #include "llvm/ADT/StringSwitch.h" #include "llvm/Analysis/TargetLibraryInfo.h" #include "llvm/BinaryFormat/ELF.h" +#include "llvm/Frontend/Offloading/OffloadWrapper.h" #include "llvm/IR/AttributeMask.h" #include "llvm/IR/CallingConv.h" #include "llvm/IR/DataLayout.h" @@ -84,7 +85,6 @@ #include "llvm/Transforms/Utils/BuildLibCalls.h" #include "llvm/Transforms/Utils/KCFIHash.h" #include "llvm/Transforms/Utils/ModuleUtils.h" -#include "llvm/Frontend/Offloading/OffloadWrapper.h" #include #include @@ -1268,9 +1268,8 @@ void CodeGenModule::Release() { if (LangOpts.SYCLIsHost && !getCodeGenOpts().SYCLTargetBinaryFileName.empty()) { auto VFS = getFileSystem(); - auto BinaryOrErr = - VFS->getBufferForFile(getCodeGenOpts().SYCLTargetBinaryFileName, - -1, false); + auto BinaryOrErr = VFS->getBufferForFile( + getCodeGenOpts().SYCLTargetBinaryFileName, -1, false); if (std::error_code EC = BinaryOrErr.getError()) { getDiags().Report(diag::err_cannot_open_file) << getCodeGenOpts().SYCLTargetBinaryFileName << EC.message(); diff --git a/clang/lib/Driver/Driver.cpp b/clang/lib/Driver/Driver.cpp index 859492e8c2b22..0f7c64279af45 100644 --- a/clang/lib/Driver/Driver.cpp +++ b/clang/lib/Driver/Driver.cpp @@ -8354,8 +8354,7 @@ Driver::BuildOffloadingActions(Compilation &C, llvm::opt::DerivedArgList &Args, ActionList AL{PackagerAction}; Action *FinalizeAction = C.MakeAction(AL, types::TY_Image); - DDep.add(*FinalizeAction, - *C.getSingleOffloadToolChain(), + DDep.add(*FinalizeAction, *C.getSingleOffloadToolChain(), /*BA=*/{}, Action::OFK_SYCL); } else { // RDC (default): package raw device bitcode to be embedded in the host diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index 0fc8040058043..0eff397a85607 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -12171,9 +12171,9 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA, // The default behaviour is rdc mode ON, which requires no special flags. // In order to enable non-rdc mode, we pass --no-sycl-rdc to the linker // wrapper. Note: -f[no-]sycl-rdc is an alias of [no-]gpu_rdc. - bool IsSYCLNoRDC = !Args.hasFlag(options::OPT_fgpu_rdc, - options::OPT_fno_gpu_rdc, - /*default=*/true); + bool IsSYCLNoRDC = + !Args.hasFlag(options::OPT_fgpu_rdc, options::OPT_fno_gpu_rdc, + /*default=*/true); if (IsSYCLNoRDC) { CmdArgs.push_back("--no-sycl-rdc"); // When invoked as a per-TU device finalizer at compile time From d8ece47b3f8d08bf44e65b9b85deee0e9b08a926 Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Thu, 30 Jul 2026 18:31:17 -0700 Subject: [PATCH 07/25] [SYCL] Address Copilot review comments on PR #22833 --- clang/lib/CodeGen/CodeGenModule.cpp | 2 +- clang/lib/Driver/ToolChains/Clang.cpp | 8 +++++--- sycl/doc/design/OffloadDesign.md | 11 ++++++----- sycl/test-e2e/AOT/early_aot.cpp | 6 +++--- 4 files changed, 15 insertions(+), 12 deletions(-) diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 74a242d4762f5..2d9452d83dd80 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -1278,7 +1278,7 @@ void CodeGenModule::Release() { (*BinaryOrErr)->getBufferSize()); if (llvm::Error E = llvm::offloading::wrapSYCLBinaries( getModule(), Buffer, llvm::offloading::SYCLJITOptions{})) - getDiags().Report(diag::err_cannot_open_file) + getDiags().Report(diag::err_fe_linking_module) << getCodeGenOpts().SYCLTargetBinaryFileName << llvm::toString(std::move(E)); } diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index 0eff397a85607..1916d1efd5cfb 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -12179,10 +12179,12 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA, // When invoked as a per-TU device finalizer at compile time // (-fno-sycl-rdc -c), the linker wrapper must run in device-link-only // mode: execute the full SYCL pipeline but skip the host linker. - // TY_Image output identifies this as the per-TU finalize action - // (BuildOffloadingActions) rather than the final link-time invocation. + // isDeviceOffloading(OFK_SYCL) distinguishes the per-TU finalizer + // (OFK_SYCL, from BuildOffloadingActions) from the final link-time + // invocation (OFK_Host) -- both produce TY_Image so output type alone + // is not a reliable discriminator. if (!Args.hasArg(options::OPT_fsycl_link_EQ) && - JA.getType() == types::TY_Image) + JA.isDeviceOffloading(Action::OFK_SYCL)) CmdArgs.push_back("--sycl-device-link"); } diff --git a/sycl/doc/design/OffloadDesign.md b/sycl/doc/design/OffloadDesign.md index 887d995a49f94..29d12b373eb26 100644 --- a/sycl/doc/design/OffloadDesign.md +++ b/sycl/doc/design/OffloadDesign.md @@ -263,11 +263,12 @@ clang++ --offload-new-driver -fsycl -fsycl-targets=T1,T2 input2.cpp -c -o object clang++ --offload-new-driver -fsycl -fsycl-targets=T1,T2 -fno-sycl-rdc object1.o object2.o -o a.out # -fno-sycl-rdc is specified ``` -A follow-up patch will add support for specifying `-fno-sycl-rdc` at the compile step -(i.e. `clang++ --offload-new-driver -fsycl -fno-sycl-rdc -c`), matching the old offload -model's usage pattern. This will be implemented by invoking `clang-linker-wrapper ---sycl-device-link --no-sycl-rdc` per translation unit at compile time to finalize each -TU's device code independently, embedding the result directly into the host object. +Specifying `-fno-sycl-rdc` at the compile step +(i.e. `clang++ --offload-new-driver -fsycl -fno-sycl-rdc -c`) is also supported, matching +the old offload model's usage pattern. In this case `clang-linker-wrapper +--sycl-device-link --no-sycl-rdc` is invoked per translation unit at compile time to +finalize each TU's device code independently, embedding the result directly into the host +object. #### Format of the --device-compiler Option The `--device-compiler` option uses the format `--device-compiler=[:][=]` where: diff --git a/sycl/test-e2e/AOT/early_aot.cpp b/sycl/test-e2e/AOT/early_aot.cpp index c7eddb56832a6..bc08b15d76b91 100644 --- a/sycl/test-e2e/AOT/early_aot.cpp +++ b/sycl/test-e2e/AOT/early_aot.cpp @@ -2,9 +2,9 @@ // REQUIRES: ocloc, gpu, target-spir -// Note: Temporary. New Offload Model requires -fno-sycl-rdc to be specified -// at linking step but Old Offload Model requires -fno-sycl-rdc to be -// specified at compilation step. +// Note: The Old Offload Model requires -fno-sycl-rdc at the compilation step. +// The New Offload Model accepts -fno-sycl-rdc at both the compile and link +// steps; the RUN lines below pass it at the link step for compatibility. // 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 From 97ecb2f9d7f8d5e321039a6b7f6091762272360f Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Fri, 31 Jul 2026 16:36:07 -0700 Subject: [PATCH 08/25] [SYCL] Move no-RDC binary embedding logic to CGSYCLRuntime Move embedSYCLNoRDCBinary out of CodeGenModule::Release() into a free function in CGSYCLRuntime.cpp to keep SYCL-specific logic out of common code. CodeGenModule::Release() retains a single call site. Co-Authored-By: Claude Sonnet 4.6 --- clang/lib/CodeGen/CGSYCLRuntime.cpp | 24 ++++++++++++++++++++++++ clang/lib/CodeGen/CGSYCLRuntime.h | 4 ++++ clang/lib/CodeGen/CodeGenModule.cpp | 27 +-------------------------- 3 files changed, 29 insertions(+), 26 deletions(-) diff --git a/clang/lib/CodeGen/CGSYCLRuntime.cpp b/clang/lib/CodeGen/CGSYCLRuntime.cpp index 6b1abe409c1d1..6ebbd650fbb6a 100644 --- a/clang/lib/CodeGen/CGSYCLRuntime.cpp +++ b/clang/lib/CodeGen/CGSYCLRuntime.cpp @@ -14,9 +14,12 @@ #include "CodeGenFunction.h" #include "clang/AST/Attr.h" #include "clang/AST/Decl.h" +#include "clang/Basic/DiagnosticFrontend.h" #include "clang/Basic/SourceLocation.h" #include "llvm/Analysis/OptimizationRemarkEmitter.h" +#include "llvm/Frontend/Offloading/OffloadWrapper.h" #include "llvm/IR/Instructions.h" +#include "llvm/Support/VirtualFileSystem.h" #include using namespace clang; @@ -131,6 +134,27 @@ bool CGSYCLRuntime::actOnGlobalVarEmit(CodeGenModule &CGM, const VarDecl &D, return true; } +void clang::CodeGen::embedSYCLNoRDCBinary(CodeGenModule &CGM) { + if (!CGM.getLangOpts().SYCLIsHost || + CGM.getCodeGenOpts().SYCLTargetBinaryFileName.empty()) + return; + auto BinaryOrErr = CGM.getFileSystem()->getBufferForFile( + CGM.getCodeGenOpts().SYCLTargetBinaryFileName, /*MaxSize=*/-1, + /*RequiresNullTerminator=*/false); + if (std::error_code EC = BinaryOrErr.getError()) { + CGM.getDiags().Report(diag::err_cannot_open_file) + << CGM.getCodeGenOpts().SYCLTargetBinaryFileName << EC.message(); + return; + } + llvm::ArrayRef Buffer((*BinaryOrErr)->getBufferStart(), + (*BinaryOrErr)->getBufferSize()); + if (llvm::Error E = llvm::offloading::wrapSYCLBinaries( + CGM.getModule(), Buffer, llvm::offloading::SYCLJITOptions{})) + CGM.getDiags().Report(diag::err_fe_linking_module) + << CGM.getCodeGenOpts().SYCLTargetBinaryFileName + << llvm::toString(std::move(E)); +} + bool Util::matchQualifiedTypeName(const CXXRecordDecl *RecTy, ArrayRef Scopes) { // The idea: check the declaration context chain starting from the type diff --git a/clang/lib/CodeGen/CGSYCLRuntime.h b/clang/lib/CodeGen/CGSYCLRuntime.h index 85a18e4bb590e..ed750153a6fe5 100644 --- a/clang/lib/CodeGen/CGSYCLRuntime.h +++ b/clang/lib/CodeGen/CGSYCLRuntime.h @@ -51,6 +51,10 @@ class CGSYCLRuntime { llvm::Value *Addr); }; +/// Embeds the per-TU finalized SYCL device binary into the host module when +/// -fsycl-include-target-binary is set (-fno-sycl-rdc compile step). +void embedSYCLNoRDCBinary(CodeGenModule &CGM); + } // namespace CodeGen } // namespace clang diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 2d9452d83dd80..c249b028ac44b 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -60,7 +60,6 @@ #include "llvm/ADT/StringSwitch.h" #include "llvm/Analysis/TargetLibraryInfo.h" #include "llvm/BinaryFormat/ELF.h" -#include "llvm/Frontend/Offloading/OffloadWrapper.h" #include "llvm/IR/AttributeMask.h" #include "llvm/IR/CallingConv.h" #include "llvm/IR/DataLayout.h" @@ -1258,31 +1257,7 @@ void CodeGenModule::Release() { if (llvm::Function *CudaCtorFunction = CUDARuntime->finalizeModule()) AddGlobalCtor(CudaCtorFunction); } - // SYCL -fno-sycl-rdc: the per-TU device image was finalized at compile time - // by clang-linker-wrapper --sycl-device-link. Read it here and call - // wrapSYCLBinaries to embed the binary and emit __sycl_register_lib / - // __sycl_unregister_lib into the host module, mirroring what CUDA does via - // CGCUDANV::makeModuleCtorFunction for -fno-gpu-rdc. - // SYCLJITOptions is left default (empty) because device code is already - // fully finalized -- there are no JIT compile/link options to encode. - if (LangOpts.SYCLIsHost && - !getCodeGenOpts().SYCLTargetBinaryFileName.empty()) { - auto VFS = getFileSystem(); - auto BinaryOrErr = VFS->getBufferForFile( - getCodeGenOpts().SYCLTargetBinaryFileName, -1, false); - if (std::error_code EC = BinaryOrErr.getError()) { - getDiags().Report(diag::err_cannot_open_file) - << getCodeGenOpts().SYCLTargetBinaryFileName << EC.message(); - } else { - llvm::ArrayRef Buffer((*BinaryOrErr)->getBufferStart(), - (*BinaryOrErr)->getBufferSize()); - if (llvm::Error E = llvm::offloading::wrapSYCLBinaries( - getModule(), Buffer, llvm::offloading::SYCLJITOptions{})) - getDiags().Report(diag::err_fe_linking_module) - << getCodeGenOpts().SYCLTargetBinaryFileName - << llvm::toString(std::move(E)); - } - } + embedSYCLNoRDCBinary(*this); if (OpenMPRuntime) { OpenMPRuntime->createOffloadEntriesAndInfoMetadata(); OpenMPRuntime->clear(); From 7a9c78205ed485771d62f0be8fee0a2105efb6e1 Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Fri, 31 Jul 2026 16:58:02 -0700 Subject: [PATCH 09/25] [SYCL][Driver] Use getFinalPhase to guard --sycl-device-link in linker wrapper --- clang/lib/Driver/ToolChains/Clang.cpp | 5 +---- 1 file changed, 1 insertion(+), 4 deletions(-) diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index 1916d1efd5cfb..08b8c930b5721 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -12176,14 +12176,11 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA, /*default=*/true); if (IsSYCLNoRDC) { CmdArgs.push_back("--no-sycl-rdc"); - // When invoked as a per-TU device finalizer at compile time - // (-fno-sycl-rdc -c), the linker wrapper must run in device-link-only - // mode: execute the full SYCL pipeline but skip the host linker. // isDeviceOffloading(OFK_SYCL) distinguishes the per-TU finalizer // (OFK_SYCL, from BuildOffloadingActions) from the final link-time // invocation (OFK_Host) -- both produce TY_Image so output type alone // is not a reliable discriminator. - if (!Args.hasArg(options::OPT_fsycl_link_EQ) && + if (C.getDriver().getFinalPhase(C.getArgs()) != phases::Link && JA.isDeviceOffloading(Action::OFK_SYCL)) CmdArgs.push_back("--sycl-device-link"); } From a45367cd94233f1c2309505853ef623809ac979a Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Fri, 31 Jul 2026 17:00:01 -0700 Subject: [PATCH 10/25] [SYCL][Driver] Trim verbose comment in SYCL no-RDC host offload block --- clang/lib/Driver/ToolChains/Clang.cpp | 8 +------- 1 file changed, 1 insertion(+), 7 deletions(-) diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index 08b8c930b5721..866c953a6e3a0 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -9154,13 +9154,7 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA, CmdArgs.push_back("-fcuda-include-gpubinary"); CmdArgs.push_back(HostOffloadingInputs.front().getFilename()); } else if (IsSYCL && !IsRDCMode) { - // SYCL no-RDC (-fno-sycl-rdc): device image was finalized per-TU by - // clang-linker-wrapper --sycl-device-link. Pass it to host CodeGen via - // -fsycl-include-target-binary so wrapSYCLBinaries embeds and registers - // it at compile time. -fembed-offload-object is intentionally NOT used: - // that would defer processing to link time, defeating the purpose. - // Note: this block is only reached during the host cc1 compile step; - // at final link time HostOffloadingInputs is empty for SYCL no-RDC. + // SYCL no-RDC: per-TU finalized device image, embed at compile time. assert(HostOffloadingInputs.size() == 1 && "One finalized image per TU"); CmdArgs.push_back("-fsycl-include-target-binary"); CmdArgs.push_back(HostOffloadingInputs.front().getFilename()); From a2bc665cd92c76fe88d63f181fe6b4a60758f9e8 Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Fri, 31 Jul 2026 17:04:49 -0700 Subject: [PATCH 11/25] [SYCL][Driver] Add -fno-sycl-rdc -flto tests to sycl-no-rdc-new-driver --- clang/test/Driver/sycl-no-rdc-new-driver.cpp | 11 +++++++++++ 1 file changed, 11 insertions(+) diff --git a/clang/test/Driver/sycl-no-rdc-new-driver.cpp b/clang/test/Driver/sycl-no-rdc-new-driver.cpp index 15a09f46f8a73..df63550f454a3 100644 --- a/clang/test/Driver/sycl-no-rdc-new-driver.cpp +++ b/clang/test/Driver/sycl-no-rdc-new-driver.cpp @@ -26,6 +26,17 @@ // RUN: | FileCheck -check-prefix=CHK-AOT-NO-RDC %s // CHK-AOT-NO-RDC: clang-linker-wrapper{{.*}} "--no-sycl-rdc" +// -fno-sycl-rdc -flto -c: per-TU device link still happens, --sycl-device-link present. +// RUN: %clang -### --offload-new-driver -Werror --target=x86_64-unknown-linux-gnu -fsycl -fno-sycl-rdc -flto -c %t.cpp 2>&1 \ +// RUN: | FileCheck -check-prefix=CHK-NO-RDC-LTO-C %s +// CHK-NO-RDC-LTO-C: clang-linker-wrapper{{.*}} "--no-sycl-rdc"{{.*}} "--sycl-device-link" + +// -fno-sycl-rdc -flto (link step): --sycl-device-link should NOT appear. +// RUN: %clang -### --offload-new-driver -Werror --target=x86_64-unknown-linux-gnu -fsycl -fno-sycl-rdc -flto %t.cpp 2>&1 \ +// RUN: | FileCheck -check-prefix=CHK-NO-RDC-LTO %s +// CHK-NO-RDC-LTO: clang-linker-wrapper{{.*}} "--no-sycl-rdc" +// CHK-NO-RDC-LTO-NOT: --sycl-device-link + // Verify pipeline with --offload-new-driver -fno-sycl-rdc. // RUN: touch %t1.cpp // RUN: touch %t2.cpp From 0e6ce2e0cabf58cb0d1a7e0fa56cb274be3d15b3 Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Fri, 31 Jul 2026 17:09:36 -0700 Subject: [PATCH 12/25] [SYCL][Driver] Trim verbose comment blocks in sycl-no-rdc-compile-step test --- clang/test/Driver/sycl-no-rdc-compile-step.cpp | 13 +++---------- 1 file changed, 3 insertions(+), 10 deletions(-) diff --git a/clang/test/Driver/sycl-no-rdc-compile-step.cpp b/clang/test/Driver/sycl-no-rdc-compile-step.cpp index ee9fc6e874c86..964382ffe8e2f 100644 --- a/clang/test/Driver/sycl-no-rdc-compile-step.cpp +++ b/clang/test/Driver/sycl-no-rdc-compile-step.cpp @@ -5,10 +5,7 @@ // RUN: touch %t.cpp -// --- CHECK 1: -fno-sycl-rdc -c --- -// Per-TU finalize: clang-linker-wrapper must be invoked with --sycl-device-link -// and --no-sycl-rdc. Host cc1 must receive -fsycl-include-target-binary. -// -fembed-offload-object must NOT appear (device code is not deferred). +// -fno-sycl-rdc -c: per-TU linker-wrapper with --sycl-device-link, host cc1 gets -fsycl-include-target-binary. // RUN: %clang -### --offload-new-driver -Werror --target=x86_64-unknown-linux-gnu \ // RUN: -fsycl -fno-sycl-rdc -c %t.cpp 2>&1 \ // RUN: | FileCheck -check-prefix=CHK-COMPILE %s @@ -16,9 +13,7 @@ // CHK-COMPILE: "-fsycl-include-target-binary" // CHK-COMPILE-NOT: -fembed-offload-object -// --- CHECK 2: default RDC -c (no flag) --- -// Default RDC path: -fembed-offload-object must appear (raw bitcode deferred to -// link time). -fsycl-include-target-binary and --no-sycl-rdc must NOT appear. +// Default RDC -c: -fembed-offload-object appears, no -fsycl-include-target-binary or --no-sycl-rdc. // RUN: %clang -### --offload-new-driver --target=x86_64-unknown-linux-gnu \ // RUN: -fsycl -c %t.cpp 2>&1 \ // RUN: | FileCheck -check-prefix=CHK-RDC %s @@ -26,9 +21,7 @@ // CHK-RDC-NOT: -fsycl-include-target-binary // CHK-RDC-NOT: --no-sycl-rdc -// --- CHECK 3: -fno-sycl-rdc at link step (regression for PR #22832) --- -// At link time --no-sycl-rdc is forwarded to clang-linker-wrapper. -// --sycl-device-link must NOT appear for the final link invocation. +// -fno-sycl-rdc at link step: --no-sycl-rdc forwarded, --sycl-device-link must NOT appear. // RUN: touch %t.o // RUN: %clang -### --offload-new-driver -Werror --target=x86_64-unknown-linux-gnu \ // RUN: -fsycl -fno-sycl-rdc %t.o 2>&1 \ From 187c79cea77c0381d34547b811c8c49bddcc6b5f Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Fri, 31 Jul 2026 17:11:21 -0700 Subject: [PATCH 13/25] [SYCL][Driver] Use implicit-check-not for -fembed-offload-object in compile-step test --- clang/test/Driver/sycl-no-rdc-compile-step.cpp | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/clang/test/Driver/sycl-no-rdc-compile-step.cpp b/clang/test/Driver/sycl-no-rdc-compile-step.cpp index 964382ffe8e2f..1c5701207eeb2 100644 --- a/clang/test/Driver/sycl-no-rdc-compile-step.cpp +++ b/clang/test/Driver/sycl-no-rdc-compile-step.cpp @@ -8,10 +8,9 @@ // -fno-sycl-rdc -c: per-TU linker-wrapper with --sycl-device-link, host cc1 gets -fsycl-include-target-binary. // RUN: %clang -### --offload-new-driver -Werror --target=x86_64-unknown-linux-gnu \ // RUN: -fsycl -fno-sycl-rdc -c %t.cpp 2>&1 \ -// RUN: | FileCheck -check-prefix=CHK-COMPILE %s +// RUN: | FileCheck -check-prefix=CHK-COMPILE --implicit-check-not=-fembed-offload-object %s // CHK-COMPILE: clang-linker-wrapper{{.*}} "--no-sycl-rdc"{{.*}} "--sycl-device-link" // CHK-COMPILE: "-fsycl-include-target-binary" -// CHK-COMPILE-NOT: -fembed-offload-object // Default RDC -c: -fembed-offload-object appears, no -fsycl-include-target-binary or --no-sycl-rdc. // RUN: %clang -### --offload-new-driver --target=x86_64-unknown-linux-gnu \ From fede3c5ead5a3becef8d4b3227278e1d17a60450 Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Fri, 31 Jul 2026 17:12:20 -0700 Subject: [PATCH 14/25] [SYCL][Driver] Anchor -fsycl-include-target-binary check to host cc1 line --- clang/test/Driver/sycl-no-rdc-compile-step.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/test/Driver/sycl-no-rdc-compile-step.cpp b/clang/test/Driver/sycl-no-rdc-compile-step.cpp index 1c5701207eeb2..80f24616fef77 100644 --- a/clang/test/Driver/sycl-no-rdc-compile-step.cpp +++ b/clang/test/Driver/sycl-no-rdc-compile-step.cpp @@ -10,7 +10,7 @@ // RUN: -fsycl -fno-sycl-rdc -c %t.cpp 2>&1 \ // RUN: | FileCheck -check-prefix=CHK-COMPILE --implicit-check-not=-fembed-offload-object %s // CHK-COMPILE: clang-linker-wrapper{{.*}} "--no-sycl-rdc"{{.*}} "--sycl-device-link" -// CHK-COMPILE: "-fsycl-include-target-binary" +// CHK-COMPILE: "-fsycl-is-host"{{.*}} "-fsycl-include-target-binary" // Default RDC -c: -fembed-offload-object appears, no -fsycl-include-target-binary or --no-sycl-rdc. // RUN: %clang -### --offload-new-driver --target=x86_64-unknown-linux-gnu \ From 7746b3728fed197061886acbe3ec12a020958b5f Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Fri, 31 Jul 2026 17:16:16 -0700 Subject: [PATCH 15/25] [SYCL][Doc] Rework NoRDC section to describe new offload model usage --- sycl/doc/design/OffloadDesign.md | 30 ++++++++++++++---------------- 1 file changed, 14 insertions(+), 16 deletions(-) diff --git a/sycl/doc/design/OffloadDesign.md b/sycl/doc/design/OffloadDesign.md index 29d12b373eb26..7f4b0ea0dc4f0 100644 --- a/sycl/doc/design/OffloadDesign.md +++ b/sycl/doc/design/OffloadDesign.md @@ -247,28 +247,26 @@ are needed to pass along this information. For the Old Offload Model support of NoRDC Mode see [NonRelocatableDeviceCode.md](NonRelocatableDeviceCode.md). -The default compiler behavior is -fsycl-rdc, which incorporates linking of device code. If -fno-sycl-rdc is specified, the compiler skips linking of device code and performs offload processing on every module individually. +The default compiler behavior is `-fsycl-rdc`, which incorporates linking of device code. +If `-fno-sycl-rdc` is specified, the compiler skips linking of device code and performs +offload processing on every module individually. -The usage scenario for the old offload model is: +With the new offload model, `-fno-sycl-rdc` can be specified at the link step: ``` -clang++ --no-offload-new-driver -fsycl -fsycl-targets=T1,T2 input1.cpp -fno-sycl-rdc -c -o object1.o # -fno-sycl-rdc is specified -clang++ --no-offload-new-driver -fsycl -fsycl-targets=T1,T2 input2.cpp -fno-sycl-rdc -c -o object2.o # -fno-sycl-rdc is specified -clang++ --no-offload-new-driver -fsycl -fsycl-targets=T1,T2 object1.o object2.o -o a.out # -fno-sycl-rdc is NOT specified +clang++ --offload-new-driver -fsycl -fsycl-targets=T1,T2 input1.cpp -c -o object1.o +clang++ --offload-new-driver -fsycl -fsycl-targets=T1,T2 input2.cpp -c -o object2.o +clang++ --offload-new-driver -fsycl -fsycl-targets=T1,T2 -fno-sycl-rdc object1.o object2.o -o a.out ``` -Currently, SYCL offload processing resides in clang-linker-wrapper. That leads to the following usage scenario, which is inverted compared to the old offload model: +It can also be specified at the compile step: ``` -clang++ --offload-new-driver -fsycl -fsycl-targets=T1,T2 input1.cpp -c -o object1.o # -fno-sycl-rdc is NOT specified -clang++ --offload-new-driver -fsycl -fsycl-targets=T1,T2 input2.cpp -c -o object2.o # -fno-sycl-rdc is NOT specified -clang++ --offload-new-driver -fsycl -fsycl-targets=T1,T2 -fno-sycl-rdc object1.o object2.o -o a.out # -fno-sycl-rdc is specified +clang++ --offload-new-driver -fsycl -fsycl-targets=T1,T2 -fno-sycl-rdc input1.cpp -c -o object1.o +clang++ --offload-new-driver -fsycl -fsycl-targets=T1,T2 -fno-sycl-rdc input2.cpp -c -o object2.o +clang++ --offload-new-driver -fsycl -fsycl-targets=T1,T2 object1.o object2.o -o a.out ``` - -Specifying `-fno-sycl-rdc` at the compile step -(i.e. `clang++ --offload-new-driver -fsycl -fno-sycl-rdc -c`) is also supported, matching -the old offload model's usage pattern. In this case `clang-linker-wrapper ---sycl-device-link --no-sycl-rdc` is invoked per translation unit at compile time to -finalize each TU's device code independently, embedding the result directly into the host -object. +In this case `clang-linker-wrapper --sycl-device-link --no-sycl-rdc` is invoked per +translation unit at compile time to finalize each TU's device code independently, +embedding the result directly into the host object. #### Format of the --device-compiler Option The `--device-compiler` option uses the format `--device-compiler=[:][=]` where: From 2be9203c04a7992e9ca8c24a05fe3b5bb057b818 Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Fri, 31 Jul 2026 17:21:11 -0700 Subject: [PATCH 16/25] [SYCL][Test] Pass -fno-sycl-rdc unconditionally in early_aot e2e test --- sycl/test-e2e/AOT/early_aot.cpp | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/sycl/test-e2e/AOT/early_aot.cpp b/sycl/test-e2e/AOT/early_aot.cpp index bc08b15d76b91..2537d92f494ce 100644 --- a/sycl/test-e2e/AOT/early_aot.cpp +++ b/sycl/test-e2e/AOT/early_aot.cpp @@ -2,14 +2,14 @@ // REQUIRES: ocloc, gpu, target-spir -// Note: The Old Offload Model requires -fno-sycl-rdc at the compilation step. -// The New Offload Model accepts -fno-sycl-rdc at both the compile and link -// steps; the RUN lines below pass it at the link step for compatibility. +// Note: For early AOT, -fno-sycl-rdc must be specified at the compilation step +// for the Old Offload Model. The New Offload Model supports -fno-sycl-rdc at +// both the compile and link steps. // 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 %if new-offload-model %{ -fno-sycl-rdc %} -DMAIN_CPP %s %t_add.o %t_sub.o -o %t.out +// RUN: %clangxx -fsycl -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen %gpu_aot_target_opts -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 -fno-sycl-rdc -c -DSUB_CPP %s -o %t_sub.o +// RUN: %clangxx -fsycl -fno-sycl-rdc -DMAIN_CPP %s %t_add.o %t_sub.o -o %t.out // RUN: %{run} %t.out From aa3d82729feadde945e2339d9bb31c7af00bc858 Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Mon, 3 Aug 2026 18:43:56 -0700 Subject: [PATCH 17/25] [SYCL][Test] Use %s instead of touch %t.cpp in no-rdc driver tests --- clang/test/Driver/sycl-no-rdc-compile-step.cpp | 6 ++---- clang/test/Driver/sycl-no-rdc-new-driver.cpp | 14 ++++++-------- 2 files changed, 8 insertions(+), 12 deletions(-) diff --git a/clang/test/Driver/sycl-no-rdc-compile-step.cpp b/clang/test/Driver/sycl-no-rdc-compile-step.cpp index 80f24616fef77..bd1b43caec520 100644 --- a/clang/test/Driver/sycl-no-rdc-compile-step.cpp +++ b/clang/test/Driver/sycl-no-rdc-compile-step.cpp @@ -3,18 +3,16 @@ /// action and routes the result to -fsycl-include-target-binary on the host /// cc1, instead of deferring device processing to link time. -// RUN: touch %t.cpp - // -fno-sycl-rdc -c: per-TU linker-wrapper with --sycl-device-link, host cc1 gets -fsycl-include-target-binary. // RUN: %clang -### --offload-new-driver -Werror --target=x86_64-unknown-linux-gnu \ -// RUN: -fsycl -fno-sycl-rdc -c %t.cpp 2>&1 \ +// RUN: -fsycl -fno-sycl-rdc -c %s 2>&1 \ // RUN: | FileCheck -check-prefix=CHK-COMPILE --implicit-check-not=-fembed-offload-object %s // CHK-COMPILE: clang-linker-wrapper{{.*}} "--no-sycl-rdc"{{.*}} "--sycl-device-link" // CHK-COMPILE: "-fsycl-is-host"{{.*}} "-fsycl-include-target-binary" // Default RDC -c: -fembed-offload-object appears, no -fsycl-include-target-binary or --no-sycl-rdc. // RUN: %clang -### --offload-new-driver --target=x86_64-unknown-linux-gnu \ -// RUN: -fsycl -c %t.cpp 2>&1 \ +// RUN: -fsycl -c %s 2>&1 \ // RUN: | FileCheck -check-prefix=CHK-RDC %s // CHK-RDC: -fembed-offload-object // CHK-RDC-NOT: -fsycl-include-target-binary diff --git a/clang/test/Driver/sycl-no-rdc-new-driver.cpp b/clang/test/Driver/sycl-no-rdc-new-driver.cpp index df63550f454a3..b3e1576050829 100644 --- a/clang/test/Driver/sycl-no-rdc-new-driver.cpp +++ b/clang/test/Driver/sycl-no-rdc-new-driver.cpp @@ -4,35 +4,33 @@ // -fno-sycl-rdc is passed. RDC is ON by default; --no-sycl-rdc signals // RDC is OFF. -// RUN: touch %t.cpp - // Default (no flag): RDC is ON by default for SYCL, so --no-sycl-rdc should NOT appear. -// RUN: %clang -### --offload-new-driver --target=x86_64-unknown-linux-gnu -fsycl %t.cpp 2>&1 \ +// RUN: %clang -### --offload-new-driver --target=x86_64-unknown-linux-gnu -fsycl %s 2>&1 \ // RUN: | FileCheck -check-prefix=CHK-DEFAULT %s // CHK-DEFAULT-NOT: --no-sycl-rdc // -fno-sycl-rdc: --no-sycl-rdc should appear. -// RUN: %clang -### --offload-new-driver -Werror --target=x86_64-unknown-linux-gnu -fsycl -fno-sycl-rdc %t.cpp 2>&1 \ +// RUN: %clang -### --offload-new-driver -Werror --target=x86_64-unknown-linux-gnu -fsycl -fno-sycl-rdc %s 2>&1 \ // RUN: | FileCheck -check-prefix=CHK-NO-RDC %s // CHK-NO-RDC: clang-linker-wrapper{{.*}} "--no-sycl-rdc" // AOT Intel GPU target, default RDC: --no-sycl-rdc should NOT appear. -// RUN: %clang -### --offload-new-driver --target=x86_64-unknown-linux-gnu -fsycl -fsycl-targets=intel_gpu_pvc %t.cpp 2>&1 \ +// RUN: %clang -### --offload-new-driver --target=x86_64-unknown-linux-gnu -fsycl -fsycl-targets=intel_gpu_pvc %s 2>&1 \ // RUN: | FileCheck -check-prefix=CHK-AOT-RDC %s // CHK-AOT-RDC-NOT: --no-sycl-rdc // AOT Intel GPU target + -fno-sycl-rdc: --no-sycl-rdc should appear. -// RUN: %clang -### --offload-new-driver -Werror --target=x86_64-unknown-linux-gnu -fsycl -fsycl-targets=intel_gpu_pvc -fno-sycl-rdc %t.cpp 2>&1 \ +// RUN: %clang -### --offload-new-driver -Werror --target=x86_64-unknown-linux-gnu -fsycl -fsycl-targets=intel_gpu_pvc -fno-sycl-rdc %s 2>&1 \ // RUN: | FileCheck -check-prefix=CHK-AOT-NO-RDC %s // CHK-AOT-NO-RDC: clang-linker-wrapper{{.*}} "--no-sycl-rdc" // -fno-sycl-rdc -flto -c: per-TU device link still happens, --sycl-device-link present. -// RUN: %clang -### --offload-new-driver -Werror --target=x86_64-unknown-linux-gnu -fsycl -fno-sycl-rdc -flto -c %t.cpp 2>&1 \ +// RUN: %clang -### --offload-new-driver -Werror --target=x86_64-unknown-linux-gnu -fsycl -fno-sycl-rdc -flto -c %s 2>&1 \ // RUN: | FileCheck -check-prefix=CHK-NO-RDC-LTO-C %s // CHK-NO-RDC-LTO-C: clang-linker-wrapper{{.*}} "--no-sycl-rdc"{{.*}} "--sycl-device-link" // -fno-sycl-rdc -flto (link step): --sycl-device-link should NOT appear. -// RUN: %clang -### --offload-new-driver -Werror --target=x86_64-unknown-linux-gnu -fsycl -fno-sycl-rdc -flto %t.cpp 2>&1 \ +// RUN: %clang -### --offload-new-driver -Werror --target=x86_64-unknown-linux-gnu -fsycl -fno-sycl-rdc -flto %s 2>&1 \ // RUN: | FileCheck -check-prefix=CHK-NO-RDC-LTO %s // CHK-NO-RDC-LTO: clang-linker-wrapper{{.*}} "--no-sycl-rdc" // CHK-NO-RDC-LTO-NOT: --sycl-device-link From 8946294d91806454049d8672c1f9329d56c07061 Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Tue, 4 Aug 2026 09:54:13 -0700 Subject: [PATCH 18/25] [SYCL][Driver] Trim verbose comment in BuildOffloadingActions SYCL no-RDC block --- clang/lib/Driver/Driver.cpp | 10 +++------- 1 file changed, 3 insertions(+), 7 deletions(-) diff --git a/clang/lib/Driver/Driver.cpp b/clang/lib/Driver/Driver.cpp index 0f7c64279af45..c9e50f9390f22 100644 --- a/clang/lib/Driver/Driver.cpp +++ b/clang/lib/Driver/Driver.cpp @@ -8342,13 +8342,9 @@ Driver::BuildOffloadingActions(Compilation &C, llvm::opt::DerivedArgList &Args, } else if (C.isOffloadingHostKind(Action::OFK_SYCL) && tools::SYCL::shouldDoPerObjectFileLinking(C) && !isa(HostAction)) { - // SYCL -fno-sycl-rdc at compile time (-c): finalize this TU's device code - // immediately via clang-linker-wrapper --sycl-device-link --no-sycl-rdc, - // producing a self-contained device image. The image is passed to the host - // cc1 via -fsycl-include-target-binary and embedded+registered at compile - // time. The final link step does no SYCL device work. This mirrors the - // CUDA/HIP -fno-gpu-rdc per-TU finalize model; downstream uses - // clang-linker-wrapper instead of clang-sycl-linker. + // SYCL -fno-sycl-rdc at compile time: finalize each TU's device code + // immediately via clang-linker-wrapper --sycl-device-link, embedding the + // result into the host object via -fsycl-include-target-binary. Action *PackagerAction = C.MakeAction(OffloadActions, types::TY_Image); ActionList AL{PackagerAction}; From 687055842e541ed18254ccaa7a3c9a2d4db81b1e Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Tue, 4 Aug 2026 10:00:48 -0700 Subject: [PATCH 19/25] [SYCL][Driver] Add --sysroot to sycl-no-rdc-compile-step driver tests --- clang/test/Driver/sycl-no-rdc-compile-step.cpp | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/clang/test/Driver/sycl-no-rdc-compile-step.cpp b/clang/test/Driver/sycl-no-rdc-compile-step.cpp index bd1b43caec520..ce2e5bc59bd83 100644 --- a/clang/test/Driver/sycl-no-rdc-compile-step.cpp +++ b/clang/test/Driver/sycl-no-rdc-compile-step.cpp @@ -5,6 +5,7 @@ // -fno-sycl-rdc -c: per-TU linker-wrapper with --sycl-device-link, host cc1 gets -fsycl-include-target-binary. // RUN: %clang -### --offload-new-driver -Werror --target=x86_64-unknown-linux-gnu \ +// RUN: --sysroot=%S/Inputs/SYCL --no-offloadlib -fno-sycl-instrument-device-code \ // RUN: -fsycl -fno-sycl-rdc -c %s 2>&1 \ // RUN: | FileCheck -check-prefix=CHK-COMPILE --implicit-check-not=-fembed-offload-object %s // CHK-COMPILE: clang-linker-wrapper{{.*}} "--no-sycl-rdc"{{.*}} "--sycl-device-link" @@ -12,6 +13,7 @@ // Default RDC -c: -fembed-offload-object appears, no -fsycl-include-target-binary or --no-sycl-rdc. // RUN: %clang -### --offload-new-driver --target=x86_64-unknown-linux-gnu \ +// RUN: --sysroot=%S/Inputs/SYCL --no-offloadlib -fno-sycl-instrument-device-code \ // RUN: -fsycl -c %s 2>&1 \ // RUN: | FileCheck -check-prefix=CHK-RDC %s // CHK-RDC: -fembed-offload-object @@ -21,7 +23,7 @@ // -fno-sycl-rdc at link step: --no-sycl-rdc forwarded, --sycl-device-link must NOT appear. // RUN: touch %t.o // RUN: %clang -### --offload-new-driver -Werror --target=x86_64-unknown-linux-gnu \ -// RUN: -fsycl -fno-sycl-rdc %t.o 2>&1 \ +// RUN: --sysroot=%S/Inputs/SYCL -fsycl -fno-sycl-rdc %t.o 2>&1 \ // RUN: | FileCheck -check-prefix=CHK-LINK %s // CHK-LINK: clang-linker-wrapper{{.*}} "--no-sycl-rdc" // CHK-LINK-NOT: clang-linker-wrapper{{.*}} "--sycl-device-link" From e48cdcfcdc233ff5b88b37c2070203329a454f0e Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Tue, 4 Aug 2026 11:00:22 -0700 Subject: [PATCH 20/25] [SYCL][E2E] Fix early_aot link step - remove -fno-sycl-rdc after compile-step finalization --- sycl/test-e2e/AOT/early_aot.cpp | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/sycl/test-e2e/AOT/early_aot.cpp b/sycl/test-e2e/AOT/early_aot.cpp index 2537d92f494ce..cb89347631a26 100644 --- a/sycl/test-e2e/AOT/early_aot.cpp +++ b/sycl/test-e2e/AOT/early_aot.cpp @@ -2,14 +2,14 @@ // REQUIRES: ocloc, gpu, target-spir -// Note: For early AOT, -fno-sycl-rdc must be specified at the compilation step -// for the Old Offload Model. The New Offload Model supports -fno-sycl-rdc at -// both the compile and link steps. +// Note: For early AOT, -fno-sycl-rdc is specified at the compilation step for +// both offload models. The New Offload Model finalizes device code per-TU at +// compile time, so -fno-sycl-rdc is not needed at the link step. // Build the early AOT device binaries // RUN: %clangxx -fsycl -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen %gpu_aot_target_opts -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 -fno-sycl-rdc -c -DSUB_CPP %s -o %t_sub.o -// RUN: %clangxx -fsycl -fno-sycl-rdc -DMAIN_CPP %s %t_add.o %t_sub.o -o %t.out +// RUN: %clangxx -fsycl -DMAIN_CPP %s %t_add.o %t_sub.o -o %t.out // RUN: %{run} %t.out From e8b78144b53edda6f4a6d4d3f1d741d55c36298a Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Tue, 4 Aug 2026 11:49:17 -0700 Subject: [PATCH 21/25] [SYCL][E2E] Add new offload model RUN lines to early_aot --- sycl/test-e2e/AOT/early_aot.cpp | 13 +++++++------ 1 file changed, 7 insertions(+), 6 deletions(-) diff --git a/sycl/test-e2e/AOT/early_aot.cpp b/sycl/test-e2e/AOT/early_aot.cpp index cb89347631a26..fa9fd97aab098 100644 --- a/sycl/test-e2e/AOT/early_aot.cpp +++ b/sycl/test-e2e/AOT/early_aot.cpp @@ -2,17 +2,18 @@ // REQUIRES: ocloc, gpu, target-spir -// Note: For early AOT, -fno-sycl-rdc is specified at the compilation step for -// both offload models. The New Offload Model finalizes device code per-TU at -// compile time, so -fno-sycl-rdc is not needed at the link step. - -// Build the early AOT device binaries +// Old offload model: -fno-sycl-rdc at compile step only. // RUN: %clangxx -fsycl -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen %gpu_aot_target_opts -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 -fno-sycl-rdc -c -DSUB_CPP %s -o %t_sub.o // RUN: %clangxx -fsycl -DMAIN_CPP %s %t_add.o %t_sub.o -o %t.out - // RUN: %{run} %t.out +// New offload model: -fno-sycl-rdc at compile step only. +// RUN: %clangxx --offload-new-driver -fsycl -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen %gpu_aot_target_opts -fno-sycl-rdc -c -DADD_CPP %s -o %t_add_new.o +// RUN: %clangxx --offload-new-driver -fsycl -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen %gpu_aot_target_opts -fno-sycl-rdc -c -DSUB_CPP %s -o %t_sub_new.o +// RUN: %clangxx --offload-new-driver -fsycl -DMAIN_CPP %s %t_add_new.o %t_sub_new.o -o %t_new.out +// RUN: %{run} %t_new.out + #ifdef MAIN_CPP // main.cpp From da4844cfe458ddb520e8cda1861fe888e0bb4c7c Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Tue, 4 Aug 2026 15:17:29 -0700 Subject: [PATCH 22/25] [SYCL][CodeGen] Use separate section for no-RDC per-TU device image embedding --- clang/lib/CodeGen/CGSYCLRuntime.cpp | 5 ++++- llvm/include/llvm/Frontend/Offloading/OffloadWrapper.h | 7 ++++++- llvm/lib/Frontend/Offloading/OffloadWrapper.cpp | 10 ++++++---- 3 files changed, 16 insertions(+), 6 deletions(-) diff --git a/clang/lib/CodeGen/CGSYCLRuntime.cpp b/clang/lib/CodeGen/CGSYCLRuntime.cpp index 6ebbd650fbb6a..3f119abd5a7ae 100644 --- a/clang/lib/CodeGen/CGSYCLRuntime.cpp +++ b/clang/lib/CodeGen/CGSYCLRuntime.cpp @@ -148,8 +148,11 @@ void clang::CodeGen::embedSYCLNoRDCBinary(CodeGenModule &CGM) { } llvm::ArrayRef Buffer((*BinaryOrErr)->getBufferStart(), (*BinaryOrErr)->getBufferSize()); + // Use a non-.llvm.offloading section - for sycl-no-rdc at compile time, the + // device image is already finalized and must not be re-processed at link time. if (llvm::Error E = llvm::offloading::wrapSYCLBinaries( - CGM.getModule(), Buffer, llvm::offloading::SYCLJITOptions{})) + CGM.getModule(), Buffer, llvm::offloading::SYCLJITOptions{}, + ".sycl_offloading.device_image")) CGM.getDiags().Report(diag::err_fe_linking_module) << CGM.getCodeGenOpts().SYCLTargetBinaryFileName << llvm::toString(std::move(E)); diff --git a/llvm/include/llvm/Frontend/Offloading/OffloadWrapper.h b/llvm/include/llvm/Frontend/Offloading/OffloadWrapper.h index dc99e29964afc..3ece1ea1458af 100644 --- a/llvm/include/llvm/Frontend/Offloading/OffloadWrapper.h +++ b/llvm/include/llvm/Frontend/Offloading/OffloadWrapper.h @@ -68,9 +68,14 @@ struct SYCLJITOptions { /// as global symbols and registers the images with the SYCL Runtime. /// \param Options Compiler and linker options to be encoded for the later /// use by a runtime for JIT compilation. Not used for AOT. +/// \param SectionName ELF section to store the binary in. Defaults to +/// ".llvm.offloading" so the linker-wrapper picks it up at link time. +/// Pass a different name to prevent link-time re-processing (e.g. for +/// per-TU finalized no-RDC images that are already fully compiled). LLVM_ABI llvm::Error wrapSYCLBinaries(llvm::Module &M, llvm::ArrayRef Buffer, - SYCLJITOptions Options = SYCLJITOptions()); + SYCLJITOptions Options = SYCLJITOptions(), + llvm::StringRef SectionName = ".llvm.offloading"); } // namespace offloading } // namespace llvm diff --git a/llvm/lib/Frontend/Offloading/OffloadWrapper.cpp b/llvm/lib/Frontend/Offloading/OffloadWrapper.cpp index ded603a1e00e3..a3acb1926885f 100644 --- a/llvm/lib/Frontend/Offloading/OffloadWrapper.cpp +++ b/llvm/lib/Frontend/Offloading/OffloadWrapper.cpp @@ -641,13 +641,14 @@ class SYCLWrapper { /// Embeds \p Buffer (a raw OffloadBinary) as a global constant and returns /// a pair of (Start, Size), where Start points to the beginning of the /// embedded data and Size is its length in bytes. - std::pair embedBinary(ArrayRef Buffer) { + std::pair embedBinary(ArrayRef Buffer, + StringRef SectionName) { Constant *Arr = ConstantDataArray::get(C, Buffer); GlobalVariable *BinaryGV = new GlobalVariable( M, Arr->getType(), /*isConstant=*/true, GlobalValue::InternalLinkage, Arr, ".sycl_offloading.binary"); BinaryGV->setUnnamedAddr(GlobalValue::UnnamedAddr::Global); - BinaryGV->setSection(".llvm.offloading"); + BinaryGV->setSection(SectionName); IntegerType *Int64Ty = Type::getInt64Ty(C); Constant *Zero = ConstantInt::get(Int64Ty, 0); @@ -749,9 +750,10 @@ Error offloading::wrapHIPBinary(Module &M, ArrayRef Image, } Error llvm::offloading::wrapSYCLBinaries(llvm::Module &M, ArrayRef Buffer, - SYCLJITOptions Options) { + SYCLJITOptions Options, + StringRef SectionName) { SYCLWrapper W(M, Options); - auto [Start, Size] = W.embedBinary(Buffer); + auto [Start, Size] = W.embedBinary(Buffer, SectionName); W.createRegisterFatbinFunction(Start, Size); W.createUnregisterFunction(Start, Size); return Error::success(); From 5eb41d95306f7bcfe55b9f0f9bf07e2883168fbe Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Tue, 4 Aug 2026 15:42:13 -0700 Subject: [PATCH 23/25] [NFC] Fix clang-format issues in no-RDC section name change --- clang/lib/CodeGen/CGSYCLRuntime.cpp | 3 ++- llvm/lib/Frontend/Offloading/OffloadWrapper.cpp | 2 +- 2 files changed, 3 insertions(+), 2 deletions(-) diff --git a/clang/lib/CodeGen/CGSYCLRuntime.cpp b/clang/lib/CodeGen/CGSYCLRuntime.cpp index 3f119abd5a7ae..77c3826770b32 100644 --- a/clang/lib/CodeGen/CGSYCLRuntime.cpp +++ b/clang/lib/CodeGen/CGSYCLRuntime.cpp @@ -149,7 +149,8 @@ void clang::CodeGen::embedSYCLNoRDCBinary(CodeGenModule &CGM) { llvm::ArrayRef Buffer((*BinaryOrErr)->getBufferStart(), (*BinaryOrErr)->getBufferSize()); // Use a non-.llvm.offloading section - for sycl-no-rdc at compile time, the - // device image is already finalized and must not be re-processed at link time. + // device image is already finalized and must not be re-processed at link + // time. if (llvm::Error E = llvm::offloading::wrapSYCLBinaries( CGM.getModule(), Buffer, llvm::offloading::SYCLJITOptions{}, ".sycl_offloading.device_image")) diff --git a/llvm/lib/Frontend/Offloading/OffloadWrapper.cpp b/llvm/lib/Frontend/Offloading/OffloadWrapper.cpp index a3acb1926885f..b5a6891d0a013 100644 --- a/llvm/lib/Frontend/Offloading/OffloadWrapper.cpp +++ b/llvm/lib/Frontend/Offloading/OffloadWrapper.cpp @@ -642,7 +642,7 @@ class SYCLWrapper { /// a pair of (Start, Size), where Start points to the beginning of the /// embedded data and Size is its length in bytes. std::pair embedBinary(ArrayRef Buffer, - StringRef SectionName) { + StringRef SectionName) { Constant *Arr = ConstantDataArray::get(C, Buffer); GlobalVariable *BinaryGV = new GlobalVariable( M, Arr->getType(), /*isConstant=*/true, GlobalValue::InternalLinkage, From b8b9fab63678f8f7b22d1c468475cbc6782fb287 Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Tue, 4 Aug 2026 17:45:59 -0700 Subject: [PATCH 24/25] [SYCL][CodeGen] Fix no-RDC per-TU device image embedding via bitcode link Replace raw-bytes wrapSYCLBinaries path with parseBitcodeFile+linkModules: clang-linker-wrapper --sycl-device-link now outputs the wrapper .bc directly (skipping runCompile), and embedSYCLNoRDCBinary links it into the host module. This ensures the correct __sycl.tgt_bin_desc descriptor and one-arg __sycl_register_lib constructor reach the SYCL runtime. Also reverts the SectionName parameter added to wrapSYCLBinaries as it is no longer needed. --- clang/lib/CodeGen/CGSYCLRuntime.cpp | 34 ++++++++++++++----- .../ClangLinkerWrapper.cpp | 12 +++++++ .../llvm/Frontend/Offloading/OffloadWrapper.h | 7 +--- .../Frontend/Offloading/OffloadWrapper.cpp | 10 +++--- 4 files changed, 42 insertions(+), 21 deletions(-) diff --git a/clang/lib/CodeGen/CGSYCLRuntime.cpp b/clang/lib/CodeGen/CGSYCLRuntime.cpp index 77c3826770b32..ef0dc36b189bd 100644 --- a/clang/lib/CodeGen/CGSYCLRuntime.cpp +++ b/clang/lib/CodeGen/CGSYCLRuntime.cpp @@ -17,8 +17,10 @@ #include "clang/Basic/DiagnosticFrontend.h" #include "clang/Basic/SourceLocation.h" #include "llvm/Analysis/OptimizationRemarkEmitter.h" +#include "llvm/Bitcode/BitcodeReader.h" #include "llvm/Frontend/Offloading/OffloadWrapper.h" #include "llvm/IR/Instructions.h" +#include "llvm/Linker/Linker.h" #include "llvm/Support/VirtualFileSystem.h" #include @@ -146,17 +148,31 @@ void clang::CodeGen::embedSYCLNoRDCBinary(CodeGenModule &CGM) { << CGM.getCodeGenOpts().SYCLTargetBinaryFileName << EC.message(); return; } - llvm::ArrayRef Buffer((*BinaryOrErr)->getBufferStart(), - (*BinaryOrErr)->getBufferSize()); - // Use a non-.llvm.offloading section - for sycl-no-rdc at compile time, the - // device image is already finalized and must not be re-processed at link - // time. - if (llvm::Error E = llvm::offloading::wrapSYCLBinaries( - CGM.getModule(), Buffer, llvm::offloading::SYCLJITOptions{}, - ".sycl_offloading.device_image")) + // The input is a wrapper bitcode module produced by clang-linker-wrapper + // --sycl-device-link. It already contains the correct __sycl.tgt_bin_desc + // structure and __sycl_register_lib constructor. Link it directly into the + // host module so the SYCL runtime finds the device image at program startup. + llvm::LLVMContext &Ctx = CGM.getModule().getContext(); + llvm::Expected> DevModOrErr = + llvm::parseBitcodeFile((*BinaryOrErr)->getMemBufferRef(), Ctx); + if (!DevModOrErr) { CGM.getDiags().Report(diag::err_fe_linking_module) << CGM.getCodeGenOpts().SYCLTargetBinaryFileName - << llvm::toString(std::move(E)); + << llvm::toString(DevModOrErr.takeError()); + return; + } + // Propagate the host data layout so the linker doesn't warn about mismatched + // layouts (the wrapper .bc is created without a data layout set). + (*DevModOrErr)->setDataLayout(CGM.getModule().getDataLayout()); + // Install a diagnostic handler that routes DK_Linker diagnostics through + // clang's diagnostic engine. The default handler asserts CurLinkModule is + // set, which is only true inside BackendConsumer's LinkModules loop. + Ctx.setDiagnosticHandler(std::make_unique()); + if (llvm::Linker::linkModules(CGM.getModule(), std::move(*DevModOrErr))) + CGM.getDiags().Report(diag::err_fe_linking_module) + << CGM.getCodeGenOpts().SYCLTargetBinaryFileName + << "linking wrapper bitcode into host module failed"; + Ctx.setDiagnosticHandler(nullptr); } bool Util::matchQualifiedTypeName(const CXXRecordDecl *RecTy, diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index 59aa81ed94a74..8d2d79e42a987 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -2790,6 +2790,18 @@ linkAndWrapDeviceFiles(ArrayRef> LinkerInputFiles, MD.SplitModules = std::move(Modules); std::scoped_lock Guard(SYCLBINModulesMtx); SYCLBINModules.emplace_back(std::move(MD)); + } else if (Args.hasArg(OPT_sycl_device_link)) { + // For per-TU -fno-sycl-rdc compile-step finalization, output the + // wrapper bitcode directly. The host cc1 will link this .bc into the + // host module via -fsycl-include-target-binary, so the SYCL + // registration IR (constructors calling __sycl_register_lib) is + // emitted correctly without compiling to an intermediate host object. + Expected OutputFile = + sycl::wrapSYCLBinariesFromFile(Modules, LinkerArgs, + /*IsEmbeddedIR=*/false); + if (!OutputFile) + return OutputFile.takeError(); + AppendImageToWrapperOutput(*OutputFile); } else { // TODO(NOM7): Remove this call and use community flow for bundle/wrap Expected OutputFile = diff --git a/llvm/include/llvm/Frontend/Offloading/OffloadWrapper.h b/llvm/include/llvm/Frontend/Offloading/OffloadWrapper.h index 3ece1ea1458af..dc99e29964afc 100644 --- a/llvm/include/llvm/Frontend/Offloading/OffloadWrapper.h +++ b/llvm/include/llvm/Frontend/Offloading/OffloadWrapper.h @@ -68,14 +68,9 @@ struct SYCLJITOptions { /// as global symbols and registers the images with the SYCL Runtime. /// \param Options Compiler and linker options to be encoded for the later /// use by a runtime for JIT compilation. Not used for AOT. -/// \param SectionName ELF section to store the binary in. Defaults to -/// ".llvm.offloading" so the linker-wrapper picks it up at link time. -/// Pass a different name to prevent link-time re-processing (e.g. for -/// per-TU finalized no-RDC images that are already fully compiled). LLVM_ABI llvm::Error wrapSYCLBinaries(llvm::Module &M, llvm::ArrayRef Buffer, - SYCLJITOptions Options = SYCLJITOptions(), - llvm::StringRef SectionName = ".llvm.offloading"); + SYCLJITOptions Options = SYCLJITOptions()); } // namespace offloading } // namespace llvm diff --git a/llvm/lib/Frontend/Offloading/OffloadWrapper.cpp b/llvm/lib/Frontend/Offloading/OffloadWrapper.cpp index b5a6891d0a013..ded603a1e00e3 100644 --- a/llvm/lib/Frontend/Offloading/OffloadWrapper.cpp +++ b/llvm/lib/Frontend/Offloading/OffloadWrapper.cpp @@ -641,14 +641,13 @@ class SYCLWrapper { /// Embeds \p Buffer (a raw OffloadBinary) as a global constant and returns /// a pair of (Start, Size), where Start points to the beginning of the /// embedded data and Size is its length in bytes. - std::pair embedBinary(ArrayRef Buffer, - StringRef SectionName) { + std::pair embedBinary(ArrayRef Buffer) { Constant *Arr = ConstantDataArray::get(C, Buffer); GlobalVariable *BinaryGV = new GlobalVariable( M, Arr->getType(), /*isConstant=*/true, GlobalValue::InternalLinkage, Arr, ".sycl_offloading.binary"); BinaryGV->setUnnamedAddr(GlobalValue::UnnamedAddr::Global); - BinaryGV->setSection(SectionName); + BinaryGV->setSection(".llvm.offloading"); IntegerType *Int64Ty = Type::getInt64Ty(C); Constant *Zero = ConstantInt::get(Int64Ty, 0); @@ -750,10 +749,9 @@ Error offloading::wrapHIPBinary(Module &M, ArrayRef Image, } Error llvm::offloading::wrapSYCLBinaries(llvm::Module &M, ArrayRef Buffer, - SYCLJITOptions Options, - StringRef SectionName) { + SYCLJITOptions Options) { SYCLWrapper W(M, Options); - auto [Start, Size] = W.embedBinary(Buffer, SectionName); + auto [Start, Size] = W.embedBinary(Buffer); W.createRegisterFatbinFunction(Start, Size); W.createUnregisterFunction(Start, Size); return Error::success(); From 2fcf277255b4c22b9aabc8701b386f56a0c399d2 Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Tue, 4 Aug 2026 19:14:27 -0700 Subject: [PATCH 25/25] [SYCL][CodeGen] Guard no-RDC .bc output path with --no-sycl-rdc check --sycl-device-link is also passed for -fsycl-link (device finalization at link step), which must still produce an object file via runWrapperAndCompile. The wrapper .bc output path should only activate when both --sycl-device-link and --no-sycl-rdc are present (per-TU -fno-sycl-rdc compile-step case). --- 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 8d2d79e42a987..c544d9d78f1ac 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -2790,7 +2790,8 @@ linkAndWrapDeviceFiles(ArrayRef> LinkerInputFiles, MD.SplitModules = std::move(Modules); std::scoped_lock Guard(SYCLBINModulesMtx); SYCLBINModules.emplace_back(std::move(MD)); - } else if (Args.hasArg(OPT_sycl_device_link)) { + } else if (Args.hasArg(OPT_sycl_device_link) && + Args.hasArg(OPT_no_sycl_rdc)) { // For per-TU -fno-sycl-rdc compile-step finalization, output the // wrapper bitcode directly. The host cc1 will link this .bc into the // host module via -fsycl-include-target-binary, so the SYCL