Skip to content

Commit 524da07

Browse files
maksimsabmaksimsab
committed
[SYCL] Add offload wrapping for SYCL kind.
1 parent 68309ad commit 524da07

File tree

11 files changed

+675
-5
lines changed

11 files changed

+675
-5
lines changed

‎clang/test/Driver/linker-wrapper-image.c‎

Lines changed: 35 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1,6 +1,7 @@
11
// REQUIRES: x86-registered-target
22
// REQUIRES: nvptx-registered-target
33
// REQUIRES: amdgpu-registered-target
4+
// REQUIRES: spirv-registered-target
45

56
// RUN: %clang -cc1 %s -triple x86_64-unknown-linux-gnu -emit-obj -o %t.elf.o
67

@@ -263,3 +264,37 @@
263264
// HIP: while.end:
264265
// HIP-NEXT: ret void
265266
// HIP-NEXT: }
267+
268+
// RUN: clang-offload-packager -o %t.out --image=file=%t.elf.o,kind=sycl,triple=spirv64-unknown-unknown,arch=generic
269+
// RUN: %clang -cc1 %s -triple x86_64-unknown-linux-gnu -emit-obj -o %t.o \
270+
// RUN: -fembed-offload-object=%t.out
271+
// RUN: clang-linker-wrapper --print-wrapped-module --dry-run --host-triple=x86_64-unknown-linux-gnu \
272+
// RUN: --linker-path=/usr/bin/ld %t.o -o a.out 2>&1 | FileCheck %s --check-prefixes=SYCL
273+
// RUN: clang-linker-wrapper --print-wrapped-module --dry-run --host-triple=x86_64-unknown-linux-gnu -r \
274+
// RUN: --linker-path=/usr/bin/ld %t.o -o a.out 2>&1 | FileCheck %s --check-prefixes=SYCL
275+
276+
// SYCL: %__sycl.tgt_device_image = type { i16, i8, i8, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr }
277+
// SYCL-NEXT: %__sycl.tgt_bin_desc = type { i16, i16, ptr, ptr, ptr }
278+
279+
// SYCL: @.sycl_offloading.target.0 = internal unnamed_addr constant [1 x i8] zeroinitializer
280+
// SYCL-NEXT: @.sycl_offloading.opts.compile.0 = internal unnamed_addr constant [1 x i8] zeroinitializer
281+
// SYCL-NEXT: @.sycl_offloading.opts.link.0 = internal unnamed_addr constant [1 x i8] zeroinitializer
282+
// SYCL-NEXT: @.sycl_offloading.0.data = internal unnamed_addr constant [0 x i8] zeroinitializer
283+
// SYCL-NEXT: @.sycl_offloading.0.info = internal local_unnamed_addr constant [2 x i64] [i64 ptrtoint (ptr @.sycl_offloading.0.data to i64), i64 0], section ".tgtimg", align 16
284+
// SYCL-NEXT: @llvm.used = appending global [1 x ptr] [ptr @.sycl_offloading.0.info], section "llvm.metadata"
285+
// SYCL-NEXT: @.sycl_offloading.device_images = internal unnamed_addr constant [1 x %__sycl.tgt_device_image] [%__sycl.tgt_device_image { i16 3, i8 8, i8 0, ptr @.sycl_offloading.target.0, ptr @.sycl_offloading.opts.compile.0, ptr @.sycl_offloading.opts.link.0, ptr @.sycl_offloading.0.data, ptr @.sycl_offloading.0.data, ptr null, ptr null, ptr null, ptr null }]
286+
// SYCL-NEXT: @.sycl_offloading.descriptor = internal constant %__sycl.tgt_bin_desc { i16 1, i16 1, ptr @.sycl_offloading.device_images, ptr null, ptr null }
287+
// SYCL-NEXT: @llvm.global_ctors = appending global [1 x { i32, ptr, ptr }] [{ i32, ptr, ptr } { i32 1, ptr @sycl.descriptor_reg, ptr null }]
288+
// SYCL-NEXT: @llvm.global_dtors = appending global [1 x { i32, ptr, ptr }] [{ i32, ptr, ptr } { i32 1, ptr @sycl.descriptor_unreg, ptr null }]
289+
290+
// SYCL: define internal void @sycl.descriptor_reg() section ".text.startup" {
291+
// SYCL-NEXT: entry:
292+
// SYCL-NEXT: call void @__sycl_register_lib(ptr @.sycl_offloading.descriptor)
293+
// SYCL-NEXT: ret void
294+
// SYCL-NEXT: }
295+
296+
// SYCL: define internal void @sycl.descriptor_unreg() section ".text.startup" {
297+
// SYCL-NEXT: entry:
298+
// SYCL-NEXT: call void @__sycl_unregister_lib(ptr @.sycl_offloading.descriptor)
299+
// SYCL-NEXT: ret void
300+
// SYCL-NEXT: }

