Skip to content

Commit 2abe0d6

Browse files
author
Konrad Kusiak
authored
[SYCL][Graph] Implement dynamic_work_group_memory for SYCL-Graphs (#17314)
Implements dynamic_work_group_memory for SYCL-Graphs: #16712 With this PR we're able to update work_group_memory size on subsequent graph executions. We're also now able to use dynamic_work_group_memory with both lambdas and free function kernels.
1 parent 2533518 commit 2abe0d6

30 files changed

+871
-19
lines changed

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

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1606,12 +1606,12 @@ def SYCLType: InheritableAttr {
16061606
let Subjects = SubjectList<[CXXRecord, Enum], ErrorDiag>;
16071607
let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost];
16081608
let Args = [EnumArgument<"Type", "SYCLType", /*is_string=*/true,
1609-
["accessor", "local_accessor", "work_group_memory",
1609+
["accessor", "local_accessor", "work_group_memory", "dynamic_work_group_memory",
16101610
"specialization_id", "kernel_handler", "buffer_location",
16111611
"no_alias", "accessor_property_list", "group",
16121612
"private_memory", "aspect", "annotated_ptr", "annotated_arg",
16131613
"stream", "sampler", "host_pipe", "multi_ptr"],
1614-
["accessor", "local_accessor", "work_group_memory",
1614+
["accessor", "local_accessor", "work_group_memory", "dynamic_work_group_memory",
16151615
"specialization_id", "kernel_handler", "buffer_location",
16161616
"no_alias", "accessor_property_list", "group",
16171617
"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
@@ -63,7 +63,8 @@ class SYCLIntegrationHeader {
6363
kind_specialization_constants_buffer,
6464
kind_stream,
6565
kind_work_group_memory,
66-
kind_last = kind_work_group_memory
66+
kind_dynamic_work_group_memory,
67+
kind_last = kind_dynamic_work_group_memory
6768
};
6869

6970
public:

‎clang/lib/Sema/SemaSYCL.cpp‎

Lines changed: 24 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -2090,7 +2090,9 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler {
20902090
}
20912091

20922092
bool handleSyclSpecialType(ParmVarDecl *PD, QualType ParamTy) final {
2093-
if (!SemaSYCL::isSyclType(ParamTy, SYCLTypeAttr::work_group_memory)) {
2093+
if (!SemaSYCL::isSyclType(ParamTy, SYCLTypeAttr::work_group_memory) &&
2094+
!SemaSYCL::isSyclType(ParamTy,
2095+
SYCLTypeAttr::dynamic_work_group_memory)) {
20942096
Diag.Report(PD->getLocation(), diag::err_bad_kernel_param_type)
20952097
<< ParamTy;
20962098
IsInvalid = true;
@@ -2246,7 +2248,8 @@ class SyclKernelUnionChecker : public SyclKernelFieldHandler {
22462248
}
22472249

22482250
bool handleSyclSpecialType(ParmVarDecl *PD, QualType ParamTy) final {
2249-
if (!SemaSYCL::isSyclType(ParamTy, SYCLTypeAttr::work_group_memory))
2251+
if (!SemaSYCL::isSyclType(ParamTy, SYCLTypeAttr::work_group_memory) &&
2252+
!SemaSYCL::isSyclType(ParamTy, SYCLTypeAttr::dynamic_work_group_memory))
22502253
unsupportedFreeFunctionParamType(); // TODO
22512254
return true;
22522255
}
@@ -3032,7 +3035,9 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
30323035
}
30333036

30343037
bool handleSyclSpecialType(ParmVarDecl *PD, QualType ParamTy) final {
3035-
if (SemaSYCL::isSyclType(ParamTy, SYCLTypeAttr::work_group_memory)) {
3038+
if (SemaSYCL::isSyclType(ParamTy, SYCLTypeAttr::work_group_memory) ||
3039+
SemaSYCL::isSyclType(ParamTy,
3040+
SYCLTypeAttr::dynamic_work_group_memory)) {
30363041
const auto *RecordDecl = ParamTy->getAsCXXRecordDecl();
30373042
assert(RecordDecl && "The type must be a RecordDecl");
30383043
CXXMethodDecl *InitMethod = getMethodByName(RecordDecl, InitMethodName);
@@ -4544,7 +4549,9 @@ class FreeFunctionKernelBodyCreator : public SyclKernelFieldHandler {
45444549
// TODO: Revisit this approach once https://github.com/intel/llvm/issues/16061
45454550
// is closed.
45464551
bool handleSyclSpecialType(ParmVarDecl *PD, QualType ParamTy) final {
4547-
if (SemaSYCL::isSyclType(ParamTy, SYCLTypeAttr::work_group_memory)) {
4552+
if (SemaSYCL::isSyclType(ParamTy, SYCLTypeAttr::work_group_memory) ||
4553+
SemaSYCL::isSyclType(ParamTy,
4554+
SYCLTypeAttr::dynamic_work_group_memory)) {
45484555
const auto *RecordDecl = ParamTy->getAsCXXRecordDecl();
45494556
AccessSpecifier DefaultConstructorAccess;
45504557
auto DefaultConstructor =
@@ -4823,6 +4830,10 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
48234830
} else if (SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::work_group_memory)) {
48244831
addParam(FieldTy, SYCLIntegrationHeader::kind_work_group_memory,
48254832
offsetOf(RD, BC.getType()->getAsCXXRecordDecl()));
4833+
} else if (SemaSYCL::isSyclType(FieldTy,
4834+
SYCLTypeAttr::dynamic_work_group_memory)) {
4835+
addParam(FieldTy, SYCLIntegrationHeader::kind_dynamic_work_group_memory,
4836+
offsetOf(RD, BC.getType()->getAsCXXRecordDecl()));
48264837
}
48274838
return true;
48284839
}
@@ -4846,6 +4857,10 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
48464857
} else if (SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::work_group_memory)) {
48474858
addParam(FieldTy, SYCLIntegrationHeader::kind_work_group_memory,
48484859
offsetOf(FD, FieldTy));
4860+
} else if (SemaSYCL::isSyclType(FieldTy,
4861+
SYCLTypeAttr::dynamic_work_group_memory)) {
4862+
addParam(FieldTy, SYCLIntegrationHeader::kind_dynamic_work_group_memory,
4863+
offsetOf(FD, FieldTy));
48494864
} else if (SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::sampler) ||
48504865
SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::annotated_ptr) ||
48514866
SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::annotated_arg)) {
@@ -4870,6 +4885,10 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
48704885
bool handleSyclSpecialType(ParmVarDecl *PD, QualType ParamTy) final {
48714886
if (SemaSYCL::isSyclType(ParamTy, SYCLTypeAttr::work_group_memory))
48724887
addParam(PD, ParamTy, SYCLIntegrationHeader::kind_work_group_memory);
4888+
else if (SemaSYCL::isSyclType(ParamTy,
4889+
SYCLTypeAttr::dynamic_work_group_memory))
4890+
addParam(PD, ParamTy,
4891+
SYCLIntegrationHeader::kind_dynamic_work_group_memory);
48734892
else
48744893
unsupportedFreeFunctionParamType(); // TODO
48754894
return true;
@@ -5993,6 +6012,7 @@ static const char *paramKind2Str(KernelParamKind K) {
59936012
CASE(specialization_constants_buffer);
59946013
CASE(pointer);
59956014
CASE(work_group_memory);
6015+
CASE(dynamic_work_group_memory);
59966016
}
59976017
return "<ERROR>";
59986018

‎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/Inputs/sycl/detail/kernel_desc.hpp‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -19,6 +19,7 @@ namespace detail {
1919
kind_specialization_constants_buffer = 4,
2020
kind_stream = 5,
2121
kind_work_group_memory = 6,
22+
kind_dynamic_work_group_memory = 7,
2223
kind_invalid = 0xf, // not a valid kernel kind
2324
};
2425

‎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)