163 lines · plain
1; RUN: llc < %s -mtriple=nvptx -mcpu=sm_20 | FileCheck %s --check-prefixes=ALL,G32,LS322; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 | FileCheck %s --check-prefixes=ALL,G64,LS643; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 --nvptx-short-ptr | FileCheck %s --check-prefixes=G64,LS324; RUN: %if ptxas-ptr32 %{ 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;; i89; ALL-LABEL: st_global_i810define void @st_global_i8(ptr addrspace(1) %ptr, i8 %a) {11; G32: st.global.b8 [%r{{[0-9]+}}], %rs{{[0-9]+}}12; G64: st.global.b8 [%rd{{[0-9]+}}], %rs{{[0-9]+}}13; ALL: ret14 store i8 %a, ptr addrspace(1) %ptr15 ret void16}17; ALL-LABEL: st_shared_i818define void @st_shared_i8(ptr addrspace(3) %ptr, i8 %a) {19; LS32: st.shared.b8 [%r{{[0-9]+}}], %rs{{[0-9]+}}20; LS64: st.shared.b8 [%rd{{[0-9]+}}], %rs{{[0-9]+}}21; ALL: ret22 store i8 %a, ptr addrspace(3) %ptr23 ret void24}25; ALL-LABEL: st_local_i826define void @st_local_i8(ptr addrspace(5) %ptr, i8 %a) {27; LS32: st.local.b8 [%r{{[0-9]+}}], %rs{{[0-9]+}}28; LS64: st.local.b8 [%rd{{[0-9]+}}], %rs{{[0-9]+}}29; ALL: ret30 store i8 %a, ptr addrspace(5) %ptr31 ret void32}33 34;; i1635; ALL-LABEL: st_global_i1636define void @st_global_i16(ptr addrspace(1) %ptr, i16 %a) {37; G32: st.global.b16 [%r{{[0-9]+}}], %rs{{[0-9]+}}38; G64: st.global.b16 [%rd{{[0-9]+}}], %rs{{[0-9]+}}39; ALL: ret40 store i16 %a, ptr addrspace(1) %ptr41 ret void42}43; ALL-LABEL: st_shared_i1644define void @st_shared_i16(ptr addrspace(3) %ptr, i16 %a) {45; LS32: st.shared.b16 [%r{{[0-9]+}}], %rs{{[0-9]+}}46; LS64: st.shared.b16 [%rd{{[0-9]+}}], %rs{{[0-9]+}}47; ALL: ret48 store i16 %a, ptr addrspace(3) %ptr49 ret void50}51; ALL-LABEL: st_local_i1652define void @st_local_i16(ptr addrspace(5) %ptr, i16 %a) {53; LS32: st.local.b16 [%r{{[0-9]+}}], %rs{{[0-9]+}}54; LS64: st.local.b16 [%rd{{[0-9]+}}], %rs{{[0-9]+}}55; ALL: ret56 store i16 %a, ptr addrspace(5) %ptr57 ret void58}59 60;; i3261; ALL-LABEL: st_global_i3262define void @st_global_i32(ptr addrspace(1) %ptr, i32 %a) {63; G32: st.global.b32 [%r{{[0-9]+}}], %r{{[0-9]+}}64; G64: st.global.b32 [%rd{{[0-9]+}}], %r{{[0-9]+}}65; ALL: ret66 store i32 %a, ptr addrspace(1) %ptr67 ret void68}69; ALL-LABEL: st_shared_i3270define void @st_shared_i32(ptr addrspace(3) %ptr, i32 %a) {71; LS32: st.shared.b32 [%r{{[0-9]+}}], %r{{[0-9]+}}72; LS64: st.shared.b32 [%rd{{[0-9]+}}], %r{{[0-9]+}}73; PTX64: ret74 store i32 %a, ptr addrspace(3) %ptr75 ret void76}77; ALL-LABEL: st_local_i3278define void @st_local_i32(ptr addrspace(5) %ptr, i32 %a) {79; LS32: st.local.b32 [%r{{[0-9]+}}], %r{{[0-9]+}}80; LS64: st.local.b32 [%rd{{[0-9]+}}], %r{{[0-9]+}}81; ALL: ret82 store i32 %a, ptr addrspace(5) %ptr83 ret void84}85 86;; i6487; ALL-LABEL: st_global_i6488define void @st_global_i64(ptr addrspace(1) %ptr, i64 %a) {89; G32: st.global.b64 [%r{{[0-9]+}}], %rd{{[0-9]+}}90; G64: st.global.b64 [%rd{{[0-9]+}}], %rd{{[0-9]+}}91; ALL: ret92 store i64 %a, ptr addrspace(1) %ptr93 ret void94}95; ALL-LABEL: st_shared_i6496define void @st_shared_i64(ptr addrspace(3) %ptr, i64 %a) {97; LS32: st.shared.b64 [%r{{[0-9]+}}], %rd{{[0-9]+}}98; LS64: st.shared.b64 [%rd{{[0-9]+}}], %rd{{[0-9]+}}99; ALL: ret100 store i64 %a, ptr addrspace(3) %ptr101 ret void102}103; ALL-LABEL: st_local_i64104define void @st_local_i64(ptr addrspace(5) %ptr, i64 %a) {105; LS32: st.local.b64 [%r{{[0-9]+}}], %rd{{[0-9]+}}106; LS64: st.local.b64 [%rd{{[0-9]+}}], %rd{{[0-9]+}}107; ALL: ret108 store i64 %a, ptr addrspace(5) %ptr109 ret void110}111 112;; f32113; ALL-LABEL: st_global_f32114define void @st_global_f32(ptr addrspace(1) %ptr, float %a) {115; G32: st.global.b32 [%r{{[0-9]+}}], %r{{[0-9]+}}116; G64: st.global.b32 [%rd{{[0-9]+}}], %r{{[0-9]+}}117; ALL: ret118 store float %a, ptr addrspace(1) %ptr119 ret void120}121; ALL-LABEL: st_shared_f32122define void @st_shared_f32(ptr addrspace(3) %ptr, float %a) {123; LS32: st.shared.b32 [%r{{[0-9]+}}], %r{{[0-9]+}}124; LS64: st.shared.b32 [%rd{{[0-9]+}}], %r{{[0-9]+}}125; ALL: ret126 store float %a, ptr addrspace(3) %ptr127 ret void128}129; ALL-LABEL: st_local_f32130define void @st_local_f32(ptr addrspace(5) %ptr, float %a) {131; LS32: st.local.b32 [%r{{[0-9]+}}], %r{{[0-9]+}}132; LS64: st.local.b32 [%rd{{[0-9]+}}], %r{{[0-9]+}}133; ALL: ret134 store float %a, ptr addrspace(5) %ptr135 ret void136}137 138;; f64139; ALL-LABEL: st_global_f64140define void @st_global_f64(ptr addrspace(1) %ptr, double %a) {141; G32: st.global.b64 [%r{{[0-9]+}}], %rd{{[0-9]+}}142; G64: st.global.b64 [%rd{{[0-9]+}}], %rd{{[0-9]+}}143; ALL: ret144 store double %a, ptr addrspace(1) %ptr145 ret void146}147; ALL-LABEL: st_shared_f64148define void @st_shared_f64(ptr addrspace(3) %ptr, double %a) {149; LS32: st.shared.b64 [%r{{[0-9]+}}], %rd{{[0-9]+}}150; LS64: st.shared.b64 [%rd{{[0-9]+}}], %rd{{[0-9]+}}151; ALL: ret152 store double %a, ptr addrspace(3) %ptr153 ret void154}155; ALL-LABEL: st_local_f64156define void @st_local_f64(ptr addrspace(5) %ptr, double %a) {157; LS32: st.local.b64 [%r{{[0-9]+}}], %rd{{[0-9]+}}158; LS64: st.local.b64 [%rd{{[0-9]+}}], %rd{{[0-9]+}}159; ALL: ret160 store double %a, ptr addrspace(5) %ptr161 ret void162}163