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 3 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] 7 ; 8 ; We expect SeparateConstOffsetFromGEP to transform it to 9 ; 10 ; float *base = &a[x][y]; 11 ; *output = base[0] + base[1] + base[32] + base[33]; 12 ; 13 ; so the backend can emit PTX that uses fewer virtual registers. 14 15 target datalayout = "e-i64:64-v16:16-v32:32-n16:32:64" 16 target triple = "nvptx64-unknown-unknown" 17 18 @array = internal addrspace(3) constant [32 x [32 x float]] zeroinitializer, align 4 19 20 define void @sum_of_array(i32 %x, i32 %y, float* nocapture %output) { 21 .preheader: 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 28 %6 = add i32 %y, 1 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 34 %12 = add i32 %x, 1 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 45 ret void 46 } 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{{\]}} 52 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 58 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) { 66 .preheader: 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 73 %6 = add i64 %0, 1 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 78 %11 = add i64 %1, 1 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 88 ret void 89 } 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{{\]}} 95 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 101 102 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)]. 108 ; 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 112 ; may be invalid. 113 define void @sum_of_array3(i32 %x, i32 %y, float* nocapture %output) { 114 .preheader: 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 138 ret void 139 } 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{{\]}} 145 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 151 152 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]. 158 ; 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) { 162 .preheader: 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 169 %6 = add i64 %0, 1 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 174 %11 = add i64 %1, 1 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 184 ret void 185 } 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{{\]}} 191 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 197 198 199 ; The source code is: 200 ; p0 = &input[sext(x + y)]; 201 ; p1 = &input[sext(x + (y + 5))]; 202 ; 203 ; Without reuniting extensions, SeparateConstOffsetFromGEP would emit 204 ; p0 = &input[sext(x + y)]; 205 ; t1 = &input[sext(x) + sext(y)]; 206 ; p1 = &t1[5]; 207 ; 208 ; With reuniting extensions, it merges p0 and t1 and thus emits 209 ; p0 = &input[sext(x + y)]; 210 ; p1 = &p0[5]; 211 define void @reunion(i32 %x, i32 %y, float* %input) { 212 ; IR-LABEL: @reunion( 213 ; PTX-LABEL: reunion( 214 entry: 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) 221 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) 230 231 ret void 232 } 233 234 declare void @use(float) 235
