[clang] [llvm] [clang][SYCL] Add -fno-sycl-rdc to finalize device code per translation unit (PR #218089)
Yury Plyakhin via cfe-commits
cfe-commits at lists.llvm.org
Fri Aug 21 19:16:57 PDT 2026
https://github.com/YuriPlyakhin created https://github.com/llvm/llvm-project/pull/218089
Activate non-relocatable SYCL device code compilation for a single device architecture.
Driver:
- Add -fsycl-rdc / -fno-sycl-rdc as aliases of -fgpu-rdc / -fno-gpu-rdc, for compatibility with the spelling SYCL users already have.
- In BuildOffloadingActions, when SYCL non-RDC is selected, package the translation unit's device images and hand them to clang-linker-wrapper as a LinkerWrapperJobAction producing the new TY_SYCL_FATBIN type. The wrapper finalizes the images and emits the fat binary directly.
- Include that fat binary into the host compilation via -foffload-include-binary, so it is embedded and registered at compile time and the final clang-linker-wrapper performs no SYCL device link.
- Generalize the --emit-fatbin-only handling in tools::LinkerWrapper to the new type.
- Give SYCLToolChain a device linker.
- -fno-sycl-rdc with more than one device triple is now diagnosed in the driver
Because the device link now happens while compiling, -fsycl-device-image-split= reaches clang-sycl-linker in non-RDC mode.
More than one device architecture is not supported yet.
Add a driver test asserting the section-level contract that keeps the two modes apart: the non-RDC host object carries its finalized image in ".sycl_fatbin" and no ".llvm.offloading", while the RDC host object is the reverse.
Add a driver test that builds a two translation unit program, checks that each object holds an offload binary wrapping a finalized SPIR-V module that offers that translation unit's kernels, then links and runs it against its own __sycl_register_lib and __sycl_unregister_lib (-nolibsycl, so no SYCL runtime or device is needed) to confirm that both images are registered before the other static initializers of the program run and unregistered after main returns. An RDC build of the same sources registers a single image instead, which is what tells the two modes apart at run time.
Add end-to-end test. libsycl/test/basic/no_rdc_multi_tu.cpp runs the equivalent program on a real device in both modes and checks the results computed by the kernels of both translation units.
Both tests also build with -fsycl-device-image-split=kernel, which is how a single fat binary holding several device images is exercised.
SYCL defaults to relocatable device code (RDC); RDC compilation is unchanged.
co-authored by claude
>From e6c93eddc11f0c209c2720df1e68d722ca9a17f7 Mon Sep 17 00:00:00 2001
From: "Plyakhin, Yury" <yury.plyakhin at intel.com>
Date: Thu, 20 Aug 2026 17:39:23 -0700
Subject: [PATCH] [clang][SYCL] Add -fno-sycl-rdc to finalize device code per
translation unit
Activate non-relocatable SYCL device code compilation for a single device
architecture, following the shape HIP non-RDC uses.
Driver:
- Add -fsycl-rdc / -fno-sycl-rdc as aliases of -fgpu-rdc / -fno-gpu-rdc, for
compatibility with the spelling SYCL users already have. -fgpu-rdc stays
the option the driver keys on, so nothing below has to know the SYCL
spellings exist and mixed spellings compose as usual.
- In BuildOffloadingActions, when SYCL non-RDC is selected, package the
translation unit's device images and hand them to clang-linker-wrapper as
a LinkerWrapperJobAction producing the new TY_SYCL_FATBIN type, instead of
packaging unlinked device code for a cross-TU device link. The wrapper
finalizes the images and emits the fat binary directly, which is what HIP
non-RDC does with TY_HIP_FATBIN.
- Include that fat binary into the host compilation via
-foffload-include-binary (rather than -fembed-offload-object), so it is
embedded and registered at compile time and the final clang-linker-wrapper
performs no SYCL device link. The host routing reuses IsRDCMode (whose
default is now SYCL-aware) instead of a separate rdc query.
- Generalize the --emit-fatbin-only handling in tools::LinkerWrapper to the
new type, and with it the verbose-output suppression that keeps a direct
fat binary link from confusing CMake's implicit linker argument parsing.
- Give SYCLToolChain a device linker. Under this shape the job the driver
builds for the device link is discarded by tools::LinkerWrapper apart from
its executable, which it reports as its own --linker-path=, so all
SPIRV::Linker has to do for TY_SYCL_FATBIN is name clang-sycl-linker.
Because the device link now happens while compiling, -fsycl-device-image-split=
reaches clang-sycl-linker in non-RDC mode, including with -c, where it used to
be reported as unused. Splitting the device code of a translation unit into
several images therefore works: clang-sycl-linker packs the images it produces
for one architecture into a single offload binary itself.
More than one device architecture is not supported yet. Each of them is
finalized by an invocation of its own, so each arrives as an offload binary of
its own, and clang-linker-wrapper reports "Expected a single bundled image for
direct fat binary output" rather than merging them into the one fat binary the
translation unit is to carry. Teaching it to merge them, and the runtime to
choose between the images of several architectures, is left to a follow-up.
More than one device triple runs into the same limit, and additionally has no
single device toolchain to bind the finalize job to, so it is diagnosed in the
driver instead of being allowed to reach that point.
Add a driver test asserting the section-level contract that keeps the two
modes apart: the non-RDC host object carries its finalized image in
".sycl_fatbin" and no ".llvm.offloading", while the RDC host object is the
reverse. clang-linker-wrapper scans ".llvm.offloading*" for device code
still needing a link, so a finalized image landing there is fed back to
clang-sycl-linker and the link fails; this test catches that without
requiring a GPU. Both object formats are checked, because the choice is made
by name and does not depend on the format.
Add end-to-end tests. clang/test/Driver/sycl-nordc-registration.cpp builds a
two translation unit program, checks that each object holds an offload binary
wrapping a finalized SPIR-V module that offers only that translation unit's
kernels, then links and runs it against its own __sycl_register_lib and
__sycl_unregister_lib (-nolibsycl, so no SYCL runtime or device is needed) to
confirm that both images are registered before the other static initializers
of the program run and unregistered after main returns. An RDC build of the
same sources registers a single image instead, which is what tells the two
modes apart at run time. libsycl/test/basic/no_rdc_multi_tu.cpp runs the
equivalent program on a real device in both modes and checks the results
computed by the kernels of both translation units. Both also build with
-fsycl-device-image-split=kernel, which is how a single fat binary holding
several device images is exercised without an ahead-of-time compiler for a
real device.
SYCL defaults to relocatable device code (RDC); RDC compilation is
unchanged.
Co-Authored-By: Claude <noreply at anthropic.com>
---
.../clang/Basic/DiagnosticDriverKinds.td | 3 +
clang/include/clang/Driver/Types.def | 1 +
clang/include/clang/Options/Options.td | 9 ++
clang/lib/Driver/Driver.cpp | 35 ++++
clang/lib/Driver/ToolChains/Clang.cpp | 26 +--
clang/lib/Driver/ToolChains/SPIRV.cpp | 13 ++
clang/lib/Driver/ToolChains/SYCL.cpp | 6 +
clang/lib/Driver/ToolChains/SYCL.h | 3 +
clang/lib/Driver/Types.cpp | 5 +-
.../sycl-nordc-registration-second-tu.cpp | 18 +++
.../test/Driver/sycl-nordc-fatbin-section.cpp | 27 ++++
clang/test/Driver/sycl-nordc-registration.cpp | 150 ++++++++++++++++++
clang/test/Driver/sycl-offload-jit.cpp | 89 +++++++++++
.../basic/Inputs/no_rdc_multi_tu_second.cpp | 8 +
.../basic/Inputs/no_rdc_multi_tu_second.hpp | 6 +
libsycl/test/basic/no_rdc_multi_tu.cpp | 68 ++++++++
16 files changed, 454 insertions(+), 13 deletions(-)
create mode 100644 clang/test/Driver/Inputs/sycl-nordc-registration-second-tu.cpp
create mode 100644 clang/test/Driver/sycl-nordc-fatbin-section.cpp
create mode 100644 clang/test/Driver/sycl-nordc-registration.cpp
create mode 100644 libsycl/test/basic/Inputs/no_rdc_multi_tu_second.cpp
create mode 100644 libsycl/test/basic/Inputs/no_rdc_multi_tu_second.hpp
create mode 100644 libsycl/test/basic/no_rdc_multi_tu.cpp
diff --git a/clang/include/clang/Basic/DiagnosticDriverKinds.td b/clang/include/clang/Basic/DiagnosticDriverKinds.td
index 929421f915a4e..7799fa7c75715 100644
--- a/clang/include/clang/Basic/DiagnosticDriverKinds.td
+++ b/clang/include/clang/Basic/DiagnosticDriverKinds.td
@@ -95,6 +95,9 @@ def err_drv_hipspv_no_hip_path : Error<
"'--hip-path' must be specified when offloading to SPIR-V unless '-nogpuinc' "
"is given">;
+def err_drv_sycl_no_rdc_multiple_targets : Error<
+ "'%0' is not supported with more than one SYCL offloading target">;
+
// SYCL Windows-specific diagnostics
def err_drv_sycl_requires_dynamic_crt : Error<
"SYCL requires dynamic C++ runtime (/MD or /MDd); "
diff --git a/clang/include/clang/Driver/Types.def b/clang/include/clang/Driver/Types.def
index 74563ff835179..38ccdaaebe717 100644
--- a/clang/include/clang/Driver/Types.def
+++ b/clang/include/clang/Driver/Types.def
@@ -122,6 +122,7 @@ TYPE("dSYM", dSYM, INVALID, "dSYM", phases
TYPE("dependencies", Dependencies, INVALID, "d", phases::Compile, phases::Backend, phases::Assemble, phases::Link)
TYPE("cuda-fatbin", CUDA_FATBIN, INVALID, "fatbin", phases::Compile, phases::Backend, phases::Assemble, phases::Link)
TYPE("hip-fatbin", HIP_FATBIN, INVALID, "hipfb", phases::Compile, phases::Backend, phases::Assemble, phases::Link)
+TYPE("sycl-fatbin", SYCL_FATBIN, INVALID, "syclfb", phases::Compile, phases::Backend, phases::Assemble, phases::Link)
TYPE("api-information", API_INFO, INVALID, "json", phases::Precompile)
TYPE("dx-container", DX_CONTAINER, INVALID, "dxo", phases::Compile, phases::Backend)
TYPE("none", Nothing, INVALID, nullptr, phases::Compile, phases::Backend, phases::Assemble, phases::Link)
diff --git a/clang/include/clang/Options/Options.td b/clang/include/clang/Options/Options.td
index a5be26fdf0dce..ed176e541dfed 100644
--- a/clang/include/clang/Options/Options.td
+++ b/clang/include/clang/Options/Options.td
@@ -7735,6 +7735,15 @@ def fsycl_host_only : Flag<["-"], "fsycl-host-only">,
def sycl_link : Flag<["--"], "sycl-link">, Flags<[HelpHidden]>,
HelpText<"Perform link through clang-sycl-linker via the target "
"offloading toolchain.">;
+def : Flag<["-"], "fsycl-rdc">, Alias<fgpu_rdc>;
+def : Flag<["-"], "fno-sycl-rdc">, Alias<fno_gpu_rdc>,
+ HelpText<"Generate non-relocatable device code during SYCL offload target "
+ "compilation. Use of this option in combination with '-c' will "
+ "produce final device binaries within the generated fat object. "
+ "When using this option, each kernel must be self-contained within "
+ "its translation unit (source file); a device function defined in "
+ "another translation unit cannot be called, even if declared "
+ "SYCL_EXTERNAL.">;
def fsycl_device_image_split_EQ : Joined<["-"], "fsycl-device-image-split=">,
HelpText<"Select the granularity at which SYCL device code is grouped into "
"device images: kernel (one device image per kernel) | "
diff --git a/clang/lib/Driver/Driver.cpp b/clang/lib/Driver/Driver.cpp
index 2b79cb9d12c2a..83b3a55ea494d 100644
--- a/clang/lib/Driver/Driver.cpp
+++ b/clang/lib/Driver/Driver.cpp
@@ -1184,6 +1184,20 @@ void Driver::CreateOffloadingDeviceToolChains(Compilation &C,
C.addOffloadDeviceToolChain(&TC, Kind);
}
}
+
+ // Non-RDC SYCL device code is finalized by a clang-linker-wrapper job
+ // bound to one device toolchain, so only one SYCL target can be requested
+ // now.
+ if (Kinds.contains(Action::OFK_SYCL)) {
+ const Arg *RDCArg = C.getInputArgs().getLastArg(options::OPT_fgpu_rdc,
+ options::OPT_fno_gpu_rdc);
+ if (RDCArg && RDCArg->getOption().matches(options::OPT_fno_gpu_rdc)) {
+ auto TCRange = C.getOffloadToolChains<Action::OFK_SYCL>();
+ if (std::distance(TCRange.first, TCRange.second) > 1)
+ Diag(clang::diag::err_drv_sycl_no_rdc_multiple_targets)
+ << RDCArg->getAsString(C.getInputArgs());
+ }
+ }
}
bool Driver::loadZOSCustomizationFile(llvm::cl::ExpansionContext &ExpCtx) {
@@ -5106,6 +5120,12 @@ Driver::BuildOffloadingActions(Compilation &C, llvm::opt::DerivedArgList &Args,
C.isOffloadingHostKind(Action::OFK_HIP) &&
!Args.hasFlag(options::OPT_fgpu_rdc, options::OPT_fno_gpu_rdc, false);
+ // SYCL defaults to relocatable device code.
+ bool SYCLNoRDC =
+ C.isOffloadingHostKind(Action::OFK_SYCL) &&
+ !Args.hasFlag(options::OPT_fgpu_rdc, options::OPT_fno_gpu_rdc,
+ /*Default=*/true);
+
bool HIPRelocatableObj =
C.isOffloadingHostKind(Action::OFK_HIP) &&
Args.hasFlag(options::OPT_fhip_emit_relocatable,
@@ -5318,6 +5338,21 @@ Driver::BuildOffloadingActions(Compilation &C, llvm::opt::DerivedArgList &Args,
DDep.add(*PackagerAction,
*C.getOffloadToolChains<Action::OFK_HIP>().first->second,
/*BA=*/{}, Action::OFK_HIP);
+ } else if (SYCLNoRDC) {
+ // Package all the offloading actions into a single output that can be
+ // embedded in the host and linked.
+ Action *PackagerAction =
+ C.MakeAction<OffloadPackagerJobAction>(OffloadActions, types::TY_Image);
+
+ // Wrap the device binary with linker wrapper before bundling with host
+ // code. Do not bind a specific arch here, as the packaged binary may
+ // contain entries for multiple architectures.
+ ActionList AL{PackagerAction};
+ PackagerAction =
+ C.MakeAction<LinkerWrapperJobAction>(AL, types::TY_SYCL_FATBIN);
+ DDep.add(*PackagerAction,
+ *C.getOffloadToolChains<Action::OFK_SYCL>().first->second,
+ /*BA=*/{}, Action::OFK_SYCL);
} else {
// Package all the offloading actions into a single output that can be
// embedded in the host and linked.
diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp
index dabc8c8d964d6..e97a6fed06caa 100644
--- a/clang/lib/Driver/ToolChains/Clang.cpp
+++ b/clang/lib/Driver/ToolChains/Clang.cpp
@@ -8342,9 +8342,10 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
CmdArgs.push_back("-foffload-include-binary");
CmdArgs.push_back(CudaDeviceInput->getFilename());
} else if (!HostOffloadingInputs.empty()) {
- if ((IsCuda || IsHIP) &&
- (!IsRDCMode || Args.hasArg(options::OPT_cuda_emit_nvcc_abi)) &&
- !UsesLLVMOffloading) {
+ if (((IsCuda || IsHIP) &&
+ (!IsRDCMode || Args.hasArg(options::OPT_cuda_emit_nvcc_abi)) &&
+ !UsesLLVMOffloading) ||
+ (IsSYCL && !IsRDCMode)) {
assert(HostOffloadingInputs.size() == 1 && "Only one input expected");
CmdArgs.push_back("-foffload-include-binary");
CmdArgs.push_back(HostOffloadingInputs.front().getFilename());
@@ -9866,11 +9867,12 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA,
OPT_fno_slp_vectorize,
OPT_hipstdpar};
const llvm::DenseSet<unsigned> LinkerOptions{OPT_mllvm, OPT_Zlinker_input};
- // Suppress verbose output for HIP non-RDC fat binaries because it confuses
+ // Suppress verbose output for non-RDC fat binaries because it confuses
// CMake implicit linker argument parsing.
- bool SuppressHIPNoRDCVerbose =
- JA.getType() == types::TY_HIP_FATBIN &&
- !Args.hasFlag(options::OPT_fgpu_rdc, options::OPT_fno_gpu_rdc, false);
+ bool SuppressNoRDCVerbose =
+ JA.getType() == types::TY_SYCL_FATBIN ||
+ (JA.getType() == types::TY_HIP_FATBIN &&
+ !Args.hasFlag(options::OPT_fgpu_rdc, options::OPT_fno_gpu_rdc, false));
auto ToolChainHasRT = [&](const ToolChain &TC, StringRef Name) {
return TC.getVFS().exists(
TC.getCompilerRT(Args, Name, ToolChain::FT_Static));
@@ -9892,7 +9894,7 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA,
};
auto ShouldForward = [&](const llvm::DenseSet<unsigned> &Set, Arg *A,
const ToolChain &TC) {
- if (A->getOption().matches(OPT_v) && SuppressHIPNoRDCVerbose)
+ if (A->getOption().matches(OPT_v) && SuppressNoRDCVerbose)
return false;
return (Set.contains(A->getOption().getID()) ||
(A->getOption().getGroup().isValid() &&
@@ -10010,7 +10012,7 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA,
CmdArgs.push_back(Args.MakeArgString("--host-triple=" +
getToolChain().getTripleString()));
- if (Args.hasArg(options::OPT_v) && !SuppressHIPNoRDCVerbose)
+ if (Args.hasArg(options::OPT_v) && !SuppressNoRDCVerbose)
CmdArgs.push_back("--wrapper-verbose");
if (Arg *A = Args.getLastArg(options::OPT_cuda_path_EQ)) {
CmdArgs.push_back(
@@ -10098,10 +10100,12 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA,
// We use action type to differentiate two use cases of the linker wrapper.
// TY_Image for normal linker wrapper work.
- // TY_HIP_FATBIN for HIP device-only links emitting a fat binary directly.
+ // TY_HIP_FATBIN and TY_SYCL_FATBIN for device-only links emitting a fat
+ // binary directly.
assert(JA.getType() == types::TY_HIP_FATBIN ||
+ JA.getType() == types::TY_SYCL_FATBIN ||
JA.getType() == types::TY_Image);
- if (JA.getType() == types::TY_HIP_FATBIN) {
+ if (JA.getType() != types::TY_Image) {
CmdArgs.push_back("--emit-fatbin-only");
CmdArgs.append({"-o", Output.getFilename()});
for (auto Input : Inputs)
diff --git a/clang/lib/Driver/ToolChains/SPIRV.cpp b/clang/lib/Driver/ToolChains/SPIRV.cpp
index 2dd6af3e0a4c9..a2c50ac90fd44 100644
--- a/clang/lib/Driver/ToolChains/SPIRV.cpp
+++ b/clang/lib/Driver/ToolChains/SPIRV.cpp
@@ -162,6 +162,19 @@ void SPIRV::Linker::ConstructJob(Compilation &C, const JobAction &JA,
std::string Linker = ToolChain.GetProgramPath(getShortName());
ArgStringList CmdArgs;
+ // A SYCL fat binary is produced by clang-linker-wrapper, which drives the
+ // device link itself by re-invoking the driver with --sycl-link once per
+ // device image. It discards this command apart from its executable, which it
+ // reports as its own --linker-path=, so all this job has to name is the tool
+ // that will finalize the images.
+ if (JA.getType() == types::TY_SYCL_FATBIN) {
+ C.addCommand(std::make_unique<Command>(
+ JA, *this, ResponseFileSupport::None(),
+ Args.MakeArgString(ToolChain.GetProgramPath("clang-sycl-linker")),
+ CmdArgs, Inputs, Output));
+ return;
+ }
+
// clang-sycl-linker needs the device target triple and architecture to
// finalize a device image. Emit the values derived from --target/-march=
// before the linker inputs, so that -triple=/-arch= passed through
diff --git a/clang/lib/Driver/ToolChains/SYCL.cpp b/clang/lib/Driver/ToolChains/SYCL.cpp
index 34aa99320473e..b0c07dd1be2ee 100644
--- a/clang/lib/Driver/ToolChains/SYCL.cpp
+++ b/clang/lib/Driver/ToolChains/SYCL.cpp
@@ -6,6 +6,7 @@
//
//===----------------------------------------------------------------------===//
#include "SYCL.h"
+#include "SPIRV.h"
#include "clang/Driver/CommonArgs.h"
#include "llvm/Support/VirtualFileSystem.h"
@@ -181,6 +182,11 @@ void SYCLToolChain::addClangWarningOptions(ArgStringList &CC1Args) const {
HostTC.addClangWarningOptions(CC1Args);
}
+Tool *SYCLToolChain::buildLinker() const {
+ // The SPIR-V linker dispatches to clang-sycl-linker for SYCL offloading.
+ return new tools::SPIRV::Linker(*this);
+}
+
ToolChain::CXXStdlibType
SYCLToolChain::GetCXXStdlibType(const ArgList &Args) const {
return HostTC.GetCXXStdlibType(Args);
diff --git a/clang/lib/Driver/ToolChains/SYCL.h b/clang/lib/Driver/ToolChains/SYCL.h
index 48f5d986c3e6e..969199ef24949 100644
--- a/clang/lib/Driver/ToolChains/SYCL.h
+++ b/clang/lib/Driver/ToolChains/SYCL.h
@@ -58,6 +58,9 @@ class LLVM_LIBRARY_VISIBILITY SYCLToolChain : public ToolChain {
computeMSVCVersion(const Driver *D,
const llvm::opt::ArgList &Args) const override;
+protected:
+ Tool *buildLinker() const override;
+
private:
const ToolChain &HostTC;
SYCLInstallationDetector SYCLInstallation;
diff --git a/clang/lib/Driver/Types.cpp b/clang/lib/Driver/Types.cpp
index 7cf8af1d1af92..4cca8acd515d7 100644
--- a/clang/lib/Driver/Types.cpp
+++ b/clang/lib/Driver/Types.cpp
@@ -123,13 +123,14 @@ bool types::canTypeBeUserSpecified(ID Id) {
TY_dSYM,
TY_Dependencies,
TY_CUDA_FATBIN,
- TY_HIP_FATBIN};
+ TY_HIP_FATBIN,
+ TY_SYCL_FATBIN};
return !llvm::is_contained(kStaticLangageTypes, Id);
}
bool types::appendSuffixForType(ID Id) {
return Id == TY_PCH || Id == TY_dSYM || Id == TY_CUDA_FATBIN ||
- Id == TY_HIP_FATBIN;
+ Id == TY_HIP_FATBIN || Id == TY_SYCL_FATBIN;
}
bool types::canLipoType(ID Id) {
diff --git a/clang/test/Driver/Inputs/sycl-nordc-registration-second-tu.cpp b/clang/test/Driver/Inputs/sycl-nordc-registration-second-tu.cpp
new file mode 100644
index 0000000000000..f35b1f71dd9fd
--- /dev/null
+++ b/clang/test/Driver/Inputs/sycl-nordc-registration-second-tu.cpp
@@ -0,0 +1,18 @@
+// Second translation unit for sycl-nordc-registration.cpp. It contributes a
+// kernel of its own so that a non-RDC build has to finalize and register two
+// independent device images.
+
+template <typename KernelName, typename... Ts>
+void sycl_kernel_launch(const char *, Ts...) {}
+
+struct second_tu_kernel_name;
+struct second_tu_kernel {
+ void operator()() const {}
+};
+
+[[clang::sycl_kernel_entry_point(second_tu_kernel_name)]]
+void launch_second_tu_kernel(second_tu_kernel KernelFunc) {
+ KernelFunc();
+}
+
+void call_second_tu() { launch_second_tu_kernel(second_tu_kernel{}); }
diff --git a/clang/test/Driver/sycl-nordc-fatbin-section.cpp b/clang/test/Driver/sycl-nordc-fatbin-section.cpp
new file mode 100644
index 0000000000000..d4f30b6098cd5
--- /dev/null
+++ b/clang/test/Driver/sycl-nordc-fatbin-section.cpp
@@ -0,0 +1,27 @@
+/// Verify which section carries the SYCL device image in the host object.
+// REQUIRES: spirv-registered-target, x86-registered-target
+
+// RUN: %clangxx --target=x86_64-unknown-linux-gnu -fsycl -fno-sycl-rdc \
+// RUN: -c %s -o %t.nordc.o
+// RUN: llvm-readelf -S %t.nordc.o \
+// RUN: | FileCheck -check-prefix=NORDC %s --implicit-check-not='.llvm.offloading'
+// NORDC: .sycl_fatbin
+
+// RUN: %clangxx --target=x86_64-unknown-linux-gnu -fsycl -fsycl-rdc \
+// RUN: -c %s -o %t.rdc.o
+// RUN: llvm-readelf -S %t.rdc.o \
+// RUN: | FileCheck -check-prefix=RDC %s --implicit-check-not='.sycl_fatbin'
+// RDC: .llvm.offloading
+
+/// The choice does not depend on the object format.
+// RUN: %clangxx --target=x86_64-pc-windows-msvc -fsycl -fno-sycl-rdc \
+// RUN: -c %s -o %t.nordc.obj
+// RUN: llvm-readobj --sections %t.nordc.obj \
+// RUN: | FileCheck -check-prefix=NORDC %s --implicit-check-not='.llvm.offloading'
+
+// RUN: %clangxx --target=x86_64-pc-windows-msvc -fsycl -fsycl-rdc \
+// RUN: -c %s -o %t.rdc.obj
+// RUN: llvm-readobj --sections %t.rdc.obj \
+// RUN: | FileCheck -check-prefix=RDC %s --implicit-check-not='.sycl_fatbin'
+
+void f() {}
diff --git a/clang/test/Driver/sycl-nordc-registration.cpp b/clang/test/Driver/sycl-nordc-registration.cpp
new file mode 100644
index 0000000000000..12e2921e2a8a7
--- /dev/null
+++ b/clang/test/Driver/sycl-nordc-registration.cpp
@@ -0,0 +1,150 @@
+// Check what a -fno-sycl-rdc object carries and when it is registered: build a
+// two translation unit program, check that each object holds its own finalized
+// device image, then link and run the program to verify that both images reach
+// the SYCL runtime.
+// The program provides its own __sycl_register_lib/__sycl_unregister_lib, so
+// -nolibsycl is used and no SYCL runtime or offload device is needed.
+
+// REQUIRES: spirv-registered-target, system-linux, native
+
+// RUN: rm -rf %t && mkdir -p %t
+// RUN: %clangxx -fsycl -fno-sycl-rdc -c %s -o %t/main.o
+// RUN: %clangxx -fsycl -fno-sycl-rdc -c \
+// RUN: %S/Inputs/sycl-nordc-registration-second-tu.cpp -o %t/second.o
+
+// The image is an offload binary (magic 0x10FF10AD) wrapping a finalized SPIR-V
+// module (magic 0x07230203), both shown little endian by the hex dump.
+// RUN: llvm-readelf --hex-dump=.sycl_fatbin %t/main.o | FileCheck %s \
+// RUN: --check-prefix=IMAGE
+// RUN: llvm-readelf --hex-dump=.sycl_fatbin %t/second.o | FileCheck %s \
+// RUN: --check-prefix=IMAGE
+// IMAGE: 10ff10ad
+// IMAGE: 03022307
+
+// Each image offers the kernels of that translation unit.
+// RUN: llvm-objcopy --dump-section=.sycl_fatbin=%t/main.image %t/main.o /dev/null
+// RUN: llvm-objcopy --dump-section=.sycl_fatbin=%t/second.image %t/second.o /dev/null
+// RUN: FileCheck %s --check-prefix=MAIN-IMAGE --input-file=%t/main.image
+// RUN: FileCheck %s --check-prefix=SECOND-IMAGE --input-file=%t/second.image
+// MAIN-IMAGE-DAG: main_tu_kernel_name
+// MAIN-IMAGE-DAG: main_tu_other_kernel_name
+// SECOND-IMAGE-DAG: second_tu_kernel_name
+
+// The finalized image is one offload binary holding one entry.
+// RUN: llvm-objdump --offloading %t/main.image | FileCheck %s \
+// RUN: --check-prefix=MAIN-ENTRIES
+// MAIN-ENTRIES: OFFLOADING IMAGE [0]:
+// MAIN-ENTRIES: kind{{ *}}spir-v
+// MAIN-ENTRIES: arch{{ *$}}
+// MAIN-ENTRIES-NOT: OFFLOADING IMAGE
+
+// Splitting a translation unit by kernel gives it an image
+// per kernel, but they are merged into a single offload binary.
+// RUN: %clangxx -fsycl -fno-sycl-rdc -fsycl-device-image-split=kernel -c %s \
+// RUN: -o %t/main.split.o
+// RUN: llvm-objcopy --dump-section=.sycl_fatbin=%t/main.split.image \
+// RUN: %t/main.split.o /dev/null
+// RUN: llvm-objdump --offloading %t/main.split.image | FileCheck %s \
+// RUN: --check-prefix=SPLIT-ENTRIES
+// SPLIT-ENTRIES: OFFLOADING IMAGE [0]:
+// SPLIT-ENTRIES: OFFLOADING IMAGE [1]:
+// SPLIT-ENTRIES-NOT: OFFLOADING IMAGE
+
+// Between them the two images still offer both kernels of the translation unit,
+// and the split program links and runs as one image did.
+// RUN: FileCheck %s --check-prefix=MAIN-IMAGE --input-file=%t/main.split.image
+// RUN: %clangxx -fsycl -fno-sycl-rdc -nolibsycl %t/main.split.o %t/second.o \
+// RUN: -o %t/nordc-split
+// RUN: %t/nordc-split | FileCheck %s --check-prefix=NORDC-OUT
+
+// Both images are registered before the other static initializers of the
+// program run, and unregistered after main returns.
+// RUN: %clangxx -fsycl -fno-sycl-rdc -nolibsycl %t/main.o %t/second.o -o %t/nordc
+// RUN: %t/nordc | FileCheck %s --check-prefix=NORDC-OUT
+// NORDC-OUT: registered image 1
+// NORDC-OUT-NEXT: registered image 2
+// NORDC-OUT-NEXT: static initializer sees 2
+// NORDC-OUT-NEXT: main sees 2
+// NORDC-OUT-NEXT: unregistered image 1
+// NORDC-OUT-NEXT: unregistered image 2
+
+// An RDC build of the same sources links the device code together, so a single
+// image is registered instead of one per translation unit.
+// RUN: %clangxx -fsycl -fsycl-rdc -c %s -o %t/main.rdc.o
+// RUN: %clangxx -fsycl -fsycl-rdc -c \
+// RUN: %S/Inputs/sycl-nordc-registration-second-tu.cpp -o %t/second.rdc.o
+// RUN: %clangxx -fsycl -fsycl-rdc -nolibsycl %t/main.rdc.o %t/second.rdc.o \
+// RUN: -o %t/rdc
+// RUN: %t/rdc | FileCheck %s --check-prefix=RDC-OUT
+// RDC-OUT: registered image 1
+// RDC-OUT-NEXT: static initializer sees 1
+// RDC-OUT-NEXT: main sees 1
+// RDC-OUT-NEXT: unregistered image 1
+
+#include <cstdio>
+#include <cstdlib>
+#include <cstring>
+
+template <typename KernelName, typename... Ts>
+void sycl_kernel_launch(const char *, Ts...) {}
+
+struct main_tu_kernel_name;
+struct main_tu_kernel {
+ void operator()() const {}
+};
+
+[[clang::sycl_kernel_entry_point(main_tu_kernel_name)]]
+void launch_main_tu_kernel(main_tu_kernel KernelFunc) {
+ KernelFunc();
+}
+
+// A second kernel of the same translation unit to test that splitting works
+// with no-rdc.
+struct main_tu_other_kernel_name;
+struct main_tu_other_kernel {
+ void operator()() const {}
+};
+
+[[clang::sycl_kernel_entry_point(main_tu_other_kernel_name)]]
+void launch_main_tu_other_kernel(main_tu_other_kernel KernelFunc) {
+ KernelFunc();
+}
+
+void call_second_tu();
+
+static int Registered = 0;
+static int Unregistered = 0;
+
+// Stand in for the SYCL runtime entry points the registration constructors and
+// destructors call. Reject anything that is not an offload binary so that a
+// malformed image is not mistaken for a passing run.
+extern "C" void __sycl_register_lib(void *Image, size_t Size) {
+ static const unsigned char OffloadBinaryMagic[] = {0x10, 0xFF, 0x10, 0xAD};
+ if (Size <= sizeof(OffloadBinaryMagic) ||
+ std::memcmp(Image, OffloadBinaryMagic, sizeof(OffloadBinaryMagic)) != 0) {
+ std::printf("registered something that is not an offload binary\n");
+ std::exit(1);
+ }
+ std::printf("registered image %d\n", ++Registered);
+}
+
+extern "C" void __sycl_unregister_lib(void *Image, size_t Size) {
+ std::printf("unregistered image %d\n", ++Unregistered);
+}
+
+// A dynamic initializer of its own, so that the registration constructors have
+// to be merged into a constructor list the translation unit already has.
+struct RegistrationObserver {
+ RegistrationObserver() {
+ std::printf("static initializer sees %d\n", Registered);
+ }
+};
+static RegistrationObserver Observer;
+
+int main() {
+ launch_main_tu_kernel(main_tu_kernel{});
+ launch_main_tu_other_kernel(main_tu_other_kernel{});
+ call_second_tu();
+ std::printf("main sees %d\n", Registered);
+ return 0;
+}
diff --git a/clang/test/Driver/sycl-offload-jit.cpp b/clang/test/Driver/sycl-offload-jit.cpp
index 6720e8d250381..64982ee5cb75d 100644
--- a/clang/test/Driver/sycl-offload-jit.cpp
+++ b/clang/test/Driver/sycl-offload-jit.cpp
@@ -60,6 +60,89 @@
// CHK-SYCL-RDC-HOST: "-cc1"{{.*}} "-fsycl-is-host" {{.*}} "-fgpu-rdc"
// CHK-SYCL-NORDC-NOT: "-fgpu-rdc"
+/// Check that -fsycl-rdc / -fno-sycl-rdc are accepted as aliases of
+/// -fgpu-rdc / -fno-gpu-rdc for the SYCL device compilation.
+// RUN: %clang -### -fsycl -fsycl-rdc -c %s 2>&1 \
+// RUN: | FileCheck -check-prefix=CHK-SYCL-RDC %s
+// RUN: %clang -### -fsycl -fno-sycl-rdc -c %s 2>&1 \
+// RUN: | FileCheck -check-prefix=CHK-SYCL-NORDC %s
+
+/// Check the phases graph in non-RDC mode.
+// RUN: %clang -ccc-print-phases --target=x86_64-unknown-linux-gnu -fsycl -fno-sycl-rdc %s 2>&1 \
+// RUN: | FileCheck -check-prefixes=CHK-PHASES-NORDC %s
+// CHK-PHASES-NORDC: 6: backend, {5}, ir, (device-sycl)
+// CHK-PHASES-NORDC-NEXT: 7: offload, "device-sycl (spirv64-unknown-unknown)" {6}, ir
+// CHK-PHASES-NORDC-NEXT: 8: llvm-offload-binary, {7}, image, (device-sycl)
+// CHK-PHASES-NORDC-NEXT: 9: clang-linker-wrapper, {8}, sycl-fatbin, (device-sycl)
+// CHK-PHASES-NORDC-NEXT: 10: offload, "host-sycl (x86_64{{.*}})" {2}, "device-sycl (spirv64{{.*}})" {9}, ir
+// CHK-PHASES-NORDC-NEXT: 11: backend, {10}, assembler, (host-sycl)
+// CHK-PHASES-NORDC-NEXT: 12: assembler, {11}, object, (host-sycl)
+// CHK-PHASES-NORDC-NEXT: 13: clang-linker-wrapper, {12}, image, (host-sycl)
+
+/// With more than one architecture the packaged image holds an entry per
+/// architecture, and a single fat binary is expected to reach the host.
+// RUN: %clang -ccc-print-phases --target=x86_64-unknown-linux-gnu -fsycl -fno-sycl-rdc \
+// RUN: --offload-targets=spirv64-unknown-unknown --offload-arch=generic --offload-arch=bmg_g21 \
+// RUN: -c %s 2>&1 | FileCheck -check-prefixes=CHK-PHASES-NORDC-ARCHS %s
+// CHK-PHASES-NORDC-ARCHS: 7: offload, "device-sycl (spirv64-unknown-unknown:bmg_g21)" {6}, ir
+// CHK-PHASES-NORDC-ARCHS: 12: offload, "device-sycl (spirv64-unknown-unknown:generic)" {11}, ir
+// CHK-PHASES-NORDC-ARCHS-NEXT: 13: llvm-offload-binary, {7, 12}, image, (device-sycl)
+// CHK-PHASES-NORDC-ARCHS-NEXT: 14: clang-linker-wrapper, {13}, sycl-fatbin, (device-sycl)
+
+/// More than one device triple is not supported today.
+// RUN: not %clang -### --target=x86_64-unknown-linux-gnu -fsycl -fno-sycl-rdc \
+// RUN: --offload-targets=spirv64-unknown-unknown,spirv32-unknown-unknown -c %s 2>&1 \
+// RUN: | FileCheck -check-prefix=CHK-NORDC-MULTI-TRIPLE %s
+// RUN: not %clang -### --target=x86_64-unknown-linux-gnu -fsycl -fno-gpu-rdc \
+// RUN: --offload-targets=spirv64-unknown-unknown,spirv32-unknown-unknown -c %s 2>&1 \
+// RUN: | FileCheck -check-prefix=CHK-NORDC-MULTI-TRIPLE-GPU %s
+// CHK-NORDC-MULTI-TRIPLE: error: '-fno-sycl-rdc' is not supported with more than one SYCL offloading target
+// CHK-NORDC-MULTI-TRIPLE-GPU: error: '-fno-gpu-rdc' is not supported with more than one SYCL offloading target
+
+/// A single target repeated is one target.
+// RUN: %clang -### --target=x86_64-unknown-linux-gnu -fsycl -fno-sycl-rdc \
+// RUN: --offload-targets=spirv64-unknown-unknown,spirv64-unknown-unknown -c %s 2>&1 \
+// RUN: | FileCheck -check-prefix=CHK-NORDC-DUP-TRIPLE %s
+// CHK-NORDC-DUP-TRIPLE-NOT: error:
+// CHK-NORDC-DUP-TRIPLE: clang-linker-wrapper{{.*}} "--emit-fatbin-only"
+
+/// The same two targets in RDC mode are supported.
+// RUN: %clang -### --target=x86_64-unknown-linux-gnu -fsycl -fsycl-rdc \
+// RUN: --offload-targets=spirv64-unknown-unknown,spirv32-unknown-unknown -c %s 2>&1 \
+// RUN: | FileCheck -check-prefix=CHK-RDC-MULTI-TRIPLE %s
+// CHK-RDC-MULTI-TRIPLE-NOT: error:
+// CHK-RDC-MULTI-TRIPLE: "-cc1" "-triple" "spirv32-unknown-unknown"{{.*}} "-fsycl-is-device"
+// CHK-RDC-MULTI-TRIPLE: "-cc1" "-triple" "spirv64-unknown-unknown"{{.*}} "-fsycl-is-device"
+
+/// Check that in non-RDC mode clang-linker-wrapper finalizes the packaged
+/// device images into a fat binary rather than a host object, and that the
+/// binary is included into the host compilation via -foffload-include-binary.
+// RUN: %clang -### --target=x86_64-unknown-linux-gnu -fsycl -fno-sycl-rdc -c %s 2>&1 \
+// RUN: | FileCheck -check-prefix=CHK-NORDC-INCLUDE %s \
+// RUN: --implicit-check-not='"-fembed-offload-object='
+// CHK-NORDC-INCLUDE: clang-linker-wrapper{{.*}} "--linker-path={{.*}}clang-sycl-linker" "--emit-fatbin-only" "-o" "[[FB:.*]].syclfb"
+// CHK-NORDC-INCLUDE: "-cc1"{{.*}} "-fsycl-is-host"{{.*}} "-foffload-include-binary" "[[FB]].syclfb"
+
+/// Conversely, RDC mode embeds unlinked device code via -fembed-offload-object.
+// RUN: %clang -### --target=x86_64-unknown-linux-gnu -fsycl -fsycl-rdc %s 2>&1 \
+// RUN: | FileCheck -check-prefix=CHK-RDC-EMBED %s \
+// RUN: --implicit-check-not='"-foffload-include-binary"' \
+// RUN: --implicit-check-not='"--emit-fatbin-only"'
+// CHK-RDC-EMBED: "-cc1"{{.*}} "-fsycl-is-host"{{.*}} "-fembed-offload-object=
+
+/// -flto on a SYCL command line requests *host* LTO. It must not divert the
+/// per-TU device finalize to llvm-lto, which would write bitcode where a
+/// finalized device image is expected; the device link is unaffected and the
+/// image is still included at compile time.
+// RUN: %clang -### --target=x86_64-unknown-linux-gnu -fsycl -fno-sycl-rdc \
+// RUN: -flto -c %s 2>&1 | FileCheck -check-prefix=CHK-NORDC-LTO %s \
+// RUN: --implicit-check-not=llvm-lto
+// RUN: %clang -### --target=x86_64-unknown-linux-gnu -fsycl -fno-sycl-rdc \
+// RUN: -flto %s 2>&1 | FileCheck -check-prefix=CHK-NORDC-LTO %s \
+// RUN: --implicit-check-not=llvm-lto
+// CHK-NORDC-LTO: clang-linker-wrapper{{.*}} "--linker-path={{.*}}clang-sycl-linker" "--emit-fatbin-only"
+// CHK-NORDC-LTO: "-cc1"{{.*}} "-fsycl-is-host"{{.*}} "-foffload-include-binary"
+
// Check that --allow-partial-linkage and --create-library are not passed to
// clang-linker-wrapper for SYCL (they are spirv-link flags, not clang-sycl-linker flags).
// RUN: %clang -### --target=x86_64-unknown-linux-gnu -fsycl %s 2>&1 \
@@ -98,6 +181,12 @@
// RUN: | FileCheck -check-prefix=CHK-SPLIT-UNUSED %s
// CHK-SPLIT-UNUSED: warning: argument unused during compilation: '-fsycl-device-image-split=kernel'
+/// In non-RDC mode the split does happen while compiling.
+// RUN: %clang -### -c --target=x86_64-unknown-linux-gnu -fsycl -fno-sycl-rdc \
+// RUN: -fsycl-device-image-split=kernel %s 2>&1 \
+// RUN: | FileCheck -check-prefix=CHK-SPLIT-KERNEL %s \
+// RUN: --implicit-check-not='argument unused during compilation'
+
/// Check for option incompatibility with -fsycl
// RUN: not %clang -### -fsycl -ffreestanding %s 2>&1 \
// RUN: | FileCheck -check-prefix=CHK-INCOMPATIBILITY %s -DINCOMPATOPT=-ffreestanding
diff --git a/libsycl/test/basic/Inputs/no_rdc_multi_tu_second.cpp b/libsycl/test/basic/Inputs/no_rdc_multi_tu_second.cpp
new file mode 100644
index 0000000000000..b835f7da0fc64
--- /dev/null
+++ b/libsycl/test/basic/Inputs/no_rdc_multi_tu_second.cpp
@@ -0,0 +1,8 @@
+// Second translation unit for basic/no_rdc_multi_tu.cpp.
+
+#include "no_rdc_multi_tu_second.hpp"
+
+void runSecondTuKernel(sycl::queue &Q, int *Data, std::size_t N) {
+ Q.parallel_for<class SecondTuKernel>(N,
+ [=](sycl::item<1> I) { Data[I] *= 3; });
+}
diff --git a/libsycl/test/basic/Inputs/no_rdc_multi_tu_second.hpp b/libsycl/test/basic/Inputs/no_rdc_multi_tu_second.hpp
new file mode 100644
index 0000000000000..9de650788e4af
--- /dev/null
+++ b/libsycl/test/basic/Inputs/no_rdc_multi_tu_second.hpp
@@ -0,0 +1,6 @@
+#pragma once
+
+#include <cstddef>
+#include <sycl/sycl.hpp>
+
+void runSecondTuKernel(sycl::queue &Q, int *Data, std::size_t N);
diff --git a/libsycl/test/basic/no_rdc_multi_tu.cpp b/libsycl/test/basic/no_rdc_multi_tu.cpp
new file mode 100644
index 0000000000000..8ee1e6cc7f367
--- /dev/null
+++ b/libsycl/test/basic/no_rdc_multi_tu.cpp
@@ -0,0 +1,68 @@
+// REQUIRES: any-device
+
+// Run a program whose kernels come from two translation units. With
+// -fno-sycl-rdc each translation unit is finalized into a device binary of its
+// own, so both binaries have to be registered and both kernels have to be found
+// at run time.
+// RUN: %clangxx -fsycl -fno-sycl-rdc -c %s -o %t.nordc.main.o
+// RUN: %clangxx -fsycl -fno-sycl-rdc -c %S/Inputs/no_rdc_multi_tu_second.cpp \
+// RUN: -o %t.nordc.second.o
+// RUN: %clangxx -fsycl -fno-sycl-rdc %t.nordc.main.o %t.nordc.second.o \
+// RUN: -o %t.nordc.out
+// RUN: %t.nordc.out
+
+// Splitting the device code of a translation unit by kernel gives it more than
+// one device image, all of which end up in the fat binary of that translation
+// unit. The runtime is handed one binary and has to find the kernels across all
+// of its images.
+// RUN: %clangxx -fsycl -fno-sycl-rdc -fsycl-device-image-split=kernel -c %s \
+// RUN: -o %t.split.main.o
+// RUN: %clangxx -fsycl -fno-sycl-rdc -fsycl-device-image-split=kernel -c \
+// RUN: %S/Inputs/no_rdc_multi_tu_second.cpp -o %t.split.second.o
+// RUN: %clangxx -fsycl -fno-sycl-rdc %t.split.main.o %t.split.second.o \
+// RUN: -o %t.split.out
+// RUN: %t.split.out
+
+// An RDC build of the same sources has to give the same result.
+// RUN: %clangxx -fsycl -fsycl-rdc -c %s -o %t.rdc.main.o
+// RUN: %clangxx -fsycl -fsycl-rdc -c %S/Inputs/no_rdc_multi_tu_second.cpp \
+// RUN: -o %t.rdc.second.o
+// RUN: %clangxx -fsycl -fsycl-rdc %t.rdc.main.o %t.rdc.second.o -o %t.rdc.out
+// RUN: %t.rdc.out
+
+#include "Inputs/no_rdc_multi_tu_second.hpp"
+
+#include <cassert>
+#include <cstddef>
+
+static void runFirstTuKernel(sycl::queue &Q, int *Data, std::size_t N) {
+ Q.parallel_for<class FirstTuKernel>(N,
+ [=](sycl::item<1> I) { Data[I] += 1; });
+}
+
+static void runFirstTuOtherKernel(sycl::queue &Q, int *Data, std::size_t N) {
+ Q.parallel_for<class FirstTuOtherKernel>(
+ N, [=](sycl::item<1> I) { Data[I] += 2; });
+}
+
+int main() {
+ constexpr std::size_t N = 16;
+
+ sycl::queue Q;
+ int *Data = sycl::malloc_shared<int>(N, Q);
+ for (std::size_t I = 0; I < N; ++I)
+ Data[I] = static_cast<int>(I);
+
+ runFirstTuKernel(Q, Data, N);
+ Q.wait();
+ runFirstTuOtherKernel(Q, Data, N);
+ Q.wait();
+ runSecondTuKernel(Q, Data, N);
+ Q.wait();
+
+ for (std::size_t I = 0; I < N; ++I)
+ assert(Data[I] == (static_cast<int>(I) + 3) * 3);
+
+ sycl::free(Data, Q);
+ return 0;
+}
More information about the cfe-commits
mailing list