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