test/CodeGen/NVPTX/st-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
   9 define void @st_global_i8(i8 addrspace(1)* %ptr, i8 %a) {
  10 ; PTX32: st.global.u8 [%r{{[0-9]+}}], %rc{{[0-9]+}}
  11 ; PTX32: ret
  12 ; PTX64: st.global.u8 [%rl{{[0-9]+}}], %rc{{[0-9]+}}
  13 ; PTX64: ret
  14   store i8 %a, i8 addrspace(1)* %ptr
  15   ret void
  16 }
  17
  18 define void @st_shared_i8(i8 addrspace(3)* %ptr, i8 %a) {
  19 ; PTX32: st.shared.u8 [%r{{[0-9]+}}], %rc{{[0-9]+}}
  20 ; PTX32: ret
  21 ; PTX64: st.shared.u8 [%rl{{[0-9]+}}], %rc{{[0-9]+}}
  22 ; PTX64: ret
  23   store i8 %a, i8 addrspace(3)* %ptr
  24   ret void
  25 }
  26
  27 define void @st_local_i8(i8 addrspace(5)* %ptr, i8 %a) {
  28 ; PTX32: st.local.u8 [%r{{[0-9]+}}], %rc{{[0-9]+}}
  29 ; PTX32: ret
  30 ; PTX64: st.local.u8 [%rl{{[0-9]+}}], %rc{{[0-9]+}}
  31 ; PTX64: ret
  32   store i8 %a, i8 addrspace(5)* %ptr
  33   ret void
  34 }
  35
  36 ;; i16
  37
  38 define void @st_global_i16(i16 addrspace(1)* %ptr, i16 %a) {
  39 ; PTX32: st.global.u16 [%r{{[0-9]+}}], %rs{{[0-9]+}}
  40 ; PTX32: ret
  41 ; PTX64: st.global.u16 [%rl{{[0-9]+}}], %rs{{[0-9]+}}
  42 ; PTX64: ret
  43   store i16 %a, i16 addrspace(1)* %ptr
  44   ret void
  45 }
  46
  47 define void @st_shared_i16(i16 addrspace(3)* %ptr, i16 %a) {
  48 ; PTX32: st.shared.u16 [%r{{[0-9]+}}], %rs{{[0-9]+}}
  49 ; PTX32: ret
  50 ; PTX64: st.shared.u16 [%rl{{[0-9]+}}], %rs{{[0-9]+}}
  51 ; PTX64: ret
  52   store i16 %a, i16 addrspace(3)* %ptr
  53   ret void
  54 }
  55
  56 define void @st_local_i16(i16 addrspace(5)* %ptr, i16 %a) {
  57 ; PTX32: st.local.u16 [%r{{[0-9]+}}], %rs{{[0-9]+}}
  58 ; PTX32: ret
  59 ; PTX64: st.local.u16 [%rl{{[0-9]+}}], %rs{{[0-9]+}}
  60 ; PTX64: ret
  61   store i16 %a, i16 addrspace(5)* %ptr
  62   ret void
  63 }
  64
  65 ;; i32
  66
  67 define void @st_global_i32(i32 addrspace(1)* %ptr, i32 %a) {
  68 ; PTX32: st.global.u32 [%r{{[0-9]+}}], %r{{[0-9]+}}
  69 ; PTX32: ret
  70 ; PTX64: st.global.u32 [%rl{{[0-9]+}}], %r{{[0-9]+}}
  71 ; PTX64: ret
  72   store i32 %a, i32 addrspace(1)* %ptr
  73   ret void
  74 }
  75
  76 define void @st_shared_i32(i32 addrspace(3)* %ptr, i32 %a) {
  77 ; PTX32: st.shared.u32 [%r{{[0-9]+}}], %r{{[0-9]+}}
  78 ; PTX32: ret
  79 ; PTX64: st.shared.u32 [%rl{{[0-9]+}}], %r{{[0-9]+}}
  80 ; PTX64: ret
  81   store i32 %a, i32 addrspace(3)* %ptr
  82   ret void
  83 }
  84
  85 define void @st_local_i32(i32 addrspace(5)* %ptr, i32 %a) {
  86 ; PTX32: st.local.u32 [%r{{[0-9]+}}], %r{{[0-9]+}}
  87 ; PTX32: ret
  88 ; PTX64: st.local.u32 [%rl{{[0-9]+}}], %r{{[0-9]+}}
  89 ; PTX64: ret
  90   store i32 %a, i32 addrspace(5)* %ptr
  91   ret void
  92 }
  93
  94 ;; i64
  95
  96 define void @st_global_i64(i64 addrspace(1)* %ptr, i64 %a) {
  97 ; PTX32: st.global.u64 [%r{{[0-9]+}}], %rl{{[0-9]+}}
  98 ; PTX32: ret
  99 ; PTX64: st.global.u64 [%rl{{[0-9]+}}], %rl{{[0-9]+}}
 100 ; PTX64: ret
 101   store i64 %a, i64 addrspace(1)* %ptr
 102   ret void
 103 }
 104
 105 define void @st_shared_i64(i64 addrspace(3)* %ptr, i64 %a) {
 106 ; PTX32: st.shared.u64 [%r{{[0-9]+}}], %rl{{[0-9]+}}
 107 ; PTX32: ret
 108 ; PTX64: st.shared.u64 [%rl{{[0-9]+}}], %rl{{[0-9]+}}
 109 ; PTX64: ret
 110   store i64 %a, i64 addrspace(3)* %ptr
 111   ret void
 112 }
 113
 114 define void @st_local_i64(i64 addrspace(5)* %ptr, i64 %a) {
 115 ; PTX32: st.local.u64 [%r{{[0-9]+}}], %rl{{[0-9]+}}
 116 ; PTX32: ret
 117 ; PTX64: st.local.u64 [%rl{{[0-9]+}}], %rl{{[0-9]+}}
 118 ; PTX64: ret
 119   store i64 %a, i64 addrspace(5)* %ptr
 120   ret void
 121 }
 122
 123 ;; f32
 124
 125 define void @st_global_f32(float addrspace(1)* %ptr, float %a) {
 126 ; PTX32: st.global.f32 [%r{{[0-9]+}}], %f{{[0-9]+}}
 127 ; PTX32: ret
 128 ; PTX64: st.global.f32 [%rl{{[0-9]+}}], %f{{[0-9]+}}
 129 ; PTX64: ret
 130   store float %a, float addrspace(1)* %ptr
 131   ret void
 132 }
 133
 134 define void @st_shared_f32(float addrspace(3)* %ptr, float %a) {
 135 ; PTX32: st.shared.f32 [%r{{[0-9]+}}], %f{{[0-9]+}}
 136 ; PTX32: ret
 137 ; PTX64: st.shared.f32 [%rl{{[0-9]+}}], %f{{[0-9]+}}
 138 ; PTX64: ret
 139   store float %a, float addrspace(3)* %ptr
 140   ret void
 141 }
 142
 143 define void @st_local_f32(float addrspace(5)* %ptr, float %a) {
 144 ; PTX32: st.local.f32 [%r{{[0-9]+}}], %f{{[0-9]+}}
 145 ; PTX32: ret
 146 ; PTX64: st.local.f32 [%rl{{[0-9]+}}], %f{{[0-9]+}}
 147 ; PTX64: ret
 148   store float %a, float addrspace(5)* %ptr
 149   ret void
 150 }
 151
 152 ;; f64
 153
 154 define void @st_global_f64(double addrspace(1)* %ptr, double %a) {
 155 ; PTX32: st.global.f64 [%r{{[0-9]+}}], %fl{{[0-9]+}}
 156 ; PTX32: ret
 157 ; PTX64: st.global.f64 [%rl{{[0-9]+}}], %fl{{[0-9]+}}
 158 ; PTX64: ret
 159   store double %a, double addrspace(1)* %ptr
 160   ret void
 161 }
 162
 163 define void @st_shared_f64(double addrspace(3)* %ptr, double %a) {
 164 ; PTX32: st.shared.f64 [%r{{[0-9]+}}], %fl{{[0-9]+}}
 165 ; PTX32: ret
 166 ; PTX64: st.shared.f64 [%rl{{[0-9]+}}], %fl{{[0-9]+}}
 167 ; PTX64: ret
 168   store double %a, double addrspace(3)* %ptr
 169   ret void
 170 }
 171
 172 define void @st_local_f64(double addrspace(5)* %ptr, double %a) {
 173 ; PTX32: st.local.f64 [%r{{[0-9]+}}], %fl{{[0-9]+}}
 174 ; PTX32: ret
 175 ; PTX64: st.local.f64 [%rl{{[0-9]+}}], %fl{{[0-9]+}}
 176 ; PTX64: ret
 177   store double %a, double addrspace(5)* %ptr
 178   ret void
 179 }