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

   1 ; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s --check-prefix=PTX
   2 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s --check-prefix=PTX
   3 ; RUN: opt < %s -S -separate-const-offset-from-gep -gvn -dce | FileCheck %s --check-prefix=IR
   4
   5 ; Verifies the SeparateConstOffsetFromGEP pass.
   6 ; The following code computes
   7 ; *output = array[x][y] + array[x][y+1] + array[x+1][y] + array[x+1][y+1]
   8 ;
   9 ; We expect SeparateConstOffsetFromGEP to transform it to
  10 ;
  11 ; float *base = &a[x][y];
  12 ; *output = base[0] + base[1] + base[32] + base[33];
  13 ;
  14 ; so the backend can emit PTX that uses fewer virtual registers.
  15
  16 target datalayout = "e-i64:64-v16:16-v32:32-n16:32:64"
  17 target triple = "nvptx64-unknown-unknown"
  18
  19 @array = internal addrspace(3) constant [32 x [32 x float]] zeroinitializer, align 4
  20
  21 define void @sum_of_array(i32 %x, i32 %y, float* nocapture %output) {
  22 .preheader:
  23   %0 = zext i32 %y to i64
  24   %1 = zext i32 %x to i64
  25   %2 = getelementptr inbounds [32 x [32 x float]] addrspace(3)* @array, i64 0, i64 %1, i64 %0
  26   %3 = addrspacecast float addrspace(3)* %2 to float*
  27   %4 = load float* %3, align 4
  28   %5 = fadd float %4, 0.000000e+00
  29   %6 = add i32 %y, 1
  30   %7 = zext i32 %6 to i64
  31   %8 = getelementptr inbounds [32 x [32 x float]] addrspace(3)* @array, i64 0, i64 %1, i64 %7
  32   %9 = addrspacecast float addrspace(3)* %8 to float*
  33   %10 = load float* %9, align 4
  34   %11 = fadd float %5, %10
  35   %12 = add i32 %x, 1
  36   %13 = zext i32 %12 to i64
  37   %14 = getelementptr inbounds [32 x [32 x float]] addrspace(3)* @array, i64 0, i64 %13, i64 %0
  38   %15 = addrspacecast float addrspace(3)* %14 to float*
  39   %16 = load float* %15, align 4
  40   %17 = fadd float %11, %16
  41   %18 = getelementptr inbounds [32 x [32 x float]] addrspace(3)* @array, i64 0, i64 %13, i64 %7
  42   %19 = addrspacecast float addrspace(3)* %18 to float*
  43   %20 = load float* %19, align 4
  44   %21 = fadd float %17, %20
  45   store float %21, float* %output, align 4
  46   ret void
  47 }
  48
  49 ; PTX-LABEL: sum_of_array(
  50 ; PTX: ld.shared.f32 {{%f[0-9]+}}, {{\[}}[[BASE_REG:%(rl|r)[0-9]+]]{{\]}}
  51 ; PTX: ld.shared.f32 {{%f[0-9]+}}, {{\[}}[[BASE_REG]]+4{{\]}}
  52 ; PTX: ld.shared.f32 {{%f[0-9]+}}, {{\[}}[[BASE_REG]]+128{{\]}}
  53 ; PTX: ld.shared.f32 {{%f[0-9]+}}, {{\[}}[[BASE_REG]]+132{{\]}}
  54
  55 ; IR-LABEL: @sum_of_array(
  56 ; IR: [[BASE_PTR:%[0-9]+]] = getelementptr inbounds [32 x [32 x float]] addrspace(3)* @array, i64 0, i32 %x, i32 %y
  57 ; IR: getelementptr float addrspace(3)* [[BASE_PTR]], i64 1
  58 ; IR: getelementptr float addrspace(3)* [[BASE_PTR]], i64 32
  59 ; IR: getelementptr float addrspace(3)* [[BASE_PTR]], i64 33