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 9define 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 18define 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 27define 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 38define 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 47define 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 56define 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 67define 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 76define 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 85define 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 96define 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 105define 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 114define 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 125define 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 134define 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 143define 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 154define 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 163define 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 172define 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} 180