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

   1 ; NOTE: Assertions have been autogenerated by utils/update_test_checks.py UTC_ARGS: --version 4
   2 ; RUN: opt < %s -mtriple=nvptx64-nvidia-cuda -S -passes=separate-const-offset-from-gep,gvn \
   3 ; RUN:       -reassociate-geps-verify-no-dead-code \
   4 ; RUN:     | FileCheck %s --check-prefix=IR
   5 ; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda -mcpu=sm_20 \
   6 ; RUN:     | FileCheck %s --check-prefix=PTX
   7
   8 ; Verifies the SeparateConstOffsetFromGEP pass.
   9 ; The following code computes
  10 ; *output = array[x][y] + array[x][y+1] + array[x+1][y] + array[x+1][y+1]
  11 ;
  12 ; We expect SeparateConstOffsetFromGEP to transform it to
  13 ;
  14 ; ptr base = &a[x][y];
  15 ; *output = base[0] + base[1] + base[32] + base[33];
  16 ;
  17 ; so the backend can emit PTX that uses fewer virtual registers.
  18
  19 @array = internal addrspace(3) global [32 x [32 x float]] zeroinitializer, align 4
  20
  21 define void @sum_of_array(i32 %x, i32 %y, ptr nocapture %output) {
  22 ; IR-LABEL: define void @sum_of_array(
  23 ; IR-SAME: i32 [[X:%.*]], i32 [[Y:%.*]], ptr nocapture [[OUTPUT:%.*]]) {
  24 ; IR-NEXT:  .preheader:
  25 ; IR-NEXT:    [[TMP0:%.*]] = sext i32 [[Y]] to i64
  26 ; IR-NEXT:    [[TMP1:%.*]] = sext i32 [[X]] to i64
  27 ; IR-NEXT:    [[TMP2:%.*]] = getelementptr [32 x [32 x float]], ptr addrspace(3) @array, i64 0, i64 [[TMP1]], i64 [[TMP0]]
  28 ; IR-NEXT:    [[TMP3:%.*]] = addrspacecast ptr addrspace(3) [[TMP2]] to ptr
  29 ; IR-NEXT:    [[TMP4:%.*]] = load float, ptr [[TMP3]], align 4
  30 ; IR-NEXT:    [[TMP5:%.*]] = fadd float [[TMP4]], 0.000000e+00
  31 ; IR-NEXT:    [[TMP6:%.*]] = getelementptr inbounds i8, ptr addrspace(3) [[TMP2]], i64 4
  32 ; IR-NEXT:    [[TMP7:%.*]] = addrspacecast ptr addrspace(3) [[TMP6]] to ptr
  33 ; IR-NEXT:    [[TMP8:%.*]] = load float, ptr [[TMP7]], align 4
  34 ; IR-NEXT:    [[TMP9:%.*]] = fadd float [[TMP5]], [[TMP8]]
  35 ; IR-NEXT:    [[TMP10:%.*]] = getelementptr inbounds i8, ptr addrspace(3) [[TMP2]], i64 128
  36 ; IR-NEXT:    [[TMP11:%.*]] = addrspacecast ptr addrspace(3) [[TMP10]] to ptr
  37 ; IR-NEXT:    [[TMP12:%.*]] = load float, ptr [[TMP11]], align 4
  38 ; IR-NEXT:    [[TMP13:%.*]] = fadd float [[TMP9]], [[TMP12]]
  39 ; IR-NEXT:    [[TMP14:%.*]] = getelementptr inbounds i8, ptr addrspace(3) [[TMP2]], i64 132
  40 ; IR-NEXT:    [[TMP15:%.*]] = addrspacecast ptr addrspace(3) [[TMP14]] to ptr
  41 ; IR-NEXT:    [[TMP16:%.*]] = load float, ptr [[TMP15]], align 4
  42 ; IR-NEXT:    [[TMP17:%.*]] = fadd float [[TMP13]], [[TMP16]]
  43 ; IR-NEXT:    store float [[TMP17]], ptr [[OUTPUT]], align 4
  44 ; IR-NEXT:    ret void
  45 ;
  46 .preheader:
  47   %0 = sext i32 %y to i64
  48   %1 = sext i32 %x to i64
  49   %2 = getelementptr inbounds [32 x [32 x float]], ptr addrspace(3) @array, i64 0, i64 %1, i64 %0
  50   %3 = addrspacecast ptr addrspace(3) %2 to ptr
  51   %4 = load float, ptr %3, align 4
  52   %5 = fadd float %4, 0.000000e+00
  53   %6 = add i32 %y, 1
  54   %7 = sext i32 %6 to i64
  55   %8 = getelementptr inbounds [32 x [32 x float]], ptr addrspace(3) @array, i64 0, i64 %1, i64 %7
  56   %9 = addrspacecast ptr addrspace(3) %8 to ptr
  57   %10 = load float, ptr %9, align 4
  58   %11 = fadd float %5, %10
  59   %12 = add i32 %x, 1
  60   %13 = sext i32 %12 to i64
  61   %14 = getelementptr inbounds [32 x [32 x float]], ptr addrspace(3) @array, i64 0, i64 %13, i64 %0
  62   %15 = addrspacecast ptr addrspace(3) %14 to ptr
  63   %16 = load float, ptr %15, align 4
  64   %17 = fadd float %11, %16
  65   %18 = getelementptr inbounds [32 x [32 x float]], ptr addrspace(3) @array, i64 0, i64 %13, i64 %7
  66   %19 = addrspacecast ptr addrspace(3) %18 to ptr
  67   %20 = load float, ptr %19, align 4
  68   %21 = fadd float %17, %20
  69   store float %21, ptr %output, align 4
  70   ret void
  71 }
  72 ; PTX-LABEL: sum_of_array(
  73 ; PTX-DAG: ld.shared.f32 {{%f[0-9]+}}, [[[BASE_REG:%(rd|r)[0-9]+]]]
  74 ; PTX-DAG: ld.shared.f32 {{%f[0-9]+}}, [[[BASE_REG]]+4]
  75 ; PTX-DAG: ld.shared.f32 {{%f[0-9]+}}, [[[BASE_REG]]+128]
  76 ; PTX-DAG: ld.shared.f32 {{%f[0-9]+}}, [[[BASE_REG]]+132]
  77
  78 ; TODO: GVN is unable to preserve the "inbounds" keyword on the first GEP. Need
  79 ; some infrastructure changes to enable such optimizations.
  80
  81 ; @sum_of_array2 is very similar to @sum_of_array. The only difference is in
  82 ; the order of "sext" and "add" when computing the array indices. @sum_of_array
  83 ; computes add before sext, e.g., array[sext(x + 1)][sext(y + 1)], while
  84 ; @sum_of_array2 computes sext before add,
  85 ; e.g., array[sext(x) + 1][sext(y) + 1]. SeparateConstOffsetFromGEP should be
  86 ; able to extract constant offsets from both forms.
  87 define void @sum_of_array2(i32 %x, i32 %y, ptr nocapture %output) {
  88 ; IR-LABEL: define void @sum_of_array2(
  89 ; IR-SAME: i32 [[X:%.*]], i32 [[Y:%.*]], ptr nocapture [[OUTPUT:%.*]]) {
  90 ; IR-NEXT:  .preheader:
  91 ; IR-NEXT:    [[TMP0:%.*]] = sext i32 [[Y]] to i64
  92 ; IR-NEXT:    [[TMP1:%.*]] = sext i32 [[X]] to i64
  93 ; IR-NEXT:    [[TMP2:%.*]] = getelementptr [32 x [32 x float]], ptr addrspace(3) @array, i64 0, i64 [[TMP1]], i64 [[TMP0]]
  94 ; IR-NEXT:    [[TMP3:%.*]] = addrspacecast ptr addrspace(3) [[TMP2]] to ptr
  95 ; IR-NEXT:    [[TMP4:%.*]] = load float, ptr [[TMP3]], align 4
  96 ; IR-NEXT:    [[TMP5:%.*]] = fadd float [[TMP4]], 0.000000e+00
  97 ; IR-NEXT:    [[TMP6:%.*]] = getelementptr inbounds i8, ptr addrspace(3) [[TMP2]], i64 4
  98 ; IR-NEXT:    [[TMP7:%.*]] = addrspacecast ptr addrspace(3) [[TMP6]] to ptr
  99 ; IR-NEXT:    [[TMP8:%.*]] = load float, ptr [[TMP7]], align 4
 100 ; IR-NEXT:    [[TMP9:%.*]] = fadd float [[TMP5]], [[TMP8]]
 101 ; IR-NEXT:    [[TMP10:%.*]] = getelementptr inbounds i8, ptr addrspace(3) [[TMP2]], i64 128
 102 ; IR-NEXT:    [[TMP11:%.*]] = addrspacecast ptr addrspace(3) [[TMP10]] to ptr
 103 ; IR-NEXT:    [[TMP12:%.*]] = load float, ptr [[TMP11]], align 4
 104 ; IR-NEXT:    [[TMP13:%.*]] = fadd float [[TMP9]], [[TMP12]]
 105 ; IR-NEXT:    [[TMP14:%.*]] = getelementptr inbounds i8, ptr addrspace(3) [[TMP2]], i64 132
 106 ; IR-NEXT:    [[TMP15:%.*]] = addrspacecast ptr addrspace(3) [[TMP14]] to ptr
 107 ; IR-NEXT:    [[TMP16:%.*]] = load float, ptr [[TMP15]], align 4
 108 ; IR-NEXT:    [[TMP17:%.*]] = fadd float [[TMP13]], [[TMP16]]
 109 ; IR-NEXT:    store float [[TMP17]], ptr [[OUTPUT]], align 4
 110 ; IR-NEXT:    ret void
 111 ;
 112 .preheader:
 113   %0 = sext i32 %y to i64
 114   %1 = sext i32 %x to i64
 115   %2 = getelementptr inbounds [32 x [32 x float]], ptr addrspace(3) @array, i64 0, i64 %1, i64 %0
 116   %3 = addrspacecast ptr addrspace(3) %2 to ptr
 117   %4 = load float, ptr %3, align 4
 118   %5 = fadd float %4, 0.000000e+00
 119   %6 = add i64 %0, 1
 120   %7 = getelementptr inbounds [32 x [32 x float]], ptr addrspace(3) @array, i64 0, i64 %1, i64 %6
 121   %8 = addrspacecast ptr addrspace(3) %7 to ptr
 122   %9 = load float, ptr %8, align 4
 123   %10 = fadd float %5, %9
 124   %11 = add i64 %1, 1
 125   %12 = getelementptr inbounds [32 x [32 x float]], ptr addrspace(3) @array, i64 0, i64 %11, i64 %0
 126   %13 = addrspacecast ptr addrspace(3) %12 to ptr
 127   %14 = load float, ptr %13, align 4
 128   %15 = fadd float %10, %14
 129   %16 = getelementptr inbounds [32 x [32 x float]], ptr addrspace(3) @array, i64 0, i64 %11, i64 %6
 130   %17 = addrspacecast ptr addrspace(3) %16 to ptr
 131   %18 = load float, ptr %17, align 4
 132   %19 = fadd float %15, %18
 133   store float %19, ptr %output, align 4
 134   ret void
 135 }
 136 ; PTX-LABEL: sum_of_array2(
 137 ; PTX-DAG: ld.shared.f32 {{%f[0-9]+}}, [[[BASE_REG:%(rd|r)[0-9]+]]]
 138 ; PTX-DAG: ld.shared.f32 {{%f[0-9]+}}, [[[BASE_REG]]+4]
 139 ; PTX-DAG: ld.shared.f32 {{%f[0-9]+}}, [[[BASE_REG]]+128]
 140 ; PTX-DAG: ld.shared.f32 {{%f[0-9]+}}, [[[BASE_REG]]+132]
 141
 142
 143
 144 ; This function loads
 145 ;   array[zext(x)][zext(y)]
 146 ;   array[zext(x)][zext(y +nuw 1)]
 147 ;   array[zext(x +nuw 1)][zext(y)]
 148 ;   array[zext(x +nuw 1)][zext(y +nuw 1)].
 149 ;
 150 ; This function is similar to @sum_of_array, but it
 151 ; 1) extends array indices using zext instead of sext;
 152 ; 2) annotates the addition with "nuw"; otherwise, zext(x + 1) => zext(x) + 1
 153 ;    may be invalid.
 154 define void @sum_of_array3(i32 %x, i32 %y, ptr nocapture %output) {
 155 ; IR-LABEL: define void @sum_of_array3(
 156 ; IR-SAME: i32 [[X:%.*]], i32 [[Y:%.*]], ptr nocapture [[OUTPUT:%.*]]) {
 157 ; IR-NEXT:  .preheader:
 158 ; IR-NEXT:    [[TMP0:%.*]] = zext i32 [[Y]] to i64
 159 ; IR-NEXT:    [[TMP1:%.*]] = zext i32 [[X]] to i64
 160 ; IR-NEXT:    [[TMP2:%.*]] = getelementptr [32 x [32 x float]], ptr addrspace(3) @array, i64 0, i64 [[TMP1]], i64 [[TMP0]]
 161 ; IR-NEXT:    [[TMP3:%.*]] = addrspacecast ptr addrspace(3) [[TMP2]] to ptr
 162 ; IR-NEXT:    [[TMP4:%.*]] = load float, ptr [[TMP3]], align 4
 163 ; IR-NEXT:    [[TMP5:%.*]] = fadd float [[TMP4]], 0.000000e+00
 164 ; IR-NEXT:    [[TMP6:%.*]] = getelementptr inbounds i8, ptr addrspace(3) [[TMP2]], i64 4
 165 ; IR-NEXT:    [[TMP7:%.*]] = addrspacecast ptr addrspace(3) [[TMP6]] to ptr
 166 ; IR-NEXT:    [[TMP8:%.*]] = load float, ptr [[TMP7]], align 4
 167 ; IR-NEXT:    [[TMP9:%.*]] = fadd float [[TMP5]], [[TMP8]]
 168 ; IR-NEXT:    [[TMP10:%.*]] = getelementptr inbounds i8, ptr addrspace(3) [[TMP2]], i64 128
 169 ; IR-NEXT:    [[TMP11:%.*]] = addrspacecast ptr addrspace(3) [[TMP10]] to ptr
 170 ; IR-NEXT:    [[TMP12:%.*]] = load float, ptr [[TMP11]], align 4
 171 ; IR-NEXT:    [[TMP13:%.*]] = fadd float [[TMP9]], [[TMP12]]
 172 ; IR-NEXT:    [[TMP14:%.*]] = getelementptr inbounds i8, ptr addrspace(3) [[TMP2]], i64 132
 173 ; IR-NEXT:    [[TMP15:%.*]] = addrspacecast ptr addrspace(3) [[TMP14]] to ptr
 174 ; IR-NEXT:    [[TMP16:%.*]] = load float, ptr [[TMP15]], align 4
 175 ; IR-NEXT:    [[TMP17:%.*]] = fadd float [[TMP13]], [[TMP16]]
 176 ; IR-NEXT:    store float [[TMP17]], ptr [[OUTPUT]], align 4
 177 ; IR-NEXT:    ret void
 178 ;
 179 .preheader:
 180   %0 = zext i32 %y to i64
 181   %1 = zext i32 %x to i64
 182   %2 = getelementptr inbounds [32 x [32 x float]], ptr addrspace(3) @array, i64 0, i64 %1, i64 %0
 183   %3 = addrspacecast ptr addrspace(3) %2 to ptr
 184   %4 = load float, ptr %3, align 4
 185   %5 = fadd float %4, 0.000000e+00
 186   %6 = add nuw i32 %y, 1
 187   %7 = zext i32 %6 to i64
 188   %8 = getelementptr inbounds [32 x [32 x float]], ptr addrspace(3) @array, i64 0, i64 %1, i64 %7
 189   %9 = addrspacecast ptr addrspace(3) %8 to ptr
 190   %10 = load float, ptr %9, align 4
 191   %11 = fadd float %5, %10
 192   %12 = add nuw i32 %x, 1
 193   %13 = zext i32 %12 to i64
 194   %14 = getelementptr inbounds [32 x [32 x float]], ptr addrspace(3) @array, i64 0, i64 %13, i64 %0
 195   %15 = addrspacecast ptr addrspace(3) %14 to ptr
 196   %16 = load float, ptr %15, align 4
 197   %17 = fadd float %11, %16
 198   %18 = getelementptr inbounds [32 x [32 x float]], ptr addrspace(3) @array, i64 0, i64 %13, i64 %7
 199   %19 = addrspacecast ptr addrspace(3) %18 to ptr
 200   %20 = load float, ptr %19, align 4
 201   %21 = fadd float %17, %20
 202   store float %21, ptr %output, align 4
 203   ret void
 204 }
 205 ; PTX-LABEL: sum_of_array3(
 206 ; PTX-DAG: ld.shared.f32 {{%f[0-9]+}}, [[[BASE_REG:%(rd|r)[0-9]+]]]
 207 ; PTX-DAG: ld.shared.f32 {{%f[0-9]+}}, [[[BASE_REG]]+4]
 208 ; PTX-DAG: ld.shared.f32 {{%f[0-9]+}}, [[[BASE_REG]]+128]
 209 ; PTX-DAG: ld.shared.f32 {{%f[0-9]+}}, [[[BASE_REG]]+132]
 210
 211
 212
 213 ; This function loads
 214 ;   array[zext(x)][zext(y)]
 215 ;   array[zext(x)][zext(y)]
 216 ;   array[zext(x) + 1][zext(y) + 1]
 217 ;   array[zext(x) + 1][zext(y) + 1].
 218 ;
 219 ; We expect the generated code to reuse the computation of
 220 ; &array[zext(x)][zext(y)]. See the expected IR and PTX for details.
 221 define void @sum_of_array4(i32 %x, i32 %y, ptr nocapture %output) {
 222 ; IR-LABEL: define void @sum_of_array4(
 223 ; IR-SAME: i32 [[X:%.*]], i32 [[Y:%.*]], ptr nocapture [[OUTPUT:%.*]]) {
 224 ; IR-NEXT:  .preheader:
 225 ; IR-NEXT:    [[TMP0:%.*]] = zext i32 [[Y]] to i64
 226 ; IR-NEXT:    [[TMP1:%.*]] = zext i32 [[X]] to i64
 227 ; IR-NEXT:    [[TMP2:%.*]] = getelementptr [32 x [32 x float]], ptr addrspace(3) @array, i64 0, i64 [[TMP1]], i64 [[TMP0]]
 228 ; IR-NEXT:    [[TMP3:%.*]] = addrspacecast ptr addrspace(3) [[TMP2]] to ptr
 229 ; IR-NEXT:    [[TMP4:%.*]] = load float, ptr [[TMP3]], align 4
 230 ; IR-NEXT:    [[TMP5:%.*]] = fadd float [[TMP4]], 0.000000e+00
 231 ; IR-NEXT:    [[TMP6:%.*]] = getelementptr inbounds i8, ptr addrspace(3) [[TMP2]], i64 4
 232 ; IR-NEXT:    [[TMP7:%.*]] = addrspacecast ptr addrspace(3) [[TMP6]] to ptr
 233 ; IR-NEXT:    [[TMP8:%.*]] = load float, ptr [[TMP7]], align 4
 234 ; IR-NEXT:    [[TMP9:%.*]] = fadd float [[TMP5]], [[TMP8]]
 235 ; IR-NEXT:    [[TMP10:%.*]] = getelementptr inbounds i8, ptr addrspace(3) [[TMP2]], i64 128
 236 ; IR-NEXT:    [[TMP11:%.*]] = addrspacecast ptr addrspace(3) [[TMP10]] to ptr
 237 ; IR-NEXT:    [[TMP12:%.*]] = load float, ptr [[TMP11]], align 4
 238 ; IR-NEXT:    [[TMP13:%.*]] = fadd float [[TMP9]], [[TMP12]]
 239 ; IR-NEXT:    [[TMP14:%.*]] = getelementptr inbounds i8, ptr addrspace(3) [[TMP2]], i64 132
 240 ; IR-NEXT:    [[TMP15:%.*]] = addrspacecast ptr addrspace(3) [[TMP14]] to ptr
 241 ; IR-NEXT:    [[TMP16:%.*]] = load float, ptr [[TMP15]], align 4
 242 ; IR-NEXT:    [[TMP17:%.*]] = fadd float [[TMP13]], [[TMP16]]
 243 ; IR-NEXT:    store float [[TMP17]], ptr [[OUTPUT]], align 4
 244 ; IR-NEXT:    ret void
 245 ;
 246 .preheader:
 247   %0 = zext i32 %y to i64
 248   %1 = zext i32 %x to i64
 249   %2 = getelementptr inbounds [32 x [32 x float]], ptr addrspace(3) @array, i64 0, i64 %1, i64 %0
 250   %3 = addrspacecast ptr addrspace(3) %2 to ptr
 251   %4 = load float, ptr %3, align 4
 252   %5 = fadd float %4, 0.000000e+00
 253   %6 = add i64 %0, 1
 254   %7 = getelementptr inbounds [32 x [32 x float]], ptr addrspace(3) @array, i64 0, i64 %1, i64 %6
 255   %8 = addrspacecast ptr addrspace(3) %7 to ptr
 256   %9 = load float, ptr %8, align 4
 257   %10 = fadd float %5, %9
 258   %11 = add i64 %1, 1
 259   %12 = getelementptr inbounds [32 x [32 x float]], ptr addrspace(3) @array, i64 0, i64 %11, i64 %0
 260   %13 = addrspacecast ptr addrspace(3) %12 to ptr
 261   %14 = load float, ptr %13, align 4
 262   %15 = fadd float %10, %14
 263   %16 = getelementptr inbounds [32 x [32 x float]], ptr addrspace(3) @array, i64 0, i64 %11, i64 %6
 264   %17 = addrspacecast ptr addrspace(3) %16 to ptr
 265   %18 = load float, ptr %17, align 4
 266   %19 = fadd float %15, %18
 267   store float %19, ptr %output, align 4
 268   ret void
 269 }
 270 ; PTX-LABEL: sum_of_array4(
 271 ; PTX-DAG: ld.shared.f32 {{%f[0-9]+}}, [[[BASE_REG:%(rd|r)[0-9]+]]]
 272 ; PTX-DAG: ld.shared.f32 {{%f[0-9]+}}, [[[BASE_REG]]+4]
 273 ; PTX-DAG: ld.shared.f32 {{%f[0-9]+}}, [[[BASE_REG]]+128]
 274 ; PTX-DAG: ld.shared.f32 {{%f[0-9]+}}, [[[BASE_REG]]+132]
 275
 276
 277
 278 ; The source code is:
 279 ;   p0 = &input[sext(x + y)];
 280 ;   p1 = &input[sext(x + (y + 5))];
 281 ;
 282 ; Without reuniting extensions, SeparateConstOffsetFromGEP would emit
 283 ;   p0 = &input[sext(x + y)];
 284 ;   t1 = &input[sext(x) + sext(y)];
 285 ;   p1 = &t1[5];
 286 ;
 287 ; With reuniting extensions, it merges p0 and t1 and thus emits
 288 ;   p0 = &input[sext(x + y)];
 289 ;   p1 = &p0[5];
 290 define void @reunion(i32 %x, i32 %y, ptr %input) {
 291 ; IR-LABEL: define void @reunion(
 292 ; IR-SAME: i32 [[X:%.*]], i32 [[Y:%.*]], ptr [[INPUT:%.*]]) {
 293 ; IR-NEXT:  entry:
 294 ; IR-NEXT:    [[XY:%.*]] = add nsw i32 [[X]], [[Y]]
 295 ; IR-NEXT:    [[TMP0:%.*]] = sext i32 [[XY]] to i64
 296 ; IR-NEXT:    [[P0:%.*]] = getelementptr float, ptr [[INPUT]], i64 [[TMP0]]
 297 ; IR-NEXT:    [[V0:%.*]] = load float, ptr [[P0]], align 4
 298 ; IR-NEXT:    call void @use(float [[V0]])
 299 ; IR-NEXT:    [[P13:%.*]] = getelementptr inbounds i8, ptr [[P0]], i64 20
 300 ; IR-NEXT:    [[V1:%.*]] = load float, ptr [[P13]], align 4
 301 ; IR-NEXT:    call void @use(float [[V1]])
 302 ; IR-NEXT:    ret void
 303 ;
 304 ; PTX-LABEL: reunion(
 305 entry:
 306   %xy = add nsw i32 %x, %y
 307   %0 = sext i32 %xy to i64
 308   %p0 = getelementptr inbounds float, ptr %input, i64 %0
 309   %v0 = load float, ptr %p0, align 4
 310 ; PTX: ld.f32 %f{{[0-9]+}}, [[[p0:%rd[0-9]+]]]
 311   call void @use(float %v0)
 312
 313   %y5 = add nsw i32 %y, 5
 314   %xy5 = add nsw i32 %x, %y5
 315   %1 = sext i32 %xy5 to i64
 316   %p1 = getelementptr inbounds float, ptr %input, i64 %1
 317   %v1 = load float, ptr %p1, align 4
 318 ; PTX: ld.f32 %f{{[0-9]+}}, [[[p0]]+20]
 319   call void @use(float %v1)
 320
 321   ret void
 322 }
 323
 324 declare void @use(float)