test/CodeGen/NVPTX/ld-addrspace.ll

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