[SYCL] Avoid runtime registration for empty device pre-link outputs

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.

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 <typename Name, typename Func>
+__attribute__((sycl_kernel)) void kernel(const Func &func) {
+  func();
+}
 
+void f() { kernel<class TestKernel>([]() {}); }
+#else
 void f() {}
+#endif

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 <typename t, typename Func>
 __attribute__((sycl_kernel)) void kernel(const Func &func) {
   func();
@@ -52,6 +80,9 @@ void __sycl_unregister_lib(void *) {}
 int main() {
   kernel<class fake_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"

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='
 
 // /<arch> 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 <typename Name, typename Func>
+__attribute__((sycl_kernel)) void kernel(const Func &func) {
+  func();
+}
+
+void f() { kernel<class TestKernel>([]() {}); }

clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp

@@ -1244,7 +1244,7 @@ wrapSYCLBinariesFromFile(ArrayRef<module_split::SplitModule> 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."));

clang/tools/clang-offload-wrapper/ClangOffloadWrapper.cpp

@@ -1014,6 +1014,7 @@ class BinaryWrapper {
     // Create initializer for the images array.
     SmallVector<Constant *, 4u> 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<GlobalVariable *> DescOrErr = createBinDesc(Kind, *Pack);
       if (!DescOrErr)
         return DescOrErr.takeError();
+      if (!*DescOrErr)
+        continue;
 
       if (EmitRegFuncs) {
         GlobalVariable *Desc = *DescOrErr;

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<SYCLImage> &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)

llvm/lib/SYCLPostLink/SYCLPostLink.cpp

@@ -192,6 +192,10 @@ llvm::sycl_post_link::performPostLinkProcessing(
       Modules.push_back(std::move(MD));
 
     for (std::unique_ptr<module_split::ModuleDesc> &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 =

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 = !{}

llvm/tools/sycl-post-link/sycl-post-link.cpp

@@ -347,7 +347,9 @@ Error saveModule(
               SplitMode, static_cast<int>(IdQueriesRange.getValue())))
         return E;
     }
-    addTableRow(*Table, CopyTriple);
+    if (!MD.getModule().global_values().empty() ||
+        !MD.getModule().getModuleInlineAsm().empty())
+      addTableRow(*Table, CopyTriple);
   }
 
   return Error::success();