Skip to content

Commit 1884bd4

Browse files
mlychkovAlexander Batashev
mlychkov
authored and
Alexander Batashev
committed
[SYCL][FPGA] Align clang with new spec of accessor_property_list
Signed-off-by: Mikhail Lychkov <[email protected]>
1 parent 0474429 commit 1884bd4

14 files changed

+141
-93
lines changed

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

Lines changed: 4 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -11033,9 +11033,10 @@ def err_sycl_compiletime_property_duplication : Error<
1103311033
def err_sycl_invalid_property_list_param_number : Error<
1103411034
"%0 must have exactly one template parameter">;
1103511035
def err_sycl_invalid_accessor_property_template_param : Error<
11036-
"Fifth template parameter of the accessor must be of a property_list type">;
11037-
def err_sycl_invalid_property_list_template_param : Error<
11038-
"%select{property_list|property_list pack argument|buffer_location}0 "
11036+
"Sixth template parameter of the accessor must be of accessor_property_list "
11037+
"or property_list type">;
11038+
def err_sycl_invalid_accessor_property_list_template_param : Error<
11039+
"%select{accessor_property_list|accessor_property_list pack argument|buffer_location}0 "
1103911040
"template parameter must be a "
1104011041
"%select{parameter pack|type|non-negative integer}1">;
1104111042
def warn_sycl_pass_by_value_deprecated

‎clang/lib/Sema/SemaSYCL.cpp

