Skip to content

[SYCL] Refactor SYCL kernel object handling in hierarchical parallelism #6212

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
bader merged 15 commits into from
Jun 15, 2022
Merged

[SYCL] Refactor SYCL kernel object handling in hierarchical parallelism #6212

Show file tree
Hide file tree
Changes from 1 commit
Commits
Show all changes
15 commits
Select commit Hold shift + click to select a range
eec1da1
Revert the logic committed by 16f64b8ff1f5fe44ff13de6c5d8fbd3818f81d43
bader May 26, 2022
6c60344
Update pass tests.
bader May 29, 2022
3530e1d
Improve debug infomation for LowerWGScope pass.
bader May 29, 2022
0e13ee9
Mark auto-generated kernel with work-group metadata.
bader May 29, 2022
e1511a7
Fix formatting.
bader May 29, 2022
eb6d2d7
Add TODO comment for a future improvement.
bader May 31, 2022
969b468
Revert "Mark auto-generated kernel with work-group metadata."
bader May 31, 2022
ca26f57
Add work-group scope attribute to the SYCL kernel
bader Jun 7, 2022
1a61c46
clang-format
bader Jun 7, 2022
b970a77
Set anonymous struct flag for lambda types only. Small refactoring.
bader Jun 12, 2022
3886d8d
Remove unused function parameter.
bader Jun 12, 2022
959e1aa
Apply code review suggestion.
bader Jun 13, 2022
2d987ac
Give a name to SYCL kernel object to make mangle-able static version …
bader Jun 14, 2022
fdeaab6
Apply code review comments.
bader Jun 14, 2022
0a5834b
Change getOwn to get.
bader Jun 14, 2022
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
12 changes: 8 additions & 4 deletions clang/lib/Sema/SemaSYCL.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2772,11 +2772,15 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
const CXXRecordDecl *KernelObj) {
TypeSourceInfo *TSInfo =
KernelObj->isLambda() ? KernelObj->getLambdaTypeInfo() : nullptr;
auto Type = QualType(KernelObj->getTypeForDecl(), 0);
Type->getAsRecordDecl()->setAnonymousStructOrUnion(KernelObj->isLambda());
// Generate a name kernel object variable.
Twine Name = Twine("__SYCLKernel");
SmallString<12> Buffer;
StringRef NameRef = Name.toStringRef(Buffer);
auto &Ident = Ctx.Idents.getOwn(NameRef);

VarDecl *VD = VarDecl::Create(
Ctx, DC, KernelObj->getLocation(), KernelObj->getLocation(),
KernelObj->getIdentifier(), Type, TSInfo, SC_None);
Ctx, DC, KernelObj->getLocation(), KernelObj->getLocation(), &Ident,
QualType(KernelObj->getTypeForDecl(), 0), TSInfo, SC_None);
return VD;
}

