Skip to content

Commit 16f83ba

Browse files
author
Konrad Kusiak
committed
Added clang tests
1 parent 70fee11 commit 16f83ba

File tree

6 files changed

+142
-7
lines changed

6 files changed

+142
-7
lines changed

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

Lines changed: 13 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -667,6 +667,19 @@ __SYCL_TYPE(work_group_memory) work_group_memory {
667667
__attribute((opencl_local)) DataT *Ptr;
668668
};
669669

670+
template <typename DataT>
671+
class __attribute__((sycl_special_class))
672+
__SYCL_TYPE(dynamic_work_group_memory) dynamic_work_group_memory {
673+
public:
674+
dynamic_work_group_memory() = default;
675+
676+
void __init(__attribute((opencl_local)) DataT *Ptr) { this->LocalMem.__init(Ptr); }
677+
work_group_memory<DataT> get() const { return LocalMem; }
678+
679+
private:
680+
work_group_memory<DataT> LocalMem;
681+
};
682+
670683
template <typename T, int dimensions = 1,
671684
typename AllocatorT = int /*fake type as AllocatorT is not used*/>
672685
class buffer {

‎clang/test/CodeGenSYCL/dynamic_work_group_memory.cpp‎

Lines changed: 39 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,39 @@
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 dynamic_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_dynamic_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+
sycl::dynamic_work_group_memory<int> dynMem;
32+
Q.submit([&](sycl::handler &CGH) {
33+
sycl::range<1> ndr;
34+
CGH.parallel_for(ndr, [=](sycl::item<1> it) {
35+
auto localMem = dynMem.get();
36+
int *ptr = &localMem; });
37+
});
38+
return 0;
39+
}

‎clang/test/CodeGenSYCL/free_function_int_header.cpp‎

Lines changed: 39 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -2,7 +2,8 @@
22
// RUN: FileCheck -input-file=%t.h %s
33
//
44
// This test checks integration header contents for free functions with scalar,
5-
// pointer, non-decomposed struct parameters and work group memory parameters.
5+
// pointer, non-decomposed struct parameters, work group memory parameters and
6+
// dynamic work group memory parameters.
67

78
#include "mock_properties.hpp"
89
#include "sycl.hpp"
@@ -101,6 +102,11 @@ __attribute__((sycl_device))
101102
void ff_8(sycl::work_group_memory<int>) {
102103
}
103104

105+
__attribute__((sycl_device))
106+
[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 0)]]
107+
void ff_9(sycl::dynamic_work_group_memory<int>) {
108+
}
109+
104110

