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 15target datalayout = "e-i64:64-v16:16-v32:32-n16:32:64" 16target triple = "nvptx64-unknown-unknown" 17 18@array = internal addrspace(3) constant [32 x [32 x float]] zeroinitializer, align 4 19 20define 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. 65define 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. 113define 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. 161define 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]; 211define void @reunion(i32 %x, i32 %y, float* %input) { 212; IR-LABEL: @reunion( 213; PTX-LABEL: reunion( 214entry: 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 234declare void @use(float) 235