Expand Down
2 changes: 1 addition & 1 deletion clang/test/CodeGenSYCL/accessor_inheritance.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -50,7 +50,7 @@ int main() {
// CHECK: [[ARG_C]].addr.ascast = addrspacecast ptr [[ARG_C]].addr to ptr addrspace(4)
//
// Lambda object alloca
// CHECK: [[KERNEL_OBJ:%[a-zA-Z0-9_]+]] = addrspacecast ptr [[KERNEL]] to ptr addrspace(4)
// CHECK: [[KERNEL_OBJ:%[a-zA-Z0-9_.]+]] = addrspacecast ptr [[KERNEL]] to ptr addrspace(4)
//
// Kernel argument stores
// CHECK: store i32 [[ARG_A]], ptr addrspace(4) [[ARG_A]].addr.ascast
Expand Down
4 changes: 2 additions & 2 deletions clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -27,12 +27,12 @@ int main() {
// Check alloca for pointer argument
// CHECK: [[MEM_ARG]].addr = alloca ptr addrspace(1)
// Check lambda object alloca
// CHECK: [[ANONALLOCA:%[0-9]+]] = alloca %class.anon
// CHECK: [[ANONALLOCA:%[a-zA-Z0-9_]+]] = alloca %class.anon
// Check allocas for ranges
// CHECK: [[ARANGEA:%agg.tmp.*]] = alloca %"struct.cl::sycl::range"
// CHECK: [[MRANGEA:%agg.tmp.*]] = alloca %"struct.cl::sycl::range"
// CHECK: [[OIDA:%agg.tmp.*]] = alloca %"struct.cl::sycl::id"
// CHECK: [[ANON:%[0-9]+]] = addrspacecast ptr [[ANONALLOCA]] to ptr addrspace(4)
// CHECK: [[ANON:%[a-zA-Z0-9_.]+]] = addrspacecast ptr [[ANONALLOCA]] to ptr addrspace(4)
// CHECK: [[ARANGET:%agg.tmp.*]] = addrspacecast ptr [[ARANGEA]] to ptr addrspace(4)
// CHECK: [[MRANGET:%agg.tmp.*]] = addrspacecast ptr [[MRANGEA]] to ptr addrspace(4)
// CHECK: [[OIDT:%agg.tmp.*]] = addrspacecast ptr [[OIDA]] to ptr addrspace(4)
Expand Down
2 changes: 1 addition & 1 deletion clang/test/CodeGenSYCL/debug-info-srcpos-kernel.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -25,7 +25,7 @@ int main() {
}

// CHECK: define{{.*}} spir_kernel {{.*}}19use_kernel_for_test({{.*}}){{.*}} !dbg [[KERNEL:![0-9]+]] {{.*}}{
// CHECK: getelementptr inbounds %class.anon, %class.anon addrspace(4)* %{{[0-9]+}}, i32 0, i32 0, !dbg [[LINE_A0:![0-9]+]]
// CHECK: getelementptr inbounds %class.anon, %class.anon addrspace(4)* %{{.*}}, i32 0, i32 0, !dbg [[LINE_A0:![0-9]+]]
// CHECK: call spir_func void {{.*}}6__init{{.*}} !dbg [[LINE_A0]]
// CHECK: call spir_func void @_ZZ4mainENKUlvE_clEv{{.*}} !dbg [[LINE_B0:![0-9]+]]
// CHECK: ret void, !dbg [[LINE_C0:![0-9]+]]
Expand Down
2 changes: 1 addition & 1 deletion clang/test/CodeGenSYCL/kernel-param-acc-array.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -39,7 +39,7 @@ int main() {
// CHECK: [[MEM_ARG2:%[a-zA-Z0-9_.]+]] = alloca ptr addrspace(1), align 8

// CHECK lambda object alloca
// CHECK: [[LOCAL_OBJECTA:%0]] = alloca %class.anon, align 4
// CHECK: [[LOCAL_OBJECTA:%__SYCLKernel]] = alloca %class.anon, align 4

// CHECK allocas for ranges
// CHECK: [[ACC_RANGE1A:%[a-zA-Z0-9_.]+]] = alloca %"struct.cl::sycl::range"
Expand Down
4 changes: 2 additions & 2 deletions clang/test/CodeGenSYCL/kernel-param-member-acc-array.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -42,7 +42,7 @@ int main() {
// CHECK: [[MEM_ARG1]].addr{{[0-9]*}} = alloca ptr addrspace(1), align 8

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

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

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

// Check addrspacecast for ranges
// CHECK: [[ACC_RANGE1AS:%.*]] = addrspacecast ptr [[ACC_RANGE1A]] to ptr addrspace(4)
Expand Down
12 changes: 6 additions & 6 deletions clang/test/CodeGenSYCL/kernel-param-pod-array.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -49,8 +49,8 @@ int main() {
// CHECK-SAME:(ptr noundef byval(%struct{{.*}}.__wrapper_class) align 4 %[[ARR_ARG:.*]])

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

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

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

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

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

// Check for Array init loop
// CHECK: %[[LAMBDA_PTR:.+]] = getelementptr inbounds %class{{.*}}.anon{{.*}}, ptr addrspace(4) %[[LOCAL_OBJECT]], i32 0, i32 0
Expand Down
2 changes: 1 addition & 1 deletion clang/test/CodeGenSYCL/no_opaque_accessor_inheritance.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -50,7 +50,7 @@ int main() {
// CHECK: [[ARG_C]].addr.ascast = addrspacecast i32* [[ARG_C]].addr to i32 addrspace(4)*
//
// Lambda object alloca
// CHECK: [[KERNEL_OBJ:%[a-zA-Z0-9_]+]] = addrspacecast %class{{.*}}.anon* [[KERNEL]] to %class{{.*}}.anon addrspace(4)*
// CHECK: [[KERNEL_OBJ:%[a-zA-Z0-9_.]+]] = addrspacecast %class{{.*}}.anon* [[KERNEL]] to %class{{.*}}.anon addrspace(4)*
//
// Kernel argument stores
// CHECK: store i32 [[ARG_A]], i32 addrspace(4)* [[ARG_A]].addr.ascast
Expand Down
4 changes: 2 additions & 2 deletions clang/test/CodeGenSYCL/no_opaque_basic-kernel-wrapper.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -27,12 +27,12 @@ int main() {
// Check alloca for pointer argument
// CHECK: [[MEM_ARG]].addr = alloca i32 addrspace(1)*
// Check lambda object alloca
// CHECK: [[ANONALLOCA:%[0-9]+]] = alloca %class.anon
// CHECK: [[ANONALLOCA:%[a-zA-Z0-9_]+]] = alloca %class.anon
// Check allocas for ranges
// CHECK: [[ARANGEA:%agg.tmp.*]] = alloca %"struct.cl::sycl::range"
// CHECK: [[MRANGEA:%agg.tmp.*]] = alloca %"struct.cl::sycl::range"
// CHECK: [[OIDA:%agg.tmp.*]] = alloca %"struct.cl::sycl::id"
// CHECK: [[ANON:%[0-9]+]] = addrspacecast %class.anon* [[ANONALLOCA]] to %class.anon addrspace(4)*
// CHECK: [[ANON:%[a-zA-Z0-9_.]+]] = addrspacecast %class.anon* [[ANONALLOCA]] to %class.anon addrspace(4)*
// CHECK: [[ARANGET:%agg.tmp.*]] = addrspacecast %"struct.cl::sycl::range"* [[ARANGEA]] to %"struct.cl::sycl::range" addrspace(4)*
// CHECK: [[MRANGET:%agg.tmp.*]] = addrspacecast %"struct.cl::sycl::range"* [[MRANGEA]] to %"struct.cl::sycl::range" addrspace(4)*
// CHECK: [[OIDT:%agg.tmp.*]] = addrspacecast %"struct.cl::sycl::id"* [[OIDA]] to %"struct.cl::sycl::id" addrspace(4)*
Expand Down
2 changes: 1 addition & 1 deletion clang/test/CodeGenSYCL/no_opaque_kernel-param-acc-array.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -39,7 +39,7 @@ int main() {
// CHECK: [[MEM_ARG2:%[a-zA-Z0-9_.]+]] = alloca i32 addrspace(1)*, align 8

// CHECK lambda object alloca
// CHECK: [[LOCAL_OBJECTA:%0]] = alloca %class.anon, align 4
// CHECK: [[LOCAL_OBJECTA:%__SYCLKernel]] = alloca %class.anon, align 4

// CHECK allocas for ranges
// CHECK: [[ACC_RANGE1A:%[a-zA-Z0-9_.]+]] = alloca %"struct.cl::sycl::range"
Expand Down
4 changes: 2 additions & 2 deletions clang/test/CodeGenSYCL/no_opaque_kernel-param-member-acc-array.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -42,7 +42,7 @@ int main() {
// CHECK: [[MEM_ARG1]].addr{{[0-9]*}} = alloca i32 addrspace(1)*, align 8

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

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

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

// Check addrspacecast for ranges
// CHECK: [[ACC_RANGE1AS:%.*]] = addrspacecast %"struct.cl::sycl::range"* [[ACC_RANGE1A]] to %"struct.cl::sycl::range" addrspace(4)*
Expand Down
12 changes: 6 additions & 6 deletions clang/test/CodeGenSYCL/no_opaque_kernel-param-pod-array.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -49,8 +49,8 @@ int main() {
// CHECK-SAME:(%struct{{.*}}.__wrapper_class* noundef byval(%struct{{.*}}.__wrapper_class) align 4 %[[ARR_ARG:.*]])

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

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

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

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

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

// Check for Array init loop
// CHECK: %[[LAMBDA_PTR:.+]] = getelementptr inbounds %class{{.*}}.anon{{.*}}, %class{{.*}}.anon{{.*}} addrspace(4)* %[[LOCAL_OBJECT]], i32 0, i32 0
Expand Down
8 changes: 4 additions & 4 deletions clang/test/CodeGenSYCL/no_opaque_sampler.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2,8 +2,8 @@
// CHECK: define {{.*}}spir_kernel void @{{[a-zA-Z0-9_]+}}(%opencl.sampler_t addrspace(2)* [[SAMPLER_ARG:%[a-zA-Z0-9_]+]])
// CHECK-NEXT: entry:
// CHECK-NEXT: [[SAMPLER_ARG]].addr = alloca %opencl.sampler_t addrspace(2)*, align 8
// CHECK: [[ANON:%[0-9]+]] = alloca %class.anon, align 8
// CHECK: [[ANONCAST:%[0-9]+]] = addrspacecast %class.anon* [[ANON]] to %class.anon addrspace(4)*
// CHECK: [[ANON:%[a-zA-Z0-9_]+]] = alloca %class.anon, align 8
// CHECK: [[ANONCAST:%[a-zA-Z0-9_.]+]] = addrspacecast %class.anon* [[ANON]] to %class.anon addrspace(4)*
// CHECK: store %opencl.sampler_t addrspace(2)* [[SAMPLER_ARG]], %opencl.sampler_t addrspace(2)* addrspace(4)* [[SAMPLER_ARG]].addr.ascast, align 8
// CHECK-NEXT: [[BITCAST:%[0-9]+]] = bitcast %class.anon* [[ANON]] to i8*
// CHECK-NEXT: call void @llvm.lifetime.start.p0i8(i64 8, i8* [[BITCAST]]) #4
Expand All @@ -17,8 +17,8 @@
// Check alloca
// CHECK: [[SAMPLER_ARG_WRAPPED]].addr = alloca %opencl.sampler_t addrspace(2)*, align 8
// CHECK: [[ARG_A]].addr = alloca i32, align 4
// CHECK: [[LAMBDAA:%[0-9]+]] = alloca %class.anon.0, align 8
// CHECK: [[LAMBDA:%[0-9]+]] = addrspacecast %class.anon.0* [[LAMBDAA]] to %class.anon.0 addrspace(4)*
// CHECK: [[LAMBDAA:%[a-zA-Z0-9_]+]] = alloca %class.anon.0, align 8
// CHECK: [[LAMBDA:%[a-zA-Z0-9_.]+]] = addrspacecast %class.anon.0* [[LAMBDAA]] to %class.anon.0 addrspace(4)*

// Check argument store
// CHECK: store %opencl.sampler_t addrspace(2)* [[SAMPLER_ARG_WRAPPED]], %opencl.sampler_t addrspace(2)* addrspace(4)* [[SAMPLER_ARG_WRAPPED]].addr.ascast, align 8
Expand Down
2 changes: 1 addition & 1 deletion clang/test/CodeGenSYCL/no_opaque_union-kernel-param.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -31,7 +31,7 @@ int main() {
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_A(%union.MyUnion* noundef byval(%union.MyUnion) align 4 [[MEM_ARG:%[a-zA-Z0-9_]+]])

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

// CHECK: [[LOCAL_OBJECTAS:%.*]] = addrspacecast %class.anon* [[LOCAL_OBJECT]] to %class.anon addrspace(4)*
// CHECK: [[MEM_ARGAS:%.*]] = addrspacecast %union.MyUnion* [[MEM_ARG]] to %union.MyUnion addrspace(4)*
Expand Down
8 changes: 4 additions & 4 deletions clang/test/CodeGenSYCL/sampler.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2,8 +2,8 @@
// CHECK: define {{.*}}spir_kernel void @{{[a-zA-Z0-9_]+}}(ptr addrspace(2) [[SAMPLER_ARG:%[a-zA-Z0-9_]+]])
// CHECK-NEXT: entry:
// CHECK-NEXT: [[SAMPLER_ARG]].addr = alloca ptr addrspace(2), align 8
// CHECK: [[ANON:%[0-9]+]] = alloca %class.anon, align 8
// CHECK: [[ANONCAST:%[0-9]+]] = addrspacecast ptr [[ANON]] to ptr addrspace(4)
// CHECK: [[ANON:%[a-zA-Z0-9_]+]] = alloca %class.anon, align 8
// CHECK: [[ANONCAST:%[a-zA-Z0-9_.]+]] = addrspacecast ptr [[ANON]] to ptr addrspace(4)
// CHECK: store ptr addrspace(2) [[SAMPLER_ARG]], ptr addrspace(4) [[SAMPLER_ARG]].addr.ascast, align 8
// CHECK-NEXT: call void @llvm.lifetime.start.p0(i64 8, ptr [[ANON]]) #4
// CHECK-NEXT: [[GEP:%[a-zA-z0-9]+]] = getelementptr inbounds %class.anon, ptr addrspace(4) [[ANONCAST]], i32 0, i32 0
Expand All @@ -16,8 +16,8 @@
// Check alloca
// CHECK: [[SAMPLER_ARG_WRAPPED]].addr = alloca ptr addrspace(2), align 8
// CHECK: [[ARG_A]].addr = alloca i32, align 4
// CHECK: [[LAMBDAA:%[0-9]+]] = alloca %class.anon.0, align 8
// CHECK: [[LAMBDA:%[0-9]+]] = addrspacecast ptr [[LAMBDAA]] to ptr addrspace(4)
// CHECK: [[LAMBDAA:%[a-zA-Z0-9_]+]] = alloca %class.anon.0, align 8
// CHECK: [[LAMBDA:%[a-zA-Z0-9_.]+]] = addrspacecast ptr [[LAMBDAA]] to ptr addrspace(4)

// Check argument store
// CHECK: store ptr addrspace(2) [[SAMPLER_ARG_WRAPPED]], ptr addrspace(4) [[SAMPLER_ARG_WRAPPED]].addr.ascast, align 8
Expand Down
6 changes: 1 addition & 5 deletions clang/test/CodeGenSYCL/union-kernel-param.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2,10 +2,6 @@

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

#include "Inputs/sycl.hpp"

using namespace cl::sycl;

union MyUnion {
int FldInt;
char FldChar;
Expand All @@ -31,7 +27,7 @@ int main() {
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_A(ptr noundef byval(%union.MyUnion) align 4 [[MEM_ARG:%[a-zA-Z0-9_]+]])

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

// CHECK: [[LOCAL_OBJECTAS:%.*]] = addrspacecast ptr [[LOCAL_OBJECT]] to ptr addrspace(4)
// CHECK: [[MEM_ARGAS:%.*]] = addrspacecast ptr [[MEM_ARG]] to ptr addrspace(4)
Expand Down
2 changes: 1 addition & 1 deletion clang/test/SemaSYCL/accessor_inheritance.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -44,7 +44,7 @@ int main() {
// CHECK: ParmVarDecl{{.*}} used _arg_C 'int'

// Check lambda initialization
// CHECK: VarDecl {{.*}} used '(lambda at {{.*}}accessor_inheritance.cpp
// CHECK: VarDecl {{.*}} used __SYCLKernel '(lambda at {{.*}}accessor_inheritance.cpp
// CHECK-NEXT: InitListExpr {{.*}}
// CHECK-NEXT: InitListExpr {{.*}} 'AccessorDerived'
// CHECK-NEXT: InitListExpr {{.*}} 'AccessorBase'
Expand Down
6 changes: 3 additions & 3 deletions clang/test/SemaSYCL/array-kernel-param.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -178,7 +178,7 @@ int main() {
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'sycl::id<1>'
// CHECK-NEXT: CompoundStmt
// CHECK-NEXT: DeclStmt
// CHECK-NEXT: VarDecl {{.*}} used '(lambda at {{.*}}array-kernel-param.cpp{{.*}})' cinit
// CHECK-NEXT: VarDecl {{.*}} used __SYCLKernel '(lambda at {{.*}}array-kernel-param.cpp{{.*}})' cinit
// CHECK-NEXT: InitListExpr {{.*}} '(lambda at {{.*}}array-kernel-param.cpp{{.*}})'
// CHECK-NEXT: InitListExpr {{.*}} 'StructWithAccessors'
// CHECK-NEXT: InitListExpr {{.*}} 'Accessor[2]'
Expand Down Expand Up @@ -219,7 +219,7 @@ int main() {
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_c 'int'
// CHECK-NEXT: CompoundStmt
// CHECK-NEXT: DeclStmt
// CHECK-NEXT: VarDecl {{.*}} used '(lambda at {{.*}}array-kernel-param.cpp{{.*}})' cinit
// CHECK-NEXT: VarDecl {{.*}} used __SYCLKernel '(lambda at {{.*}}array-kernel-param.cpp{{.*}})' cinit
// CHECK-NEXT: InitListExpr {{.*}} '(lambda at {{.*}}array-kernel-param.cpp{{.*}})'

// Initializer for struct array i.e. DecomposedStruct DecompStructArray[2]
Expand Down Expand Up @@ -330,7 +330,7 @@ int main() {
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_s 'S<int>':'S<int>'
// CHECK-NEXT: CompoundStmt
// CHECK-NEXT: DeclStmt
// CHECK-NEXT: VarDecl {{.*}} used '(lambda at {{.*}}array-kernel-param.cpp{{.*}})' cinit
// CHECK-NEXT: VarDecl {{.*}} used __SYCLKernel '(lambda at {{.*}}array-kernel-param.cpp{{.*}})' cinit
// CHECK-NEXT: InitListExpr {{.*}} '(lambda at {{.*}}array-kernel-param.cpp{{.*}})'
// CHECK-NEXT: CXXConstructExpr {{.*}} 'S<int>':'S<int>' 'void (const S<int> &) noexcept'
// CHECK-NEXT: ImplicitCastExpr
Expand Down
2 changes: 1 addition & 1 deletion clang/test/SemaSYCL/basic-kernel-wrapper.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -34,7 +34,7 @@ int main() {
// Check lambda declaration inside the wrapper

// CHECK: DeclStmt
// CHECK-NEXT: VarDecl {{.*}} used '(lambda at {{.*}}basic-kernel-wrapper.cpp{{.*}})'
// CHECK-NEXT: VarDecl {{.*}} used __SYCLKernel '(lambda at {{.*}}basic-kernel-wrapper.cpp{{.*}})'

// Check accessor initialization

Expand Down
2 changes: 1 addition & 1 deletion clang/test/SemaSYCL/inheritance.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -50,7 +50,7 @@ int main() {
// Check initializers for derived and base classes.
// Each class has it's own initializer list
// Base classes should be initialized first.
// CHECK: VarDecl {{.*}} derived 'derived' cinit
// CHECK: VarDecl {{.*}} used __SYCLKernel 'derived' cinit
// CHECK-NEXT: InitListExpr {{.*}} 'derived'
// CHECK-NEXT: CXXConstructExpr {{.*}} 'base' 'void (const base &) noexcept'
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const base' lvalue <NoOp>
Expand Down