‎clang/test/Driver/linker-wrapper.c‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -54,7 +54,7 @@ __attribute__((visibility("protected"), used)) int x;
5454
// RUN: clang-offload-packager -o %t.out \
5555
// RUN: --image=file=%t.spirv.bc,kind=sycl,triple=spirv64-unknown-unknown,arch=generic
5656
// RUN: %clang -cc1 %s -triple x86_64-unknown-linux-gnu -emit-obj -o %t.o -fembed-offload-object=%t.out
57-
// RUN: not clang-linker-wrapper --host-triple=x86_64-unknown-linux-gnu --dry-run \
57+
// RUN: clang-linker-wrapper --host-triple=x86_64-unknown-linux-gnu --dry-run \
5858
// RUN: --linker-path=/usr/bin/ld %t.o -o a.out 2>&1 | FileCheck %s --check-prefix=SPIRV-LINK
5959

6060
// SPIRV-LINK: clang{{.*}} -o {{.*}}.img --target=spirv64-unknown-unknown {{.*}}.o --sycl-link -Xlinker -triple=spirv64-unknown-unknown -Xlinker -arch=

‎clang/tools/clang-linker-wrapper/CMakeLists.txt‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -16,6 +16,7 @@ set(LLVM_LINK_COMPONENTS
1616
CodeGen
1717
LTO
1818
FrontendOffloading
19+
FrontendSYCL
1920
)
2021

2122
set(LLVM_TARGET_DEFINITIONS LinkerWrapperOpts.td)

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

Lines changed: 40 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -22,6 +22,7 @@
2222
#include "llvm/CodeGen/CommandFlags.h"
2323
#include "llvm/Frontend/Offloading/OffloadWrapper.h"
2424
#include "llvm/Frontend/Offloading/Utility.h"
25+
#include "llvm/Frontend/SYCL/OffloadWrapper.h"
2526
#include "llvm/IR/Constants.h"
2627
#include "llvm/IR/DiagnosticPrinter.h"
2728
#include "llvm/IR/Module.h"
@@ -711,6 +712,13 @@ wrapDeviceImages(ArrayRef<std::unique_ptr<MemoryBuffer>> Buffers,
711712
M, BuffersToWrap.front(), offloading::getOffloadEntryArray(M)))
712713
return std::move(Err);
713714
break;
715+
case OFK_SYCL: {
716+
offloading::sycl::SYCLWrappingOptions WrappingOptions;
717+
if (Error Err = offloading::sycl::wrapSYCLBinaries(M, BuffersToWrap,
718+
WrappingOptions))
719+
return Err;
720+
break;
721+
}
714722
default:
715723
return createStringError(getOffloadKindName(Kind) +
716724
" wrapping is not supported");
@@ -748,6 +756,36 @@ bundleOpenMP(ArrayRef<OffloadingImage> Images) {
748756
return std::move(Buffers);
749757
}
750758

