Skip to content

Commit 0c7a1e1

Browse files
baderbader
authored
[SYCL] Refactor SYCL kernel object handling in hierarchical parallelism (#6212)
This patch refactors #1455 to avoid uses of deprecated `getPointerElementType` function. #1455 introduces the code that uses pointer type information to create a shadow copy of SYCL kernel object. The same can be achieved by applying `work-group` scope attribute the SYCL kernel object. Compiler allocates such object in local address space, so object is shared among all work-items in the work-group.
1 parent 4043dda commit 0c7a1e1

24 files changed

+186
-213
lines changed

‎clang/lib/Sema/SemaSYCL.cpp

Lines changed: 45 additions & 37 deletions
Original file line numberDiff line numberDiff line change
@@ -2505,38 +2505,44 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
25052505
return CompoundStmt::Create(SemaRef.getASTContext(), BodyStmts, {}, {});
25062506
}
25072507

2508-
void markParallelWorkItemCalls() {
2509-
if (getKernelInvocationKind(KernelCallerFunc) ==
2510-
InvokeParallelForWorkGroup) {
2511-
// Fetch the kernel object and the associated call operator
2512-
// (of either the lambda or the function object).
2513-
CXXRecordDecl *KernelObj =
2514-
GetSYCLKernelObjectType(KernelCallerFunc)->getAsCXXRecordDecl();
2515-
CXXMethodDecl *WGLambdaFn = nullptr;
2516-
if (KernelObj->isLambda())
2517-
WGLambdaFn = KernelObj->getLambdaCallOperator();
2518-
else
2519-
WGLambdaFn = getOperatorParens(KernelObj);
2520-
assert(WGLambdaFn && "non callable object is passed as kernel obj");
2521-
// Mark the function that it "works" in a work group scope:
2522-
// NOTE: In case of parallel_for_work_item the marker call itself is
2523-
// marked with work item scope attribute, here the '()' operator of the
2524-
// object passed as parameter is marked. This is an optimization -
2525-
// there are a lot of locals created at parallel_for_work_group
2526-
// scope before calling the lambda - it is more efficient to have
2527-
// all of them in the private address space rather then sharing via
2528-
// the local AS. See parallel_for_work_group implementation in the
2529-
// SYCL headers.
2530-
if (!WGLambdaFn->hasAttr<SYCLScopeAttr>()) {
2531-
WGLambdaFn->addAttr(SYCLScopeAttr::CreateImplicit(
2532-
SemaRef.getASTContext(), SYCLScopeAttr::Level::WorkGroup));
2533-
// Search and mark parallel_for_work_item calls:
2534-
MarkWIScopeFnVisitor MarkWIScope(SemaRef.getASTContext());
2535-
MarkWIScope.TraverseDecl(WGLambdaFn);
2536-
// Now mark local variables declared in the PFWG lambda with work group
2537-
// scope attribute
2538-
addScopeAttrToLocalVars(*WGLambdaFn);
2539-
}
2508+
void annotateHierarchicalParallelismAPICalls() {
2509+
// Is this a hierarchical parallelism kernel invocation?
2510+
if (getKernelInvocationKind(KernelCallerFunc) != InvokeParallelForWorkGroup)
2511+
return;
2512+
2513+
// Mark kernel object with work-group scope attribute to avoid work-item
2514+
// scope memory allocation.
2515+
KernelObjClone->addAttr(SYCLScopeAttr::CreateImplicit(
2516+
SemaRef.getASTContext(), SYCLScopeAttr::Level::WorkGroup));
2517+
2518+
// Fetch the kernel object and the associated call operator
2519+
// (of either the lambda or the function object).
2520+
CXXRecordDecl *KernelObj =
2521+
GetSYCLKernelObjectType(KernelCallerFunc)->getAsCXXRecordDecl();
2522+
CXXMethodDecl *WGLambdaFn = nullptr;
2523+
if (KernelObj->isLambda())
2524+
WGLambdaFn = KernelObj->getLambdaCallOperator();
2525+
else
2526+
WGLambdaFn = getOperatorParens(KernelObj);
2527+
assert(WGLambdaFn && "non callable object is passed as kernel obj");
2528+
// Mark the function that it "works" in a work group scope:
2529+
// NOTE: In case of parallel_for_work_item the marker call itself is
2530+
// marked with work item scope attribute, here the '()' operator of the
2531+
// object passed as parameter is marked. This is an optimization -
2532+
// there are a lot of locals created at parallel_for_work_group
2533+
// scope before calling the lambda - it is more efficient to have
2534+
// all of them in the private address space rather then sharing via
2535+
// the local AS. See parallel_for_work_group implementation in the
2536+
// SYCL headers.
2537+
if (!WGLambdaFn->hasAttr<SYCLScopeAttr>()) {
2538+
WGLambdaFn->addAttr(SYCLScopeAttr::CreateImplicit(
2539+
SemaRef.getASTContext(), SYCLScopeAttr::Level::WorkGroup));
2540+
// Search and mark parallel_for_work_item calls:
2541+
MarkWIScopeFnVisitor MarkWIScope(SemaRef.getASTContext());
2542+
MarkWIScope.TraverseDecl(WGLambdaFn);
2543+
// Now mark local variables declared in the PFWG lambda with work group
2544+
// scope attribute
2545+
addScopeAttrToLocalVars(*WGLambdaFn);
25402546
}
25412547
}
25422548

@@ -2766,11 +2772,13 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
27662772
const CXXRecordDecl *KernelObj) {
27672773
TypeSourceInfo *TSInfo =
27682774
KernelObj->isLambda() ? KernelObj->getLambdaTypeInfo() : nullptr;
2769-
VarDecl *VD = VarDecl::Create(
2770-
Ctx, DC, KernelObj->getLocation(), KernelObj->getLocation(),
2771-
KernelObj->getIdentifier(), QualType(KernelObj->getTypeForDecl(), 0),
2772-
TSInfo, SC_None);
2775+
IdentifierInfo *Ident = KernelObj->getIdentifier();
2776+
if (!Ident)
2777+
Ident = &Ctx.Idents.get("__SYCLKernel");
27732778

2779+
VarDecl *VD = VarDecl::Create(
2780+
Ctx, DC, KernelObj->getLocation(), KernelObj->getLocation(), Ident,
2781+
QualType(KernelObj->getTypeForDecl(), 0), TSInfo, SC_None);
27742782
return VD;
27752783
}
27762784

@@ -2851,7 +2859,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
28512859
KernelObj(KernelObj), KernelCallerFunc(KernelCallerFunc),
28522860
KernelCallerSrcLoc(KernelCallerFunc->getLocation()) {
28532861
CollectionInitExprs.push_back(createInitListExpr(KernelObj));
2854-
markParallelWorkItemCalls();
2862+
annotateHierarchicalParallelismAPICalls();
28552863

28562864
Stmt *DS = new (S.Context) DeclStmt(DeclGroupRef(KernelObjClone),
28572865
KernelCallerSrcLoc, KernelCallerSrcLoc);

‎clang/test/CodeGenSYCL/accessor_inheritance.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -50,7 +50,7 @@ int main() {
5050
// CHECK: [[ARG_C]].addr.ascast = addrspacecast ptr [[ARG_C]].addr to ptr addrspace(4)
5151
//
5252
// Lambda object alloca
53-
// CHECK: [[KERNEL_OBJ:%[a-zA-Z0-9_]+]] = addrspacecast ptr [[KERNEL]] to ptr addrspace(4)
53+
// CHECK: [[KERNEL_OBJ:%[a-zA-Z0-9_.]+]] = addrspacecast ptr [[KERNEL]] to ptr addrspace(4)
5454
//
5555
// Kernel argument stores
5656
// CHECK: store i32 [[ARG_A]], ptr addrspace(4) [[ARG_A]].addr.ascast

‎clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -27,12 +27,12 @@ int main() {
2727
// Check alloca for pointer argument
2828
// CHECK: [[MEM_ARG]].addr = alloca ptr addrspace(1)
2929
// Check lambda object alloca
30-
// CHECK: [[ANONALLOCA:%[0-9]+]] = alloca %class.anon
30+
// CHECK: [[ANONALLOCA:%[a-zA-Z0-9_]+]] = alloca %class.anon
3131
// Check allocas for ranges
3232
// CHECK: [[ARANGEA:%agg.tmp.*]] = alloca %"struct.cl::sycl::range"
3333
// CHECK: [[MRANGEA:%agg.tmp.*]] = alloca %"struct.cl::sycl::range"
3434
// CHECK: [[OIDA:%agg.tmp.*]] = alloca %"struct.cl::sycl::id"
35-
// CHECK: [[ANON:%[0-9]+]] = addrspacecast ptr [[ANONALLOCA]] to ptr addrspace(4)
35+
// CHECK: [[ANON:%[a-zA-Z0-9_.]+]] = addrspacecast ptr [[ANONALLOCA]] to ptr addrspace(4)
3636
// CHECK: [[ARANGET:%agg.tmp.*]] = addrspacecast ptr [[ARANGEA]] to ptr addrspace(4)
3737
// CHECK: [[MRANGET:%agg.tmp.*]] = addrspacecast ptr [[MRANGEA]] to ptr addrspace(4)
3838
// CHECK: [[OIDT:%agg.tmp.*]] = addrspacecast ptr [[OIDA]] to ptr addrspace(4)

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

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -25,7 +25,7 @@ int main() {
2525
}
2626

2727
// CHECK: define{{.*}} spir_kernel {{.*}}19use_kernel_for_test({{.*}}){{.*}} !dbg [[KERNEL:![0-9]+]] {{.*}}{
28-
// CHECK: getelementptr inbounds %class.anon, %class.anon addrspace(4)* %{{[0-9]+}}, i32 0, i32 0, !dbg [[LINE_A0:![0-9]+]]
28+
// CHECK: getelementptr inbounds %class.anon, %class.anon addrspace(4)* %{{.*}}, i32 0, i32 0, !dbg [[LINE_A0:![0-9]+]]
2929
// CHECK: call spir_func void {{.*}}6__init{{.*}} !dbg [[LINE_A0]]
3030
// CHECK: call spir_func void @_ZZ4mainENKUlvE_clEv{{.*}} !dbg [[LINE_B0:![0-9]+]]
3131
// CHECK: ret void, !dbg [[LINE_C0:![0-9]+]]

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

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -39,7 +39,7 @@ int main() {
3939
// CHECK: [[MEM_ARG2:%[a-zA-Z0-9_.]+]] = alloca ptr addrspace(1), align 8
4040

4141
// CHECK lambda object alloca
42-
// CHECK: [[LOCAL_OBJECTA:%0]] = alloca %class.anon, align 4
42+
// CHECK: [[LOCAL_OBJECTA:%__SYCLKernel]] = alloca %class.anon, align 4
4343

4444
// CHECK allocas for ranges
4545
// CHECK: [[ACC_RANGE1A:%[a-zA-Z0-9_.]+]] = alloca %"struct.cl::sycl::range"

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

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -42,7 +42,7 @@ int main() {
4242
// CHECK: [[MEM_ARG1]].addr{{[0-9]*}} = alloca ptr addrspace(1), align 8
4343

4444
// Check lambda object alloca
45-
// CHECK: [[LOCAL_OBJECTA:%0]] = alloca %class{{.*}}.anon, align 4
45+
// CHECK: [[LOCAL_OBJECTA:%__SYCLKernel]] = alloca %class{{.*}}.anon, align 4
4646

4747
// Check allocas for ranges
4848
// CHECK: [[ACC_RANGE1A:%[a-zA-Z0-9_.]+]] = alloca %"struct.cl::sycl::range"
@@ -53,7 +53,7 @@ int main() {
5353
// CHECK: [[OFFSET2A:%[a-zA-Z0-9_.]+]] = alloca %"struct.cl::sycl::id"
5454

5555
// Check lambda object addrspacecast
56-
// CHECK: [[LOCAL_OBJECT:%.*]] = addrspacecast ptr %0 to ptr addrspace(4)
56+
// CHECK: [[LOCAL_OBJECT:%.*]] = addrspacecast ptr %__SYCLKernel to ptr addrspace(4)
5757

5858
// Check addrspacecast for ranges
5959
// CHECK: [[ACC_RANGE1AS:%.*]] = addrspacecast ptr [[ACC_RANGE1A]] to ptr addrspace(4)

‎clang/test/CodeGenSYCL/kernel-param-pod-array.cpp

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -49,8 +49,8 @@ int main() {
4949
// CHECK-SAME:(ptr noundef byval(%struct{{.*}}.__wrapper_class) align 4 %[[ARR_ARG:.*]])
5050

5151
// Check local lambda object alloca
52-
// CHECK: %[[LOCAL_OBJECTA:[0-9]+]] = alloca %class{{.*}}.anon, align 4
53-
// CHECK: %[[LOCAL_OBJECT:[0-9]+]] = addrspacecast ptr %[[LOCAL_OBJECTA]] to ptr addrspace(4)
52+
// CHECK: %[[LOCAL_OBJECTA:[a-zA-Z0-9_]+]] = alloca %class{{.*}}.anon, align 4
53+
// CHECK: %[[LOCAL_OBJECT:[a-zA-Z0-9_.]+]] = addrspacecast ptr %[[LOCAL_OBJECTA]] to ptr addrspace(4)
5454

5555
// Check for Array init loop
5656
// CHECK: %[[LAMBDA_PTR:.+]] = getelementptr inbounds %class{{.*}}.anon, ptr addrspace(4) %[[LOCAL_OBJECT]], i32 0, i32 0
@@ -74,8 +74,8 @@ int main() {
7474
// CHECK-SAME:(ptr noundef byval(%struct{{.*}}.__wrapper_class{{.*}}) align 4 %[[ARR_ARG:.*]])
7575

7676
// Check local lambda object alloca
77-
// CHECK: %[[LOCAL_OBJECTA:[0-9]+]] = alloca %class{{.*}}.anon{{.*}}, align 4
78-
// CHECK: %[[LOCAL_OBJECT:[0-9]+]] = addrspacecast ptr %[[LOCAL_OBJECTA]] to ptr addrspace(4)
77+
// CHECK: %[[LOCAL_OBJECTA:[a-zA-Z0-9_]+]] = alloca %class{{.*}}.anon{{.*}}, align 4
78+
// CHECK: %[[LOCAL_OBJECT:[a-zA-Z0-9_.]+]] = addrspacecast ptr %[[LOCAL_OBJECTA]] to ptr addrspace(4)
7979

8080
// Check for Array init loop
8181
// CHECK: %[[LAMBDA_PTR:.+]] = getelementptr inbounds %class{{.*}}.anon{{.*}}, ptr addrspace(4) %[[LOCAL_OBJECT]], i32 0, i32 0
@@ -98,8 +98,8 @@ int main() {
9898
// CHECK-SAME:(ptr noundef byval(%struct{{.*}}.__wrapper_class{{.*}}) align 4 %[[ARR_ARG:.*]])
9999

100100
// Check local lambda object alloca
101-
// CHECK: %[[LOCAL_OBJECTA:[0-9]+]] = alloca %class{{.*}}.anon{{.*}}, align 4
102-
// CHECK: %[[LOCAL_OBJECT:[0-9]+]] = addrspacecast ptr %[[LOCAL_OBJECTA]] to ptr addrspace(4)
101+
// CHECK: %[[LOCAL_OBJECTA:[a-zA-Z0-9_]+]] = alloca %class{{.*}}.anon{{.*}}, align 4
102+
// CHECK: %[[LOCAL_OBJECT:[a-zA-Z0-9_.]+]] = addrspacecast ptr %[[LOCAL_OBJECTA]] to ptr addrspace(4)
103103

104104
// Check for Array init loop
105105
// CHECK: %[[LAMBDA_PTR:.+]] = getelementptr inbounds %class{{.*}}.anon{{.*}}, ptr addrspace(4) %[[LOCAL_OBJECT]], i32 0, i32 0

‎clang/test/CodeGenSYCL/no_opaque_accessor_inheritance.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -50,7 +50,7 @@ int main() {
5050
// CHECK: [[ARG_C]].addr.ascast = addrspacecast i32* [[ARG_C]].addr to i32 addrspace(4)*
5151
//
5252
// Lambda object alloca
53-
// CHECK: [[KERNEL_OBJ:%[a-zA-Z0-9_]+]] = addrspacecast %class{{.*}}.anon* [[KERNEL]] to %class{{.*}}.anon addrspace(4)*
53+
// CHECK: [[KERNEL_OBJ:%[a-zA-Z0-9_.]+]] = addrspacecast %class{{.*}}.anon* [[KERNEL]] to %class{{.*}}.anon addrspace(4)*
5454
//
5555
// Kernel argument stores
5656
// CHECK: store i32 [[ARG_A]], i32 addrspace(4)* [[ARG_A]].addr.ascast

‎clang/test/CodeGenSYCL/no_opaque_basic-kernel-wrapper.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -27,12 +27,12 @@ int main() {
2727
// Check alloca for pointer argument
2828
// CHECK: [[MEM_ARG]].addr = alloca i32 addrspace(1)*
2929
// Check lambda object alloca
30-
// CHECK: [[ANONALLOCA:%[0-9]+]] = alloca %class.anon
30+
// CHECK: [[ANONALLOCA:%[a-zA-Z0-9_]+]] = alloca %class.anon
3131
// Check allocas for ranges
3232
// CHECK: [[ARANGEA:%agg.tmp.*]] = alloca %"struct.cl::sycl::range"
3333
// CHECK: [[MRANGEA:%agg.tmp.*]] = alloca %"struct.cl::sycl::range"
3434
// CHECK: [[OIDA:%agg.tmp.*]] = alloca %"struct.cl::sycl::id"
35-
// CHECK: [[ANON:%[0-9]+]] = addrspacecast %class.anon* [[ANONALLOCA]] to %class.anon addrspace(4)*
35+
// CHECK: [[ANON:%[a-zA-Z0-9_.]+]] = addrspacecast %class.anon* [[ANONALLOCA]] to %class.anon addrspace(4)*
3636
// CHECK: [[ARANGET:%agg.tmp.*]] = addrspacecast %"struct.cl::sycl::range"* [[ARANGEA]] to %"struct.cl::sycl::range" addrspace(4)*
3737
// CHECK: [[MRANGET:%agg.tmp.*]] = addrspacecast %"struct.cl::sycl::range"* [[MRANGEA]] to %"struct.cl::sycl::range" addrspace(4)*
3838
// CHECK: [[OIDT:%agg.tmp.*]] = addrspacecast %"struct.cl::sycl::id"* [[OIDA]] to %"struct.cl::sycl::id" addrspace(4)*

‎clang/test/CodeGenSYCL/no_opaque_kernel-param-acc-array.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -39,7 +39,7 @@ int main() {
3939
// CHECK: [[MEM_ARG2:%[a-zA-Z0-9_.]+]] = alloca i32 addrspace(1)*, align 8
4040

4141
// CHECK lambda object alloca
42-
// CHECK: [[LOCAL_OBJECTA:%0]] = alloca %class.anon, align 4
42+
// CHECK: [[LOCAL_OBJECTA:%__SYCLKernel]] = alloca %class.anon, align 4
4343

4444
// CHECK allocas for ranges
4545
// CHECK: [[ACC_RANGE1A:%[a-zA-Z0-9_.]+]] = alloca %"struct.cl::sycl::range"

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

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -42,7 +42,7 @@ int main() {
4242
// CHECK: [[MEM_ARG1]].addr{{[0-9]*}} = alloca i32 addrspace(1)*, align 8
4343

4444
// Check lambda object alloca
45-
// CHECK: [[LOCAL_OBJECTA:%0]] = alloca %class{{.*}}.anon, align 4
45+
// CHECK: [[LOCAL_OBJECTA:%__SYCLKernel]] = alloca %class{{.*}}.anon, align 4
4646

4747
// Check allocas for ranges
4848
// CHECK: [[ACC_RANGE1A:%[a-zA-Z0-9_.]+]] = alloca %"struct.cl::sycl::range"
@@ -53,7 +53,7 @@ int main() {
5353
// CHECK: [[OFFSET2A:%[a-zA-Z0-9_.]+]] = alloca %"struct.cl::sycl::id"
5454

5555
// Check lambda object addrspacecast
56-
// CHECK: [[LOCAL_OBJECT:%.*]] = addrspacecast %class{{.*}}.anon* %0 to %class{{.*}}.anon addrspace(4)*
56+
// CHECK: [[LOCAL_OBJECT:%.*]] = addrspacecast %class{{.*}}.anon* %__SYCLKernel to %class{{.*}}.anon addrspace(4)*
5757

5858
// Check addrspacecast for ranges
5959
// CHECK: [[ACC_RANGE1AS:%.*]] = addrspacecast %"struct.cl::sycl::range"* [[ACC_RANGE1A]] to %"struct.cl::sycl::range" addrspace(4)*

‎clang/test/CodeGenSYCL/no_opaque_kernel-param-pod-array.cpp

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -49,8 +49,8 @@ int main() {
4949
// CHECK-SAME:(%struct{{.*}}.__wrapper_class* noundef byval(%struct{{.*}}.__wrapper_class) align 4 %[[ARR_ARG:.*]])
5050

5151
// Check local lambda object alloca
52-
// CHECK: %[[LOCAL_OBJECTA:[0-9]+]] = alloca %class{{.*}}.anon, align 4
53-
// CHECK: %[[LOCAL_OBJECT:[0-9]+]] = addrspacecast %class{{.*}}.anon* %[[LOCAL_OBJECTA]] to %class{{.*}}.anon addrspace(4)*
52+
// CHECK: %[[LOCAL_OBJECTA:[a-zA-Z0-9_]+]] = alloca %class{{.*}}.anon, align 4
53+
// CHECK: %[[LOCAL_OBJECT:[a-zA-Z0-9_.]+]] = addrspacecast %class{{.*}}.anon* %[[LOCAL_OBJECTA]] to %class{{.*}}.anon addrspace(4)*
5454

5555
// Check for Array init loop
5656
// CHECK: %[[LAMBDA_PTR:.+]] = getelementptr inbounds %class{{.*}}.anon, %class{{.*}}.anon addrspace(4)* %[[LOCAL_OBJECT]], i32 0, i32 0
@@ -74,8 +74,8 @@ int main() {
7474
// CHECK-SAME:(%struct{{.*}}.__wrapper_class{{.*}}* noundef byval(%struct{{.*}}.__wrapper_class{{.*}}) align 4 %[[ARR_ARG:.*]])
7575

7676
// Check local lambda object alloca
77-
// CHECK: %[[LOCAL_OBJECTA:[0-9]+]] = alloca %class{{.*}}.anon{{.*}}, align 4
78-
// CHECK: %[[LOCAL_OBJECT:[0-9]+]] = addrspacecast %class{{.*}}.anon{{.*}}* %[[LOCAL_OBJECTA]] to %class{{.*}}.anon{{.*}} addrspace(4)*
77+
// CHECK: %[[LOCAL_OBJECTA:[a-zA-Z0-9_]+]] = alloca %class{{.*}}.anon{{.*}}, align 4
78+
// CHECK: %[[LOCAL_OBJECT:[a-zA-Z0-9_.]+]] = addrspacecast %class{{.*}}.anon{{.*}}* %[[LOCAL_OBJECTA]] to %class{{.*}}.anon{{.*}} addrspace(4)*
7979

8080
// Check for Array init loop
8181
// CHECK: %[[LAMBDA_PTR:.+]] = getelementptr inbounds %class{{.*}}.anon{{.*}}, %class{{.*}}.anon{{.*}} addrspace(4)* %[[LOCAL_OBJECT]], i32 0, i32 0
@@ -100,8 +100,8 @@ int main() {
100100
// CHECK-SAME:(%struct{{.*}}.__wrapper_class{{.*}}* noundef byval(%struct{{.*}}.__wrapper_class{{.*}}) align 4 %[[ARR_ARG:.*]])
101101

102102
// Check local lambda object alloca
103-
// CHECK: %[[LOCAL_OBJECTA:[0-9]+]] = alloca %class{{.*}}.anon{{.*}}, align 4
104-
// CHECK: %[[LOCAL_OBJECT:[0-9]+]] = addrspacecast %class{{.*}}.anon{{.*}}* %[[LOCAL_OBJECTA]] to %class{{.*}}.anon{{.*}} addrspace(4)*
103+
// CHECK: %[[LOCAL_OBJECTA:[a-zA-Z0-9_]+]] = alloca %class{{.*}}.anon{{.*}}, align 4
104+
// CHECK: %[[LOCAL_OBJECT:[a-zA-Z0-9_.]+]] = addrspacecast %class{{.*}}.anon{{.*}}* %[[LOCAL_OBJECTA]] to %class{{.*}}.anon{{.*}} addrspace(4)*
105105

106106
// Check for Array init loop
107107
// CHECK: %[[LAMBDA_PTR:.+]] = getelementptr inbounds %class{{.*}}.anon{{.*}}, %class{{.*}}.anon{{.*}} addrspace(4)* %[[LOCAL_OBJECT]], i32 0, i32 0

‎clang/test/CodeGenSYCL/no_opaque_sampler.cpp

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -2,8 +2,8 @@
22
// CHECK: define {{.*}}spir_kernel void @{{[a-zA-Z0-9_]+}}(%opencl.sampler_t addrspace(2)* [[SAMPLER_ARG:%[a-zA-Z0-9_]+]])
33
// CHECK-NEXT: entry:
44
// CHECK-NEXT: [[SAMPLER_ARG]].addr = alloca %opencl.sampler_t addrspace(2)*, align 8
5-
// CHECK: [[ANON:%[0-9]+]] = alloca %class.anon, align 8
6-
// CHECK: [[ANONCAST:%[0-9]+]] = addrspacecast %class.anon* [[ANON]] to %class.anon addrspace(4)*
5+
// CHECK: [[ANON:%[a-zA-Z0-9_]+]] = alloca %class.anon, align 8
6+
// CHECK: [[ANONCAST:%[a-zA-Z0-9_.]+]] = addrspacecast %class.anon* [[ANON]] to %class.anon addrspace(4)*
77
// CHECK: store %opencl.sampler_t addrspace(2)* [[SAMPLER_ARG]], %opencl.sampler_t addrspace(2)* addrspace(4)* [[SAMPLER_ARG]].addr.ascast, align 8
88
// CHECK-NEXT: [[BITCAST:%[0-9]+]] = bitcast %class.anon* [[ANON]] to i8*
99
// CHECK-NEXT: call void @llvm.lifetime.start.p0i8(i64 8, i8* [[BITCAST]]) #4
@@ -17,8 +17,8 @@
1717
// Check alloca
1818
// CHECK: [[SAMPLER_ARG_WRAPPED]].addr = alloca %opencl.sampler_t addrspace(2)*, align 8
1919
// CHECK: [[ARG_A]].addr = alloca i32, align 4
20-
// CHECK: [[LAMBDAA:%[0-9]+]] = alloca %class.anon.0, align 8
21-
// CHECK: [[LAMBDA:%[0-9]+]] = addrspacecast %class.anon.0* [[LAMBDAA]] to %class.anon.0 addrspace(4)*
20+
// CHECK: [[LAMBDAA:%[a-zA-Z0-9_]+]] = alloca %class.anon.0, align 8
21+
// CHECK: [[LAMBDA:%[a-zA-Z0-9_.]+]] = addrspacecast %class.anon.0* [[LAMBDAA]] to %class.anon.0 addrspace(4)*
2222

2323
// Check argument store
2424
// CHECK: store %opencl.sampler_t addrspace(2)* [[SAMPLER_ARG_WRAPPED]], %opencl.sampler_t addrspace(2)* addrspace(4)* [[SAMPLER_ARG_WRAPPED]].addr.ascast, align 8

‎clang/test/CodeGenSYCL/no_opaque_union-kernel-param.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -31,7 +31,7 @@ int main() {
3131
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_A(%union.MyUnion* noundef byval(%union.MyUnion) align 4 [[MEM_ARG:%[a-zA-Z0-9_]+]])
3232

3333
// Check lambda object alloca
34-
// CHECK: [[LOCAL_OBJECT:%0]] = alloca %class.anon, align 4
34+
// CHECK: [[LOCAL_OBJECT:%__SYCLKernel]] = alloca %class.anon, align 4
3535

3636
// CHECK: [[LOCAL_OBJECTAS:%.*]] = addrspacecast %class.anon* [[LOCAL_OBJECT]] to %class.anon addrspace(4)*
3737
// CHECK: [[MEM_ARGAS:%.*]] = addrspacecast %union.MyUnion* [[MEM_ARG]] to %union.MyUnion addrspace(4)*

‎clang/test/CodeGenSYCL/sampler.cpp

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -2,8 +2,8 @@
22
// CHECK: define {{.*}}spir_kernel void @{{[a-zA-Z0-9_]+}}(ptr addrspace(2) [[SAMPLER_ARG:%[a-zA-Z0-9_]+]])
33
// CHECK-NEXT: entry:
44
// CHECK-NEXT: [[SAMPLER_ARG]].addr = alloca ptr addrspace(2), align 8
5-
// CHECK: [[ANON:%[0-9]+]] = alloca %class.anon, align 8
6-
// CHECK: [[ANONCAST:%[0-9]+]] = addrspacecast ptr [[ANON]] to ptr addrspace(4)
5+
// CHECK: [[ANON:%[a-zA-Z0-9_]+]] = alloca %class.anon, align 8
6+
// CHECK: [[ANONCAST:%[a-zA-Z0-9_.]+]] = addrspacecast ptr [[ANON]] to ptr addrspace(4)
77
// CHECK: store ptr addrspace(2) [[SAMPLER_ARG]], ptr addrspace(4) [[SAMPLER_ARG]].addr.ascast, align 8
88
// CHECK-NEXT: call void @llvm.lifetime.start.p0(i64 8, ptr [[ANON]]) #4
99
// CHECK-NEXT: [[GEP:%[a-zA-z0-9]+]] = getelementptr inbounds %class.anon, ptr addrspace(4) [[ANONCAST]], i32 0, i32 0
@@ -16,8 +16,8 @@
1616
// Check alloca
1717
// CHECK: [[SAMPLER_ARG_WRAPPED]].addr = alloca ptr addrspace(2), align 8
1818
// CHECK: [[ARG_A]].addr = alloca i32, align 4
19-
// CHECK: [[LAMBDAA:%[0-9]+]] = alloca %class.anon.0, align 8
20-
// CHECK: [[LAMBDA:%[0-9]+]] = addrspacecast ptr [[LAMBDAA]] to ptr addrspace(4)
19+
// CHECK: [[LAMBDAA:%[a-zA-Z0-9_]+]] = alloca %class.anon.0, align 8
20+
// CHECK: [[LAMBDA:%[a-zA-Z0-9_.]+]] = addrspacecast ptr [[LAMBDAA]] to ptr addrspace(4)
2121

2222
// Check argument store
2323
// CHECK: store ptr addrspace(2) [[SAMPLER_ARG_WRAPPED]], ptr addrspace(4) [[SAMPLER_ARG_WRAPPED]].addr.ascast, align 8

‎clang/test/CodeGenSYCL/union-kernel-param.cpp

Lines changed: 1 addition & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -2,10 +2,6 @@
22

33
// This test checks a kernel argument that is union with both array and non-array fields.
44

5-
#include "Inputs/sycl.hpp"
6-
7-
using namespace cl::sycl;
8-
95
union MyUnion {
106
int FldInt;
117
char FldChar;
@@ -31,7 +27,7 @@ int main() {
3127
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_A(ptr noundef byval(%union.MyUnion) align 4 [[MEM_ARG:%[a-zA-Z0-9_]+]])
3228

3329
// Check lambda object alloca
34-
// CHECK: [[LOCAL_OBJECT:%0]] = alloca %class.anon, align 4
30+
// CHECK: [[LOCAL_OBJECT:%__SYCLKernel]] = alloca %class.anon, align 4
3531

3632
// CHECK: [[LOCAL_OBJECTAS:%.*]] = addrspacecast ptr [[LOCAL_OBJECT]] to ptr addrspace(4)
3733
// CHECK: [[MEM_ARGAS:%.*]] = addrspacecast ptr [[MEM_ARG]] to ptr addrspace(4)

0 commit comments

Comments
 (0)