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