Skip to content

Commit 52cf2fe

Browse files
jinge90jinge90
committed
Merge remote-tracking branch 'origin/sycl' into link_devicelib_by_default_in_aot_style
2 parents d184db2 + 0474429 commit 52cf2fe

24 files changed

+701
-85
lines changed

‎clang/lib/Sema/SemaSYCL.cpp

+29-22
Original file line numberDiff line numberDiff line change
@@ -1752,6 +1752,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
17521752
const CXXRecordDecl *KernelObj;
17531753
llvm::SmallVector<Expr *, 16> MemberExprBases;
17541754
FunctionDecl *KernelCallerFunc;
1755+
SourceLocation KernelCallerSrcLoc; // KernelCallerFunc source location.
17551756
// Contains a count of how many containers we're in. This is used by the
17561757
// pointer-struct-wrapping code to ensure that we don't try to wrap
17571758
// non-top-level pointers.
@@ -1821,7 +1822,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
18211822

18221823
QualType ParamType = KernelParameter->getOriginalType();
18231824
Expr *DRE = SemaRef.BuildDeclRefExpr(KernelParameter, ParamType, VK_LValue,
1824-
SourceLocation());
1825+
KernelCallerSrcLoc);
18251826
return DRE;
18261827
}
18271828

@@ -1833,7 +1834,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
18331834

18341835
QualType ParamType = KernelParameter->getOriginalType();
18351836
Expr *DRE = SemaRef.BuildDeclRefExpr(KernelParameter, ParamType, VK_LValue,
1836-
SourceLocation());
1837+
KernelCallerSrcLoc);
18371838

18381839
// Struct Type kernel arguments are decomposed. The pointer fields are
18391840
// then wrapped inside a compiler generated struct. Therefore when
@@ -1846,6 +1847,10 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
18461847
ParamType = Pointer->getType();
18471848
}
18481849

1850+
DRE =
1851+
ImplicitCastExpr::Create(SemaRef.Context, ParamType, CK_LValueToRValue,
1852+
DRE, /*BasePath=*/nullptr, VK_RValue);
1853+
18491854
if (PointerTy->getPointeeType().getAddressSpace() !=
18501855
ParamType->getPointeeType().getAddressSpace())
18511856
DRE = ImplicitCastExpr::Create(SemaRef.Context, PointerTy,
@@ -1876,7 +1881,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
18761881

18771882
void addFieldInit(FieldDecl *FD, QualType Ty, MultiExprArg ParamRef) {
18781883
InitializationKind InitKind =
1879-
InitializationKind::CreateCopy(SourceLocation(), SourceLocation());
1884+
InitializationKind::CreateCopy(KernelCallerSrcLoc, KernelCallerSrcLoc);
18801885
addFieldInit(FD, Ty, ParamRef, InitKind);
18811886
}
18821887

@@ -1913,10 +1918,10 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
19131918
MemberExpr *buildMemberExpr(Expr *Base, ValueDecl *Member) {
19141919
DeclAccessPair MemberDAP = DeclAccessPair::make(Member, AS_none);
19151920
MemberExpr *Result = SemaRef.BuildMemberExpr(
1916-
Base, /*IsArrow */ false, SourceLocation(), NestedNameSpecifierLoc(),
1917-
SourceLocation(), Member, MemberDAP,
1921+
Base, /*IsArrow */ false, KernelCallerSrcLoc, NestedNameSpecifierLoc(),
1922+
KernelCallerSrcLoc, Member, MemberDAP,
19181923
/*HadMultipleCandidates*/ false,
1919-
DeclarationNameInfo(Member->getDeclName(), SourceLocation()),
1924+
DeclarationNameInfo(Member->getDeclName(), KernelCallerSrcLoc),
19201925
Member->getType(), VK_LValue, OK_Ordinary);
19211926
return Result;
19221927
}
@@ -1944,7 +1949,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
19441949
for (size_t I = 0; I < NumParams; ++I) {
19451950
QualType ParamType = KernelParameters[I]->getOriginalType();
19461951
ParamDREs[I] = SemaRef.BuildDeclRefExpr(KernelParameters[I], ParamType,
1947-
VK_LValue, SourceLocation());
1952+
VK_LValue, KernelCallerSrcLoc);
19481953
}
19491954

