1; RUN: llc < %s -mtriple=nvptx -mcpu=sm_20 | FileCheck %s --check-prefixes=ALL,G32,LS32 2; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 | FileCheck %s --check-prefixes=ALL,G64,LS64 3; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 --nvptx-short-ptr | FileCheck %s --check-prefixes=G64,LS32 4; RUN: %if ptxas && !ptxas-12.0 %{ llc < %s -mtriple=nvptx -mcpu=sm_20 | %ptxas-verify %} 5; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 | %ptxas-verify %} 6; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 --nvptx-short-ptr | %ptxas-verify %} 7 8;; i8 9; ALL-LABEL: st_global_i8 10define void @st_global_i8(ptr 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, ptr addrspace(1) %ptr 15 ret void 16} 17; ALL-LABEL: st_shared_i8 18define void @st_shared_i8(ptr 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, ptr addrspace(3) %ptr 23 ret void 24} 25; ALL-LABEL: st_local_i8 26define void @st_local_i8(ptr 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, ptr addrspace(5) %ptr 31 ret void 32} 33 34;; i16 35; ALL-LABEL: st_global_i16 36define void @st_global_i16(ptr 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, ptr addrspace(1) %ptr 41 ret void 42} 43; ALL-LABEL: st_shared_i16 44define void @st_shared_i16(ptr 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, ptr addrspace(3) %ptr 49 ret void 50} 51; ALL-LABEL: st_local_i16 52define void @st_local_i16(ptr 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, ptr addrspace(5) %ptr 57 ret void 58} 59 60;; i32 61; ALL-LABEL: st_global_i32 62define void @st_global_i32(ptr 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, ptr addrspace(1) %ptr 67 ret void 68} 69; ALL-LABEL: st_shared_i32 70define void @st_shared_i32(ptr 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, ptr addrspace(3) %ptr 75 ret void 76} 77; ALL-LABEL: st_local_i32 78define void @st_local_i32(ptr 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, ptr addrspace(5) %ptr 83 ret void 84} 85 86;; i64 87; ALL-LABEL: st_global_i64 88define void @st_global_i64(ptr 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, ptr addrspace(1) %ptr 93 ret void 94} 95; ALL-LABEL: st_shared_i64 96define void @st_shared_i64(ptr 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, ptr addrspace(3) %ptr 101 ret void 102} 103; ALL-LABEL: st_local_i64 104define void @st_local_i64(ptr 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, ptr addrspace(5) %ptr 109 ret void 110} 111 112;; f32 113; ALL-LABEL: st_global_f32 114define void @st_global_f32(ptr 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, ptr addrspace(1) %ptr 119 ret void 120} 121; ALL-LABEL: st_shared_f32 122define void @st_shared_f32(ptr 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, ptr addrspace(3) %ptr 127 ret void 128} 129; ALL-LABEL: st_local_f32 130define void @st_local_f32(ptr 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, ptr addrspace(5) %ptr 135 ret void 136} 137 138;; f64 139; ALL-LABEL: st_global_f64 140define void @st_global_f64(ptr 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, ptr addrspace(1) %ptr 145 ret void 146} 147; ALL-LABEL: st_shared_f64 148define void @st_shared_f64(ptr 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, ptr addrspace(3) %ptr 153 ret void 154} 155; ALL-LABEL: st_local_f64 156define void @st_local_f64(ptr 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, ptr addrspace(5) %ptr 161 ret void 162} 163