Skip to content

Commit 4cb702b

Browse files
AlexVlxAlexVlx
committed
Update old test, add new test.
1 parent 966029b commit 4cb702b

File tree

2 files changed

+331
-239
lines changed

2 files changed

+331
-239
lines changed

‎clang/test/CodeGenOpenCL/address-of-automatic-variable.cl‎

Lines changed: 135 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,135 @@
1+
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 5
2+
// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -finclude-default-header -fdeclare-opencl-builtins -triple x86_64-unknown-linux-gnu -emit-llvm %s -o - | FileCheck %s
3+
// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -finclude-default-header -fdeclare-opencl-builtins -triple amdgcn-amd-amdhsa -emit-llvm %s -o - | FileCheck -check-prefixes=AMDGCN %s
4+
5+
// CHECK-LABEL: define dso_local zeroext i1 @helperFunction(
6+
// CHECK-SAME: ptr noundef [[PPPP:%.*]]) #[[ATTR0:[0-9]+]] {
7+
// CHECK-NEXT: [[ENTRY:.*:]]
8+
// CHECK-NEXT: [[RETVAL:%.*]] = alloca i1, align 1
9+
// CHECK-NEXT: [[PPPP_ADDR:%.*]] = alloca ptr, align 8
10+
// CHECK-NEXT: store ptr [[PPPP]], ptr [[PPPP_ADDR]], align 8
11+
// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[PPPP_ADDR]], align 8
12+
// CHECK-NEXT: [[TMP1:%.*]] = call ptr @__to_private(ptr [[TMP0]])
13+
// CHECK-NEXT: [[CMP:%.*]] = icmp eq ptr [[TMP1]], null
14+
// CHECK-NEXT: br i1 [[CMP]], label %[[IF_THEN:.*]], label %[[IF_END:.*]]
15+
// CHECK: [[IF_THEN]]:
16+
// CHECK-NEXT: store i1 false, ptr [[RETVAL]], align 1
17+
// CHECK-NEXT: br label %[[RETURN:.*]]
18+
// CHECK: [[IF_END]]:
19+
// CHECK-NEXT: [[TMP2:%.*]] = load ptr, ptr [[PPPP_ADDR]], align 8
20+
// CHECK-NEXT: [[TMP3:%.*]] = load i64, ptr [[TMP2]], align 8
21+
// CHECK-NEXT: [[CMP1:%.*]] = icmp ne i64 [[TMP3]], 5
22+
// CHECK-NEXT: br i1 [[CMP1]], label %[[IF_THEN2:.*]], label %[[IF_END3:.*]]
23+
// CHECK: [[IF_THEN2]]:
24+
// CHECK-NEXT: store i1 false, ptr [[RETVAL]], align 1
25+
// CHECK-NEXT: br label %[[RETURN]]
26+
// CHECK: [[IF_END3]]:
27+
// CHECK-NEXT: store i1 true, ptr [[RETVAL]], align 1
28+
// CHECK-NEXT: br label %[[RETURN]]
29+
// CHECK: [[RETURN]]:
30+
// CHECK-NEXT: [[TMP4:%.*]] = load i1, ptr [[RETVAL]], align 1
31+
// CHECK-NEXT: ret i1 [[TMP4]]
32+
//
33+
// AMDGCN-LABEL: define dso_local zeroext i1 @helperFunction(
34+
// AMDGCN-SAME: ptr noundef [[PPPP:%.*]]) #[[ATTR0:[0-9]+]] {
35+
// AMDGCN-NEXT: [[ENTRY:.*:]]
36+
// AMDGCN-NEXT: [[RETVAL:%.*]] = alloca i1, align 1, addrspace(5)
37+
// AMDGCN-NEXT: [[PPPP_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
38+
// AMDGCN-NEXT: store ptr [[PPPP]], ptr addrspace(5) [[PPPP_ADDR]], align 8
39+
// AMDGCN-NEXT: [[TMP0:%.*]] = load ptr, ptr addrspace(5) [[PPPP_ADDR]], align 8
40+
// AMDGCN-NEXT: [[TMP1:%.*]] = call ptr addrspace(5) @__to_private(ptr [[TMP0]])
41+
// AMDGCN-NEXT: [[CMP:%.*]] = icmp eq ptr addrspace(5) [[TMP1]], addrspacecast (ptr null to ptr addrspace(5))
42+
// AMDGCN-NEXT: br i1 [[CMP]], label %[[IF_THEN:.*]], label %[[IF_END:.*]]
43+
// AMDGCN: [[IF_THEN]]:
44+
// AMDGCN-NEXT: store i1 false, ptr addrspace(5) [[RETVAL]], align 1
45+
// AMDGCN-NEXT: br label %[[RETURN:.*]]
46+
// AMDGCN: [[IF_END]]:
47+
// AMDGCN-NEXT: [[TMP2:%.*]] = load ptr, ptr addrspace(5) [[PPPP_ADDR]], align 8
48+
// AMDGCN-NEXT: [[TMP3:%.*]] = load i64, ptr [[TMP2]], align 8
49+
// AMDGCN-NEXT: [[CMP1:%.*]] = icmp ne i64 [[TMP3]], 5
50+
// AMDGCN-NEXT: br i1 [[CMP1]], label %[[IF_THEN2:.*]], label %[[IF_END3:.*]]
51+
// AMDGCN: [[IF_THEN2]]:
52+
// AMDGCN-NEXT: store i1 false, ptr addrspace(5) [[RETVAL]], align 1
53+
// AMDGCN-NEXT: br label %[[RETURN]]
54+
// AMDGCN: [[IF_END3]]:
55+
// AMDGCN-NEXT: store i1 true, ptr addrspace(5) [[RETVAL]], align 1
56+
// AMDGCN-NEXT: br label %[[RETURN]]
57+
// AMDGCN: [[RETURN]]:
58+
// AMDGCN-NEXT: [[TMP4:%.*]] = load i1, ptr addrspace(5) [[RETVAL]], align 1
59+
// AMDGCN-NEXT: ret i1 [[TMP4]]
60+
//
61+
bool helperFunction(long *pppp) {
62+
if (to_private(pppp) == NULL) {
63+
return false;
64+
}
65+
if (*pppp != 5) {
66+
return false;
67+
}
68+
return true;
69+
}
70+
71+
// CHECK-LABEL: define dso_local spir_kernel void @testKernel(
72+
// CHECK-SAME: ptr noundef align 4 [[RESULTS:%.*]]) #[[ATTR1:[0-9]+]] !kernel_arg_addr_space [[META3:![0-9]+]] !kernel_arg_access_qual [[META4:![0-9]+]] !kernel_arg_type [[META5:![0-9]+]] !kernel_arg_base_type [[META5]] !kernel_arg_type_qual [[META6:![0-9]+]] {
73+
// CHECK-NEXT: [[ENTRY:.*:]]
74+
// CHECK-NEXT: [[RESULTS_ADDR:%.*]] = alloca ptr, align 8
75+
// CHECK-NEXT: [[VVVV:%.*]] = alloca i64, align 8
76+
// CHECK-NEXT: [[PPPP:%.*]] = alloca ptr, align 8
77+
// CHECK-NEXT: [[TID:%.*]] = alloca i32, align 4
78+
// CHECK-NEXT: store ptr [[RESULTS]], ptr [[RESULTS_ADDR]], align 8
79+
// CHECK-NEXT: store i64 5, ptr [[VVVV]], align 8
80+
// CHECK-NEXT: store ptr [[VVVV]], ptr [[PPPP]], align 8
81+
// CHECK-NEXT: [[CALL:%.*]] = call i64 @_Z13get_global_idj(i32 noundef 0) #[[ATTR3:[0-9]+]]
82+
// CHECK-NEXT: [[CONV:%.*]] = trunc i64 [[CALL]] to i32
83+
// CHECK-NEXT: store i32 [[CONV]], ptr [[TID]], align 4
84+
// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[PPPP]], align 8
85+
// CHECK-NEXT: [[CALL1:%.*]] = call zeroext i1 @helperFunction(ptr noundef [[TMP0]]) #[[ATTR4:[0-9]+]]
86+
// CHECK-NEXT: [[CONV2:%.*]] = zext i1 [[CALL1]] to i32
87+
// CHECK-NEXT: [[TMP1:%.*]] = load ptr, ptr [[RESULTS_ADDR]], align 8
88+
// CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr [[TID]], align 4
89+
// CHECK-NEXT: [[IDXPROM:%.*]] = zext i32 [[TMP2]] to i64
90+
// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds nuw i32, ptr [[TMP1]], i64 [[IDXPROM]]
91+
// CHECK-NEXT: store i32 [[CONV2]], ptr [[ARRAYIDX]], align 4
92+
// CHECK-NEXT: ret void
93+
//
94+
// AMDGCN-LABEL: define dso_local amdgpu_kernel void @testKernel(
95+
// AMDGCN-SAME: ptr addrspace(1) noundef align 4 [[RESULTS:%.*]]) #[[ATTR1:[0-9]+]] !kernel_arg_addr_space [[META4:![0-9]+]] !kernel_arg_access_qual [[META5:![0-9]+]] !kernel_arg_type [[META6:![0-9]+]] !kernel_arg_base_type [[META6]] !kernel_arg_type_qual [[META7:![0-9]+]] {
96+
// AMDGCN-NEXT: [[ENTRY:.*:]]
97+
// AMDGCN-NEXT: [[RESULTS_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5)
98+
// AMDGCN-NEXT: [[VVVV:%.*]] = alloca i64, align 8, addrspace(5)
99+
// AMDGCN-NEXT: [[PPPP:%.*]] = alloca ptr addrspace(5), align 4, addrspace(5)
100+
// AMDGCN-NEXT: [[TID:%.*]] = alloca i32, align 4, addrspace(5)
101+
// AMDGCN-NEXT: store ptr addrspace(1) [[RESULTS]], ptr addrspace(5) [[RESULTS_ADDR]], align 8
102+
// AMDGCN-NEXT: store i64 5, ptr addrspace(5) [[VVVV]], align 8
103+
// AMDGCN-NEXT: store ptr addrspace(5) [[VVVV]], ptr addrspace(5) [[PPPP]], align 4
104+
// AMDGCN-NEXT: [[CALL:%.*]] = call i64 @_Z13get_global_idj(i32 noundef 0) #[[ATTR3:[0-9]+]]
105+
// AMDGCN-NEXT: [[CONV:%.*]] = trunc i64 [[CALL]] to i32
106+
// AMDGCN-NEXT: store i32 [[CONV]], ptr addrspace(5) [[TID]], align 4
107+
// AMDGCN-NEXT: [[TMP0:%.*]] = load ptr addrspace(5), ptr addrspace(5) [[PPPP]], align 4
108+
// AMDGCN-NEXT: [[TMP1:%.*]] = addrspacecast ptr addrspace(5) [[TMP0]] to ptr
109+
// AMDGCN-NEXT: [[CALL1:%.*]] = call zeroext i1 @helperFunction(ptr noundef [[TMP1]]) #[[ATTR4:[0-9]+]]
110+
// AMDGCN-NEXT: [[CONV2:%.*]] = zext i1 [[CALL1]] to i32
111+
// AMDGCN-NEXT: [[TMP2:%.*]] = load ptr addrspace(1), ptr addrspace(5) [[RESULTS_ADDR]], align 8
112+
// AMDGCN-NEXT: [[TMP3:%.*]] = load i32, ptr addrspace(5) [[TID]], align 4
113+
// AMDGCN-NEXT: [[IDXPROM:%.*]] = zext i32 [[TMP3]] to i64
114+
// AMDGCN-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds nuw i32, ptr addrspace(1) [[TMP2]], i64 [[IDXPROM]]
115+
// AMDGCN-NEXT: store i32 [[CONV2]], ptr addrspace(1) [[ARRAYIDX]], align 4
116+
// AMDGCN-NEXT: ret void
117+
//
118+
__kernel void testKernel(__global uint *results) {
119+
long vvvv = 5;
120+
__private long *pppp = &vvvv;
121+
122+
uint tid = get_global_id(0);
123+
results[tid] = helperFunction(pppp);
124+
}
125+
//.
126+
// CHECK: [[META3]] = !{i32 1}
127+
// CHECK: [[META4]] = !{!"none"}
128+
// CHECK: [[META5]] = !{!"uint*"}
129+
// CHECK: [[META6]] = !{!""}
130+
//.
131+
// AMDGCN: [[META4]] = !{i32 1}
132+
// AMDGCN: [[META5]] = !{!"none"}
133+
// AMDGCN: [[META6]] = !{!"uint*"}
134+
// AMDGCN: [[META7]] = !{!""}
135+
//.

0 commit comments

Comments
 (0)