19501955
MemberExpr *MethodME = buildMemberExpr(MemberExprBases.back(), Method);
@@ -1954,12 +1959,12 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
19541959
ResultTy = ResultTy.getNonLValueExprType(SemaRef.Context);
19551960
llvm::SmallVector<Expr *, 4> ParamStmts;
19561961
const auto *Proto = cast<FunctionProtoType>(Method->getType());
1957-
SemaRef.GatherArgumentsForCall(SourceLocation(), Method, Proto, 0,
1962+
SemaRef.GatherArgumentsForCall(KernelCallerSrcLoc, Method, Proto, 0,
19581963
ParamDREs, ParamStmts);
19591964
// [kernel_obj or wrapper object].accessor.__init(_ValueType*,
19601965
// range<int>, range<int>, id<int>)
19611966
AddTo.push_back(CXXMemberCallExpr::Create(
1962-
SemaRef.Context, MethodME, ParamStmts, ResultTy, VK, SourceLocation(),
1967+
SemaRef.Context, MethodME, ParamStmts, ResultTy, VK, KernelCallerSrcLoc,
19631968
FPOptionsOverride()));
19641969
}
19651970

@@ -1981,7 +1986,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
19811986

19821987
InitListExpr *createInitListExpr(QualType InitTy, uint64_t NumChildInits) {
19831988
InitListExpr *ILE = new (SemaRef.getASTContext()) InitListExpr(
1984-
SemaRef.getASTContext(), SourceLocation(), {}, SourceLocation());
1989+
SemaRef.getASTContext(), KernelCallerSrcLoc, {}, KernelCallerSrcLoc);
19851990
ILE->reserveInits(SemaRef.getASTContext(), NumChildInits);
19861991
ILE->setType(InitTy);
19871992

@@ -2007,16 +2012,17 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
20072012
TypeSourceInfo *TSInfo =
20082013
KernelObj->isLambda() ? KernelObj->getLambdaTypeInfo() : nullptr;
20092014
VarDecl *VD = VarDecl::Create(
2010-
Ctx, DC, SourceLocation(), SourceLocation(), KernelObj->getIdentifier(),
2011-
QualType(KernelObj->getTypeForDecl(), 0), TSInfo, SC_None);
2015+
Ctx, DC, KernelObj->getLocation(), KernelObj->getLocation(),
2016+
KernelObj->getIdentifier(), QualType(KernelObj->getTypeForDecl(), 0),
2017+
TSInfo, SC_None);
20122018

20132019
return VD;
20142020
}
20152021

20162022
// Default inits the type, then calls the init-method in the body.
20172023
bool handleSpecialType(FieldDecl *FD, QualType Ty) {
20182024
addFieldInit(FD, Ty, None,
2019-
InitializationKind::CreateDefault(SourceLocation()));
2025+
InitializationKind::CreateDefault(KernelCallerSrcLoc));
20202026

20212027
addFieldMemberExpr(FD, Ty);
20222028

@@ -2030,7 +2036,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
20302036

