1 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s --check-prefix=PTX 2 ; RUN: opt < %s -S -separate-const-offset-from-gep -gvn -dce | 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]] addrspace(3)* @array, i64 0, i64 %1, i64 %0 25 %3 = addrspacecast float addrspace(3)* %2 to float* 26 %4 = load 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]] addrspace(3)* @array, i64 0, i64 %1, i64 %7 31 %9 = addrspacecast float addrspace(3)* %8 to float* 32 %10 = load 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]] addrspace(3)* @array, i64 0, i64 %13, i64 %0 37 %15 = addrspacecast float addrspace(3)* %14 to float* 38 %16 = load float* %15, align 4 39 %17 = fadd float %11, %16 40 %18 = getelementptr inbounds [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* %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:%(rl|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]] addrspace(3)* @array, i64 0, i64 %{{[a-zA-Z0-9]+}}, i64 %{{[a-zA-Z0-9]+}} 55 ; IR: getelementptr float addrspace(3)* [[BASE_PTR]], i64 1 56 ; IR: getelementptr float addrspace(3)* [[BASE_PTR]], i64 32 57 ; IR: getelementptr 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]] addrspace(3)* @array, i64 0, i64 %1, i64 %0 70 %3 = addrspacecast float addrspace(3)* %2 to float* 71 %4 = load 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]] addrspace(3)* @array, i64 0, i64 %1, i64 %6 75 %8 = addrspacecast float addrspace(3)* %7 to float* 76 %9 = load 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]] addrspace(3)* @array, i64 0, i64 %11, i64 %0 80 %13 = addrspacecast float addrspace(3)* %12 to float* 81 %14 = load float* %13, align 4 82 %15 = fadd float %10, %14 83 %16 = getelementptr inbounds [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* %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:%(rl|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]] addrspace(3)* @array, i64 0, i64 %{{[a-zA-Z0-9]+}}, i64 %{{[a-zA-Z0-9]+}} 98 ; IR: getelementptr float addrspace(3)* [[BASE_PTR]], i64 1 99 ; IR: getelementptr float addrspace(3)* [[BASE_PTR]], i64 32 100 ; IR: getelementptr float addrspace(3)* [[BASE_PTR]], i64 33 101 102 ; Similar to @sum_of_array3, but extends array indices using zext instead of 103 ; sext. e.g., array[zext(x + 1)][zext(y + 1)]. 104 define void @sum_of_array3(i32 %x, i32 %y, float* nocapture %output) { 105 .preheader: 106 %0 = zext i32 %y to i64 107 %1 = zext i32 %x to i64 108 %2 = getelementptr inbounds [32 x [32 x float]] addrspace(3)* @array, i64 0, i64 %1, i64 %0 109 %3 = addrspacecast float addrspace(3)* %2 to float* 110 %4 = load float* %3, align 4 111 %5 = fadd float %4, 0.000000e+00 112 %6 = add i32 %y, 1 113 %7 = zext i32 %6 to i64 114 %8 = getelementptr inbounds [32 x [32 x float]] addrspace(3)* @array, i64 0, i64 %1, i64 %7 115 %9 = addrspacecast float addrspace(3)* %8 to float* 116 %10 = load float* %9, align 4 117 %11 = fadd float %5, %10 118 %12 = add i32 %x, 1 119 %13 = zext i32 %12 to i64 120 %14 = getelementptr inbounds [32 x [32 x float]] addrspace(3)* @array, i64 0, i64 %13, i64 %0 121 %15 = addrspacecast float addrspace(3)* %14 to float* 122 %16 = load float* %15, align 4 123 %17 = fadd float %11, %16 124 %18 = getelementptr inbounds [32 x [32 x float]] addrspace(3)* @array, i64 0, i64 %13, i64 %7 125 %19 = addrspacecast float addrspace(3)* %18 to float* 126 %20 = load float* %19, align 4 127 %21 = fadd float %17, %20 128 store float %21, float* %output, align 4 129 ret void 130 } 131 ; PTX-LABEL: sum_of_array3( 132 ; PTX: ld.shared.f32 {{%f[0-9]+}}, {{\[}}[[BASE_REG:%(rl|r)[0-9]+]]{{\]}} 133 ; PTX: ld.shared.f32 {{%f[0-9]+}}, {{\[}}[[BASE_REG]]+4{{\]}} 134 ; PTX: ld.shared.f32 {{%f[0-9]+}}, {{\[}}[[BASE_REG]]+128{{\]}} 135 ; PTX: ld.shared.f32 {{%f[0-9]+}}, {{\[}}[[BASE_REG]]+132{{\]}} 136 137 ; IR-LABEL: @sum_of_array3( 138 ; IR: [[BASE_PTR:%[a-zA-Z0-9]+]] = getelementptr inbounds [32 x [32 x float]] addrspace(3)* @array, i64 0, i64 %{{[a-zA-Z0-9]+}}, i64 %{{[a-zA-Z0-9]+}} 139 ; IR: getelementptr float addrspace(3)* [[BASE_PTR]], i64 1 140 ; IR: getelementptr float addrspace(3)* [[BASE_PTR]], i64 32 141 ; IR: getelementptr float addrspace(3)* [[BASE_PTR]], i64 33 142
