Skip to content

Commit 998c97f

Browse files
author
Alexander Batashev
authored
[SYCL][FPGA] Align clang with new spec of accessor_property_list (#2447)
Signed-off-by: Mikhail Lychkov <[email protected]>
1 parent f253851 commit 998c97f

14 files changed

+125
-103
lines changed

‎clang/include/clang/Basic/DiagnosticSemaKinds.td‎

Lines changed: 4 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -11036,9 +11036,10 @@ def err_sycl_compiletime_property_duplication : Error<
1103611036
def err_sycl_invalid_property_list_param_number : Error<
1103711037
"%0 must have exactly one template parameter">;
1103811038
def err_sycl_invalid_accessor_property_template_param : Error<
11039-
"Fifth template parameter of the accessor must be of a property_list type">;
11040-
def err_sycl_invalid_property_list_template_param : Error<
11041-
"%select{property_list|property_list pack argument|buffer_location}0 "
11039+
"sixth template parameter of the accessor must be of accessor_property_list "
11040+
"type">;
11041+
def err_sycl_invalid_accessor_property_list_template_param : Error<
11042+
"%select{accessor_property_list|accessor_property_list pack argument|buffer_location}0 "
1104211043
"template parameter must be a "
1104311044
"%select{parameter pack|type|non-negative integer}1">;
1104411045
def warn_sycl_pass_by_value_deprecated

‎clang/lib/Sema/SemaSYCL.cpp‎

Lines changed: 40 additions & 30 deletions
Original file line numberDiff line numberDiff line change
@@ -82,8 +82,8 @@ class Util {
8282
static bool isSyclHalfType(const QualType &Ty);
8383

8484
/// Checks whether given clang type is a full specialization of the SYCL
85-
/// property_list class.
86-
static bool isPropertyListType(const QualType &Ty);
85+
/// accessor_property_list class.
86+
static bool isAccessorPropertyListType(const QualType &Ty);
8787

8888
/// Checks whether given clang type is a full specialization of the SYCL
8989
/// buffer_location class.
@@ -1194,29 +1194,31 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler {
11941194
return;
11951195
}
11961196
QualType PropListTy = PropList.getAsType();
1197-
if (!Util::isPropertyListType(PropListTy)) {
1197+
if (!Util::isAccessorPropertyListType(PropListTy)) {
11981198
SemaRef.Diag(Loc,
11991199
diag::err_sycl_invalid_accessor_property_template_param);
12001200
return;
12011201
}
1202-
const auto *PropListDecl =
1202+
const auto *AccPropListDecl =
12031203
cast<ClassTemplateSpecializationDecl>(PropListTy->getAsRecordDecl());
1204-
if (PropListDecl->getTemplateArgs().size() != 1) {
1204+
if (AccPropListDecl->getTemplateArgs().size() != 1) {
12051205
SemaRef.Diag(Loc, diag::err_sycl_invalid_property_list_param_number)
1206-
<< "property_list";
1206+
<< "accessor_property_list";
12071207
return;
12081208
}
1209-
const auto TemplArg = PropListDecl->getTemplateArgs()[0];
1209+
const auto TemplArg = AccPropListDecl->getTemplateArgs()[0];
12101210
if (TemplArg.getKind() != TemplateArgument::ArgKind::Pack) {
1211-
SemaRef.Diag(Loc, diag::err_sycl_invalid_property_list_template_param)
1212-
<< /*property_list*/ 0 << /*parameter pack*/ 0;
1211+
SemaRef.Diag(Loc,
1212+
diag::err_sycl_invalid_accessor_property_list_template_param)
1213+
<< /*accessor_property_list*/ 0 << /*parameter pack*/ 0;
12131214
return;
12141215
}
12151216
for (TemplateArgument::pack_iterator Prop = TemplArg.pack_begin();
12161217
Prop != TemplArg.pack_end(); ++Prop) {
12171218
if (Prop->getKind() != TemplateArgument::ArgKind::Type) {
1218-
SemaRef.Diag(Loc, diag::err_sycl_invalid_property_list_template_param)
1219-
<< /*property_list pack argument*/ 1 << /*type*/ 1;
1219+
SemaRef.Diag(
1220+
Loc, diag::err_sycl_invalid_accessor_property_list_template_param)
1221+
<< /*accessor_property_list pack argument*/ 1 << /*type*/ 1;
12201222
return;
12211223
}
12221224
QualType PropTy = Prop->getAsType();
@@ -1235,13 +1237,15 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler {
12351237
}
12361238
const auto BufferLoc = PropDecl->getTemplateArgs()[0];
12371239
if (BufferLoc.getKind() != TemplateArgument::ArgKind::Integral) {
1238-
SemaRef.Diag(Loc, diag::err_sycl_invalid_property_list_template_param)
1240+
SemaRef.Diag(Loc,
1241+
diag::err_sycl_invalid_accessor_property_list_template_param)
12391242
<< /*buffer_location*/ 2 << /*non-negative integer*/ 2;
12401243
return;
12411244
}
12421245
int LocationID = static_cast<int>(BufferLoc.getAsIntegral().getExtValue());
12431246
if (LocationID < 0) {
1244-
SemaRef.Diag(Loc, diag::err_sycl_invalid_property_list_template_param)
1247+
SemaRef.Diag(Loc,
1248+
diag::err_sycl_invalid_accessor_property_list_template_param)
12451249
<< /*buffer_location*/ 2 << /*non-negative integer*/ 2;
12461250
return;
12471251
}
@@ -1414,19 +1418,18 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
14141418
}
14151419

14161420
// Handle accessor properties. If any properties were found in
1417-
// the property_list - add the appropriate attributes to ParmVarDecl.
1421+
// the accessor_property_list - add the appropriate attributes to ParmVarDecl.
14181422
void handleAccessorPropertyList(ParmVarDecl *Param,
14191423
const CXXRecordDecl *RecordDecl,
14201424
SourceLocation Loc) {
14211425
const auto *AccTy = cast<ClassTemplateSpecializationDecl>(RecordDecl);
1422-
// TODO: when SYCL headers' part is ready - replace this 'if' with an error
14231426
if (AccTy->getTemplateArgs().size() < 6)
14241427
return;
14251428
const auto PropList = cast<TemplateArgument>(AccTy->getTemplateArgs()[5]);
14261429
QualType PropListTy = PropList.getAsType();
1427-
const auto *PropListDecl =
1430+
const auto *AccPropListDecl =
14281431
cast<ClassTemplateSpecializationDecl>(PropListTy->getAsRecordDecl());
1429-
const auto TemplArg = PropListDecl->getTemplateArgs()[0];
1432+
const auto TemplArg = AccPropListDecl->getTemplateArgs()[0];
14301433
// Move through TemplateArgs list of a property list and search for
14311434
// properties. If found - apply the appropriate attribute to ParmVarDecl.
14321435
for (TemplateArgument::pack_iterator Prop = TemplArg.pack_begin();
@@ -3455,20 +3458,17 @@ bool Util::isSyclSpecConstantType(const QualType &Ty) {
34553458
return matchQualifiedTypeName(Ty, Scopes);
34563459
}
34573460

3458-
bool Util::isPropertyListType(const QualType &Ty) {
3459-
return isSyclType(Ty, "property_list", true /*Tmpl*/);
3460-
}
3461-
34623461
bool Util::isSyclBufferLocationType(const QualType &Ty) {
3463-
const StringRef &Name = "buffer_location";
3464-
std::array<DeclContextDesc, 4> Scopes = {
3465-
Util::DeclContextDesc{clang::Decl::Kind::Namespace, "cl"},
3466-
Util::DeclContextDesc{clang::Decl::Kind::Namespace, "sycl"},
3467-
// TODO: this doesn't belong to property namespace, instead it shall be
3468-
// in its own namespace. Change it, when the actual implementation in SYCL
3469-
// headers is ready
3470-
Util::DeclContextDesc{clang::Decl::Kind::Namespace, "property"},
3471-
Util::DeclContextDesc{Decl::Kind::ClassTemplateSpecialization, Name}};
3462+
const StringRef &PropertyName = "buffer_location";
3463+
const StringRef &InstanceName = "instance";
3464+
std::array<DeclContextDesc, 6> Scopes = {
3465+
Util::DeclContextDesc{Decl::Kind::Namespace, "cl"},
3466+
Util::DeclContextDesc{Decl::Kind::Namespace, "sycl"},
3467+
Util::DeclContextDesc{Decl::Kind::Namespace, "INTEL"},
3468+
Util::DeclContextDesc{Decl::Kind::Namespace, "property"},
3469+
Util::DeclContextDesc{Decl::Kind::CXXRecord, PropertyName},
3470+
Util::DeclContextDesc{Decl::Kind::ClassTemplateSpecialization,
3471+
InstanceName}};
34723472
return matchQualifiedTypeName(Ty, Scopes);
34733473
}
34743474

@@ -3482,6 +3482,16 @@ bool Util::isSyclType(const QualType &Ty, StringRef Name, bool Tmpl) {
34823482
return matchQualifiedTypeName(Ty, Scopes);
34833483
}
34843484

3485+
bool Util::isAccessorPropertyListType(const QualType &Ty) {
3486+
const StringRef &Name = "accessor_property_list";
3487+
std::array<DeclContextDesc, 4> Scopes = {
3488+
Util::DeclContextDesc{clang::Decl::Kind::Namespace, "cl"},
3489+
Util::DeclContextDesc{clang::Decl::Kind::Namespace, "sycl"},
3490+
Util::DeclContextDesc{clang::Decl::Kind::Namespace, "ONEAPI"},
3491+
Util::DeclContextDesc{Decl::Kind::ClassTemplateSpecialization, Name}};
3492+
return matchQualifiedTypeName(Ty, Scopes);
3493+
}
3494+
34853495
bool Util::matchQualifiedTypeName(const QualType &Ty,
34863496
ArrayRef<Util::DeclContextDesc> Scopes) {
34873497
// The idea: check the declaration context chain starting from the type

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

Lines changed: 21 additions & 18 deletions
Original file line numberDiff line numberDiff line change
@@ -71,19 +71,11 @@ enum prop_type {
7171
base_prop
7272
};
7373

74-
// Compile time known accessor property
75-
// TODO: this doesn't belong to property namespace, instead it shall be in its
76-
// own namespace. Change it, when the actual implementation in SYCL headers is
77-
// ready
78-
template <int>
79-
class buffer_location {};
80-
8174
struct property_base {
8275
virtual prop_type type() const = 0;
8376
};
8477
} // namespace property
8578

86-
template <typename... properties>
8779
class property_list {
8880
public:
8981
template <typename... propertiesTN>
@@ -102,6 +94,20 @@ class property_list {
10294
bool operator!=(const property_list &rhs) const { return false; }
10395
};
10496

97+
namespace INTEL {
98+
namespace property {
99+
// Compile time known accessor property
100+
struct buffer_location {
101+
template <int> class instance {};
102+
};
103+
} // namespace property
104+
} // namespace INTEL
105+
106+
namespace ONEAPI {
107+
template <typename... properties>
108+
class accessor_property_list {};
109+
} // namespace ONEAPI
110+
105111
template <int dim>
106112
struct id {
107113
template <typename... T>
@@ -136,7 +142,7 @@ struct _ImplT {
136142
template <typename dataT, int dimensions, access::mode accessmode,
137143
access::target accessTarget = access::target::global_buffer,
138144
access::placeholder isPlaceholder = access::placeholder::false_t,
139-
typename propertyListT = property_list<>>
145+
typename propertyListT = ONEAPI::accessor_property_list<>>
140146
class accessor {
141147

142148
public:
@@ -150,8 +156,6 @@ class accessor {
150156
private:
151157
void __init(__attribute__((opencl_global)) dataT *Ptr, range<dimensions> AccessRange,
152158
range<dimensions> MemRange, id<dimensions> Offset) {}
153-
154-
propertyListT prop_list;
155159
};
156160

157161
template <int dimensions, access::mode accessmode, access::target accesstarget>
@@ -339,8 +343,7 @@ const stream& operator<<(const stream &S, T&&) {
339343
}
340344

341345
template <typename T, int dimensions = 1,
342-
typename AllocatorT = int /*fake type as AllocatorT is not used*/,
343-
typename... properties>
346+
typename AllocatorT = int /*fake type as AllocatorT is not used*/>
344347
class buffer {
345348
public:
346349
using value_type = T;
@@ -352,13 +355,13 @@ class buffer {
352355
buffer(ParamTypes... args) {} // fake constructor
353356

354357
buffer(const range<dimensions> &bufferRange,
355-
const property_list<properties...> &propList = {}) {}
358+
const property_list &propList = {}) {}
356359

357360
buffer(T *hostData, const range<dimensions> &bufferRange,
358-
const property_list<properties...> &propList = {}) {}
361+
const property_list &propList = {}) {}
359362

360363
buffer(const T *hostData, const range<dimensions> &bufferRange,
361-
const property_list<properties...> &propList = {}) {}
364+
const property_list &propList = {}) {}
362365

363366
buffer(const buffer &rhs) = default;
364367

@@ -426,12 +429,12 @@ enum class image_channel_type : unsigned int {
426429
fp32
427430
};
428431

429-
template <int dimensions = 1, typename AllocatorT = int, typename... properties>
432+
template <int dimensions = 1, typename AllocatorT = int>
430433
class image {
431434
public:
432435
image(image_channel_order Order, image_channel_type Type,
433436
const range<dimensions> &Range,
434-
const property_list<properties...> &PropList = {}) {}
437+
const property_list &PropList = {}) {}
435438

436439
/* -- common interface members -- */
437440

‎clang/test/CodeGenSYCL/accessor_inheritance.cpp‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -67,13 +67,13 @@ int main() {
6767
// CHECK: [[ACC_FIELD:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, %struct{{.*}}Base* [[BITCAST]], i32 0, i32 2
6868
// CHECK: [[ACC1_AS_CAST:%[a-zA-Z0-9_]+]] = addrspacecast %"class{{.*}}cl::sycl::accessor"* [[ACC_FIELD]] to %"class{{.*}}cl::sycl::accessor" addrspace(4)*
6969
// Default constructor call
70-
// CHECK: call spir_func void @_ZN2cl4sycl8accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_13property_listIJEEEEC1Ev(%"class{{.*}}cl::sycl::accessor" addrspace(4)* [[ACC1_AS_CAST]])
70+
// CHECK: call spir_func void @_ZN2cl4sycl8accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEC1Ev(%"class{{.*}}cl::sycl::accessor" addrspace(4)* [[ACC1_AS_CAST]])
7171
// CHECK: [[BITCAST1:%[a-zA-Z0-9_]+]] = bitcast %struct{{.*}}Captured* [[GEP]] to i8*
72-
// CHECK: [[GEP1:%[a-zA-Z0-9_]+]] = getelementptr inbounds i8, i8* [[BITCAST1]], i64 24
72+
// CHECK: [[GEP1:%[a-zA-Z0-9_]+]] = getelementptr inbounds i8, i8* [[BITCAST1]], i64 20
7373
// CHECK: [[BITCAST2:%[a-zA-Z0-9_]+]] = bitcast i8* [[GEP1]] to %"class{{.*}}cl::sycl::accessor"*
7474
// CHECK: [[ACC2_AS_CAST:%[a-zA-Z0-9_]+]] = addrspacecast %"class{{.*}}cl::sycl::accessor"* [[BITCAST2]] to %"class{{.*}}cl::sycl::accessor" addrspace(4)*
7575
// Default constructor call
76-
// CHECK: call spir_func void @_ZN2cl4sycl8accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_13property_listIJEEEEC2Ev(%"class{{.*}}cl::sycl::accessor" addrspace(4)* [[ACC2_AS_CAST]])
76+
// CHECK: call spir_func void @_ZN2cl4sycl8accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEC2Ev(%"class{{.*}}cl::sycl::accessor" addrspace(4)* [[ACC2_AS_CAST]])
7777

7878
// CHECK C field initialization
7979
// CHECK: [[FIELD_C:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Captured, %struct{{.*}}Captured* [[GEP]], i32 0, i32 2

‎clang/test/CodeGenSYCL/buffer_location.cpp‎

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -10,17 +10,17 @@ struct Base {
1010
cl::sycl::accessor<char, 1, cl::sycl::access::mode::read,
1111
cl::sycl::access::target::global_buffer,
1212
cl::sycl::access::placeholder::false_t,
13-
cl::sycl::property_list<
14-
cl::sycl::property::buffer_location<2>>>
13+
cl::sycl::ONEAPI::accessor_property_list<
14+
cl::sycl::INTEL::property::buffer_location::instance<2>>>
1515
AccField;
1616
};
1717

1818
struct Captured : Base,
1919
cl::sycl::accessor<char, 1, cl::sycl::access::mode::read,
2020
cl::sycl::access::target::global_buffer,
2121
cl::sycl::access::placeholder::false_t,
22-
cl::sycl::property_list<
23-
cl::sycl::property::buffer_location<2>>> {
22+
cl::sycl::ONEAPI::accessor_property_list<
23+
cl::sycl::INTEL::property::buffer_location::instance<2>>> {
2424
int C;
2525
};
2626

@@ -29,8 +29,8 @@ int main() {
2929
cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write,
3030
cl::sycl::access::target::global_buffer,
3131
cl::sycl::access::placeholder::false_t,
32-
cl::sycl::property_list<
33-
cl::sycl::property::buffer_location<3>>>
32+
cl::sycl::ONEAPI::accessor_property_list<
33+
cl::sycl::INTEL::property::buffer_location::instance<3>>>
3434
accessorA;
3535
cl::sycl::kernel_single_task<class kernel_function>(
3636
[=]() {

‎clang/test/CodeGenSYCL/integration_header.cpp‎

Lines changed: 9 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -31,18 +31,18 @@
3131
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 1, 4 },
3232
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 8 },
3333
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 12 },
34-
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 6112, 28 },
35-
// CHECK-NEXT: { kernel_param_kind_t::kind_sampler, 8, 48 },
34+
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 6112, 24 },
35+
// CHECK-NEXT: { kernel_param_kind_t::kind_sampler, 8, 40 },
3636
// CHECK-EMPTY:
3737
// CHECK-NEXT: //--- _ZTSN16second_namespace13second_kernelIcEE
3838
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 0 },
3939
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 6112, 4 },
40-
// CHECK-NEXT: { kernel_param_kind_t::kind_sampler, 8, 24 },
40+
// CHECK-NEXT: { kernel_param_kind_t::kind_sampler, 8, 16 },
4141
// CHECK-EMPTY:
4242
// CHECK-NEXT: //--- _ZTS12third_kernelILi1Ei5pointIZ4mainE1XEE
4343
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 0 },
4444
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 6112, 4 },
45-
// CHECK-NEXT: { kernel_param_kind_t::kind_sampler, 8, 24 },
45+
// CHECK-NEXT: { kernel_param_kind_t::kind_sampler, 8, 16 },
4646
// CHECK-EMPTY:
4747
// CHECK-NEXT: //--- _ZTS13fourth_kernelIJN15template_arg_ns14namespaced_argILi1EEEEE
4848
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 0 },
@@ -52,11 +52,11 @@
5252
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 0 },
5353
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 4 },
5454
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 8 },
55-
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 24 },
56-
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 28 },
57-
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 44 },
58-
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 48 },
59-
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 64 },
55+
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 20 },
56+
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 24 },
57+
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 36 },
58+
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 40 },
59+
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 52 },
6060
// CHECK-EMPTY:
6161
// CHECK-NEXT: };
6262
//

