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
9 define void @st_global_i8(i8 addrspace(1)* %ptr, i8 %a) {
10 ; PTX32: st.global.u8 [%r{{[0-9]+}}], %rc{{[0-9]+}}
12 ; PTX64: st.global.u8 [%rl{{[0-9]+}}], %rc{{[0-9]+}}
14 store i8 %a, i8 addrspace(1)* %ptr
18 define void @st_shared_i8(i8 addrspace(3)* %ptr, i8 %a) {
19 ; PTX32: st.shared.u8 [%r{{[0-9]+}}], %rc{{[0-9]+}}
21 ; PTX64: st.shared.u8 [%rl{{[0-9]+}}], %rc{{[0-9]+}}
23 store i8 %a, i8 addrspace(3)* %ptr
27 define void @st_local_i8(i8 addrspace(5)* %ptr, i8 %a) {
28 ; PTX32: st.local.u8 [%r{{[0-9]+}}], %rc{{[0-9]+}}
30 ; PTX64: st.local.u8 [%rl{{[0-9]+}}], %rc{{[0-9]+}}
32 store i8 %a, i8 addrspace(5)* %ptr
38 define void @st_global_i16(i16 addrspace(1)* %ptr, i16 %a) {
39 ; PTX32: st.global.u16 [%r{{[0-9]+}}], %rs{{[0-9]+}}
41 ; PTX64: st.global.u16 [%rl{{[0-9]+}}], %rs{{[0-9]+}}
43 store i16 %a, i16 addrspace(1)* %ptr
47 define void @st_shared_i16(i16 addrspace(3)* %ptr, i16 %a) {
48 ; PTX32: st.shared.u16 [%r{{[0-9]+}}], %rs{{[0-9]+}}
50 ; PTX64: st.shared.u16 [%rl{{[0-9]+}}], %rs{{[0-9]+}}
52 store i16 %a, i16 addrspace(3)* %ptr
56 define void @st_local_i16(i16 addrspace(5)* %ptr, i16 %a) {
57 ; PTX32: st.local.u16 [%r{{[0-9]+}}], %rs{{[0-9]+}}
59 ; PTX64: st.local.u16 [%rl{{[0-9]+}}], %rs{{[0-9]+}}
61 store i16 %a, i16 addrspace(5)* %ptr
67 define void @st_global_i32(i32 addrspace(1)* %ptr, i32 %a) {
68 ; PTX32: st.global.u32 [%r{{[0-9]+}}], %r{{[0-9]+}}
70 ; PTX64: st.global.u32 [%rl{{[0-9]+}}], %r{{[0-9]+}}
72 store i32 %a, i32 addrspace(1)* %ptr
76 define void @st_shared_i32(i32 addrspace(3)* %ptr, i32 %a) {
77 ; PTX32: st.shared.u32 [%r{{[0-9]+}}], %r{{[0-9]+}}
79 ; PTX64: st.shared.u32 [%rl{{[0-9]+}}], %r{{[0-9]+}}
81 store i32 %a, i32 addrspace(3)* %ptr
85 define void @st_local_i32(i32 addrspace(5)* %ptr, i32 %a) {
86 ; PTX32: st.local.u32 [%r{{[0-9]+}}], %r{{[0-9]+}}
88 ; PTX64: st.local.u32 [%rl{{[0-9]+}}], %r{{[0-9]+}}
90 store i32 %a, i32 addrspace(5)* %ptr
96 define void @st_global_i64(i64 addrspace(1)* %ptr, i64 %a) {
97 ; PTX32: st.global.u64 [%r{{[0-9]+}}], %rl{{[0-9]+}}
99 ; PTX64: st.global.u64 [%rl{{[0-9]+}}], %rl{{[0-9]+}}
101 store i64 %a, i64 addrspace(1)* %ptr
105 define void @st_shared_i64(i64 addrspace(3)* %ptr, i64 %a) {
106 ; PTX32: st.shared.u64 [%r{{[0-9]+}}], %rl{{[0-9]+}}
108 ; PTX64: st.shared.u64 [%rl{{[0-9]+}}], %rl{{[0-9]+}}
110 store i64 %a, i64 addrspace(3)* %ptr
114 define void @st_local_i64(i64 addrspace(5)* %ptr, i64 %a) {
115 ; PTX32: st.local.u64 [%r{{[0-9]+}}], %rl{{[0-9]+}}
117 ; PTX64: st.local.u64 [%rl{{[0-9]+}}], %rl{{[0-9]+}}
119 store i64 %a, i64 addrspace(5)* %ptr
125 define void @st_global_f32(float addrspace(1)* %ptr, float %a) {
126 ; PTX32: st.global.f32 [%r{{[0-9]+}}], %f{{[0-9]+}}
128 ; PTX64: st.global.f32 [%rl{{[0-9]+}}], %f{{[0-9]+}}
130 store float %a, float addrspace(1)* %ptr
134 define void @st_shared_f32(float addrspace(3)* %ptr, float %a) {
135 ; PTX32: st.shared.f32 [%r{{[0-9]+}}], %f{{[0-9]+}}
137 ; PTX64: st.shared.f32 [%rl{{[0-9]+}}], %f{{[0-9]+}}
139 store float %a, float addrspace(3)* %ptr
143 define void @st_local_f32(float addrspace(5)* %ptr, float %a) {
144 ; PTX32: st.local.f32 [%r{{[0-9]+}}], %f{{[0-9]+}}
146 ; PTX64: st.local.f32 [%rl{{[0-9]+}}], %f{{[0-9]+}}
148 store float %a, float addrspace(5)* %ptr
154 define void @st_global_f64(double addrspace(1)* %ptr, double %a) {
155 ; PTX32: st.global.f64 [%r{{[0-9]+}}], %fl{{[0-9]+}}
157 ; PTX64: st.global.f64 [%rl{{[0-9]+}}], %fl{{[0-9]+}}
159 store double %a, double addrspace(1)* %ptr
163 define void @st_shared_f64(double addrspace(3)* %ptr, double %a) {
164 ; PTX32: st.shared.f64 [%r{{[0-9]+}}], %fl{{[0-9]+}}
166 ; PTX64: st.shared.f64 [%rl{{[0-9]+}}], %fl{{[0-9]+}}
168 store double %a, double addrspace(3)* %ptr
172 define void @st_local_f64(double addrspace(5)* %ptr, double %a) {
173 ; PTX32: st.local.f64 [%r{{[0-9]+}}], %fl{{[0-9]+}}
175 ; PTX64: st.local.f64 [%rl{{[0-9]+}}], %fl{{[0-9]+}}
177 store double %a, double addrspace(5)* %ptr