105111
// CHECK: const char* const kernel_names[] = {
106112
// CHECK-NEXT: {{.*}}__sycl_kernel_ff_2Piii
@@ -112,6 +118,7 @@ void ff_8(sycl::work_group_memory<int>) {
112118
// CHECK-NEXT: {{.*}}__sycl_kernel_ff_6I3Agg7DerivedEvT_T0_i
113119
// CHECK-NEXT: {{.*}}__sycl_kernel_ff_7ILi3EEv16KArgWithPtrArrayIXT_EE
114120
// CHECK-NEXT: {{.*}}__sycl_kernel_ff_8N4sycl3_V117work_group_memoryIiEE
121+
// CHECK-NEXT: {{.*}}__sycl_kernel_ff_9N4sycl3_V125dynamic_work_group_memoryIiEE
115122
// CHECK-NEXT: ""
116123
// CHECK-NEXT: };
117124

@@ -158,6 +165,9 @@ void ff_8(sycl::work_group_memory<int>) {
158165
// CHECK: //--- _Z18__sycl_kernel_ff_8N4sycl3_V117work_group_memoryIiEE
159166
// CHECK-NEXT: { kernel_param_kind_t::kind_work_group_memory, 8, 0 },
160167

168+
// CHECK: //--- _Z18__sycl_kernel_ff_9N4sycl3_V125dynamic_work_group_memoryIiEE
169+
// CHECK-NEXT: { kernel_param_kind_t::kind_dynamic_work_group_memory, 8, 0 },
170+
161171
// CHECK: { kernel_param_kind_t::kind_invalid, -987654321, -987654321 },
162172
// CHECK-NEXT: };
163173

@@ -324,6 +334,26 @@ void ff_8(sycl::work_group_memory<int>) {
324334
// CHECK-NEXT: };
325335
// CHECK-NEXT: }
326336

337+
// CHECK: // Definition of _Z18__sycl_kernel_ff_9N4sycl3_V125dynamic_work_group_memoryIiEE as a free function kernel
338+
//
339+
// CHECK: Forward declarations of kernel and its argument types:
340+
// CHECK: template <typename DataT> class dynamic_work_group_memory;
341+
342+
// CHECK: void ff_9(sycl::dynamic_work_group_memory<int>);
343+
// CHECK-NEXT: static constexpr auto __sycl_shim10() {
344+
// CHECK-NEXT: return (void (*)(class sycl::dynamic_work_group_memory<int>))ff_9;
345+
// CHECK-NEXT: }
346+
// CHECK-NEXT: namespace sycl {
347+
// CHECK-NEXT: template <>
348+
// CHECK-NEXT: struct ext::oneapi::experimental::is_kernel<__sycl_shim10()> {
349+
// CHECK-NEXT: static constexpr bool value = true;
350+
// CHECK-NEXT: };
351+
// CHECK-NEXT: template <>
352+
// CHECK-NEXT: struct ext::oneapi::experimental::is_single_task_kernel<__sycl_shim10()> {
353+
// CHECK-NEXT: static constexpr bool value = true;
354+
// CHECK-NEXT: };
355+
// CHECK-NEXT: }
356+
327357
// CHECK: #include <sycl/kernel_bundle.hpp>
328358

329359
// CHECK: Definition of kernel_id of _Z18__sycl_kernel_ff_2Piii
@@ -397,3 +427,11 @@ void ff_8(sycl::work_group_memory<int>) {
397427
// CHECK-NEXT: return sycl::detail::get_kernel_id_impl(std::string_view{"_Z18__sycl_kernel_ff_8N4sycl3_V117work_group_memoryIiEE"});
398428
// CHECK-NEXT: }
399429
// CHECK-NEXT: }
430+
//
431+
// CHECK: // Definition of kernel_id of _Z18__sycl_kernel_ff_9N4sycl3_V125dynamic_work_group_memoryIiEE
432+
// CHECK-NEXT: namespace sycl {
433+
// CHECK-NEXT: template <>
434+
// CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim10()>() {
435+
// CHECK-NEXT: return sycl::detail::get_kernel_id_impl(std::string_view{"_Z18__sycl_kernel_ff_9N4sycl3_V125dynamic_work_group_memoryIiEE"});
436+
// CHECK-NEXT: }
437+
// CHECK-NEXT: }

‎clang/test/CodeGenSYCL/free_function_kernel_params.cpp‎

Lines changed: 14 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,7 @@
11
// RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -triple spir64 \
22
// RUN: -emit-llvm %s -o - | FileCheck %s
33
// This test checks parameter IR generation for free functions with parameters
4-
// of non-decomposed struct type and work group memory type.
4+
// of non-decomposed struct type, work group memory type and dynamic work group memory type.
55

66
#include "sycl.hpp"
77

@@ -71,3 +71,16 @@ void ff_7(sycl::work_group_memory<int> mem) {
7171
// CHECK-NEXT: [[REGISTER:%[a-zA-Z0-9_]+]] = load ptr addrspace(3), ptr addrspace(4) %__arg_Ptr.addr.ascast, align 8
7272
// CHECK-NEXT: call spir_func void @{{.*}}work_group_memory{{.*}}__init{{.*}}(ptr addrspace(4) noundef align 8 dereferenceable_or_null(8) %mem.ascast, ptr addrspace(3) noundef [[REGISTER]])
7373

74+
__attribute__((sycl_device))
75+
[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 0)]]
76+
void ff_7(sycl::dynamic_work_group_memory<int> DynMem) {
77+
}
78+
79+
// CHECK: define dso_local spir_kernel void @{{.*}}__sycl_kernel_ff_7{{.*}}(ptr addrspace(3) noundef align 4 %__arg_Ptr)
80+
// CHECK: %__arg_Ptr.addr = alloca ptr addrspace(3), align 8
81+
// CHECK-NEXT: %DynMem = alloca %"class.sycl::_V1::dynamic_work_group_memory", align 8
82+
// CHECK: %__arg_Ptr.addr.ascast = addrspacecast ptr %__arg_Ptr.addr to ptr addrspace(4)
83+
// CHECK-NEXT: %DynMem.ascast = addrspacecast ptr %DynMem to ptr addrspace(4)
84+
// CHECK: store ptr addrspace(3) %__arg_Ptr, ptr addrspace(4) %__arg_Ptr.addr.ascast, align 8
85+
// CHECK-NEXT: [[REGISTER:%[a-zA-Z0-9_]+]] = load ptr addrspace(3), ptr addrspace(4) %__arg_Ptr.addr.ascast, align 8
86+
// CHECK-NEXT: call spir_func void @{{.*}}dynamic_work_group_memory{{.*}}__init{{.*}}(ptr addrspace(4) noundef align 8 dereferenceable_or_null(8) %DynMem.ascast, ptr addrspace(3) noundef [[REGISTER]])

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

Lines changed: 15 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -452,11 +452,9 @@ class __SYCL_TYPE(multi_ptr) multi_ptr<T, AS, access::decorated::legacy> {
452452
template <typename DataT>
453453
class __attribute__((sycl_special_class))
454454
__SYCL_TYPE(work_group_memory) work_group_memory {
455-
456-
// Default constructor for objects later initialized with __init member.
457-
work_group_memory() = default;
458-
459455
public:
456+
// Default constructor for objects later initialized with __init member.
457+
work_group_memory() = default;
460458
work_group_memory(handler &CGH) {}
461459

462460
void __init(__attribute((opencl_local)) DataT *Ptr) { this->Ptr = Ptr; }
@@ -465,6 +463,19 @@ __SYCL_TYPE(work_group_memory) work_group_memory {
465463
__attribute((opencl_local)) DataT *Ptr;
466464
};
467465

466+
template <typename DataT>
467+
class __attribute__((sycl_special_class))
468+
__SYCL_TYPE(dynamic_work_group_memory) dynamic_work_group_memory {
469+
public:
470+
dynamic_work_group_memory() = default;
471+
472+
void __init(__attribute((opencl_local)) DataT *Ptr) { this->LocalMem.__init(Ptr); }
473+
work_group_memory<DataT> get() const { return LocalMem; }
474+
475+
private:
476+
work_group_memory<DataT> LocalMem;
477+
};
478+
468479
namespace ext {
469480
namespace oneapi {
470481
namespace experimental {

‎clang/test/SemaSYCL/free_function_kernel_params.cpp‎

Lines changed: 22 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,8 @@
11
// RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -ast-dump \
22
// RUN: %s -o - | FileCheck %s
33
// This test checks parameter rewriting for free functions with parameters
4-
// of type scalar, pointer, non-decomposed struct and work group memory.
4+
// of type scalar, pointer, non-decomposed struct, work group memory and
5+
// dynamic work group memory.
56

67
#include "sycl.hpp"
78

@@ -191,3 +192,23 @@ void ff_7(sycl::work_group_memory<int> mem) {
191192
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'void (*)(sycl::work_group_memory<int>)' <FunctionToPointerDecay>
192193
// CHECK-NEXT: DeclRefExpr {{.*}} 'void (sycl::work_group_memory<int>)' lvalue Function {{.*}} 'ff_7' 'void (sycl::work_group_memory<int>)'
193194
// CHECK-NEXT: DeclRefExpr {{.*}} 'sycl::work_group_memory<int>' Var {{.*}} 'mem' 'sycl::work_group_memory<int>'
195+
196+
__attribute__((sycl_device))
197+
[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 0)]]
198+
void ff_8(sycl::dynamic_work_group_memory<int> DynMem) {
199+
}
200+
// CHECK: FunctionDecl {{.*}}__sycl_kernel{{.*}}'void (__local int *)'
201+
// CHECK-NEXT: ParmVarDecl {{.*}} used __arg_Ptr '__local int *'
202+
// CHECK-NEXT: CompoundStmt
203+
// CHECK-NEXT: DeclStmt
204+
// CHECK-NEXT: VarDecl {{.*}} used DynMem 'sycl::dynamic_work_group_memory<int>' callinit
205+
// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::dynamic_work_group_memory<int>' 'void () noexcept'
206+
// CHECK-NEXT: CXXMemberCallExpr {{.*}} 'void'
207+
// CHECK-NEXT: MemberExpr {{.*}} 'void (__local int *)' lvalue .__init
208+
// CHECK-NEXT: DeclRefExpr {{.*}} 'sycl::dynamic_work_group_memory<int>' Var {{.*}} 'DynMem' 'sycl::dynamic_work_group_memory<int>'
209+
// CHECK-NEXT: ImplicitCastExpr {{.*}} '__local int *' <LValueToRValue>
210+
// CHECK-NEXT: DeclRefExpr {{.*}} '__local int *' lvalue ParmVar {{.*}} '__arg_Ptr' '__local int *'
211+
// CHECK-NEXT: CallExpr {{.*}} 'void'
212+
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'void (*)(sycl::dynamic_work_group_memory<int>)' <FunctionToPointerDecay>
213+
// CHECK-NEXT: DeclRefExpr {{.*}} 'void (sycl::dynamic_work_group_memory<int>)' lvalue Function {{.*}} 'ff_8' 'void (sycl::dynamic_work_group_memory<int>)'
214+
// CHECK-NEXT: DeclRefExpr {{.*}} 'sycl::dynamic_work_group_memory<int>' Var {{.*}} 'DynMem' 'sycl::dynamic_work_group_memory<int>'

0 commit comments

Comments
 (0)