‎clang/test/CodeGenSYCL/kernel-param-acc-array-ih.cpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -21,7 +21,7 @@
2121
// CHECK-NEXT: const kernel_param_desc_t kernel_signatures[] = {
2222
// CHECK-NEXT: //--- _ZTSZ4mainE8kernel_A
2323
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 0 },
24-
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 16 },
24+
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 12 },
2525
// CHECK-EMPTY:
2626
// CHECK-NEXT: };
2727

‎clang/test/CodeGenSYCL/kernel-param-member-acc-array-ih.cpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -21,7 +21,7 @@
2121
// CHECK-NEXT: const kernel_param_desc_t kernel_signatures[] = {
2222
// CHECK-NEXT: //--- _ZTSZ4mainE8kernel_C
2323
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 0 },
24-
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 16 },
24+
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 12 },
2525
// CHECK-EMPTY:
2626
// CHECK-NEXT: };
2727

‎clang/test/CodeGenSYCL/struct_kernel_param.cpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -4,12 +4,12 @@
44
// CHECK: const kernel_param_desc_t kernel_signatures[] = {
55
// CHECK-NEXT: //--- _ZTSZZ5test0vENK3$_0clERN2cl4sycl7handlerEE8MyKernel
66
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 0 },
7+
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 12 },
78
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 16 },
89
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 20 },
910
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 24 },
1011
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 28 },
1112
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 32 },
12-
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 36 },
1313
// CHECK-EMPTY:
1414
// CHECK-NEXT:};
1515

0 commit comments

Comments
 (0)