20312037
bool handleSpecialType(const CXXBaseSpecifier &BS, QualType Ty) {
20322038
const auto *RecordDecl = Ty->getAsCXXRecordDecl();
2033-
addBaseInit(BS, Ty, InitializationKind::CreateDefault(SourceLocation()));
2039+
addBaseInit(BS, Ty, InitializationKind::CreateDefault(KernelCallerSrcLoc));
20342040
createSpecialMethodCall(RecordDecl, InitMethodName, BodyStmts);
20352041
return true;
20362042
}
@@ -2043,15 +2049,16 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
20432049
KernelObjClone(createKernelObjClone(S.getASTContext(),
20442050
DC.getKernelDecl(), KernelObj)),
20452051
VarEntity(InitializedEntity::InitializeVariable(KernelObjClone)),
2046-
KernelObj(KernelObj), KernelCallerFunc(KernelCallerFunc) {
2052+
KernelObj(KernelObj), KernelCallerFunc(KernelCallerFunc),
2053+
KernelCallerSrcLoc(KernelCallerFunc->getLocation()) {
20472054
CollectionInitExprs.push_back(createInitListExpr(KernelObj));
20482055
markParallelWorkItemCalls();
20492056

20502057
Stmt *DS = new (S.Context) DeclStmt(DeclGroupRef(KernelObjClone),
2051-
SourceLocation(), SourceLocation());
2058+
KernelCallerSrcLoc, KernelCallerSrcLoc);
20522059
BodyStmts.push_back(DS);
20532060
DeclRefExpr *KernelObjCloneRef = DeclRefExpr::Create(
2054-
S.Context, NestedNameSpecifierLoc(), SourceLocation(), KernelObjClone,
2061+
S.Context, NestedNameSpecifierLoc(), KernelCallerSrcLoc, KernelObjClone,
20552062
false, DeclarationNameInfo(), QualType(KernelObj->getTypeForDecl(), 0),
20562063
VK_LValue);
20572064
MemberExprBases.push_back(KernelObjCloneRef);
@@ -2099,7 +2106,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
20992106

21002107
bool handlePointerType(FieldDecl *FD, QualType FieldTy) final {
21012108
Expr *PointerRef =
2102-
createPointerParamReferenceExpr(FD->getType(), StructDepth != 0);
2109+
createPointerParamReferenceExpr(FieldTy, StructDepth != 0);
21032110
addFieldInit(FD, FieldTy, PointerRef);
21042111
return true;
21052112
}
@@ -2162,7 +2169,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
21622169
CXXCastPath BasePath;
21632170
QualType DerivedTy(RD->getTypeForDecl(), 0);
21642171
QualType BaseTy = BS.getType();
2165-
SemaRef.CheckDerivedToBaseConversion(DerivedTy, BaseTy, SourceLocation(),
2172+
SemaRef.CheckDerivedToBaseConversion(DerivedTy, BaseTy, KernelCallerSrcLoc,
21662173
SourceRange(), &BasePath,
21672174
/*IgnoreBaseAccess*/ true);
21682175
auto Cast = ImplicitCastExpr::Create(
@@ -2211,11 +2218,11 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
22112218
Index, SizeT->isSignedIntegerType()};
22122219

22132220
auto IndexLiteral = IntegerLiteral::Create(
2214-
SemaRef.getASTContext(), IndexVal, SizeT, SourceLocation());
2221+
SemaRef.getASTContext(), IndexVal, SizeT, KernelCallerSrcLoc);
22152222

22162223
ExprResult IndexExpr = SemaRef.CreateBuiltinArraySubscriptExpr(
2217-
MemberExprBases.back(), SourceLocation{}, IndexLiteral,
2218-
SourceLocation{});
2224+
MemberExprBases.back(), KernelCallerSrcLoc, IndexLiteral,
2225+
KernelCallerSrcLoc);
22192226

22202227
assert(!IndexExpr.isInvalid());
22212228
MemberExprBases.push_back(IndexExpr.get());

‎clang/test/CodeGenSYCL/debug-info-srcpos-kernel.cpp

+8-4
Original file line numberDiff line numberDiff line change
@@ -3,9 +3,10 @@
33
// Verify the SYCL kernel routine is marked artificial and has no source
44
// correlation.
55
//
6-
// The SYCL kernel should have no source correlation of its own, so it needs
7-
// to be marked artificial or it will inherit source correlation from the
8-
// surrounding code.
6+
// In order to placate the profiling tools, which can't cope with instructions
7+
// mapped to line 0, we've made the change so that the artificial code in a
8+
// SYCL kernel gets the source line info for the kernel caller function (the
9+
// 'kernel' template function on line 15 in this file).
910
//
1011

