From d58e3cb5cfe76fd381ddfa97942946ae38b4e4ef Mon Sep 17 00:00:00 2001 From: Vishal Shetagar Date: Tue, 20 Oct 2026 00:50:15 +0530 Subject: [PATCH 1/2] [SYCL] Avoid runtime registration for empty device pre-link outputs When -fsycl-link has no device images or only empty images, emit the required output object without references to __sycl_register_lib or __sycl_unregister_lib. This allows a host linker to link host-only and kernel-free configurations without the SYCL runtime. Handle empty images in both offload wrappers, omit empty modules from post-link image tables, and generate an empty object when the new driver has no device-link outputs. Preserve registration for nonempty images and NativeCPU placeholders. Extend existing regression tests for empty and mixed images, host-only inputs, and empty modules through both post-link implementations. Fixes https://github.com/intel/llvm/issues/23528 Assisted-by: OpenAI Codex --- clang/test/Driver/clang-offload-wrapper.c | 17 ++++++++++ .../linker-wrapper-sycl-image.cpp | 31 +++++++++++++++++++ .../ClangLinkerWrapper.cpp | 8 ++++- .../ClangOffloadWrapper.cpp | 8 +++++ .../Offloading/SYCLOffloadWrapper.cpp | 7 +++++ llvm/lib/SYCLPostLink/SYCLPostLink.cpp | 4 +++ .../sycl-registered-kernels-empty.ll | 5 +++ llvm/tools/sycl-post-link/sycl-post-link.cpp | 4 ++- 8 files changed, 82 insertions(+), 2 deletions(-) diff --git a/clang/test/Driver/clang-offload-wrapper.c b/clang/test/Driver/clang-offload-wrapper.c index bdbd6e761b747..139bcaee8995e 100644 --- a/clang/test/Driver/clang-offload-wrapper.c +++ b/clang/test/Driver/clang-offload-wrapper.c @@ -1,5 +1,18 @@ // REQUIRES: x86-registered-target +// Empty SYCL images must not introduce a runtime dependency. +// RUN: touch %t.empty +// RUN: echo '[Code|Properties|Symbols]' > %t.empty.table +// RUN: clang-offload-wrapper -host=x86_64-unknown-linux-gnu -kind=sycl -batch %t.empty.table -o %t.empty.bc +// RUN: llvm-dis %t.empty.bc -o - | FileCheck %s --check-prefix=EMPTY +// RUN: clang-offload-wrapper -host=x86_64-unknown-linux-gnu -kind=sycl %t.empty -o %t.empty.bc +// RUN: llvm-dis %t.empty.bc -o - | FileCheck %s --check-prefix=EMPTY +// RUN: clang-offload-wrapper -host=x86_64-pc-windows-msvc -kind=sycl %t.empty -target=spir64 %t.empty -o %t.empty.bc +// RUN: llvm-dis %t.empty.bc -o - | FileCheck %s --check-prefix=EMPTY +// EMPTY: target triple +// EMPTY-NOT: __sycl_register_lib +// EMPTY-NOT: __sycl_unregister_lib + // // Check help message. // @@ -93,6 +106,10 @@ // RUN: echo 'Content of device file1' > %t1.tgt // RUN: echo 'Content of device file2' > %t2.tgt // RUN: echo 'Content of device file3' > %t3.tgt +// RUN: clang-offload-wrapper -host=x86_64-unknown-linux-gnu -kind=sycl %t.empty -target=spir64 %t1.tgt -o %t.mixed.bc +// RUN: llvm-dis %t.mixed.bc -o - | FileCheck %s --check-prefix=MIXED +// MIXED: call void @__sycl_register_lib +// MIXED: call void @__sycl_unregister_lib // // ------- // Check bitcode produced by the wrapper tool. diff --git a/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-sycl-image.cpp b/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-sycl-image.cpp index 88eaed8082278..4c8bea9e8de5d 100644 --- a/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-sycl-image.cpp +++ b/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-sycl-image.cpp @@ -38,6 +38,34 @@ // RUN: -sycl-post-link-options=-properties --linker-path=/usr/bin/ld \ // RUN: %t.o -o %t.thin-lto.out 2>&1 | FileCheck %s --check-prefix=THIN-LTO +// Device pre-linking must produce an object without runtime references when +// there are no device inputs or when the device module is empty. +// RUN: %clang_cc1 -DEMPTY -triple=x86_64-unknown-linux-gnu %s -emit-obj -o %t.host-only.o +// RUN: clang-linker-wrapper --sycl-device-link --host-triple=x86_64-unknown-linux-gnu \ +// RUN: --linker-path=/usr/bin/ld %t.host-only.o -o %t.empty.o +// RUN: %clang %t.host-only.o %t.empty.o -o %t.host-only.out +// RUN: clang-linker-wrapper --sycl-device-link --print-wrapped-module \ +// RUN: --host-triple=x86_64-pc-windows-msvc --linker-path=/usr/bin/ld \ +// RUN: %t.host-only.o -o %t.empty.obj 2>&1 | FileCheck %s --check-prefix=EMPTY +// RUN: %clang_cc1 -DEMPTY -fsycl-is-device -triple=spir64-unknown-unknown %s -emit-llvm-bc -o %t.empty.bc +// RUN: llvm-offload-binary -o %t.empty.fat --image=file=%t.empty.bc,kind=sycl,triple=spir64-unknown-unknown +// RUN: %clang_cc1 -DEMPTY -triple=x86_64-unknown-linux-gnu %s -emit-obj -o %t.empty-input.o -fembed-offload-object=%t.empty.fat +// RUN: clang-linker-wrapper --sycl-device-link --print-wrapped-module \ +// RUN: --host-triple=x86_64-unknown-linux-gnu --linker-path=/usr/bin/ld \ +// RUN: --bitcode-library=spir64-unknown-unknown=%t.devicelib.bc \ +// RUN: -sycl-post-link-options=-split=auto -sycl-post-link-options=-symbols \ +// RUN: -sycl-post-link-options=-properties %t.empty-input.o -o %t.empty.o 2>&1 | FileCheck %s --check-prefix=EMPTY +// RUN: %clang %t.host-only.o %t.empty.o -o %t.empty.out +// RUN: clang-linker-wrapper --sycl-device-link --no-use-sycl-post-link-tool \ +// RUN: --print-wrapped-module --host-triple=x86_64-unknown-linux-gnu \ +// RUN: --linker-path=/usr/bin/ld --bitcode-library=spir64-unknown-unknown=%t.devicelib.bc \ +// RUN: %t.empty-input.o -o %t.empty.o 2>&1 | FileCheck %s --check-prefix=EMPTY +// RUN: %clang %t.host-only.o %t.empty.o -o %t.empty-library.out +// EMPTY: target triple +// EMPTY-NOT: __sycl_register_lib +// EMPTY-NOT: __sycl_unregister_lib + +#ifndef EMPTY template __attribute__((sycl_kernel)) void kernel(const Func &func) { func(); @@ -52,6 +80,9 @@ void __sycl_unregister_lib(void *) {} int main() { kernel([]() {}); } +#else +int main() {} +#endif // THIN-LTO-NOT: invalid SYCL Table file. // THIN-LTO: @.sycl_offloading.target.0 = internal unnamed_addr constant [7 x i8] c"spir64\00" diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index c6b5824e376e5..ee3d22740c588 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -1244,7 +1244,7 @@ wrapSYCLBinariesFromFile(ArrayRef SplitModules, } StringRef Target = Args.getLastArgValue(OPT_triple_EQ); - if (Target.empty()) + if (Target.empty() && !SplitModules.empty()) return createStringError( inconvertibleErrorCode(), "can't wrap SYCL image. -triple argument is missing."); @@ -3671,6 +3671,12 @@ int main(int Argc, char **Argv) { } else if (Args.hasArg(OPT_sycl_device_link)) { // Skip host linker if --sycl-device-link option is set. // Just copy the output of device linking and wrapping action. + if (FilesOrErr->empty()) { + auto OutputFile = sycl::runWrapperAndCompile({}, Args); + if (!OutputFile) + reportError(OutputFile.takeError()); + FilesOrErr->push_back(*OutputFile); + } if (FilesOrErr->size() != 1) reportError( createStringError("Expect single output from the device linker.")); diff --git a/clang/tools/clang-offload-wrapper/ClangOffloadWrapper.cpp b/clang/tools/clang-offload-wrapper/ClangOffloadWrapper.cpp index 344f483e785ab..0384bf3473ffd 100644 --- a/clang/tools/clang-offload-wrapper/ClangOffloadWrapper.cpp +++ b/clang/tools/clang-offload-wrapper/ClangOffloadWrapper.cpp @@ -1014,6 +1014,7 @@ class BinaryWrapper { // Create initializer for the images array. SmallVector ImagesInits; unsigned ImgId = 0; + bool HasNonEmptyImage = false; for (const auto &ImgPtr : Pack) { const BinaryWrapper::Image &Img = *(ImgPtr.get()); @@ -1049,6 +1050,8 @@ class BinaryWrapper { if (!BinOrErr) return BinOrErr.takeError(); MemoryBuffer *Bin = *BinOrErr; + HasNonEmptyImage |= + Bin->getBufferSize() != 0 || Triple(Img.Tgt).isNativeCPU(); if (Img.File != "-" && Kind == OldOffloadKind::OpenMP && AddOpenMPOffloadNotes) { // Adding ELF notes for STDIN is not supported yet. @@ -1141,6 +1144,9 @@ class BinaryWrapper { ImgId++; } + if (Kind == OldOffloadKind::SYCL && !HasNonEmptyImage) + return nullptr; + // Then create images array. auto *ImagesData = Kind == OldOffloadKind::SYCL @@ -1356,6 +1362,8 @@ class BinaryWrapper { Expected DescOrErr = createBinDesc(Kind, *Pack); if (!DescOrErr) return DescOrErr.takeError(); + if (!*DescOrErr) + continue; if (EmitRegFuncs) { GlobalVariable *Desc = *DescOrErr; diff --git a/llvm/lib/Frontend/Offloading/SYCLOffloadWrapper.cpp b/llvm/lib/Frontend/Offloading/SYCLOffloadWrapper.cpp index 9457b278e0346..fcf1f506d94c3 100644 --- a/llvm/lib/Frontend/Offloading/SYCLOffloadWrapper.cpp +++ b/llvm/lib/Frontend/Offloading/SYCLOffloadWrapper.cpp @@ -16,6 +16,7 @@ #include "llvm/Frontend/Offloading/SYCLOffloadWrapper.h" #include "llvm/ADT/ArrayRef.h" +#include "llvm/ADT/STLExtras.h" #include "llvm/ADT/SmallVector.h" #include "llvm/ADT/StringRef.h" #include "llvm/ADT/StringSwitch.h" @@ -773,6 +774,12 @@ struct Wrapper { Error llvm::offloading::wrapSYCLBinaries(llvm::Module &M, const SmallVector &Images, SYCLWrappingOptions Options) { + if (llvm::all_of(Images, [](const SYCLImage &Image) { + return Image.Image->getBufferSize() == 0 && + Image.Target != "native_cpu"; + })) + return Error::success(); + Wrapper W(M, Options); GlobalVariable *Desc = W.createFatbinDesc(Images); if (!Desc) diff --git a/llvm/lib/SYCLPostLink/SYCLPostLink.cpp b/llvm/lib/SYCLPostLink/SYCLPostLink.cpp index 02fa87bebd681..e913ec28c138b 100644 --- a/llvm/lib/SYCLPostLink/SYCLPostLink.cpp +++ b/llvm/lib/SYCLPostLink/SYCLPostLink.cpp @@ -192,6 +192,10 @@ llvm::sycl_post_link::performPostLinkProcessing( Modules.push_back(std::move(MD)); for (std::unique_ptr &MD : Modules) { + if (MD->getModule().global_values().empty() && + MD->getModule().getModuleInlineAsm().empty()) + continue; + size_t ID = SplitModules.size(); StringRef Suffix = MD->isESIMD() ? "_esimd" : ""; std::string OutIRFilename = diff --git a/llvm/test/tools/sycl-post-link/sycl-registered-kernels-empty.ll b/llvm/test/tools/sycl-post-link/sycl-registered-kernels-empty.ll index bb92dd8b081ac..ef4a0e67ea038 100644 --- a/llvm/test/tools/sycl-post-link/sycl-registered-kernels-empty.ll +++ b/llvm/test/tools/sycl-post-link/sycl-registered-kernels-empty.ll @@ -1,3 +1,8 @@ ; RUN: sycl-post-link -properties %s -o %t.table ; RUN: FileCheck %s -input-file=%t_0.prop --implicit-check-not="[SYCL/registered kernels]" +; RUN: FileCheck %s -input-file=%t.table --check-prefix=EMPTY +; RUN: sycl-post-link -properties -symbols -split=auto %s -o %t.table +; RUN: FileCheck %s -input-file=%t.table --check-prefix=EMPTY +; EMPTY: [Code|Properties{{(\|Symbols)?}}] +; EMPTY-NOT: {{.}} !sycl_registered_kernels = !{} diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index 0a038dc376acf..95178276ebc50 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -347,7 +347,9 @@ Error saveModule( SplitMode, static_cast(IdQueriesRange.getValue()))) return E; } - addTableRow(*Table, CopyTriple); + if (!MD.getModule().global_values().empty() || + !MD.getModule().getModuleInlineAsm().empty()) + addTableRow(*Table, CopyTriple); } return Error::success(); From d8add9c85fe7dd9af464ce9aea9fa32b1d7db399 Mon Sep 17 00:00:00 2001 From: Vishal Shetagar Date: Tue, 20 Oct 2026 02:49:53 +0530 Subject: [PATCH 2/2] [SYCL][Test] Use nonempty inputs for device image checks Give the non-RDC wrapper and linker-wrapper command tests a kernel so post-link processing retains an image. Empty sources now correctly omit images and registration, so they cannot exercise image sections or the missing-ocloc diagnostic. Disable integration-header inclusion to keep these standalone fixtures independent of installed SYCL headers and host C++ libraries. Extend the non-RDC test to check that empty ELF and COFF objects contain no embedded device image or SYCL registration references. Reproduced both original failures with the PR's CI-built compiler. The updated tests pass, along with the surrounding SYCL driver, wrapper, and post-link tests: 294 passed, 3 expected failures, 15 unsupported, and no unexpected failures. Assisted-by: OpenAI Codex --- .../test/Driver/sycl-nordc-wrapper-object.cpp | 30 ++++++++++++++++--- .../linker-wrapper-sycl.cpp | 25 ++++++++++------ 2 files changed, 42 insertions(+), 13 deletions(-) diff --git a/clang/test/Driver/sycl-nordc-wrapper-object.cpp b/clang/test/Driver/sycl-nordc-wrapper-object.cpp index a98899baa5f54..2e49457843255 100644 --- a/clang/test/Driver/sycl-nordc-wrapper-object.cpp +++ b/clang/test/Driver/sycl-nordc-wrapper-object.cpp @@ -7,33 +7,55 @@ /// and no ".llvm.offloading", so the final link does not device-link it again. /// The RDC host object is the reverse. // RUN: %clangxx --target=x86_64-unknown-linux-gnu -fsycl --offload-new-driver \ -// RUN: --no-offloadlib -fno-gpu-rdc -c %s -o %t.nordc.o +// RUN: --no-offloadlib -fno-sycl-use-header -fno-gpu-rdc -c %s -o %t.nordc.o // RUN: llvm-readelf -S %t.nordc.o \ // RUN: | FileCheck -check-prefix=NORDC %s --implicit-check-not='.llvm.offloading' // RUN: llvm-nm %t.nordc.o | FileCheck -check-prefix=REG %s // RUN: %clangxx --target=x86_64-unknown-linux-gnu -fsycl --offload-new-driver \ -// RUN: --no-offloadlib -fgpu-rdc -c %s -o %t.rdc.o +// RUN: --no-offloadlib -fno-sycl-use-header -fgpu-rdc -c %s -o %t.rdc.o // RUN: llvm-readelf -S %t.rdc.o \ // RUN: | FileCheck -check-prefix=RDC %s --implicit-check-not='.tgtimg' // RUN: llvm-nm %t.rdc.o | FileCheck -check-prefix=NOREG %s /// The choice does not depend on the object format. // RUN: %clangxx --target=x86_64-pc-windows-msvc -fsycl --offload-new-driver \ -// RUN: --no-offloadlib -fno-gpu-rdc -c %s -o %t.nordc.obj +// RUN: --no-offloadlib -fno-sycl-use-header -fno-gpu-rdc -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: llvm-nm %t.nordc.obj | FileCheck -check-prefix=REG %s // RUN: %clangxx --target=x86_64-pc-windows-msvc -fsycl --offload-new-driver \ -// RUN: --no-offloadlib -fgpu-rdc -c %s -o %t.rdc.obj +// RUN: --no-offloadlib -fno-sycl-use-header -fgpu-rdc -c %s -o %t.rdc.obj // RUN: llvm-readobj --sections %t.rdc.obj \ // RUN: | FileCheck -check-prefix=RDC %s --implicit-check-not='.tgtimg' // RUN: llvm-nm %t.rdc.obj | FileCheck -check-prefix=NOREG %s +// Empty device modules need neither a device image nor runtime registration. +// RUN: %clangxx --target=x86_64-unknown-linux-gnu -fsycl --offload-new-driver \ +// RUN: --no-offloadlib -fno-sycl-use-header -fno-gpu-rdc -DEMPTY -c %s -o %t.empty.o +// RUN: llvm-readobj --sections --symbols %t.empty.o | FileCheck -check-prefix=EMPTY %s +// RUN: %clangxx --target=x86_64-pc-windows-msvc -fsycl --offload-new-driver \ +// RUN: --no-offloadlib -fno-sycl-use-header -fno-gpu-rdc -DEMPTY -c %s -o %t.empty.obj +// RUN: llvm-readobj --sections --symbols %t.empty.obj | FileCheck -check-prefix=EMPTY %s + // NORDC: .tgtimg // RDC: .llvm.offloading // REG: U __sycl_register_lib // NOREG-NOT: __sycl_register_lib +// EMPTY: Format: +// EMPTY-NOT: .tgtimg +// EMPTY-NOT: .llvm.offloading +// EMPTY-NOT: __sycl_register_lib +// EMPTY-NOT: __sycl_unregister_lib + +#ifndef EMPTY +template +__attribute__((sycl_kernel)) void kernel(const Func &func) { + func(); +} +void f() { kernel([]() {}); } +#else void f() {} +#endif diff --git a/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-sycl.cpp b/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-sycl.cpp index 1b380896381cf..1307d2f25a492 100644 --- a/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-sycl.cpp +++ b/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-sycl.cpp @@ -4,12 +4,12 @@ // ------- // Generate .o file as linker wrapper input. // -// RUN: %clang %s -fsycl -fsycl-targets=spir64-unknown-unknown -c --offload-new-driver --no-offloadlib -fno-sycl-instrument-device-code -o %t.o -// RUN: %clang %s -fsycl -fsycl-targets=intel_gpu_pvc -c --offload-new-driver --no-offloadlib -fno-sycl-instrument-device-code -o %t_aot_gpu.o -// RUN: %clang %s -fsycl -fsycl-targets=spir64_x86_64 -c --offload-new-driver --no-offloadlib -fno-sycl-instrument-device-code -o %t_aot_cpu.o -// RUN: %clang %s -fsycl -fsycl-targets=nvptx64-nvidia-cuda --cuda-gpu-arch=sm_50 -nocudalib -fno-sycl-libspirv -c --offload-new-driver -o %t_nvptx.o -// RUN: %clang %s -fsycl -fsycl-targets=amdgcn-amd-amdhsa -Xsycl-target-backend=amdgcn-amd-amdhsa --offload-arch=gfx803 -fgpu-rdc -nogpulib -fno-sycl-libspirv -c --offload-new-driver -o %t_amdgcn.o -// RUN: %clang %s -fsycl -fsycl-targets=native_cpu -fno-sycl-libspirv -c --offload-new-driver -o %t_native_cpu.o +// RUN: %clang %s -fno-sycl-use-header -fsycl -fsycl-targets=spir64-unknown-unknown -c --offload-new-driver --no-offloadlib -fno-sycl-instrument-device-code -o %t.o +// RUN: %clang %s -fno-sycl-use-header -fsycl -fsycl-targets=intel_gpu_pvc -c --offload-new-driver --no-offloadlib -fno-sycl-instrument-device-code -o %t_aot_gpu.o +// RUN: %clang %s -fno-sycl-use-header -fsycl -fsycl-targets=spir64_x86_64 -c --offload-new-driver --no-offloadlib -fno-sycl-instrument-device-code -o %t_aot_cpu.o +// RUN: %clang %s -fno-sycl-use-header -fsycl -fsycl-targets=nvptx64-nvidia-cuda --cuda-gpu-arch=sm_50 -nocudalib -fno-sycl-libspirv -c --offload-new-driver -o %t_nvptx.o +// RUN: %clang %s -fno-sycl-use-header -fsycl -fsycl-targets=amdgcn-amd-amdhsa -Xsycl-target-backend=amdgcn-amd-amdhsa --offload-arch=gfx803 -fgpu-rdc -nogpulib -fno-sycl-libspirv -c --offload-new-driver -o %t_amdgcn.o +// RUN: %clang %s -fno-sycl-use-header -fsycl -fsycl-targets=native_cpu -fno-sycl-libspirv -c --offload-new-driver -o %t_native_cpu.o // // Generate .bc file as SYCL device library file. // @@ -216,7 +216,7 @@ // CHK-OCLOC-PATH-NOARG: no directory given for '--ocloc-path=' // / qualifier routes each value to that arch's ocloc call. -// RUN: %clang %s -fsycl -fsycl-targets=intel_gpu_skl -c --offload-new-driver --no-offloadlib -fno-sycl-instrument-device-code -o %t_aot_skl.o +// RUN: %clang %s -fno-sycl-use-header -fsycl -fsycl-targets=intel_gpu_skl -c --offload-new-driver --no-offloadlib -fno-sycl-instrument-device-code -o %t_aot_skl.o // RUN: clang-linker-wrapper \ // RUN: --device-compiler=sycl:spir64_gen-unknown-unknown/pvc=--ocloc-options=-extraopt_pvc \ // RUN: --device-compiler=sycl:spir64_gen-unknown-unknown/skl=--ocloc-options=-extraopt_skl \ @@ -437,7 +437,7 @@ // CHECK-ERROR-WITH-WRONG-SPIRV-SUBARCH: SYCL device compilation is not supported for the target: spirv64v1.0 // Check -no-sycl-rdc -// RUN: %clang %s -fsycl -fsycl-targets=spir64-unknown-unknown -c --offload-new-driver --no-offloadlib -fno-sycl-instrument-device-code -o %t_rdc.o +// RUN: %clang %s -fno-sycl-use-header -fsycl -fsycl-targets=spir64-unknown-unknown -c --offload-new-driver --no-offloadlib -fno-sycl-instrument-device-code -o %t_rdc.o // // RUN: clang-linker-wrapper --host-triple=x86_64-unknown-linux-gnu --linker-path=/usr/bin/ld -- -o /dev/null %t.o %t_rdc.o --dry-run 2>&1 | FileCheck -check-prefix=CHECK-RDC %s // CHECK-RDC: spirv-to-ir-wrapper{{.*}} -o [[FIRSTLLVMLINKIN:.*]].bc --llvm-spirv-opts --spirv-preserve-auxdata --spirv-target-env=SPV-IR --spirv-builtin-format=global @@ -500,7 +500,7 @@ // CHECK-WRAPPER-AMD: llvm-link{{.*}} {{.*}}.amd.devicelib.bc // Check that each target uses only its own libraries with multi-target input. -// RUN: %clang %s -fsycl -fsycl-targets=amdgcn-amd-amdhsa,nvptx64-nvidia-cuda \ +// RUN: %clang %s -fno-sycl-use-header -fsycl -fsycl-targets=amdgcn-amd-amdhsa,nvptx64-nvidia-cuda \ // RUN: -Xsycl-target-backend=amdgcn-amd-amdhsa --offload-arch=gfx803 \ // RUN: -fno-sycl-libspirv -Wno-unsafe-libspirv-not-linked \ // RUN: --offload-new-driver -c -o %t_multi.o -nocudalib -nogpulib -fgpu-rdc @@ -511,3 +511,10 @@ // RUN: --linker-path=/usr/bin/ld %t_multi.o -o /dev/null 2>&1 | FileCheck -check-prefix=CHECK-WRAPPER-MULTI %s // CHECK-WRAPPER-MULTI: llvm-link{{.*}} {{.*}}.amd.devicelib.bc // CHECK-WRAPPER-MULTI: llvm-link{{.*}} {{.*}}.nvptx.devicelib.bc {{.*}}.nvptx.libdummy.bc + +template +__attribute__((sycl_kernel)) void kernel(const Func &func) { + func(); +} + +void f() { kernel([]() {}); }