Skip to content

Commit 17d409e

Browse files
ayylolayylol
committed
Merge branch 'sycl' into any-device-is
2 parents 8fb587e + 439ec9e commit 17d409e

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

57 files changed

+1502
-37
lines changed

‎.github/CODEOWNERS‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -36,6 +36,7 @@ sycl/doc/extensions/ @intel/dpcpp-specification-reviewers
3636

3737
# Unified Runtime
3838
sycl/cmake/modules/FetchUnifiedRuntime.cmake @intel/unified-runtime-reviewers
39+
sycl/cmake/modules/UnifiedRuntimeTag.cmake @intel/unified-runtime-reviewers
3940
sycl/include/sycl/detail/ur.hpp @intel/unified-runtime-reviewers
4041
sycl/source/detail/posix_ur.cpp @intel/unified-runtime-reviewers
4142
sycl/source/detail/ur.cpp @intel/unified-runtime-reviewers

‎buildbot/configure.py‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -177,6 +177,8 @@ def do_configure(args):
177177
"-DLLVM_ENABLE_PROJECTS={}".format(llvm_enable_projects),
178178
"-DSYCL_BUILD_PI_HIP_PLATFORM={}".format(sycl_build_pi_hip_platform),
179179
"-DLLVM_BUILD_TOOLS=ON",
180+
"-DLLVM_ENABLE_ZSTD=ON",
181+
"-DLLVM_USE_STATIC_ZSTD=ON",
180182
"-DSYCL_ENABLE_WERROR={}".format(sycl_werror),
181183
"-DCMAKE_INSTALL_PREFIX={}".format(install_dir),
182184
"-DSYCL_INCLUDE_TESTS=ON", # Explicitly include all kinds of SYCL tests.