1112
#include <sycl.hpp>
@@ -33,8 +34,11 @@ int main() {
3334
// CHECK-SAME: scope: [[FILE]]
3435
// CHECK-SAME: file: [[FILE]]
3536
// CHECK-SAME: flags: DIFlagArtificial | DIFlagPrototyped
36-
// CHECK: [[LINE_A0]] = !DILocation(line: 0
37+
// CHECK: [[LINE_A0]] = !DILocation(line: 15,{{.*}}scope: [[KERNEL]]
3738
// CHECK: [[LINE_B0]] = !DILocation(line: 0
3839

40+
// TODO: [[LINE_B0]] should be mapped to line 15 as well. That said,
41+
// this 'line 0' assignment is less problematic as the lambda function
42+
// call would be inlined in most cases.
3943
// TODO: SYCL specific fail - analyze and enable
4044
// XFAIL: windows-msvc

‎clang/test/CodeGenSYCL/pointers-in-structs.cpp

+6-1
Original file line numberDiff line numberDiff line change
@@ -13,6 +13,7 @@ struct B {
1313
int *F1;
1414
float *F2;
1515
A F3;
16+
int *F4[2];
1617
};
1718

1819
int main() {
@@ -35,9 +36,13 @@ int main() {
3536
// CHECK: %[[WRAPPER_F1:[a-zA-Z0-9_.]+]] = type { i32 addrspace(1)* }
3637
// CHECK: %[[WRAPPER_F2:[a-zA-Z0-9_.]+]] = type { float addrspace(1)* }
3738
// CHECK: %[[WRAPPER_F:[a-zA-Z0-9_.]+]] = type { float addrspace(1)* }
39+
// CHECK: %[[WRAPPER_F4_1:[a-zA-Z0-9_.]+]] = type { i32 addrspace(1)* }
40+
// CHECK: %[[WRAPPER_F4_2:[a-zA-Z0-9_.]+]] = type { i32 addrspace(1)* }
3841
// CHECK: %[[WRAPPER_LAMBDA_PTR:[a-zA-Z0-9_.]+]] = type { float addrspace(1)* }
3942
// CHECK: define spir_kernel void @{{.*}}structs
4043
// CHECK-SAME: %[[WRAPPER_F1]]* byval(%[[WRAPPER_F1]]) align 8 %_arg_F1,
4144
// CHECK-SAME: %[[WRAPPER_F2]]* byval(%[[WRAPPER_F2]]) align 8 %_arg_F2,
42-
// CHECK-SAME: %[[WRAPPER_F]]* byval(%[[WRAPPER_F]]) align 8 %_arg_F
45+
// CHECK-SAME: %[[WRAPPER_F]]* byval(%[[WRAPPER_F]]) align 8 %_arg_F,
46+
// CHECK-SAME: %[[WRAPPER_F4_1]]* byval(%[[WRAPPER_F4_1]]) align 8 %_arg_F4
47+
// CHECK-SAME: %[[WRAPPER_F4_2]]* byval(%[[WRAPPER_F4_2]]) align 8 %_arg_F41
4348
// CHECK: define spir_kernel void @{{.*}}lambdas{{.*}}(%[[WRAPPER_LAMBDA_PTR]]* byval(%[[WRAPPER_LAMBDA_PTR]]) align 8 %_arg_)

‎clang/test/CodeGenSYCL/pointers-int-header.cpp

+10
Original file line numberDiff line numberDiff line change
@@ -10,6 +10,8 @@
1010
struct struct_with_pointer {
1111
int data_in_struct;
1212
int *ptr_in_struct;
13+
int *ptr_array_in_struct1[2];
14+
int *ptr_array_in_struct2[2][3];
1315
};
1416

1517
int main() {
@@ -27,3 +29,11 @@ int main() {
2729
// CHECK:{ kernel_param_kind_t::kind_pointer, 8, 0 },
2830
// CHECK:{ kernel_param_kind_t::kind_std_layout, 4, 8 },
2931
// CHECK:{ kernel_param_kind_t::kind_std_layout, 8, 16 },
32+
// CHECK:{ kernel_param_kind_t::kind_std_layout, 8, 24 },
33+
// CHECK:{ kernel_param_kind_t::kind_std_layout, 8, 32 },
34+
// CHECK:{ kernel_param_kind_t::kind_std_layout, 8, 40 },
35+
// CHECK:{ kernel_param_kind_t::kind_std_layout, 8, 48 },
36+
// CHECK:{ kernel_param_kind_t::kind_std_layout, 8, 56 },
37+
// CHECK:{ kernel_param_kind_t::kind_std_layout, 8, 64 },
38+
// CHECK:{ kernel_param_kind_t::kind_std_layout, 8, 72 },
39+
// CHECK:{ kernel_param_kind_t::kind_std_layout, 8, 80 },

‎clang/test/SemaSYCL/built-in-type-kernel-arg.cpp

+65-5
Original file line numberDiff line numberDiff line change
@@ -12,7 +12,9 @@ __attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
1212

1313
struct test_struct {
1414
int data;
15-
int *ptr; // Unused pointer in struct
15+
int *ptr;
16+
int *ptr_array1[2];
17+
int *ptr_array2[2][3];
1618
};
1719

1820
void test(const int some_const) {
@@ -26,6 +28,7 @@ int main() {
2628
int data = 5;
2729
int* data_addr = &data;
2830
int* new_data_addr = nullptr;
31+
int *ptr_array[2];
2932
test_struct s;
3033
s.data = data;
3134
kernel<class kernel_int>(
@@ -40,7 +43,9 @@ int main() {
4043
kernel<class kernel_pointer>(
4144
[=]() {
4245
new_data_addr[0] = data_addr[0];
46+
int *local = ptr_array[1];
4347
});
48+
4449
const int some_const = 10;
4550
test(some_const);
4651
return 0;
@@ -66,9 +71,18 @@ int main() {
6671
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_' 'int'
6772

6873
// Check kernel parameters
69-
// CHECK: {{.*}}kernel_struct{{.*}} 'void (int, __wrapper_class)'
74+
// CHECK: {{.*}}kernel_struct{{.*}} 'void (int, __wrapper_class, __wrapper_class, __wrapper_class
75+
// CHECK-SAME: __wrapper_class, __wrapper_class, __wrapper_class, __wrapper_class, __wrapper_class, __wrapper_class)'
7076
// CHECK: ParmVarDecl {{.*}} used _arg_data 'int'
7177
// CHECK: ParmVarDecl {{.*}} used _arg_ptr '__wrapper_class'
78+
// CHECK: ParmVarDecl {{.*}} used _arg_ptr_array1 '__wrapper_class'
79+
// CHECK: ParmVarDecl {{.*}} used _arg_ptr_array1 '__wrapper_class'
80+
// CHECK: ParmVarDecl {{.*}} used _arg_ptr_array2 '__wrapper_class'
81+
// CHECK: ParmVarDecl {{.*}} used _arg_ptr_array2 '__wrapper_class'
82+
// CHECK: ParmVarDecl {{.*}} used _arg_ptr_array2 '__wrapper_class'
83+
// CHECK: ParmVarDecl {{.*}} used _arg_ptr_array2 '__wrapper_class'
84+
// CHECK: ParmVarDecl {{.*}} used _arg_ptr_array2 '__wrapper_class'
85+
// CHECK: ParmVarDecl {{.*}} used _arg_ptr_array2 '__wrapper_class'
7286

7387
// Check that lambda field of struct type is initialized
7488
// CHECK: VarDecl {{.*}}'(lambda at {{.*}}built-in-type-kernel-arg.cpp{{.*}})'
@@ -77,20 +91,66 @@ int main() {
7791
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
7892
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_data' 'int'
7993
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <AddressSpaceConversion>
94+
// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' <LValueToRValue>
8095
// CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue . {{.*}}
8196
// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ptr' '__wrapper_class'
97+
// CHECK-NEXT: InitListExpr {{.*}} 'int *[2]'
98+
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <AddressSpaceConversion>
99+
// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' <LValueToRValue>
100+
// CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue . {{.*}}
101+
// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ptr_array1' '__wrapper_class'
102+
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <AddressSpaceConversion>
103+
// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' <LValueToRValue>
104+
// CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue . {{.*}}
105+
// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ptr_array1' '__wrapper_class'
106+
// CHECK-NEXT: InitListExpr {{.*}} 'int *[2][3]'
107+
// CHECK-NEXT: InitListExpr {{.*}} 'int *[3]'
108+
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <AddressSpaceConversion>
109+
// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' <LValueToRValue>
110+
// CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue . {{.*}}
111+
// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ptr_array2' '__wrapper_class'
112+
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <AddressSpaceConversion>
113+
// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' <LValueToRValue>
114+
// CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue . {{.*}}
115+
// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ptr_array2' '__wrapper_class'
116+
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <AddressSpaceConversion>
117+
// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' <LValueToRValue>
118+
// CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue . {{.*}}
119+
// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ptr_array2' '__wrapper_class'
120+
// CHECK-NEXT: InitListExpr {{.*}} 'int *[3]'
121+
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <AddressSpaceConversion>
122+
// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' <LValueToRValue>
123+
// CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue . {{.*}}
124+
// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ptr_array2' '__wrapper_class'
125+
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <AddressSpaceConversion>
126+
// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' <LValueToRValue>
127+
// CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue . {{.*}}
128+
// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ptr_array2' '__wrapper_class'
129+
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <AddressSpaceConversion>
130+
// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' <LValueToRValue>
131+
// CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue . {{.*}}
132+
// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ptr_array2' '__wrapper_class'
82133

83134
// Check kernel parameters
84-
// CHECK: {{.*}}kernel_pointer{{.*}} 'void (__global int *, __global int *)'
135+
// CHECK: {{.*}}kernel_pointer{{.*}} 'void (__global int *, __global int *, __global int *, __global int *)'
136+
// CHECK: ParmVarDecl {{.*}} used _arg_ '__global int *'
137+
// CHECK: ParmVarDecl {{.*}} used _arg_ '__global int *'
85138
// CHECK: ParmVarDecl {{.*}} used _arg_ '__global int *'
86139
// CHECK: ParmVarDecl {{.*}} used _arg_ '__global int *'
87140
// CHECK: VarDecl {{.*}}'(lambda at {{.*}}built-in-type-kernel-arg.cpp{{.*}})'
88141

89142
// Check that lambda fields of pointer types are initialized
90143
// CHECK: InitListExpr
91144
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <AddressSpaceConversion>
145+
// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' <LValueToRValue>
92146
// CHECK-NEXT: DeclRefExpr {{.*}} '__global int *' lvalue ParmVar {{.*}} '_arg_' '__global int *'
93147
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <AddressSpaceConversion>
148+
// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' <LValueToRValue>
149+
// CHECK-NEXT: DeclRefExpr {{.*}} '__global int *' lvalue ParmVar {{.*}} '_arg_' '__global int *'
150+
// CHECK: InitListExpr {{.*}} 'int *[2]'
151+
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <AddressSpaceConversion>
152+
// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' <LValueToRValue>
153+
// CHECK-NEXT: DeclRefExpr {{.*}} '__global int *' lvalue ParmVar {{.*}} '_arg_' '__global int *'
154+
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <AddressSpaceConversion>
155+
// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' <LValueToRValue>
94156
// CHECK-NEXT: DeclRefExpr {{.*}} '__global int *' lvalue ParmVar {{.*}} '_arg_' '__global int *'
95-
96-
// Check kernel parameters

‎libclc/generic/lib/workitem/get_global_offset.cl

+1-1
Original file line numberDiff line numberDiff line change
@@ -9,7 +9,7 @@
99
#include <clc/clc.h>
1010
#include <spirv/spirv.h>
1111

12-
_CLC_DEF size_t get_global_offset(uint dim) {
12+
_CLC_DEF _CLC_OVERLOAD size_t get_global_offset(uint dim) {
1313
switch (dim) {
1414
case 0: return __spirv_GlobalOffset_x();
1515
case 1: return __spirv_GlobalOffset_y();

‎libclc/generic/lib/workitem/get_group_id.cl

+1-1
Original file line numberDiff line numberDiff line change
@@ -9,7 +9,7 @@
99
#include <clc/clc.h>
1010
#include <spirv/spirv.h>
1111

12-
_CLC_DEF size_t get_group_id(uint dim) {
12+
_CLC_DEF _CLC_OVERLOAD size_t get_group_id(uint dim) {
1313
switch (dim) {
1414
case 0: return __spirv_WorkgroupId_x();
1515
case 1: return __spirv_WorkgroupId_y();

‎libclc/generic/lib/workitem/get_local_id.cl

+1-1
Original file line numberDiff line numberDiff line change
@@ -9,7 +9,7 @@
99
#include <clc/clc.h>
1010
#include <spirv/spirv.h>
1111

12-
_CLC_DEF size_t get_local_id(uint dim) {
12+
_CLC_DEF _CLC_OVERLOAD size_t get_local_id(uint dim) {
1313
switch (dim) {
1414
case 0: return __spirv_LocalInvocationId_x();
1515
case 1: return __spirv_LocalInvocationId_y();

0 commit comments

Comments
 (0)