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
8define 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
16define 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
24define 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
34define 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
42define 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
50define 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
60define 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
68define 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
76define 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
86define 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
94define 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
102define 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
112define 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
120define 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
128define 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
138define 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
146define 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
154define 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}
161