From bad7522132e141b174b6603096dd2e7df2d6016a Mon Sep 17 00:00:00 2001 From: Alexey Bader Date: Mon, 19 Oct 2026 17:32:00 -0700 Subject: [PATCH 1/4] [SYCL] Emit SYCL fat binary for host-side wrapping Align SYCL no-RDC with the upstream flow: clang-linker-wrapper --emit-fatbin-only emits finalized post-link device images in a single offload binary container, without wrapping or host linking, and the host compilation wraps it for runtime registration via -foffload-include-binary. The same container also supports device-only image extraction. Each container entry carries the image together with the SYCL runtime metadata (target, image format, kernel entries, properties and compile/link options). The host compilation reads the entries back, builds the descriptor-based registration the SYCL runtime expects and links it into the host module. Image dumping and compression are shared between the wrapped and fat binary outputs. NativeCPU and --sycl-embed-ir for NVPTX/AMDGCN, which add host objects, are diagnosed with --emit-fatbin-only. Related to CMPLRLLVM-77274 ticket. Co-Authored-By: OpenAI Codex --- clang/lib/CodeGen/CMakeLists.txt | 1 + clang/lib/CodeGen/CodeGenSYCL.cpp | 17 +- clang/lib/CodeGen/ModuleLinker.cpp | 56 +++--- .../test/CodeGen/sycl-no-rdc-embed-binary.cpp | 48 +++-- clang/test/Driver/sycl-nordc-fatbin.cpp | 40 ++-- .../test/Driver/sycl-nordc-wrapper-object.cpp | 7 +- .../linker-wrapper-sycl-fatbin.cpp | 53 +++++ .../linker-wrapper-sycl-image.cpp | 18 +- .../linker-wrapper-sycl.cpp | 13 +- .../ClangLinkerWrapper.cpp | 182 +++++++++++------- .../Frontend/Offloading/SYCLOffloadWrapper.h | 18 ++ .../Offloading/SYCLOffloadWrapper.cpp | 83 ++++++++ 12 files changed, 384 insertions(+), 152 deletions(-) create mode 100644 clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-sycl-fatbin.cpp diff --git a/clang/lib/CodeGen/CMakeLists.txt b/clang/lib/CodeGen/CMakeLists.txt index f2d2560435713..f7252f299119c 100644 --- a/clang/lib/CodeGen/CMakeLists.txt +++ b/clang/lib/CodeGen/CMakeLists.txt @@ -2,6 +2,7 @@ set(LLVM_LINK_COMPONENTS ABI AggressiveInstCombine Analysis + BinaryFormat BitReader BitWriter CodeGen diff --git a/clang/lib/CodeGen/CodeGenSYCL.cpp b/clang/lib/CodeGen/CodeGenSYCL.cpp index 793ad85b8c1f3..665de1ace5c6d 100644 --- a/clang/lib/CodeGen/CodeGenSYCL.cpp +++ b/clang/lib/CodeGen/CodeGenSYCL.cpp @@ -14,7 +14,7 @@ #include "CodeGenFunction.h" #include "CodeGenModule.h" #include "clang/Basic/DiagnosticFrontend.h" -#include "llvm/Bitcode/BitcodeReader.h" +#include "llvm/BinaryFormat/Magic.h" #include "llvm/Frontend/Offloading/OffloadWrapper.h" #include "llvm/Support/MemoryBuffer.h" #include "llvm/Support/VirtualFileSystem.h" @@ -112,14 +112,13 @@ llvm::Function *CodeGenModule::embedSYCLDeviceBinary() { return nullptr; } std::unique_ptr Buffer = std::move(BufferOrErr.get()); - // A wrapper module in bitcode form already carries the registration - // descriptor the SYCL runtime expects. It is linked into the host module by - // loadLinkModules, so there is nothing to embed here. In SYCL CUDA - // compatibility mode the file is the CUDA fat binary, which CUDA registers. - if (getLangOpts().CUDA || - llvm::isBitcode( - reinterpret_cast(Buffer->getBufferStart()), - reinterpret_cast(Buffer->getBufferEnd()))) + // The SYCL runtime registers device images through a descriptor that also + // carries their properties, kernel entries and options. loadLinkModules + // wraps an offload binary container into such a descriptor, so there is + // nothing to embed here. In SYCL CUDA compatibility mode the file is the + // CUDA fat binary, which CUDA registers. + if (getLangOpts().CUDA || llvm::identify_magic(Buffer->getBuffer()) == + llvm::file_magic::offload_binary) return nullptr; llvm::Function *RegistrationFunc = nullptr; if (llvm::Error Err = llvm::offloading::wrapSYCLBinaries( diff --git a/clang/lib/CodeGen/ModuleLinker.cpp b/clang/lib/CodeGen/ModuleLinker.cpp index 0cfd0cc32b9f4..f7f91d9319db1 100644 --- a/clang/lib/CodeGen/ModuleLinker.cpp +++ b/clang/lib/CodeGen/ModuleLinker.cpp @@ -11,7 +11,9 @@ #include "clang/Basic/CodeGenOptions.h" #include "clang/Basic/DiagnosticFrontend.h" #include "clang/Frontend/CompilerInstance.h" +#include "llvm/BinaryFormat/Magic.h" #include "llvm/Bitcode/BitcodeReader.h" +#include "llvm/Frontend/Offloading/SYCLOffloadWrapper.h" #include "llvm/IR/Module.h" using namespace clang; @@ -47,39 +49,45 @@ bool clang::loadLinkModules(CompilerInstance &CI, llvm::LLVMContext &Ctx, F.Internalize, F.LinkFlags}); } - // For SYCL no-RDC, link the per-TU wrapper module produced by - // clang-linker-wrapper --emit-fatbin-only into the host module so the SYCL - // runtime finds the device image at program startup. A raw device binary is - // embedded by CodeGenModule::embedSYCLDeviceBinary instead. - if (CI.getLangOpts().SYCLIsHost && - !CI.getCodeGenOpts().OffloadBinaryToEmbedFile.empty()) { - auto BCBuf = CI.getFileManager().getBufferForFile( - CI.getCodeGenOpts().OffloadBinaryToEmbedFile); - if (!BCBuf) { + // For SYCL no-RDC, wrap the device images of the fat binary produced by + // clang-linker-wrapper --emit-fatbin-only into a module that registers them + // with the SYCL runtime, and link it into the host module. Other binaries + // are embedded by CodeGenModule::embedSYCLDeviceBinary instead. + StringRef FatbinFile = CI.getCodeGenOpts().OffloadBinaryToEmbedFile; + if (CI.getLangOpts().SYCLIsHost && !CI.getLangOpts().CUDA && + !FatbinFile.empty()) { + auto FatbinBuf = CI.getFileManager().getBufferForFile(FatbinFile); + if (!FatbinBuf) { CI.getDiagnostics().Report(diag::err_cannot_open_file) - << CI.getCodeGenOpts().OffloadBinaryToEmbedFile - << BCBuf.getError().message(); + << FatbinFile << FatbinBuf.getError().message(); LinkModules.clear(); return true; } - if (!llvm::isBitcode( - reinterpret_cast((*BCBuf)->getBufferStart()), - reinterpret_cast((*BCBuf)->getBufferEnd()))) + if (llvm::identify_magic((*FatbinBuf)->getBuffer()) != + llvm::file_magic::offload_binary) return false; - llvm::Expected> MOrErr = - llvm::parseBitcodeFile((*BCBuf)->getMemBufferRef(), Ctx); - if (!MOrErr) { - llvm::handleAllErrors(MOrErr.takeError(), [&](llvm::ErrorInfoBase &EIB) { + + auto ReportError = [&](llvm::Error Err) { + llvm::handleAllErrors(std::move(Err), [&](llvm::ErrorInfoBase &EIB) { CI.getDiagnostics().Report(diag::err_fe_linking_module) - << CI.getCodeGenOpts().OffloadBinaryToEmbedFile << EIB.message(); + << FatbinFile << EIB.message(); }); LinkModules.clear(); return true; - } - // The wrapper .bc from clang-linker-wrapper has no data layout set. - // Stamp the host layout so linkModules doesn't warn about a mismatch. - (*MOrErr)->setDataLayout(CI.getTarget().getDataLayoutString()); - LinkModules.push_back({std::move(*MOrErr), + }; + auto ImagesOrErr = + llvm::offloading::readSYCLImages((*FatbinBuf)->getMemBufferRef()); + if (!ImagesOrErr) + return ReportError(ImagesOrErr.takeError()); + + auto WrapperM = + std::make_unique("offload.wrapper.object", Ctx); + WrapperM->setTargetTriple(CI.getTarget().getTriple()); + WrapperM->setDataLayout(CI.getTarget().getDataLayoutString()); + if (llvm::Error Err = + llvm::offloading::wrapSYCLBinaries(*WrapperM, *ImagesOrErr)) + return ReportError(std::move(Err)); + LinkModules.push_back({std::move(WrapperM), /*PropagateAttrs=*/false, /*Internalize=*/false, /*LinkFlags=*/0}); diff --git a/clang/test/CodeGen/sycl-no-rdc-embed-binary.cpp b/clang/test/CodeGen/sycl-no-rdc-embed-binary.cpp index f43301be7f0cd..79b1b12ed749b 100644 --- a/clang/test/CodeGen/sycl-no-rdc-embed-binary.cpp +++ b/clang/test/CodeGen/sycl-no-rdc-embed-binary.cpp @@ -1,5 +1,5 @@ -// Tests for -foffload-include-binary with a SYCL wrapper module in bitcode form -// (SYCL -fno-sycl-rdc compile-step device embedding via loadLinkModules). +// Tests for -foffload-include-binary with a SYCL fat binary in offload binary +// form (SYCL -fno-sycl-rdc compile-step device embedding via loadLinkModules). // Missing file: err_cannot_open_file // RUN: not %clang_cc1 -triple x86_64-unknown-linux-gnu -fsycl-is-host \ @@ -7,19 +7,42 @@ // RUN: | FileCheck --check-prefix=CHECK-NO-FILE %s // CHECK-NO-FILE: fatal error: cannot open file 'no-such-file.bc' -// Wrapper module bitcode: linked into the host module as is, so its -// registration constructor is kept and no raw device binary is embedded. +// Offload binary: each image is wrapped into the descriptor the SYCL runtime +// registers, together with its target, options and kernel entries. The raw +// container itself is not embedded. // RUN: split-file %s %t -// RUN: llvm-as %t/wrapper.ll -o %t/wrapper.bc +// RUN: echo -n 'FAKE_SYCL_DEVICE_IMAGE' > %t/image.bin +// RUN: llvm-offload-binary -o %t/fatbin.syclfb \ +// RUN: --image=file=%t/image.bin,kind=sycl,triple=spir64-unknown-unknown,arch=,sycl-target=spir64,sycl-image-format=0,symbols=_ZTS1K,compile-opts=-g,link-opts=-cl-opt // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsycl-is-host \ -// RUN: -foffload-include-binary %t/wrapper.bc -emit-llvm %t/host.cpp -o - \ +// RUN: -foffload-include-binary %t/fatbin.syclfb -emit-llvm %t/host.cpp -o - \ // RUN: | FileCheck --check-prefix=CHECK-LINK %s \ // RUN: --implicit-check-not='.sycl_fatbin' \ // RUN: --implicit-check-not='.sycl_offloading.binary' -// CHECK-LINK: @llvm.global_ctors = {{.*}} ptr @sycl.descriptor_reg +// CHECK-LINK-DAG: @.sycl_offloading.target.0 = internal unnamed_addr constant [7 x i8] c"spir64\00" +// CHECK-LINK-DAG: @.sycl_offloading.opts.compile.0 = internal unnamed_addr constant [3 x i8] c"-g\00" +// CHECK-LINK-DAG: @.sycl_offloading.opts.link.0 = internal unnamed_addr constant [8 x i8] c"-cl-opt\00" +// CHECK-LINK-DAG: @.sycl_offloading.0.data = internal unnamed_addr constant [22 x i8] c"FAKE_SYCL_DEVICE_IMAGE" +// CHECK-LINK-DAG: @__sycl_offload_entry_name = internal unnamed_addr constant [7 x i8] c"_ZTS1K\00" +// CHECK-LINK-DAG: @llvm.global_ctors = {{.*}} ptr @sycl.descriptor_reg // CHECK-LINK: define internal void @sycl.descriptor_reg() // CHECK-LINK: call void @__sycl_register_lib(ptr @.sycl_offloading.descriptor) +// Malformed runtime metadata is diagnosed. +// RUN: llvm-offload-binary -o %t/no-format.syclfb \ +// RUN: --image=file=%t/image.bin,kind=sycl,triple=spir64-unknown-unknown +// RUN: not %clang_cc1 -triple x86_64-unknown-linux-gnu -fsycl-is-host \ +// RUN: -foffload-include-binary %t/no-format.syclfb -emit-llvm %t/host.cpp \ +// RUN: -o /dev/null 2>&1 | FileCheck --check-prefix=CHECK-NO-FORMAT %s +// CHECK-NO-FORMAT: error: cannot link module '{{.*}}no-format.syclfb': invalid SYCL image format '' + +// RUN: llvm-offload-binary -o %t/openmp.syclfb \ +// RUN: --image=file=%t/image.bin,kind=openmp,triple=spir64-unknown-unknown +// RUN: not %clang_cc1 -triple x86_64-unknown-linux-gnu -fsycl-is-host \ +// RUN: -foffload-include-binary %t/openmp.syclfb -emit-llvm %t/host.cpp \ +// RUN: -o /dev/null 2>&1 | FileCheck --check-prefix=CHECK-NOT-SYCL %s +// CHECK-NOT-SYCL: error: cannot link module '{{.*}}openmp.syclfb': offload binary entry 0 is not a SYCL device image + // SYCL CUDA compatibility mode: the host compilation is also a CUDA host // compilation and the file is the CUDA fat binary. Only CUDA registers it; it // is neither linked nor embedded as a SYCL device image. @@ -33,17 +56,6 @@ // CHECK-SYCL-CUDA: section ".nv_fatbin" // CHECK-SYCL-CUDA: call ptr @__cudaRegisterFatBinary( -//--- wrapper.ll -@.sycl_offloading.descriptor = internal constant [1 x i8] zeroinitializer -@llvm.global_ctors = appending global [1 x { i32, ptr, ptr }] [{ i32, ptr, ptr } { i32 1, ptr @sycl.descriptor_reg, ptr null }] - -define internal void @sycl.descriptor_reg() { - call void @__sycl_register_lib(ptr @.sycl_offloading.descriptor) - ret void -} - -declare void @__sycl_register_lib(ptr) - //--- host.cpp void f() {} diff --git a/clang/test/Driver/sycl-nordc-fatbin.cpp b/clang/test/Driver/sycl-nordc-fatbin.cpp index 48f42a9a497b3..a7241ed2b19a7 100644 --- a/clang/test/Driver/sycl-nordc-fatbin.cpp +++ b/clang/test/Driver/sycl-nordc-fatbin.cpp @@ -1,42 +1,46 @@ // REQUIRES: spirv-registered-target, x86-registered-target -// TODO: With the intel/llvm SYCL driver the .sycl_fatbin section holds LLVM -// bitcode rather than an offload binary wrapping a finalized SPIR-V module. -// Enable once the -fno-sycl-rdc flow matches upstream. -// XFAIL: * +// REQUIRES: spirv-to-ir-wrapper, sycl-post-link /// Verify which section carries the SYCL device binary in the host object. -// RUN: %clangxx --target=x86_64-unknown-linux-gnu -fsycl -fno-gpu-rdc \ +// RUN: %clangxx --target=x86_64-unknown-linux-gnu -fsycl --offload-new-driver --no-offloadlib -fno-gpu-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' -// RUN: %clangxx --target=x86_64-unknown-linux-gnu -fsycl -fgpu-rdc \ +/// -fsycl-link also forwards --sycl-device-link to the per-TU no-RDC job. +/// Fatbin-only completion must skip the device-link copy branch. +// RUN: %clangxx --target=x86_64-unknown-linux-gnu -fsycl --offload-new-driver --no-offloadlib -fno-gpu-rdc \ +// RUN: -fsycl-link -c %s -o %t.nordc.devlink.o +// RUN: llvm-readelf -S %t.nordc.devlink.o \ +// RUN: | FileCheck -check-prefix=NORDC %s --implicit-check-not='.llvm.offloading' + +// RUN: %clangxx --target=x86_64-unknown-linux-gnu -fsycl --offload-new-driver --no-offloadlib -fgpu-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' +// RUN: | FileCheck -check-prefix=RDC %s --implicit-check-not='.tgtimg' /// The choice does not depend on the object format. -// RUN: %clangxx --target=x86_64-pc-windows-msvc -fsycl -fno-gpu-rdc \ +// RUN: %clangxx --target=x86_64-pc-windows-msvc -fsycl --offload-new-driver --no-offloadlib -fno-gpu-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 -fgpu-rdc \ +// RUN: %clangxx --target=x86_64-pc-windows-msvc -fsycl --offload-new-driver --no-offloadlib -fgpu-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' +// RUN: | FileCheck -check-prefix=RDC %s --implicit-check-not='.tgtimg' -// NORDC: .sycl_fatbin +// NORDC: .tgtimg // RDC: .llvm.offloading -/// The section holds an offload binary (magic 0x10FF10AD) whose image is a -/// finalized SPIR-V module (magic 0x07230203), both shown little endian by the -/// hex dump. -// RUN: llvm-readelf --hex-dump=.sycl_fatbin %t.nordc.o \ -// RUN: | FileCheck -check-prefix=BINARY %s +/// The host compiler wraps the fat binary into the runtime's descriptor-based +/// registration ABI rather than registering the raw container. +// RUN: %clangxx --target=x86_64-unknown-linux-gnu -fsycl --offload-new-driver --no-offloadlib -fno-gpu-rdc \ +// RUN: -S -emit-llvm %s -o - | FileCheck -check-prefix=WRAPPER %s -// BINARY: 10ff10ad -// BINARY: 03022307 +// WRAPPER: @.sycl_offloading.descriptor = internal constant %__sycl.tgt_bin_desc +// WRAPPER: call void @__sycl_register_lib(ptr @.sycl_offloading.descriptor) +// WRAPPER-NOT: call void @__sycl_register_lib(ptr {{.*}}, i64 void f() {} diff --git a/clang/test/Driver/sycl-nordc-wrapper-object.cpp b/clang/test/Driver/sycl-nordc-wrapper-object.cpp index a98899baa5f54..dbfa2f93a4d18 100644 --- a/clang/test/Driver/sycl-nordc-wrapper-object.cpp +++ b/clang/test/Driver/sycl-nordc-wrapper-object.cpp @@ -1,10 +1,9 @@ // REQUIRES: spirv-registered-target, x86-registered-target // REQUIRES: spirv-to-ir-wrapper, sycl-post-link -/// intel/llvm counterpart of LLORG's sycl-nordc-fatbin.cpp. In intel/llvm the -/// non-RDC host object carries the wrapper module produced at compile time: -/// the device image in ".tgtimg" plus the __sycl_register_lib registration, -/// and no ".llvm.offloading", so the final link does not device-link it again. +/// Verify that the non-RDC host object carries the finalized device image in +/// ".tgtimg" and a runtime descriptor registered with __sycl_register_lib. It +/// has 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 diff --git a/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-sycl-fatbin.cpp b/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-sycl-fatbin.cpp new file mode 100644 index 0000000000000..aa49a3e728fc0 --- /dev/null +++ b/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-sycl-fatbin.cpp @@ -0,0 +1,53 @@ +// REQUIRES: system-linux + +// Fatbin-only output must not generate or compile a host wrapper. +// RUN: touch %t.o +// RUN: llvm-offload-binary -o %t.pack --image=file=%t.o,triple=spir64-unknown-unknown,arch=generic,kind=sycl +// RUN: clang-linker-wrapper --dry-run --emit-fatbin-only --linker-path=/usr/bin/ld %t.pack -o %t.fatbin 2>&1 | FileCheck %s +// CHECK: sycl-post-link +// CHECK: llvm-spirv +// CHECK-NOT: offload-wrapper +// CHECK-NOT: "{{.*}}clang"{{.*}} -c +// CHECK-NOT: "/usr/bin/ld" +// RUN: llvm-objdump --offloading %t.fatbin | FileCheck %s --check-prefix=CONTAINER +// CONTAINER: OFFLOADING IMAGE [0]: +// CONTAINER: kind elf +// CONTAINER: triple spir64-unknown-unknown +// CONTAINER: producer sycl + +// Fatbin-only completion takes precedence over the device-link copy branch. +// RUN: clang-linker-wrapper --dry-run --emit-fatbin-only --sycl-device-link --linker-path=/usr/bin/ld %t.pack -o %t.devlink.fatbin 2>&1 | FileCheck %s +// RUN: llvm-objdump --offloading %t.devlink.fatbin | FileCheck %s --check-prefix=CONTAINER + +// Keep all independent no-RDC modules in a single container. +// RUN: llvm-offload-binary -o %t.multi.pack --image=file=%t.o,triple=spir64-unknown-unknown,arch=generic,kind=sycl --image=file=%t.o,triple=spir64-unknown-unknown,arch=generic,kind=sycl +// RUN: clang-linker-wrapper --dry-run --emit-fatbin-only --no-sycl-rdc --linker-path=/usr/bin/ld %t.multi.pack -o %t.multi.fatbin 2>&1 | FileCheck %s +// RUN: llvm-objdump --offloading %t.multi.fatbin | FileCheck %s --check-prefix=MULTI +// MULTI: OFFLOADING IMAGE [0]: +// MULTI: producer sycl +// MULTI: OFFLOADING IMAGE [1]: +// MULTI: producer sycl +// MULTI-NOT: OFFLOADING IMAGE + +// Reject modes that add host objects through the pipeline callback rather +// than allowing them to overwrite the device-image container. +// RUN: llvm-offload-binary -o %t.native.pack --image=file=%t.o,triple=native_cpu,arch=generic,kind=sycl +// RUN: not clang-linker-wrapper --dry-run --emit-fatbin-only --linker-path=/usr/bin/ld %t.native.pack -o %t.native.fatbin 2>&1 | FileCheck %s --check-prefix=NATIVE +// RUN: test ! -e %t.native.fatbin +// NATIVE-NOT: sycl-post-link +// NATIVE: error: --emit-fatbin-only is not supported for NativeCPU +// NATIVE-NOT: sycl-post-link + +// RUN: llvm-offload-binary -o %t.nvptx.pack --image=file=%t.o,triple=nvptx64-nvidia-cuda,arch=sm_50,kind=sycl +// RUN: not clang-linker-wrapper --dry-run --emit-fatbin-only --sycl-embed-ir --linker-path=/usr/bin/ld %t.nvptx.pack -o %t.nvptx.fatbin 2>&1 | FileCheck %s --check-prefix=EMBED +// RUN: test ! -e %t.nvptx.fatbin +// RUN: llvm-offload-binary -o %t.amdgcn.pack --image=file=%t.o,triple=amdgcn-amd-amdhsa,arch=gfx908,kind=sycl +// RUN: not clang-linker-wrapper --dry-run --emit-fatbin-only --sycl-embed-ir --linker-path=/usr/bin/ld %t.amdgcn.pack -o %t.amdgcn.fatbin 2>&1 | FileCheck %s --check-prefix=EMBED +// RUN: test ! -e %t.amdgcn.fatbin +// EMBED-NOT: sycl-post-link +// EMBED: error: --emit-fatbin-only is not supported with --sycl-embed-ir for NVPTX or AMDGCN +// EMBED-NOT: sycl-post-link + +// Embedding IR does not add host objects for SPIR targets. +// RUN: clang-linker-wrapper --dry-run --emit-fatbin-only --sycl-embed-ir --linker-path=/usr/bin/ld %t.pack -o %t.spir.fatbin 2>&1 | FileCheck %s +// RUN: llvm-objdump --offloading %t.spir.fatbin | FileCheck %s --check-prefix=CONTAINER 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..d0adc12e10d01 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 @@ -15,21 +15,21 @@ // RUN: -sycl-post-link-options=-properties --linker-path=/usr/bin/ld \ // RUN: %t.o -o %t.out 2>&1 | FileCheck %s -// With --emit-fatbin-only (-fno-sycl-rdc compile step) the same wrapper module -// is written to the output as bitcode for the host compilation to link in. +// With --emit-fatbin-only, the output is an offload container holding the +// finalized SPIR-V with its runtime metadata. The host compilation wraps it +// into the same descriptor-based registration. // RUN: clang-linker-wrapper --emit-fatbin-only --no-sycl-rdc \ // RUN: --host-triple=x86_64-unknown-linux-gnu \ // 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 --linker-path=/usr/bin/ld \ // RUN: %t.o -o %t.syclfb -// RUN: llvm-dis %t.syclfb -o - | FileCheck %s --check-prefix=NORDC -// NORDC-DAG: @SYCL_PropSetName.1 = internal unnamed_addr constant [22 x i8] c"SYCL/kernel param opt\00" -// NORDC-DAG: @.sycl_offloading.descriptor = internal constant %__sycl.tgt_bin_desc { i16 1, i16 1, ptr @.sycl_offloading.device_images, ptr null, ptr null } -// NORDC-DAG: @llvm.global_ctors = {{.*}} { i32 1, ptr @sycl.descriptor_reg, ptr null }] -// NORDC: define internal void @sycl.descriptor_reg() section ".text.startup" { -// NORDC-NEXT: entry: -// NORDC-NEXT: call void @__sycl_register_lib(ptr @.sycl_offloading.descriptor) +// RUN: llvm-objdump --offloading %t.syclfb | FileCheck %s --check-prefix=NORDC +// NORDC: OFFLOADING IMAGE [0]: +// NORDC: triple spir64-unknown-unknown +// NORDC: producer sycl +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsycl-is-host \ +// RUN: -foffload-include-binary %t.syclfb -emit-llvm %s -o - | FileCheck %s // With thin LTO, sycl-post-link emits a two-column [Code|Properties] table. // RUN: clang-linker-wrapper --print-wrapped-module --host-triple=x86_64-unknown-linux-gnu \ 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..e21e75bd9ef70 100644 --- a/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-sycl.cpp +++ b/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-sycl.cpp @@ -339,13 +339,16 @@ // CHK-DEVLINK-CMDS-NEXT: "{{.*cp|copy}}" [[CLANGOUT]] /dev/null // CHK-DEVLINK-CMDS-NOT: "{{.*}}/ld" -// Verify that with --emit-fatbin-only (-fno-sycl-rdc compile step) the wrapper -// module is copied to the output as bitcode, without compiling it or calling -// the host linker. +// Verify that --emit-fatbin-only writes a raw fatbin without generating +// wrapper bitcode, compiling it, or host linking. // RUN: clang-linker-wrapper --emit-fatbin-only --no-sycl-rdc --bitcode-library=spir64-unknown-unknown=%t.devicelib.bc -sycl-post-link-options=SYCL_POST_LINK_OPTIONS -llvm-spirv-options=LLVM_SPIRV_OPTIONS --host-triple=x86_64-unknown-linux-gnu --linker-path=/usr/bin/ld -o %t.syclfb %t.o --dry-run 2>&1 | FileCheck -check-prefix=CHK-FATBIN-ONLY %s +// CHK-FATBIN-ONLY-NOT: offload-wrapper: +// CHK-FATBIN-ONLY-NOT: "{{.*}}clang"{{.*}} -c +// CHK-FATBIN-ONLY-NOT: "{{.*cp|copy}}" // CHK-FATBIN-ONLY: llvm-spirv{{.*}} LLVM_SPIRV_OPTIONS -o {{.*}} -// CHK-FATBIN-ONLY-NEXT: offload-wrapper: output: [[WRAPPEROUT:.*]].bc, input: {{.*}} -// CHK-FATBIN-ONLY-NEXT: "{{.*cp|copy}}" [[WRAPPEROUT]].bc {{.*}}.syclfb +// CHK-FATBIN-ONLY-NOT: offload-wrapper: +// CHK-FATBIN-ONLY-NOT: "{{.*}}clang"{{.*}} -c +// CHK-FATBIN-ONLY-NOT: "{{.*cp|copy}}" // CHK-FATBIN-ONLY-NOT: "{{.*}}/ld" // Verify list of commands when syclbin is used diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index c6b5824e376e5..b26c9267fe00f 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -1212,37 +1212,12 @@ static Error compressImages(SmallVectorImpl &Images, return Error::success(); } -/// Reads device images from the given \p InputFile and wraps them -/// in one LLVM IR Module as a constant data. -/// -/// \returns A path to the LLVM Module that contains wrapped images. -Expected -wrapSYCLBinariesFromFile(ArrayRef SplitModules, - const ArgList &Args, bool IsEmbeddedIR) { - auto OutputFileOrErr = createOutputFile( - sys::path::filename(ExecutableName) + ".sycl.image.wrapper", "bc"); - if (!OutputFileOrErr) - return OutputFileOrErr.takeError(); - - StringRef OutputFilePath = *OutputFileOrErr; - if (Verbose || DryRun) { - std::string InputFiles; - SmallString<0> Msg; - for (size_t I = 0, E = SplitModules.size(); I != E; ++I) { - const module_split::SplitModule &SM = SplitModules[I]; - Msg.append(formatv(" input: {0}, compile-opts: {1}, link-opts: {2}", - SM.ModuleFilePath, SM.CompileOptions, SM.LinkOptions) - .sstr<128>()); - if (I + 1 < E) - Msg.append(","); - } - - errs() << formatv(" offload-wrapper: output: {0}, {1}\n", OutputFilePath, - Msg); - if (DryRun) - return OutputFilePath; - } - +/// Reads the post-link device images in \p SplitModules and prepares them for +/// SYCL runtime registration: dumps them if requested, assigns the runtime +/// target and compresses them if requested. +static Expected> +createSYCLImages(ArrayRef SplitModules, + const ArgList &Args, bool IsEmbeddedIR) { StringRef Target = Args.getLastArgValue(OPT_triple_EQ); if (Target.empty()) return createStringError( @@ -1261,7 +1236,7 @@ wrapSYCLBinariesFromFile(ArrayRef SplitModules, RegularTarget = "spir64"; for (const module_split::SplitModule &SI : SplitModules) { - if (!OffloadImageDumpDir.empty()) { + if (!OffloadImageDumpDir.empty() && !DryRun) { StringRef CopyFrom = SI.ModuleFilePath; SmallString<128> CopyTo = OffloadImageDumpDir; StringRef Filename = sys::path::filename(CopyFrom); @@ -1274,8 +1249,11 @@ wrapSYCLBinariesFromFile(ArrayRef SplitModules, } auto MBOrDesc = MemoryBuffer::getFile(SI.ModuleFilePath); - if (!MBOrDesc) - return createFileError(SI.ModuleFilePath, MBOrDesc.getError()); + if (std::error_code EC = MBOrDesc.getError()) { + if (!DryRun) + return createFileError(SI.ModuleFilePath, EC); + MBOrDesc = MemoryBuffer::getMemBuffer("", SI.ModuleFilePath); + } StringRef ImageTarget = IsEmbeddedIR ? StringRef(EmbeddedIRTarget) : StringRef(RegularTarget); @@ -1285,6 +1263,45 @@ wrapSYCLBinariesFromFile(ArrayRef SplitModules, if (Error E = compressImages(Images, Args)) return std::move(E); + return std::move(Images); +} + +/// Reads device images from the given \p InputFile and wraps them +/// in one LLVM IR Module as a constant data. +/// +/// \returns A path to the LLVM Module that contains wrapped images. +Expected +wrapSYCLBinariesFromFile(ArrayRef SplitModules, + const ArgList &Args, bool IsEmbeddedIR) { + auto OutputFileOrErr = createOutputFile( + sys::path::filename(ExecutableName) + ".sycl.image.wrapper", "bc"); + if (!OutputFileOrErr) + return OutputFileOrErr.takeError(); + + StringRef OutputFilePath = *OutputFileOrErr; + if (Verbose || DryRun) { + std::string InputFiles; + SmallString<0> Msg; + for (size_t I = 0, E = SplitModules.size(); I != E; ++I) { + const module_split::SplitModule &SM = SplitModules[I]; + Msg.append(formatv(" input: {0}, compile-opts: {1}, link-opts: {2}", + SM.ModuleFilePath, SM.CompileOptions, SM.LinkOptions) + .sstr<128>()); + if (I + 1 < E) + Msg.append(","); + } + + errs() << formatv(" offload-wrapper: output: {0}, {1}\n", OutputFilePath, + Msg); + if (DryRun) + return OutputFilePath; + } + + Expected> ImagesOrErr = + createSYCLImages(SplitModules, Args, IsEmbeddedIR); + if (!ImagesOrErr) + return ImagesOrErr.takeError(); + SmallVector &Images = *ImagesOrErr; LLVMContext C; Module M("offload.wrapper.object", C); @@ -2933,16 +2950,19 @@ extractSYCLCompileLinkOptions(ArrayRef OffloadFiles) { } /// Transforms all the extracted offloading input files into an image that can -/// be registered by the runtime. If NeedsWrapping is false, writes bundled +/// be registered by the runtime. If EmitFatbinOnly is true, writes bundled /// output directly without wrapping or host linking. Expected> linkAndWrapDeviceFiles(ArrayRef> LinkerInputFiles, const InputArgList &Args, char **Argv, int Argc, - bool NeedsWrapping) { + bool EmitFatbinOnly) { llvm::TimeTraceScope TimeScope("Handle all device input"); std::mutex ImageMtx; MapVector> Images; + // Owns the strings of the SYCL image entries; guarded by ImageMtx. + BumpPtrAllocator SYCLImageStringAlloc; + StringSaver SYCLImageStringSaver(SYCLImageStringAlloc); // Create a binary image of each offloading image and embed it into a new // object file. SmallVector WrappedOutput; @@ -3008,6 +3028,17 @@ linkAndWrapDeviceFiles(ArrayRef> LinkerInputFiles, // check in runSYCLOffloadingPipeline(). const llvm::Triple TargetTriple( LinkerArgs.getLastArgValue(OPT_triple_EQ)); + if (EmitFatbinOnly && !OutputSYCLBIN) { + // These modes add host objects through the pipeline callback, which + // cannot be represented in the device-image container. + if (TargetTriple.isNativeCPU()) + return createStringError( + "--emit-fatbin-only is not supported for NativeCPU"); + if ((TargetTriple.isNVPTX() || TargetTriple.isAMDGCN()) && + LinkerArgs.hasArg(OPT_sycl_embed_ir)) + return createStringError("--emit-fatbin-only is not supported with " + "--sycl-embed-ir for NVPTX or AMDGCN"); + } const bool IsAOT = TargetTriple.isSPIRAOT(); std::vector AOTDeviceArgs; auto AddDeviceOptions = [&](opt::OptSpecifier OptID, bool IsLink, @@ -3056,15 +3087,31 @@ linkAndWrapDeviceFiles(ArrayRef> LinkerInputFiles, MD.SplitModules = std::move(Modules); std::scoped_lock Guard(SYCLBINModulesMtx); SYCLBINModules.emplace_back(std::move(MD)); - } else if (!NeedsWrapping) { - // -fno-sycl-rdc compile step: emit the wrapper module as bitcode; the - // host compilation links it in via -foffload-include-binary. - Expected OutputFile = - sycl::wrapSYCLBinariesFromFile(Modules, LinkerArgs, - /*IsEmbeddedIR=*/false); - if (!OutputFile) - return OutputFile.takeError(); - AppendImageToWrapperOutput(*OutputFile); + } else if (EmitFatbinOnly) { + // Package the finalized images with their runtime metadata; the host + // compilation wraps them for registration with the SYCL runtime. + Expected> SYCLImagesOrErr = + sycl::createSYCLImages(Modules, LinkerArgs, + /*IsEmbeddedIR=*/false); + if (!SYCLImagesOrErr) + return SYCLImagesOrErr.takeError(); + + std::scoped_lock Guard(ImageMtx); + for (offloading::SYCLImage &SYCLImage : *SYCLImagesOrErr) { + OffloadingImage TheImage{}; + TheImage.TheImageKind = + Args.hasArg(OPT_embed_bitcode) ? IMG_Bitcode : IMG_Object; + TheImage.TheOffloadKind = OFK_SYCL; + TheImage.StringData["triple"] = + Args.MakeArgString(LinkerArgs.getLastArgValue(OPT_triple_EQ)); + TheImage.StringData["arch"] = + Args.MakeArgString(LinkerArgs.getLastArgValue(OPT_arch_EQ)); + offloading::addSYCLImageStrings(SYCLImage, TheImage.StringData, + SYCLImageStringSaver); + TheImage.Image = std::move(SYCLImage.Image); + + Images[OFK_SYCL].emplace_back(std::move(TheImage)); + } } else { // TODO(NOM7): Remove this call and use community flow for bundle/wrap Expected OutputFile = @@ -3145,7 +3192,7 @@ linkAndWrapDeviceFiles(ArrayRef> LinkerInputFiles, } for (auto &[Kind, Input] : Images) { - if (Kind == OFK_SYCL) + if (Kind == OFK_SYCL && !EmitFatbinOnly) continue; // We sort the entries before bundling so they appear in a deterministic // order in the final binary. @@ -3160,11 +3207,23 @@ linkAndWrapDeviceFiles(ArrayRef> LinkerInputFiles, return ArchA > ArchB; return A.TheOffloadKind < B.TheOffloadKind; }); + if (Kind == OFK_SYCL) { + // Package all post-link images into a single offload container. + SmallString<0> Buffer = OffloadBinary::write(Input); + auto FOBOrErr = FileOutputBuffer::create(ExecutableName, Buffer.size()); + if (!FOBOrErr) + return FOBOrErr.takeError(); + llvm::copy(Buffer, (*FOBOrErr)->getBufferStart()); + if (Error E = (*FOBOrErr)->commit()) + return std::move(E); + continue; + } + auto BundledImagesOrErr = bundleLinkedOutput(Input, Args, Kind); if (!BundledImagesOrErr) return BundledImagesOrErr.takeError(); - if (!NeedsWrapping) { + if (EmitFatbinOnly) { if (BundledImagesOrErr->size() != 1) return createStringError( "Expected a single bundled image for direct fat binary output"); @@ -3661,34 +3720,27 @@ int main(int Argc, char **Argv) { // Link and process the device images. The function may emit a direct fat // binary if --emit-fatbin-only is specified. auto FilesOrErr = linkAndWrapDeviceFiles(*DeviceInputFiles, Args, Argv, - Argc, !EmitFatbinOnly); + Argc, EmitFatbinOnly); if (!FilesOrErr) reportError(FilesOrErr.takeError()); if (OutputSYCLBIN) { if (Error Err = sycl::mergeSYCLBIN(*FilesOrErr, Args)) reportError(std::move(Err)); - } 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->size() != 1) - reportError( - createStringError("Expect single output from the device linker.")); - if (Error Err = sycl::copyFileToFinalExecutable((*FilesOrErr)[0], Args)) - reportError(std::move(Err)); - } else { - // Run the host linking job with the rendered arguments. - if (!EmitFatbinOnly) { - if (Error Err = runLinker(*FilesOrErr, Args)) - reportError(std::move(Err)); - } else if (!FilesOrErr->empty()) { - // Other offload kinds write the fat binary directly; the SYCL wrapper - // module still has to be copied to the requested output. + } else if (!EmitFatbinOnly) { + // Fatbin-only output has already been written by linkAndWrapDeviceFiles. + 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->size() != 1) reportError(createStringError( - "Expect single output from the SYCL device linker.")); + "Expect single output from the device linker.")); if (Error Err = sycl::copyFileToFinalExecutable((*FilesOrErr)[0], Args)) reportError(std::move(Err)); + } else { + // Run the host linking job with the rendered arguments. + if (Error Err = runLinker(*FilesOrErr, Args)) + reportError(std::move(Err)); } } } diff --git a/llvm/include/llvm/Frontend/Offloading/SYCLOffloadWrapper.h b/llvm/include/llvm/Frontend/Offloading/SYCLOffloadWrapper.h index 4e1765a56a00d..5c4f2b7f5d29e 100644 --- a/llvm/include/llvm/Frontend/Offloading/SYCLOffloadWrapper.h +++ b/llvm/include/llvm/Frontend/Offloading/SYCLOffloadWrapper.h @@ -9,11 +9,14 @@ #ifndef LLVM_CLANG_TOOLS_CLANG_LINKER_WRAPPER_SYCL_OFFLOAD_WRAPPER_H #define LLVM_CLANG_TOOLS_CLANG_LINKER_WRAPPER_SYCL_OFFLOAD_WRAPPER_H +#include "llvm/ADT/MapVector.h" #include "llvm/ADT/SmallVector.h" #include "llvm/ADT/StringRef.h" #include "llvm/IR/Module.h" +#include "llvm/Support/Error.h" #include "llvm/Support/MemoryBuffer.h" #include "llvm/Support/PropertySetIO.h" +#include "llvm/Support/StringSaver.h" #include #include @@ -88,6 +91,21 @@ llvm::Error wrapSYCLBinaries(llvm::Module &M, const llvm::SmallVector &Images, SYCLWrappingOptions Options = SYCLWrappingOptions()); +/// Records the fields of \p Image other than the image bytes in the string +/// table \p StringData of an offload binary entry, so that a fat binary can be +/// wrapped later by readSYCLImages() and wrapSYCLBinaries(). \p Saver owns the +/// serialized strings, which must outlive the offload binary serialization. +void addSYCLImageStrings( + const SYCLImage &Image, + llvm::MapVector &StringData, + llvm::StringSaver &Saver); + +/// Reads SYCL images from the offload binary container in \p Buffer, whose +/// entries carry the strings recorded by addSYCLImageStrings(). The returned +/// images reference the data of \p Buffer, which must outlive them. +llvm::Expected> +readSYCLImages(llvm::MemoryBufferRef Buffer); + } // namespace offloading } // namespace llvm diff --git a/llvm/lib/Frontend/Offloading/SYCLOffloadWrapper.cpp b/llvm/lib/Frontend/Offloading/SYCLOffloadWrapper.cpp index 9457b278e0346..b7a88ff0654af 100644 --- a/llvm/lib/Frontend/Offloading/SYCLOffloadWrapper.cpp +++ b/llvm/lib/Frontend/Offloading/SYCLOffloadWrapper.cpp @@ -28,12 +28,14 @@ #include "llvm/IR/LLVMContext.h" #include "llvm/IR/Module.h" #include "llvm/IR/Type.h" +#include "llvm/Object/OffloadBinary.h" #include "llvm/SYCLLowerIR/SYCLUtils.h" #include "llvm/SYCLLowerIR/UtilsSYCLNativeCPU.h" #include "llvm/Support/Error.h" #include "llvm/Support/ErrorHandling.h" #include "llvm/Support/LineIterator.h" #include "llvm/Support/PropertySetIO.h" +#include "llvm/Support/raw_ostream.h" #include "llvm/TargetParser/Triple.h" #include "llvm/Transforms/Utils/ModuleUtils.h" #include @@ -787,3 +789,84 @@ Error llvm::offloading::wrapSYCLBinaries(llvm::Module &M, } return Error::success(); } + +namespace { +// Offload binary string keys that carry the SYCLImage fields. +constexpr char TargetKey[] = "sycl-target"; +constexpr char FormatKey[] = "sycl-image-format"; +constexpr char EntriesKey[] = "symbols"; +constexpr char PropertiesKey[] = "properties"; +constexpr char CompileOptionsKey[] = "compile-opts"; +constexpr char LinkOptionsKey[] = "link-opts"; +} // anonymous namespace + +void llvm::offloading::addSYCLImageStrings( + const SYCLImage &Image, MapVector &StringData, + StringSaver &Saver) { + std::string Properties; + raw_string_ostream PropertiesOS(Properties); + Image.PropertyRegistry.write(PropertiesOS); + StringData[TargetKey] = Saver.save(Image.Target); + StringData[FormatKey] = Saver.save( + Twine(static_cast(binaryImageFormatToInt8(Image.Format)))); + StringData[EntriesKey] = Saver.save(Image.Entries); + StringData[PropertiesKey] = Saver.save(Properties); + StringData[CompileOptionsKey] = Saver.save(Image.CompileOptions); + StringData[LinkOptionsKey] = Saver.save(Image.LinkOptions); +} + +Expected> +llvm::offloading::readSYCLImages(MemoryBufferRef Buffer) { + auto BinariesOrErr = object::OffloadBinary::create(Buffer); + if (!BinariesOrErr) + return BinariesOrErr.takeError(); + + SmallVector Images; + for (const std::unique_ptr &Binary : *BinariesOrErr) { + if (Binary->getOffloadKind() != object::OFK_SYCL) + return createStringError("offload binary entry " + + Twine(Binary->getIndex()) + + " is not a SYCL device image"); + + std::unique_ptr PropertiesBuffer = + MemoryBuffer::getMemBufferCopy(Binary->getString(PropertiesKey)); + auto PropertiesOrErr = PropertySetRegistry::read(PropertiesBuffer.get()); + if (!PropertiesOrErr) + return PropertiesOrErr.takeError(); + + SYCLImage Image(MemoryBuffer::getMemBuffer( + Binary->getImage(), Buffer.getBufferIdentifier(), + /*RequiresNullTerminator=*/false), + **PropertiesOrErr, Binary->getString(EntriesKey), + Binary->getString(TargetKey), + Binary->getString(CompileOptionsKey), + Binary->getString(LinkOptionsKey)); + + unsigned Format = 0; + if (Binary->getString(FormatKey).getAsInteger(10, Format)) + return createStringError("invalid SYCL image format '" + + Binary->getString(FormatKey) + "'"); + switch (Format) { + case 0: + Image.Format = SYCLBinaryImageFormat::BIF_None; + break; + case 1: + Image.Format = SYCLBinaryImageFormat::BIF_Native; + break; + case 2: + Image.Format = SYCLBinaryImageFormat::BIF_SPIRV; + break; + case 3: + Image.Format = SYCLBinaryImageFormat::BIF_LLVMBC; + break; + case 4: + Image.Format = SYCLBinaryImageFormat::BIF_Compressed; + break; + default: + return createStringError("invalid SYCL image format '" + Twine(Format) + + "'"); + } + Images.emplace_back(std::move(Image)); + } + return std::move(Images); +} From 806f5c667d94a6c7c932ede029994f70608dec33 Mon Sep 17 00:00:00 2001 From: Alexey Bader Date: Mon, 19 Oct 2026 18:05:31 -0700 Subject: [PATCH 2/4] Update the comment. --- clang/test/Driver/sycl-nordc-fatbin.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/clang/test/Driver/sycl-nordc-fatbin.cpp b/clang/test/Driver/sycl-nordc-fatbin.cpp index a7241ed2b19a7..74f50b4e8adbc 100644 --- a/clang/test/Driver/sycl-nordc-fatbin.cpp +++ b/clang/test/Driver/sycl-nordc-fatbin.cpp @@ -2,7 +2,8 @@ // REQUIRES: spirv-to-ir-wrapper, sycl-post-link -/// Verify which section carries the SYCL device binary in the host object. +/// Verify that no-RDC emits image address/size records in .tgtimg, while RDC +/// embeds device code for linking in .llvm.offloading. // RUN: %clangxx --target=x86_64-unknown-linux-gnu -fsycl --offload-new-driver --no-offloadlib -fno-gpu-rdc \ // RUN: -c %s -o %t.nordc.o // RUN: llvm-readelf -S %t.nordc.o \ From 0f73afdf9836a27d6c5b177313192ebd15a8d064 Mon Sep 17 00:00:00 2001 From: Alexey Bader Date: Mon, 19 Oct 2026 20:57:28 -0700 Subject: [PATCH 3/4] [SYCL][Driver][Test] Deduplicate no-RDC wrapper coverage --- clang/test/Driver/sycl-nordc-fatbin.cpp | 9 ---- .../test/Driver/sycl-nordc-wrapper-object.cpp | 41 ++++--------------- 2 files changed, 9 insertions(+), 41 deletions(-) diff --git a/clang/test/Driver/sycl-nordc-fatbin.cpp b/clang/test/Driver/sycl-nordc-fatbin.cpp index 74f50b4e8adbc..9c07c9d76127b 100644 --- a/clang/test/Driver/sycl-nordc-fatbin.cpp +++ b/clang/test/Driver/sycl-nordc-fatbin.cpp @@ -35,13 +35,4 @@ // NORDC: .tgtimg // RDC: .llvm.offloading -/// The host compiler wraps the fat binary into the runtime's descriptor-based -/// registration ABI rather than registering the raw container. -// RUN: %clangxx --target=x86_64-unknown-linux-gnu -fsycl --offload-new-driver --no-offloadlib -fno-gpu-rdc \ -// RUN: -S -emit-llvm %s -o - | FileCheck -check-prefix=WRAPPER %s - -// WRAPPER: @.sycl_offloading.descriptor = internal constant %__sycl.tgt_bin_desc -// WRAPPER: call void @__sycl_register_lib(ptr @.sycl_offloading.descriptor) -// WRAPPER-NOT: call void @__sycl_register_lib(ptr {{.*}}, i64 - void f() {} diff --git a/clang/test/Driver/sycl-nordc-wrapper-object.cpp b/clang/test/Driver/sycl-nordc-wrapper-object.cpp index dbfa2f93a4d18..92458837df291 100644 --- a/clang/test/Driver/sycl-nordc-wrapper-object.cpp +++ b/clang/test/Driver/sycl-nordc-wrapper-object.cpp @@ -1,38 +1,15 @@ -// REQUIRES: spirv-registered-target, x86-registered-target -// REQUIRES: spirv-to-ir-wrapper, sycl-post-link +// REQUIRES: sycl-post-link -/// Verify that the non-RDC host object carries the finalized device image in -/// ".tgtimg" and a runtime descriptor registered with __sycl_register_lib. It -/// has no ".llvm.offloading", so the final link does not device-link it again. -/// The RDC host object is the reverse. +/// Verify that the no-RDC driver flow registers and unregisters the descriptor +/// expected by the Intel SYCL runtime. // 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: 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: 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: 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: --no-offloadlib -fno-gpu-rdc -S -emit-llvm %s -o - \ +// RUN: | FileCheck %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: 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 +// RUN: --no-offloadlib -fno-gpu-rdc -S -emit-llvm %s -o - \ +// RUN: | FileCheck %s -// NORDC: .tgtimg -// RDC: .llvm.offloading -// REG: U __sycl_register_lib -// NOREG-NOT: __sycl_register_lib +// CHECK: call void @__sycl_register_lib(ptr @.sycl_offloading.descriptor) +// CHECK: call void @__sycl_unregister_lib(ptr @.sycl_offloading.descriptor) void f() {} From 7540df2a763b9dcefda4fc9ea67709c5a7293a6b Mon Sep 17 00:00:00 2001 From: Alexey Bader Date: Tue, 20 Oct 2026 11:04:22 -0700 Subject: [PATCH 4/4] Apply code review suggestion. --- .../Frontend/Offloading/SYCLOffloadWrapper.h | 3 ++ .../Offloading/SYCLOffloadWrapper.cpp | 45 ++++++++++--------- 2 files changed, 27 insertions(+), 21 deletions(-) diff --git a/llvm/include/llvm/Frontend/Offloading/SYCLOffloadWrapper.h b/llvm/include/llvm/Frontend/Offloading/SYCLOffloadWrapper.h index 5c4f2b7f5d29e..ce260454694dc 100644 --- a/llvm/include/llvm/Frontend/Offloading/SYCLOffloadWrapper.h +++ b/llvm/include/llvm/Frontend/Offloading/SYCLOffloadWrapper.h @@ -42,6 +42,9 @@ enum class SYCLBinaryImageFormat { /// Returns the ABI value of \p Format shared with the SYCL runtime. int8_t binaryImageFormatToInt8(SYCLBinaryImageFormat Format); +/// Returns the SYCL image format for an ABI value, or an error if invalid. +Expected int8ToBinaryImageFormat(int8_t Format); + /// Misc property with the pre-compression format; mirrored in SYCL RT. constexpr char OriginalFormatPropName[] = "originalFormat"; diff --git a/llvm/lib/Frontend/Offloading/SYCLOffloadWrapper.cpp b/llvm/lib/Frontend/Offloading/SYCLOffloadWrapper.cpp index b7a88ff0654af..19c0f7fcd1364 100644 --- a/llvm/lib/Frontend/Offloading/SYCLOffloadWrapper.cpp +++ b/llvm/lib/Frontend/Offloading/SYCLOffloadWrapper.cpp @@ -65,6 +65,25 @@ int8_t llvm::offloading::binaryImageFormatToInt8(SYCLBinaryImageFormat Format) { } } +Expected +llvm::offloading::int8ToBinaryImageFormat(int8_t Format) { + switch (Format) { + case 0: + return SYCLBinaryImageFormat::BIF_None; + case 1: + return SYCLBinaryImageFormat::BIF_Native; + case 2: + return SYCLBinaryImageFormat::BIF_SPIRV; + case 3: + return SYCLBinaryImageFormat::BIF_LLVMBC; + case 4: + return SYCLBinaryImageFormat::BIF_Compressed; + default: + return createStringError("invalid SYCL image format '" + + Twine(static_cast(Format)) + "'"); + } +} + namespace { /// Wrapper helper class that creates all LLVM IRs wrapping given images. @@ -842,30 +861,14 @@ llvm::offloading::readSYCLImages(MemoryBufferRef Buffer) { Binary->getString(CompileOptionsKey), Binary->getString(LinkOptionsKey)); - unsigned Format = 0; + int8_t Format = 0; if (Binary->getString(FormatKey).getAsInteger(10, Format)) return createStringError("invalid SYCL image format '" + Binary->getString(FormatKey) + "'"); - switch (Format) { - case 0: - Image.Format = SYCLBinaryImageFormat::BIF_None; - break; - case 1: - Image.Format = SYCLBinaryImageFormat::BIF_Native; - break; - case 2: - Image.Format = SYCLBinaryImageFormat::BIF_SPIRV; - break; - case 3: - Image.Format = SYCLBinaryImageFormat::BIF_LLVMBC; - break; - case 4: - Image.Format = SYCLBinaryImageFormat::BIF_Compressed; - break; - default: - return createStringError("invalid SYCL image format '" + Twine(Format) + - "'"); - } + auto FormatOrErr = int8ToBinaryImageFormat(Format); + if (!FormatOrErr) + return FormatOrErr.takeError(); + Image.Format = *FormatOrErr; Images.emplace_back(std::move(Image)); } return std::move(Images);