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