llvm/test/CodeGen/NVPTX/ldg-invariant.ll

   1 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_35 -verify-machineinstrs | FileCheck %s
   2 ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_35 | %ptxas-verify %}
   3
   4 ; Check that invariant loads from the global addrspace are lowered to
   5 ; ld.global.nc.
   6
   7 ; CHECK-LABEL: @ld_global
   8 define i32 @ld_global(ptr addrspace(1) %ptr) {
   9 ; CHECK: ld.global.nc.{{[a-z]}}32
  10   %a = load i32, ptr addrspace(1) %ptr, !invariant.load !0
  11   ret i32 %a
  12 }
  13
  14 ; CHECK-LABEL: @ld_global_v2f16
  15 define half @ld_global_v2f16(ptr addrspace(1) %ptr) {
  16 ; Load of v2f16 is weird. We consider it to be a legal type, which happens to be
  17 ; loaded/stored as a 32-bit scalar.
  18 ; CHECK: ld.global.nc.u32
  19   %a = load <2 x half>, ptr addrspace(1) %ptr, !invariant.load !0
  20   %v1 = extractelement <2 x half> %a, i32 0
  21   %v2 = extractelement <2 x half> %a, i32 1
  22   %sum = fadd half %v1, %v2
  23   ret half %sum
  24 }
  25
  26 ; CHECK-LABEL: @ld_global_v4f16
  27 define half @ld_global_v4f16(ptr addrspace(1) %ptr) {
  28 ; Larger f16 vectors may be split into individual f16 elements and multiple
  29 ; loads/stores may be vectorized using f16 element type. Practically it's
  30 ; limited to v4 variant only.
  31 ; CHECK: ld.global.nc.v4.u16
  32   %a = load <4 x half>, ptr addrspace(1) %ptr, !invariant.load !0
  33   %v1 = extractelement <4 x half> %a, i32 0
  34   %v2 = extractelement <4 x half> %a, i32 1
  35   %v3 = extractelement <4 x half> %a, i32 2
  36   %v4 = extractelement <4 x half> %a, i32 3
  37   %sum1 = fadd half %v1, %v2
  38   %sum2 = fadd half %v3, %v4
  39   %sum = fadd half %sum1, %sum2
  40   ret half %sum
  41 }
  42
  43 ; CHECK-LABEL: @ld_global_v8f16
  44 define half @ld_global_v8f16(ptr addrspace(1) %ptr) {
  45 ; Larger vectors are, again, loaded as v4i32. PTX has no v8 variants of loads/stores,
  46 ; so load/store vectorizer has to convert v8f16 -> v4 x v2f16.
  47 ; CHECK: ld.global.nc.v4.u32
  48   %a = load <8 x half>, ptr addrspace(1) %ptr, !invariant.load !0
  49   %v1 = extractelement <8 x half> %a, i32 0
  50   %v2 = extractelement <8 x half> %a, i32 2
  51   %v3 = extractelement <8 x half> %a, i32 4
  52   %v4 = extractelement <8 x half> %a, i32 6
  53   %sum1 = fadd half %v1, %v2
  54   %sum2 = fadd half %v3, %v4
  55   %sum = fadd half %sum1, %sum2
  56   ret half %sum
  57 }
  58
  59 ; CHECK-LABEL: @ld_global_v2i32
  60 define i32 @ld_global_v2i32(ptr addrspace(1) %ptr) {
  61 ; CHECK: ld.global.nc.v2.{{[a-z]}}32
  62   %a = load <2 x i32>, ptr addrspace(1) %ptr, !invariant.load !0
  63   %v1 = extractelement <2 x i32> %a, i32 0
  64   %v2 = extractelement <2 x i32> %a, i32 1
  65   %sum = add i32 %v1, %v2
  66   ret i32 %sum
  67 }
  68
  69 ; CHECK-LABEL: @ld_global_v4i32
  70 define i32 @ld_global_v4i32(ptr addrspace(1) %ptr) {
  71 ; CHECK: ld.global.nc.v4.{{[a-z]}}32
  72   %a = load <4 x i32>, ptr addrspace(1) %ptr, !invariant.load !0
  73   %v1 = extractelement <4 x i32> %a, i32 0
  74   %v2 = extractelement <4 x i32> %a, i32 1
  75   %v3 = extractelement <4 x i32> %a, i32 2
  76   %v4 = extractelement <4 x i32> %a, i32 3
  77   %sum1 = add i32 %v1, %v2
  78   %sum2 = add i32 %v3, %v4
  79   %sum3 = add i32 %sum1, %sum2
  80   ret i32 %sum3
  81 }
  82
  83 ; CHECK-LABEL: @ld_not_invariant
  84 define i32 @ld_not_invariant(ptr addrspace(1) %ptr) {
  85 ; CHECK: ld.global.{{[a-z]}}32
  86   %a = load i32, ptr addrspace(1) %ptr
  87   ret i32 %a
  88 }
  89
  90 ; CHECK-LABEL: @ld_not_global_addrspace
  91 define i32 @ld_not_global_addrspace(ptr addrspace(0) %ptr) {
  92 ; CHECK: ld.{{[a-z]}}32
  93   %a = load i32, ptr addrspace(0) %ptr
  94   ret i32 %a
  95 }
  96
  97 !0 = !{}