1- // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 5
1+ // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --include-generated-funcs -- version 5
22// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s
33
44// Check there's no assertion when passing a pointer to an address space
@@ -8,55 +8,67 @@ extern void private_ptr(__private int *);
88extern void local_ptr (__local int * );
99extern void generic_ptr (__generic int * );
1010
11+
12+ void use_of_private_var ()
13+ {
14+ int x = 0 ;
15+ private_ptr (& x );
16+ generic_ptr (& x );
17+ }
18+
19+
20+ void addr_of_arg (int x )
21+ {
22+ private_ptr (& x );
23+ generic_ptr (& x );
24+ }
25+
26+
27+ __kernel void use_of_local_var ()
28+ {
29+ __local int x ;
30+ local_ptr (& x );
31+ generic_ptr (& x );
32+ }
33+
1134// CHECK-LABEL: define dso_local void @use_of_private_var(
1235// CHECK-SAME: ) #[[ATTR0:[0-9]+]] {
1336// CHECK-NEXT: [[ENTRY:.*:]]
1437// CHECK-NEXT: [[X:%.*]] = alloca i32, align 4, addrspace(5)
15- // CHECK-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[X]]) #[[ATTR4 :[0-9]+]]
38+ // CHECK-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[X]]) #[[ATTR5 :[0-9]+]]
1639// CHECK-NEXT: store i32 0, ptr addrspace(5) [[X]], align 4, !tbaa [[TBAA4:![0-9]+]]
17- // CHECK-NEXT: call void @private_ptr(ptr addrspace(5) noundef [[X]]) #[[ATTR5 :[0-9]+]]
40+ // CHECK-NEXT: call void @private_ptr(ptr addrspace(5) noundef [[X]]) #[[ATTR6 :[0-9]+]]
1841// CHECK-NEXT: [[X_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[X]] to ptr
19- // CHECK-NEXT: call void @generic_ptr(ptr noundef [[X_ASCAST]]) #[[ATTR5 ]]
20- // CHECK-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[X]]) #[[ATTR4 ]]
42+ // CHECK-NEXT: call void @generic_ptr(ptr noundef [[X_ASCAST]]) #[[ATTR6 ]]
43+ // CHECK-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[X]]) #[[ATTR5 ]]
2144// CHECK-NEXT: ret void
2245//
23- void use_of_private_var ()
24- {
25- int x = 0 ;
26- private_ptr (& x );
27- generic_ptr (& x );
28- }
29-
46+ //
3047// CHECK-LABEL: define dso_local void @addr_of_arg(
3148// CHECK-SAME: i32 noundef [[X:%.*]]) #[[ATTR0]] {
3249// CHECK-NEXT: [[ENTRY:.*:]]
3350// CHECK-NEXT: [[X_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
3451// CHECK-NEXT: store i32 [[X]], ptr addrspace(5) [[X_ADDR]], align 4, !tbaa [[TBAA4]]
35- // CHECK-NEXT: call void @private_ptr(ptr addrspace(5) noundef [[X_ADDR]]) #[[ATTR5 ]]
52+ // CHECK-NEXT: call void @private_ptr(ptr addrspace(5) noundef [[X_ADDR]]) #[[ATTR6 ]]
3653// CHECK-NEXT: [[X_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[X_ADDR]] to ptr
37- // CHECK-NEXT: call void @generic_ptr(ptr noundef [[X_ADDR_ASCAST]]) #[[ATTR5 ]]
54+ // CHECK-NEXT: call void @generic_ptr(ptr noundef [[X_ADDR_ASCAST]]) #[[ATTR6 ]]
3855// CHECK-NEXT: ret void
3956//
40- void addr_of_arg (int x )
41- {
42- private_ptr (& x );
43- generic_ptr (& x );
44- }
45-
57+ //
4658// CHECK-LABEL: define dso_local amdgpu_kernel void @use_of_local_var(
4759// CHECK-SAME: ) #[[ATTR3:[0-9]+]] !kernel_arg_addr_space [[META8:![0-9]+]] !kernel_arg_access_qual [[META8]] !kernel_arg_type [[META8]] !kernel_arg_base_type [[META8]] !kernel_arg_type_qual [[META8]] {
4860// CHECK-NEXT: [[ENTRY:.*:]]
49- // CHECK-NEXT: call void @local_ptr(ptr addrspace(3) noundef @use_of_local_var.x) #[[ATTR5]]
50- // CHECK-NEXT: call void @generic_ptr(ptr noundef addrspacecast (ptr addrspace(3) @use_of_local_var.x to ptr)) #[[ATTR5]]
61+ // CHECK-NEXT: call void @__clang_ocl_kern_imp_use_of_local_var() #[[ATTR7:[0-9]+]]
62+ // CHECK-NEXT: ret void
63+ //
64+ //
65+ // CHECK-LABEL: define dso_local void @__clang_ocl_kern_imp_use_of_local_var(
66+ // CHECK-SAME: ) #[[ATTR4:[0-9]+]] !kernel_arg_addr_space [[META8]] !kernel_arg_access_qual [[META8]] !kernel_arg_type [[META8]] !kernel_arg_base_type [[META8]] !kernel_arg_type_qual [[META8]] {
67+ // CHECK-NEXT: [[ENTRY:.*:]]
68+ // CHECK-NEXT: call void @local_ptr(ptr addrspace(3) noundef @use_of_local_var.x) #[[ATTR6]]
69+ // CHECK-NEXT: call void @generic_ptr(ptr noundef addrspacecast (ptr addrspace(3) @use_of_local_var.x to ptr)) #[[ATTR6]]
5170// CHECK-NEXT: ret void
5271//
53- __kernel void use_of_local_var ()
54- {
55- __local int x ;
56- local_ptr (& x );
57- generic_ptr (& x );
58- }
59-
6072//.
6173// CHECK: [[TBAA4]] = !{[[META5:![0-9]+]], [[META5]], i64 0}
6274// CHECK: [[META5]] = !{!"int", [[META6:![0-9]+]], i64 0}
0 commit comments