Lines changed: 60 additions & 26 deletions
Original file line numberDiff line numberDiff line change
@@ -85,6 +85,10 @@ class Util {
8585
/// property_list class.
8686
static bool isPropertyListType(const QualType &Ty);
8787

88+
/// Checks whether given clang type is a full specialization of the SYCL
89+
/// accessor_property_list class.
90+
static bool isAccessorPropertyListType(const QualType &Ty);
91+
8892
/// Checks whether given clang type is a full specialization of the SYCL
8993
/// buffer_location class.
9094
static bool isSyclBufferLocationType(const QualType &Ty);
@@ -1194,29 +1198,33 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler {
11941198
return;
11951199
}
11961200
QualType PropListTy = PropList.getAsType();
1197-
if (!Util::isPropertyListType(PropListTy)) {
1201+
if (Util::isPropertyListType(PropListTy))
1202+
return;
1203+
if (!Util::isAccessorPropertyListType(PropListTy)) {
11981204
SemaRef.Diag(Loc,
11991205
diag::err_sycl_invalid_accessor_property_template_param);
12001206
return;
12011207
}
1202-
const auto *PropListDecl =
1208+
const auto *AccPropListDecl =
12031209
cast<ClassTemplateSpecializationDecl>(PropListTy->getAsRecordDecl());
1204-
if (PropListDecl->getTemplateArgs().size() != 1) {
1210+
if (AccPropListDecl->getTemplateArgs().size() != 1) {
12051211
SemaRef.Diag(Loc, diag::err_sycl_invalid_property_list_param_number)
1206-
<< "property_list";
1212+
<< "accessor_property_list";
12071213
return;
12081214
}
1209-
const auto TemplArg = PropListDecl->getTemplateArgs()[0];
1215+
const auto TemplArg = AccPropListDecl->getTemplateArgs()[0];
12101216
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;
1217+
SemaRef.Diag(Loc,
1218+
diag::err_sycl_invalid_accessor_property_list_template_param)
1219+
<< /*accessor_property_list*/ 0 << /*parameter pack*/ 0;
12131220
return;
12141221
}
12151222
for (TemplateArgument::pack_iterator Prop = TemplArg.pack_begin();
12161223
Prop != TemplArg.pack_end(); ++Prop) {
12171224
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;
1225+
SemaRef.Diag(
1226+
Loc, diag::err_sycl_invalid_accessor_property_list_template_param)
1227+
<< /*accessor_property_list pack argument*/ 1 << /*type*/ 1;
12201228
return;
12211229
}
12221230
QualType PropTy = Prop->getAsType();
@@ -1235,13 +1243,15 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler {
12351243
}
12361244
const auto BufferLoc = PropDecl->getTemplateArgs()[0];
12371245
if (BufferLoc.getKind() != TemplateArgument::ArgKind::Integral) {
1238-
SemaRef.Diag(Loc, diag::err_sycl_invalid_property_list_template_param)
1246+
SemaRef.Diag(Loc,
1247+
diag::err_sycl_invalid_accessor_property_list_template_param)
12391248
<< /*buffer_location*/ 2 << /*non-negative integer*/ 2;
12401249
return;
12411250
}
12421251
int LocationID = static_cast<int>(BufferLoc.getAsIntegral().getExtValue());
12431252
if (LocationID < 0) {
1244-
SemaRef.Diag(Loc, diag::err_sycl_invalid_property_list_template_param)
1253+
SemaRef.Diag(Loc,
1254+
diag::err_sycl_invalid_accessor_property_list_template_param)
12451255
<< /*buffer_location*/ 2 << /*non-negative integer*/ 2;
12461256
return;
12471257
}
@@ -1251,17 +1261,29 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler {
12511261
assert(Util::isSyclAccessorType(Ty) &&
12521262
"Should only be called on SYCL accessor types.");
12531263

1254-
const RecordDecl *RecD = Ty->getAsRecordDecl();
1255-
if (const ClassTemplateSpecializationDecl *CTSD =
1256-
dyn_cast<ClassTemplateSpecializationDecl>(RecD)) {
1264+
RecordDecl *RecD = Ty->getAsRecordDecl();
1265+
if (auto *CTSD = dyn_cast<ClassTemplateSpecializationDecl>(RecD)) {
12571266
const TemplateArgumentList &TAL = CTSD->getTemplateArgs();
1267+
12581268
TemplateArgument TA = TAL.get(0);
12591269
const QualType TemplateArgTy = TA.getAsType();
1260-
1261-
if (TAL.size() > 5)
1262-
checkPropertyListType(TAL.get(5), Loc.getBegin());
12631270
llvm::DenseSet<QualType> Visited;
12641271
checkSYCLType(SemaRef, TemplateArgTy, Loc, Visited);
1272+
1273+
if (TAL.size() < 6) {
1274+
// Not enough arguments for this parameter pack.
1275+
SemaRef.Diag(Loc.getBegin(),
1276+
diag::err_template_arg_list_different_arity)
1277+
<< /*not enough args*/ 0
1278+
<< (int)SemaRef.getTemplateNameKindForDiagnostics(
1279+
TemplateName(CTSD->getSpecializedTemplate()))
1280+
<< CTSD;
1281+
SemaRef.Diag(CTSD->getLocation(), diag::note_template_decl_here)
1282+
<< CTSD->getSourceRange();
1283+
return;
1284+
}
1285+
1286+
checkPropertyListType(TAL.get(5), Loc.getBegin());
12651287
}
12661288
}
12671289

@@ -1402,19 +1424,21 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
14021424
}
14031425

14041426
// Handle accessor properties. If any properties were found in
1405-
// the property_list - add the appropriate attributes to ParmVarDecl.
1427+
// the accessor_property_list - add the appropriate attributes to ParmVarDecl.
14061428
void handleAccessorPropertyList(ParmVarDecl *Param,
14071429
const CXXRecordDecl *RecordDecl,
14081430
SourceLocation Loc) {
14091431
const auto *AccTy = cast<ClassTemplateSpecializationDecl>(RecordDecl);
1410-
// TODO: when SYCL headers' part is ready - replace this 'if' with an error
14111432
if (AccTy->getTemplateArgs().size() < 6)
14121433
return;
14131434
const auto PropList = cast<TemplateArgument>(AccTy->getTemplateArgs()[5]);
14141435
QualType PropListTy = PropList.getAsType();
1415-
const auto *PropListDecl =
1436+
// property_list contains runtime properties, it shouldn't be handled here.
1437+
if (Util::isPropertyListType(PropListTy))
1438+
return;
1439+
const auto *AccPropListDecl =
14161440
cast<ClassTemplateSpecializationDecl>(PropListTy->getAsRecordDecl());
1417-
const auto TemplArg = PropListDecl->getTemplateArgs()[0];
1441+
const auto TemplArg = AccPropListDecl->getTemplateArgs()[0];
14181442
// Move through TemplateArgs list of a property list and search for
14191443
// properties. If found - apply the appropriate attribute to ParmVarDecl.
14201444
for (TemplateArgument::pack_iterator Prop = TemplArg.pack_begin();
@@ -3444,17 +3468,16 @@ bool Util::isSyclSpecConstantType(const QualType &Ty) {
34443468
}
34453469

34463470
bool Util::isPropertyListType(const QualType &Ty) {
3447-
return isSyclType(Ty, "property_list", true /*Tmpl*/);
3471+
return isSyclType(Ty, "property_list");
34483472
}
34493473

34503474
bool Util::isSyclBufferLocationType(const QualType &Ty) {
34513475
const StringRef &Name = "buffer_location";
3452-
std::array<DeclContextDesc, 4> Scopes = {
3476+
std::array<DeclContextDesc, 6> Scopes = {
34533477
Util::DeclContextDesc{clang::Decl::Kind::Namespace, "cl"},
34543478
Util::DeclContextDesc{clang::Decl::Kind::Namespace, "sycl"},
3455-
// TODO: this doesn't belong to property namespace, instead it shall be
3456-
// in its own namespace. Change it, when the actual implementation in SYCL
3457-
// headers is ready
3479+
Util::DeclContextDesc{clang::Decl::Kind::Namespace, "ext"},
3480+
Util::DeclContextDesc{clang::Decl::Kind::Namespace, "INTEL"},
34583481
Util::DeclContextDesc{clang::Decl::Kind::Namespace, "property"},
34593482
Util::DeclContextDesc{Decl::Kind::ClassTemplateSpecialization, Name}};
34603483
return matchQualifiedTypeName(Ty, Scopes);
@@ -3470,6 +3493,17 @@ bool Util::isSyclType(const QualType &Ty, StringRef Name, bool Tmpl) {
34703493
return matchQualifiedTypeName(Ty, Scopes);
34713494
}
34723495

3496+
bool Util::isAccessorPropertyListType(const QualType &Ty) {
3497+
const StringRef &Name = "accessor_property_list";
3498+
std::array<DeclContextDesc, 5> Scopes = {
3499+
Util::DeclContextDesc{clang::Decl::Kind::Namespace, "cl"},
3500+
Util::DeclContextDesc{clang::Decl::Kind::Namespace, "sycl"},
3501+
Util::DeclContextDesc{clang::Decl::Kind::Namespace, "ext"},
3502+
Util::DeclContextDesc{clang::Decl::Kind::Namespace, "ONEAPI"},
3503+
Util::DeclContextDesc{Decl::Kind::ClassTemplateSpecialization, Name}};
3504+
return matchQualifiedTypeName(Ty, Scopes);
3505+
}
3506+
34733507
bool Util::matchQualifiedTypeName(const QualType &Ty,
34743508
ArrayRef<Util::DeclContextDesc> Scopes) {
34753509
// The idea: check the declaration context chain starting from the type

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

Lines changed: 22 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,21 @@ class property_list {
10294
bool operator!=(const property_list &rhs) const { return false; }
10395
};
10496

97+
namespace ext {
98+
namespace INTEL {
99+
namespace property {
100+
// Compile time known accessor property
101+
template <int>
102+
class buffer_location {};
103+
} // namespace property
104+
} // namespace INTEL
105+
106+
namespace ONEAPI {
107+
template <typename... properties>
108+
class accessor_property_list {};
109+
} // namespace ONEAPI
110+
} // namespace ext
111+
105112
template <int dim>
106113
struct id {
107114
template <typename... T>
@@ -136,7 +143,7 @@ struct _ImplT {
136143
template <typename dataT, int dimensions, access::mode accessmode,
137144
access::target accessTarget = access::target::global_buffer,
138145
access::placeholder isPlaceholder = access::placeholder::false_t,
139-
typename propertyListT = property_list<>>
146+
typename propertyListT = property_list>
140147
class accessor {
141148

142149
public:
@@ -150,8 +157,6 @@ class accessor {
150157
private:
151158
void __init(__attribute__((opencl_global)) dataT *Ptr, range<dimensions> AccessRange,
152159
range<dimensions> MemRange, id<dimensions> Offset) {}
153-
154-
propertyListT prop_list;
155160
};
156161

157162
template <int dimensions, access::mode accessmode, access::target accesstarget>
@@ -339,8 +344,7 @@ const stream& operator<<(const stream &S, T&&) {
339344
}
340345

341346
template <typename T, int dimensions = 1,
342-
typename AllocatorT = int /*fake type as AllocatorT is not used*/,
343-
typename... properties>
347+
typename AllocatorT = int /*fake type as AllocatorT is not used*/>
344348
class buffer {
345349
public:
346350
using value_type = T;
@@ -352,13 +356,13 @@ class buffer {
352356
buffer(ParamTypes... args) {} // fake constructor
353357

354358
buffer(const range<dimensions> &bufferRange,
355-
const property_list<properties...> &propList = {}) {}
359+
const property_list &propList = {}) {}
356360

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

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

363367
buffer(const buffer &rhs) = default;
364368

@@ -426,12 +430,12 @@ enum class image_channel_type : unsigned int {
426430
fp32
427431
};
428432

429-
template <int dimensions = 1, typename AllocatorT = int, typename... properties>
433+
template <int dimensions = 1, typename AllocatorT = int>
430434
class image {
431435
public:
432436
image(image_channel_order Order, image_channel_type Type,
433437
const range<dimensions> &Range,
434-
const property_list<properties...> &PropList = {}) {}
438+
const property_list &PropList = {}) {}
435439

436440
/* -- common interface members -- */
437441

‎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_13property_listEEC1Ev(%"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_13property_listEEC2Ev(%"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::ext::ONEAPI::accessor_property_list<
14+
cl::sycl::ext::INTEL::property::buffer_location<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::ext::ONEAPI::accessor_property_list<
23+
cl::sycl::ext::INTEL::property::buffer_location<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::ext::ONEAPI::accessor_property_list<
33+
cl::sycl::ext::INTEL::property::buffer_location<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

0 commit comments

Comments
 (0)