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 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]] 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. 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]] 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)]. 104define 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