test/Transforms/SeparateConstOffsetFromGEP/NVPTX/split-gep-and-gvn.ll

   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