759+
Expected<SmallVector<std::unique_ptr<MemoryBuffer>>>
760+
bundleSYCL(ArrayRef<OffloadingImage> Images) {
761+
SmallVector<std::unique_ptr<MemoryBuffer>> Buffers;
762+
if (DryRun) {
763+
// In dry-run mode there is an empty input which is insufficient for
764+
// the testing. Therefore, we insert a stub value.
765+
OffloadBinary::OffloadingImage Image;
766+
Image.TheOffloadKind = OffloadKind::OFK_SYCL;
767+
Image.Image = MemoryBuffer::getMemBufferCopy("");
768+
SmallString<0> SerializedImage = OffloadBinary::write(Image);
769+
Buffers.emplace_back(MemoryBuffer::getMemBufferCopy(SerializedImage));
770+
return Buffers;
771+
}
772+
773+
for (const OffloadingImage &TheImage : Images) {
774+
SmallVector<OffloadFile> OffloadBinaries;
775+
if (Error E = extractOffloadBinaries(*TheImage.Image, OffloadBinaries))
776+
return E;
777+
778+
for (const OffloadFile &File : OffloadBinaries) {
779+
const OffloadBinary &Binary = *File.getBinary();
780+
SmallString<0> SerializedImage =
781+
OffloadBinary::write(Binary.getOffloadingImage());
782+
Buffers.emplace_back(MemoryBuffer::getMemBufferCopy(SerializedImage));
783+
}
784+
}
785+
786+
return Buffers;
787+
}
788+
751789
Expected<SmallVector<std::unique_ptr<MemoryBuffer>>>
752790
bundleCuda(ArrayRef<OffloadingImage> Images, const ArgList &Args) {
753791
SmallVector<std::pair<StringRef, StringRef>, 4> InputFiles;
@@ -800,8 +838,9 @@ bundleLinkedOutput(ArrayRef<OffloadingImage> Images, const ArgList &Args,
800838
llvm::TimeTraceScope TimeScope("Bundle linked output");
801839
switch (Kind) {
802840
case OFK_OpenMP:
803-
case OFK_SYCL:
804841
return bundleOpenMP(Images);
842+
case OFK_SYCL:
843+
return bundleSYCL(Images);
805844
case OFK_Cuda:
806845
return bundleCuda(Images, Args);
807846
case OFK_HIP:

‎llvm/include/llvm/Frontend/SYCL/OffloadWrapper.h‎

Lines changed: 44 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,44 @@
1+
//===----- OffloadWrapper.h -------------------------------------*- C++ -*-===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
9+
#ifndef LLVM_FRONTEND_SYCL_OFFLOAD_WRAPPER_H
10+
#define LLVM_FRONTEND_SYCL_OFFLOAD_WRAPPER_H
11+
12+
#include "llvm/ADT/ArrayRef.h"
13+
#include "llvm/Object/OffloadBinary.h"
14+
15+
#include <string>
16+
17+
namespace llvm {
18+
19+
class Module;
20+
21+
namespace offloading {
22+
namespace sycl {
23+
24+
struct SYCLWrappingOptions {
25+
// target/compiler specific options what are suggested to use to "compile"
26+
// program at runtime.
27+
std::string CompileOptions;
28+
// Target/Compiler specific options that are suggested to use to "link"
29+
// program at runtime.
30+
std::string LinkOptions;
31+
};
32+
33+
/// Wraps OffloadBinaries in the given \p Buffers into the module \p M
34+
/// as global symbols and registers the images with the SYCL Runtime.
35+
/// \param Options Settings that allows to turn on optional data and settings.
36+
llvm::Error
37+
wrapSYCLBinaries(llvm::Module &M, llvm::ArrayRef<llvm::ArrayRef<char>> Buffers,
38+
SYCLWrappingOptions Options = SYCLWrappingOptions());
39+
40+
} // namespace sycl
41+
} // namespace offloading
42+
} // namespace llvm
43+
44+
#endif // LLVM_FRONTEND_SYCL_OFFLOAD_WRAPPER_H

‎llvm/include/llvm/Object/OffloadBinary.h‎

Lines changed: 6 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -48,6 +48,7 @@ enum ImageKind : uint16_t {
4848
IMG_Cubin,
4949
IMG_Fatbinary,
5050
IMG_PTX,
51+
IMG_SPIRV,
5152
IMG_LAST,
5253
};
5354

@@ -70,9 +71,9 @@ class OffloadBinary : public Binary {
7071

7172
/// The offloading metadata that will be serialized to a memory buffer.
7273
struct OffloadingImage {
73-
ImageKind TheImageKind;
74-
OffloadKind TheOffloadKind;
75-
uint32_t Flags;
74+
ImageKind TheImageKind = ImageKind::IMG_None;
75+
OffloadKind TheOffloadKind = OffloadKind::OFK_None;
76+
uint32_t Flags = 0;
7677
MapVector<StringRef, StringRef> StringData;
7778
std::unique_ptr<MemoryBuffer> Image;
7879
};
@@ -84,6 +85,8 @@ class OffloadBinary : public Binary {
8485
/// Serialize the contents of \p File to a binary buffer to be read later.
8586
LLVM_ABI static SmallString<0> write(const OffloadingImage &);
8687

88+
OffloadingImage getOffloadingImage() const;
89+
8790
static uint64_t getAlignment() { return 8; }
8891

8992
ImageKind getImageKind() const { return TheEntry->TheImageKind; }

‎llvm/lib/Frontend/CMakeLists.txt‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -5,3 +5,4 @@ add_subdirectory(HLSL)
55
add_subdirectory(OpenACC)
66
add_subdirectory(OpenMP)
77
add_subdirectory(Offloading)
8+
add_subdirectory(SYCL)

‎llvm/lib/Frontend/SYCL/CMakeLists.txt‎

Lines changed: 14 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,14 @@
1+
add_llvm_component_library(LLVMFrontendSYCL
2+
OffloadWrapper.cpp
3+
4+
ADDITIONAL_HEADER_DIRS
5+
${LLVM_MAIN_INCLUDE_DIR}/llvm/Frontend
6+
${LLVM_MAIN_INCLUDE_DIR}/llvm/Frontend/SYCL
7+
8+
LINK_COMPONENTS
9+
Core
10+
FrontendOffloading
11+
Object
12+
Support
13+
TransformUtils
14+
)

0 commit comments

Comments
 (0)