llvm/test/CodeGen/NVPTX/ld-addrspace.ll

   1 ; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s --check-prefixes=ALL,G32,LS32
   2 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s --check-prefixes=ALL,G64,LS64
   3 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 --nvptx-short-ptr | FileCheck %s --check-prefixes=G64,LS32
   4 ; RUN: %if ptxas && !ptxas-12.0 %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %}
   5 ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
   6 ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 --nvptx-short-ptr | %ptxas-verify %}
   7
   8
   9 ;; i8
  10 define i8 @ld_global_i8(ptr addrspace(1) %ptr) {
  11 ; ALL-LABEL: ld_global_i8
  12 ; G32: ld.global.u8 %{{.*}}, [%r{{[0-9]+}}]
  13 ; G64: ld.global.u8 %{{.*}}, [%rd{{[0-9]+}}]
  14 ; ALL: ret
  15   %a = load i8, ptr addrspace(1) %ptr
  16   ret i8 %a
  17 }
  18 define i8 @ld_shared_i8(ptr addrspace(3) %ptr) {
  19 ; ALL-LABEL: ld_shared_i8
  20 ; LS32: ld.shared.u8 %{{.*}}, [%r{{[0-9]+}}]
  21 ; LS64: ld.shared.u8 %{{.*}}, [%rd{{[0-9]+}}]
  22 ; ALL: ret
  23   %a = load i8, ptr addrspace(3) %ptr
  24   ret i8 %a
  25 }
  26 define i8 @ld_local_i8(ptr addrspace(5) %ptr) {
  27 ; ALL-LABEL: ld_local_i8
  28 ; LS32: ld.local.u8 %{{.*}}, [%r{{[0-9]+}}]
  29 ; LS64: ld.local.u8 %{{.*}}, [%rd{{[0-9]+}}]
  30 ; ALL: ret
  31   %a = load i8, ptr addrspace(5) %ptr
  32   ret i8 %a
  33 }
  34
  35 ;; i16
  36 define i16 @ld_global_i16(ptr addrspace(1) %ptr) {
  37 ; ALL-LABEL: ld_global_i16
  38 ; G32: ld.global.u16 %{{.*}}, [%r{{[0-9]+}}]
  39 ; G64: ld.global.u16 %{{.*}}, [%rd{{[0-9]+}}]
  40 ; ALL: ret
  41   %a = load i16, ptr addrspace(1) %ptr
  42   ret i16 %a
  43 }
  44 define i16 @ld_shared_i16(ptr addrspace(3) %ptr) {
  45 ; ALL-LABEL: ld_shared_i16
  46 ; LS32: ld.shared.u16 %{{.*}}, [%r{{[0-9]+}}]
  47 ; LS64: ld.shared.u16 %{{.*}}, [%rd{{[0-9]+}}]
  48 ; ALL: ret
  49   %a = load i16, ptr addrspace(3) %ptr
  50   ret i16 %a
  51 }
  52 define i16 @ld_local_i16(ptr addrspace(5) %ptr) {
  53 ; ALL-LABEL: ld_local_i16
  54 ; LS32: ld.local.u16 %{{.*}}, [%r{{[0-9]+}}]
  55 ; LS64: ld.local.u16 %{{.*}}, [%rd{{[0-9]+}}]
  56 ; ALL: ret
  57   %a = load i16, ptr addrspace(5) %ptr
  58   ret i16 %a
  59 }
  60
  61 ;; i32
  62 define i32 @ld_global_i32(ptr addrspace(1) %ptr) {
  63 ; ALL-LABEL: ld_global_i32
  64 ; G32: ld.global.u32 %{{.*}}, [%r{{[0-9]+}}]
  65 ; G64: ld.global.u32 %{{.*}}, [%rd{{[0-9]+}}]
  66 ; ALL: ret
  67   %a = load i32, ptr addrspace(1) %ptr
  68   ret i32 %a
  69 }
  70 define i32 @ld_shared_i32(ptr addrspace(3) %ptr) {
  71 ; ALL-LABEL: ld_shared_i32
  72 ; LS32: ld.shared.u32 %{{.*}}, [%r{{[0-9]+}}]
  73 ; LS64: ld.shared.u32 %{{.*}}, [%rd{{[0-9]+}}]
  74 ; PTX64: ret
  75   %a = load i32, ptr addrspace(3) %ptr
  76   ret i32 %a
  77 }
  78 define i32 @ld_local_i32(ptr addrspace(5) %ptr) {
  79 ; ALL-LABEL: ld_local_i32
  80 ; LS32: ld.local.u32 %{{.*}}, [%r{{[0-9]+}}]
  81 ; LS64: ld.local.u32 %{{.*}}, [%rd{{[0-9]+}}]
  82 ; ALL: ret
  83   %a = load i32, ptr addrspace(5) %ptr
  84   ret i32 %a
  85 }
  86
  87 ;; i64
  88 define i64 @ld_global_i64(ptr addrspace(1) %ptr) {
  89 ; ALL-LABEL: ld_global_i64
  90 ; G32: ld.global.u64 %{{.*}}, [%r{{[0-9]+}}]
  91 ; G64: ld.global.u64 %{{.*}}, [%rd{{[0-9]+}}]
  92 ; ALL: ret
  93   %a = load i64, ptr addrspace(1) %ptr
  94   ret i64 %a
  95 }
  96 define i64 @ld_shared_i64(ptr addrspace(3) %ptr) {
  97 ; ALL-LABEL: ld_shared_i64
  98 ; LS32: ld.shared.u64 %{{.*}}, [%r{{[0-9]+}}]
  99 ; LS64: ld.shared.u64 %{{.*}}, [%rd{{[0-9]+}}]
 100 ; ALL: ret
 101   %a = load i64, ptr addrspace(3) %ptr
 102   ret i64 %a
 103 }
 104 define i64 @ld_local_i64(ptr addrspace(5) %ptr) {
 105 ; ALL-LABEL: ld_local_i64
 106 ; LS32: ld.local.u64 %{{.*}}, [%r{{[0-9]+}}]
 107 ; LS64: ld.local.u64 %{{.*}}, [%rd{{[0-9]+}}]
 108 ; ALL: ret
 109   %a = load i64, ptr addrspace(5) %ptr
 110   ret i64 %a
 111 }
 112
 113 ;; f32
 114 define float @ld_global_f32(ptr addrspace(1) %ptr) {
 115 ; ALL-LABEL: ld_global_f32
 116 ; G32: ld.global.f32 %{{.*}}, [%r{{[0-9]+}}]
 117 ; G64: ld.global.f32 %{{.*}}, [%rd{{[0-9]+}}]
 118 ; ALL: ret
 119   %a = load float, ptr addrspace(1) %ptr
 120   ret float %a
 121 }
 122 define float @ld_shared_f32(ptr addrspace(3) %ptr) {
 123 ; ALL-LABEL: ld_shared_f32
 124 ; LS32: ld.shared.f32 %{{.*}}, [%r{{[0-9]+}}]
 125 ; LS64: ld.shared.f32 %{{.*}}, [%rd{{[0-9]+}}]
 126 ; ALL: ret
 127   %a = load float, ptr addrspace(3) %ptr
 128   ret float %a
 129 }
 130 define float @ld_local_f32(ptr addrspace(5) %ptr) {
 131 ; ALL-LABEL: ld_local_f32
 132 ; LS32: ld.local.f32 %{{.*}}, [%r{{[0-9]+}}]
 133 ; LS64: ld.local.f32 %{{.*}}, [%rd{{[0-9]+}}]
 134 ; ALL: ret
 135   %a = load float, ptr addrspace(5) %ptr
 136   ret float %a
 137 }
 138
 139 ;; f64
 140 define double @ld_global_f64(ptr addrspace(1) %ptr) {
 141 ; ALL-LABEL: ld_global_f64
 142 ; G32: ld.global.f64 %{{.*}}, [%r{{[0-9]+}}]
 143 ; G64: ld.global.f64 %{{.*}}, [%rd{{[0-9]+}}]
 144 ; ALL: ret
 145   %a = load double, ptr addrspace(1) %ptr
 146   ret double %a
 147 }
 148 define double @ld_shared_f64(ptr addrspace(3) %ptr) {
 149 ; ALL-LABEL: ld_shared_f64
 150 ; LS32: ld.shared.f64 %{{.*}}, [%r{{[0-9]+}}]
 151 ; LS64: ld.shared.f64 %{{.*}}, [%rd{{[0-9]+}}]
 152 ; ALL: ret
 153   %a = load double, ptr addrspace(3) %ptr
 154   ret double %a
 155 }
 156 define double @ld_local_f64(ptr addrspace(5) %ptr) {
 157 ; ALL-LABEL: ld_local_f64
 158 ; LS32: ld.local.f64 %{{.*}}, [%r{{[0-9]+}}]
 159 ; LS64: ld.local.f64 %{{.*}}, [%rd{{[0-9]+}}]
 160 ; ALL: ret
 161   %a = load double, ptr addrspace(5) %ptr
 162   ret double %a
 163 }