‎clang/include/clang/Basic/Attr.td‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1566,12 +1566,12 @@ def SYCLType: InheritableAttr {
15661566
let Subjects = SubjectList<[CXXRecord, Enum], ErrorDiag>;
15671567
let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost];
15681568
let Args = [EnumArgument<"Type", "SYCLType", /*is_string=*/true,
1569-
["accessor", "local_accessor",
1569+
["accessor", "local_accessor", "work_group_memory",
15701570
"specialization_id", "kernel_handler", "buffer_location",
15711571
"no_alias", "accessor_property_list", "group",
15721572
"private_memory", "aspect", "annotated_ptr", "annotated_arg",
15731573
"stream", "sampler", "host_pipe", "multi_ptr"],
1574-
["accessor", "local_accessor",
1574+
["accessor", "local_accessor", "work_group_memory",
15751575
"specialization_id", "kernel_handler", "buffer_location",
15761576
"no_alias", "accessor_property_list", "group",
15771577
"private_memory", "aspect", "annotated_ptr", "annotated_arg",

‎clang/include/clang/Sema/SemaSYCL.h‎

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -62,7 +62,8 @@ class SYCLIntegrationHeader {
6262
kind_pointer,
6363
kind_specialization_constants_buffer,
6464
kind_stream,
65-
kind_last = kind_stream
65+
kind_work_group_memory,
66+
kind_last = kind_work_group_memory
6667
};
6768

6869
public:

‎clang/lib/Driver/ToolChains/Clang.cpp‎

Lines changed: 13 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -10174,6 +10174,19 @@ void OffloadWrapper::ConstructJob(Compilation &C, const JobAction &JA,
1017410174
SmallString<128> TargetTripleOpt = TT.getArchName();
1017510175
bool WrapFPGADevice = false;
1017610176
bool FPGAEarly = false;
10177+
10178+
// Validate and propogate CLI options related to device image compression.
10179+
// -offload-compress
10180+
if (C.getInputArgs().getLastArg(options::OPT_offload_compress)) {
10181+
WrapperArgs.push_back(
10182+
C.getArgs().MakeArgString(Twine("-offload-compress")));
10183+
// -offload-compression-level=<>
10184+
if (Arg *A = C.getInputArgs().getLastArg(
10185+
options::OPT_offload_compression_level_EQ))
10186+
WrapperArgs.push_back(C.getArgs().MakeArgString(
10187+
Twine("-offload-compression-level=") + A->getValue()));
10188+
}
10189+
1017710190
if (Arg *A = C.getInputArgs().getLastArg(options::OPT_fsycl_link_EQ)) {
1017810191
WrapFPGADevice = true;
1017910192
FPGAEarly = (A->getValue() == StringRef("early"));

‎clang/lib/Sema/SemaSYCL.cpp‎

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -4693,6 +4693,9 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
46934693
CurOffset + offsetOf(FD, FieldTy));
46944694
} else if (SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::stream)) {
46954695
addParam(FD, FieldTy, SYCLIntegrationHeader::kind_stream);
4696+
} else if (SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::work_group_memory)) {
4697+
addParam(FieldTy, SYCLIntegrationHeader::kind_work_group_memory,
4698+
offsetOf(FD, FieldTy));
46964699
} else if (SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::sampler) ||
46974700
SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::annotated_ptr) ||
46984701
SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::annotated_arg)) {
@@ -5773,6 +5776,7 @@ static const char *paramKind2Str(KernelParamKind K) {
57735776
CASE(stream);
57745777
CASE(specialization_constants_buffer);
57755778
CASE(pointer);
5779+
CASE(work_group_memory);
57765780
}
57775781
return "<ERROR>";
57785782

‎clang/test/CodeGenSYCL/Inputs/sycl.hpp‎

Lines changed: 18 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -649,6 +649,24 @@ const stream& operator<<(const stream &S, T&&) {
649649
return S;
650650
}
651651

652+
// Dummy implementation of work_group_memory for use in CodeGenSYCL tests.
653+
template <typename DataT>
654+
class __attribute__((sycl_special_class))
655+
__SYCL_TYPE(work_group_memory) work_group_memory {
656+
public:
657+
work_group_memory(handler &CGH) {}
658+
#ifdef __SYCL_DEVICE_ONLY__
659+
// Default constructor for objects later initialized with __init member.
660+
work_group_memory() = default;
661+
#endif
662+
663+
void __init(__attribute((opencl_local)) DataT *Ptr) { this->Ptr = Ptr; }
664+
__attribute((opencl_local)) DataT *operator&() const { return Ptr; }
665+
666+
private:
667+
__attribute((opencl_local)) DataT *Ptr;
668+
};
669+
652670
template <typename T, int dimensions = 1,
653671
typename AllocatorT = int /*fake type as AllocatorT is not used*/>
654672
class buffer {

‎clang/test/CodeGenSYCL/work_group_memory.cpp‎

Lines changed: 37 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,37 @@
1+
// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm %s -o %t.ll
2+
// RUN: FileCheck < %t.ll %s --check-prefix CHECK-IR
3+
// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -fsycl-int-header=%t.h %s
4+
// RUN: FileCheck < %t.h %s --check-prefix CHECK-INT-HEADER
5+
//
6+
// Tests for work_group_memory kernel parameter using the dummy implementation in Inputs/sycl.hpp.
7+
// The first two RUN commands verify that the init call is generated with the correct arguments in LLVM IR
8+
// and the second two RUN commands verify the contents of the integration header produced by the frontend.
9+
//
10+
// CHECK-IR: define dso_local spir_kernel void @
11+
// CHECK-IR-SAME: ptr addrspace(3) noundef align 4 [[PTR:%[a-zA-Z0-9_]+]]
12+
//
13+
// CHECK-IR: [[PTR]].addr = alloca ptr addrspace(3), align 8
14+
// CHECK-IR: [[PTR]].addr.ascast = addrspacecast ptr [[PTR]].addr to ptr addrspace(4)
15+
// CHECK-IR: store ptr addrspace(3) [[PTR]], ptr addrspace(4) [[PTR]].addr.ascast, align 8
16+
// CHECK-IR: [[PTR_LOAD:%[a-zA-Z0-9_]+]] = load ptr addrspace(3), ptr addrspace(4) [[PTR]].addr.ascast, align 8
17+
//
18+
// CHECK-IR: call spir_func void @{{.*}}__init{{.*}}(ptr addrspace(4) noundef align 8 dereferenceable_or_null(8) %{{[a-zA-Z0-9_]+}}, ptr addrspace(3) noundef [[PTR_LOAD]])
19+
//
20+
// CHECK-INT-HEADER: const kernel_param_desc_t kernel_signatures[] = {
21+
// CHECK-INT-HEADER-NEXT: //--- _ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlNS0_4itemILi1EEEE_
22+
// CHECK-INT-HEADER-NEXT: { kernel_param_kind_t::kind_work_group_memory, {{[4,8]}}, 0 },
23+
// CHECK-INT-HEADER-EMPTY:
24+
// CHECK-INT-HEADER-NEXT: { kernel_param_kind_t::kind_invalid, -987654321, -987654321 },
25+
// CHECK-INT-HEADER-NEXT: };
26+
27+
#include "Inputs/sycl.hpp"
28+
29+
int main() {
30+
sycl::queue Q;
31+
Q.submit([&](sycl::handler &CGH) {
32+
sycl::work_group_memory<int> mem;
33+
sycl::range<1> ndr;
34+
CGH.parallel_for(ndr, [=](sycl::item<1> it) { int *ptr = &mem; });
35+
});
36+
return 0;
37+
}

‎clang/test/Driver/clang-offload-wrapper-zstd.c‎

Lines changed: 40 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,40 @@
1+
// REQUIRES: zstd && (system-windows || system-linux)
2+
3+
// clang-offload-wrapper compression test: checks that the wrapper can compress the device images.
4+
// Checks the '--offload-compress', '--offload-compression-level', and '--offload-compression-threshold'
5+
// CLI options.
6+
7+
// --- Prepare test data by creating the debice binary image.
8+
// RUN: echo -e -n 'device binary image1\n' > %t.bin
9+
// RUN: echo -e -n '[Category1]\nint_prop1=1|10\n[Category2]\nint_prop2=1|20\n' > %t.props
10+
// RUN: echo -e -n 'kernel1\nkernel2\n' > %t.sym
11+
// RUN: echo -e -n 'Manifest file - arbitrary data generated by the toolchain\n' > %t.mnf
12+
// RUN: echo '[Code|Properties|Symbols|Manifest]' > %t.img1
13+
// RUN: echo %t.bin"|"%t.props"|"%t.sym"|"%t.mnf >> %t.img1
14+
15+
///////////////////////////////////////////////////////
16+
// Compress the test image using clang-offload-wrapper.
17+
///////////////////////////////////////////////////////
18+
19+
// RUN: clang-offload-wrapper -kind=sycl -target=TARGET -batch %t.img1 -o %t.wrapped.bc -v \
20+
// RUN: --offload-compress --offload-compression-level=9 --offload-compression-threshold=0 \
21+
// RUN: 2>&1 | FileCheck %s --check-prefix=CHECK-COMPRESS
22+
23+
// CHECK-COMPRESS: [Compression] Original image size:
24+
// CHECK-COMPRESS: [Compression] Compressed image size:
25+
// CHECK-COMPRESS: [Compression] Compression level used: 9
26+
27+
///////////////////////////////////////////////////////////
28+
// Check that there is no compression when the threshold is set to a value higher than the image size
29+
// or '--offload-compress' is not set.
30+
///////////////////////////////////////////////////////////
31+
32+
// RUN: clang-offload-wrapper -kind=sycl -target=TARGET -batch %t.img1 -o %t.wrapped.bc -v \
33+
// RUN: --offload-compress --offload-compression-level=3 --offload-compression-threshold=1000 \
34+
// RUN: 2>&1 | FileCheck %s --check-prefix=CHECK-NO-COMPRESS
35+
36+
// RUN: clang-offload-wrapper -kind=sycl -target=TARGET -batch %t.img1 -o %t.wrapped.bc -v \
37+
// RUN: --offload-compression-level=3 --offload-compression-threshold=0 \
38+
// RUN: 2>&1 | FileCheck %s --check-prefix=CHECK-NO-COMPRESS
39+
40+
// CHECK-NO-COMPRESS-NOT: [Compression] Original image size:

‎clang/test/Driver/sycl-offload-wrapper-compression.cpp‎

Lines changed: 14 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,14 @@
1+
///
2+
/// Check if '--offload-compress' and '--offload-compression-level' CLI
3+
/// options are passed to the clang-offload-wrapper.
4+
///
5+
6+
// RUN: %clangxx -### -fsycl --offload-compress --offload-compression-level=3 %s 2>&1 | FileCheck %s --check-prefix=CHECK-COMPRESS
7+
// CHECK-COMPRESS: {{.*}}clang-offload-wrapper{{.*}}"-offload-compress"{{.*}}"-offload-compression-level=3"{{.*}}
8+
9+
// Make sure that the compression options are not passed when --offload-compress is not set.
10+
// RUN: %clangxx -### -fsycl %s 2>&1 | FileCheck %s --check-prefix=CHECK-NO-COMPRESS
11+
// RUN: %clangxx -### -fsycl --offload-compression-level=3 %s 2>&1 | FileCheck %s --check-prefix=CHECK-NO-COMPRESS
12+
13+
// CHECK-NO-COMPRESS-NOT: {{.*}}clang-offload-wrapper{{.*}}"-offload-compress"{{.*}}
14+
// CHECK-NO-COMPRESS-NOT: {{.*}}clang-offload-wrapper{{.*}}"-offload-compression-level=3"{{.*}}

0 commit comments

Comments
 (0)