1 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s --check-prefix=PTX
2 ; RUN: opt < %s -S -separate-const-offset-from-gep -reassociate-geps-verify-no-dead-code -gvn | FileCheck %s --check-prefix=IR
4 ; Verifies the SeparateConstOffsetFromGEP pass.
5 ; The following code computes
6 ; *output = array[x][y] + array[x][y+1] + array[x+1][y] + array[x+1][y+1]
8 ; We expect SeparateConstOffsetFromGEP to transform it to
10 ; float *base = &a[x][y];
11 ; *output = base[0] + base[1] + base[32] + base[33];
13 ; so the backend can emit PTX that uses fewer virtual registers.
15 target datalayout = "e-i64:64-v16:16-v32:32-n16:32:64"
16 target triple = "nvptx64-unknown-unknown"
18 @array = internal addrspace(3) constant [32 x [32 x float]] zeroinitializer, align 4
20 define void @sum_of_array(i32 %x, i32 %y, float* nocapture %output) {
22 %0 = sext i32 %y to i64
23 %1 = sext i32 %x to i64
24 %2 = getelementptr inbounds [32 x [32 x float]], [32 x [32 x float]] addrspace(3)* @array, i64 0, i64 %1, i64 %0
25 %3 = addrspacecast float addrspace(3)* %2 to float*
26 %4 = load float, float* %3, align 4
27 %5 = fadd float %4, 0.000000e+00
29 %7 = sext i32 %6 to i64
30 %8 = getelementptr inbounds [32 x [32 x float]], [32 x [32 x float]] addrspace(3)* @array, i64 0, i64 %1, i64 %7
31 %9 = addrspacecast float addrspace(3)* %8 to float*
32 %10 = load float, float* %9, align 4
33 %11 = fadd float %5, %10
35 %13 = sext i32 %12 to i64
36 %14 = getelementptr inbounds [32 x [32 x float]], [32 x [32 x float]] addrspace(3)* @array, i64 0, i64 %13, i64 %0
37 %15 = addrspacecast float addrspace(3)* %14 to float*
38 %16 = load float, float* %15, align 4
39 %17 = fadd float %11, %16
40 %18 = getelementptr inbounds [32 x [32 x float]], [32 x [32 x float]] addrspace(3)* @array, i64 0, i64 %13, i64 %7
41 %19 = addrspacecast float addrspace(3)* %18 to float*
42 %20 = load float, float* %19, align 4
43 %21 = fadd float %17, %20
44 store float %21, float* %output, align 4
47 ; PTX-LABEL: sum_of_array(
48 ; PTX: ld.shared.f32 {{%f[0-9]+}}, {{\[}}[[BASE_REG:%(rd|r)[0-9]+]]{{\]}}
49 ; PTX: ld.shared.f32 {{%f[0-9]+}}, {{\[}}[[BASE_REG]]+4{{\]}}
50 ; PTX: ld.shared.f32 {{%f[0-9]+}}, {{\[}}[[BASE_REG]]+128{{\]}}
51 ; PTX: ld.shared.f32 {{%f[0-9]+}}, {{\[}}[[BASE_REG]]+132{{\]}}
53 ; IR-LABEL: @sum_of_array(
54 ; IR: [[BASE_PTR:%[a-zA-Z0-9]+]] = getelementptr inbounds [32 x [32 x float]], [32 x [32 x float]] addrspace(3)* @array, i64 0, i64 %{{[a-zA-Z0-9]+}}, i64 %{{[a-zA-Z0-9]+}}
55 ; IR: getelementptr inbounds float, float addrspace(3)* [[BASE_PTR]], i64 1
56 ; IR: getelementptr inbounds float, float addrspace(3)* [[BASE_PTR]], i64 32
57 ; IR: getelementptr inbounds float, float addrspace(3)* [[BASE_PTR]], i64 33
59 ; @sum_of_array2 is very similar to @sum_of_array. The only difference is in
60 ; the order of "sext" and "add" when computing the array indices. @sum_of_array
61 ; computes add before sext, e.g., array[sext(x + 1)][sext(y + 1)], while
62 ; @sum_of_array2 computes sext before add,
63 ; e.g., array[sext(x) + 1][sext(y) + 1]. SeparateConstOffsetFromGEP should be
64 ; able to extract constant offsets from both forms.
65 define void @sum_of_array2(i32 %x, i32 %y, float* nocapture %output) {
67 %0 = sext i32 %y to i64
68 %1 = sext i32 %x to i64
69 %2 = getelementptr inbounds [32 x [32 x float]], [32 x [32 x float]] addrspace(3)* @array, i64 0, i64 %1, i64 %0
70 %3 = addrspacecast float addrspace(3)* %2 to float*
71 %4 = load float, float* %3, align 4
72 %5 = fadd float %4, 0.000000e+00
74 %7 = getelementptr inbounds [32 x [32 x float]], [32 x [32 x float]] addrspace(3)* @array, i64 0, i64 %1, i64 %6
75 %8 = addrspacecast float addrspace(3)* %7 to float*
76 %9 = load float, float* %8, align 4
77 %10 = fadd float %5, %9
79 %12 = getelementptr inbounds [32 x [32 x float]], [32 x [32 x float]] addrspace(3)* @array, i64 0, i64 %11, i64 %0
80 %13 = addrspacecast float addrspace(3)* %12 to float*
81 %14 = load float, float* %13, align 4
82 %15 = fadd float %10, %14
83 %16 = getelementptr inbounds [32 x [32 x float]], [32 x [32 x float]] addrspace(3)* @array, i64 0, i64 %11, i64 %6
84 %17 = addrspacecast float addrspace(3)* %16 to float*
85 %18 = load float, float* %17, align 4
86 %19 = fadd float %15, %18
87 store float %19, float* %output, align 4
90 ; PTX-LABEL: sum_of_array2(
91 ; PTX: ld.shared.f32 {{%f[0-9]+}}, {{\[}}[[BASE_REG:%(rd|r)[0-9]+]]{{\]}}
92 ; PTX: ld.shared.f32 {{%f[0-9]+}}, {{\[}}[[BASE_REG]]+4{{\]}}
93 ; PTX: ld.shared.f32 {{%f[0-9]+}}, {{\[}}[[BASE_REG]]+128{{\]}}
94 ; PTX: ld.shared.f32 {{%f[0-9]+}}, {{\[}}[[BASE_REG]]+132{{\]}}
96 ; IR-LABEL: @sum_of_array2(
97 ; IR: [[BASE_PTR:%[a-zA-Z0-9]+]] = getelementptr inbounds [32 x [32 x float]], [32 x [32 x float]] addrspace(3)* @array, i64 0, i64 %{{[a-zA-Z0-9]+}}, i64 %{{[a-zA-Z0-9]+}}
98 ; IR: getelementptr inbounds float, float addrspace(3)* [[BASE_PTR]], i64 1
99 ; IR: getelementptr inbounds float, float addrspace(3)* [[BASE_PTR]], i64 32
100 ; IR: getelementptr inbounds float, float addrspace(3)* [[BASE_PTR]], i64 33
103 ; This function loads
104 ; array[zext(x)][zext(y)]
105 ; array[zext(x)][zext(y +nuw 1)]
106 ; array[zext(x +nuw 1)][zext(y)]
107 ; array[zext(x +nuw 1)][zext(y +nuw 1)].
109 ; This function is similar to @sum_of_array, but it
110 ; 1) extends array indices using zext instead of sext;
111 ; 2) annotates the addition with "nuw"; otherwise, zext(x + 1) => zext(x) + 1
113 define void @sum_of_array3(i32 %x, i32 %y, float* nocapture %output) {
115 %0 = zext i32 %y to i64
116 %1 = zext i32 %x to i64
117 %2 = getelementptr inbounds [32 x [32 x float]], [32 x [32 x float]] addrspace(3)* @array, i64 0, i64 %1, i64 %0
118 %3 = addrspacecast float addrspace(3)* %2 to float*
119 %4 = load float, float* %3, align 4
120 %5 = fadd float %4, 0.000000e+00
121 %6 = add nuw i32 %y, 1
122 %7 = zext i32 %6 to i64
123 %8 = getelementptr inbounds [32 x [32 x float]], [32 x [32 x float]] addrspace(3)* @array, i64 0, i64 %1, i64 %7
124 %9 = addrspacecast float addrspace(3)* %8 to float*
125 %10 = load float, float* %9, align 4
126 %11 = fadd float %5, %10
127 %12 = add nuw i32 %x, 1
128 %13 = zext i32 %12 to i64
129 %14 = getelementptr inbounds [32 x [32 x float]], [32 x [32 x float]] addrspace(3)* @array, i64 0, i64 %13, i64 %0
130 %15 = addrspacecast float addrspace(3)* %14 to float*
131 %16 = load float, float* %15, align 4
132 %17 = fadd float %11, %16
133 %18 = getelementptr inbounds [32 x [32 x float]], [32 x [32 x float]] addrspace(3)* @array, i64 0, i64 %13, i64 %7
134 %19 = addrspacecast float addrspace(3)* %18 to float*
135 %20 = load float, float* %19, align 4
136 %21 = fadd float %17, %20
137 store float %21, float* %output, align 4
140 ; PTX-LABEL: sum_of_array3(
141 ; PTX: ld.shared.f32 {{%f[0-9]+}}, {{\[}}[[BASE_REG:%(rd|r)[0-9]+]]{{\]}}
142 ; PTX: ld.shared.f32 {{%f[0-9]+}}, {{\[}}[[BASE_REG]]+4{{\]}}
143 ; PTX: ld.shared.f32 {{%f[0-9]+}}, {{\[}}[[BASE_REG]]+128{{\]}}
144 ; PTX: ld.shared.f32 {{%f[0-9]+}}, {{\[}}[[BASE_REG]]+132{{\]}}
146 ; IR-LABEL: @sum_of_array3(
147 ; IR: [[BASE_PTR:%[a-zA-Z0-9]+]] = getelementptr inbounds [32 x [32 x float]], [32 x [32 x float]] addrspace(3)* @array, i64 0, i64 %{{[a-zA-Z0-9]+}}, i64 %{{[a-zA-Z0-9]+}}
148 ; IR: getelementptr inbounds float, float addrspace(3)* [[BASE_PTR]], i64 1
149 ; IR: getelementptr inbounds float, float addrspace(3)* [[BASE_PTR]], i64 32
150 ; IR: getelementptr inbounds float, float addrspace(3)* [[BASE_PTR]], i64 33
153 ; This function loads
154 ; array[zext(x)][zext(y)]
155 ; array[zext(x)][zext(y)]
156 ; array[zext(x) + 1][zext(y) + 1]
157 ; array[zext(x) + 1][zext(y) + 1].
159 ; We expect the generated code to reuse the computation of
160 ; &array[zext(x)][zext(y)]. See the expected IR and PTX for details.
161 define void @sum_of_array4(i32 %x, i32 %y, float* nocapture %output) {
163 %0 = zext i32 %y to i64
164 %1 = zext i32 %x to i64
165 %2 = getelementptr inbounds [32 x [32 x float]], [32 x [32 x float]] addrspace(3)* @array, i64 0, i64 %1, i64 %0
166 %3 = addrspacecast float addrspace(3)* %2 to float*
167 %4 = load float, float* %3, align 4
168 %5 = fadd float %4, 0.000000e+00
170 %7 = getelementptr inbounds [32 x [32 x float]], [32 x [32 x float]] addrspace(3)* @array, i64 0, i64 %1, i64 %6
171 %8 = addrspacecast float addrspace(3)* %7 to float*
172 %9 = load float, float* %8, align 4
173 %10 = fadd float %5, %9
175 %12 = getelementptr inbounds [32 x [32 x float]], [32 x [32 x float]] addrspace(3)* @array, i64 0, i64 %11, i64 %0
176 %13 = addrspacecast float addrspace(3)* %12 to float*
177 %14 = load float, float* %13, align 4
178 %15 = fadd float %10, %14
179 %16 = getelementptr inbounds [32 x [32 x float]], [32 x [32 x float]] addrspace(3)* @array, i64 0, i64 %11, i64 %6
180 %17 = addrspacecast float addrspace(3)* %16 to float*
181 %18 = load float, float* %17, align 4
182 %19 = fadd float %15, %18
183 store float %19, float* %output, align 4
186 ; PTX-LABEL: sum_of_array4(
187 ; PTX: ld.shared.f32 {{%f[0-9]+}}, {{\[}}[[BASE_REG:%(rd|r)[0-9]+]]{{\]}}
188 ; PTX: ld.shared.f32 {{%f[0-9]+}}, {{\[}}[[BASE_REG]]+4{{\]}}
189 ; PTX: ld.shared.f32 {{%f[0-9]+}}, {{\[}}[[BASE_REG]]+128{{\]}}
190 ; PTX: ld.shared.f32 {{%f[0-9]+}}, {{\[}}[[BASE_REG]]+132{{\]}}
192 ; IR-LABEL: @sum_of_array4(
193 ; IR: [[BASE_PTR:%[a-zA-Z0-9]+]] = getelementptr inbounds [32 x [32 x float]], [32 x [32 x float]] addrspace(3)* @array, i64 0, i64 %{{[a-zA-Z0-9]+}}, i64 %{{[a-zA-Z0-9]+}}
194 ; IR: getelementptr inbounds float, float addrspace(3)* [[BASE_PTR]], i64 1
195 ; IR: getelementptr inbounds float, float addrspace(3)* [[BASE_PTR]], i64 32
196 ; IR: getelementptr inbounds float, float addrspace(3)* [[BASE_PTR]], i64 33
199 ; The source code is:
200 ; p0 = &input[sext(x + y)];
201 ; p1 = &input[sext(x + (y + 5))];
203 ; Without reuniting extensions, SeparateConstOffsetFromGEP would emit
204 ; p0 = &input[sext(x + y)];
205 ; t1 = &input[sext(x) + sext(y)];
208 ; With reuniting extensions, it merges p0 and t1 and thus emits
209 ; p0 = &input[sext(x + y)];
211 define void @reunion(i32 %x, i32 %y, float* %input) {
212 ; IR-LABEL: @reunion(
213 ; PTX-LABEL: reunion(
215 %xy = add nsw i32 %x, %y
216 %0 = sext i32 %xy to i64
217 %p0 = getelementptr inbounds float, float* %input, i64 %0
218 %v0 = load float, float* %p0, align 4
219 ; PTX: ld.f32 %f{{[0-9]+}}, {{\[}}[[p0:%rd[0-9]+]]{{\]}}
220 call void @use(float %v0)
222 %y5 = add nsw i32 %y, 5
223 %xy5 = add nsw i32 %x, %y5
224 %1 = sext i32 %xy5 to i64
225 %p1 = getelementptr inbounds float, float* %input, i64 %1
226 ; IR: getelementptr inbounds float, float* %p0, i64 5
227 %v1 = load float, float* %p1, align 4
228 ; PTX: ld.f32 %f{{[0-9]+}}, {{\[}}[[p0]]+20{{\]}}
229 call void @use(float %v1)
234 declare void @use(float)