164 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 9;; i810define i8 @ld_global_i8(ptr addrspace(1) %ptr) {11; ALL-LABEL: ld_global_i812; G32: ld.global.b8 %{{.*}}, [%r{{[0-9]+}}]13; G64: ld.global.b8 %{{.*}}, [%rd{{[0-9]+}}]14; ALL: ret15 %a = load i8, ptr addrspace(1) %ptr16 ret i8 %a17}18define i8 @ld_shared_i8(ptr addrspace(3) %ptr) {19; ALL-LABEL: ld_shared_i820; LS32: ld.shared.b8 %{{.*}}, [%r{{[0-9]+}}]21; LS64: ld.shared.b8 %{{.*}}, [%rd{{[0-9]+}}]22; ALL: ret23 %a = load i8, ptr addrspace(3) %ptr24 ret i8 %a25}26define i8 @ld_local_i8(ptr addrspace(5) %ptr) {27; ALL-LABEL: ld_local_i828; LS32: ld.local.b8 %{{.*}}, [%r{{[0-9]+}}]29; LS64: ld.local.b8 %{{.*}}, [%rd{{[0-9]+}}]30; ALL: ret31 %a = load i8, ptr addrspace(5) %ptr32 ret i8 %a33}34 35;; i1636define i16 @ld_global_i16(ptr addrspace(1) %ptr) {37; ALL-LABEL: ld_global_i1638; G32: ld.global.b16 %{{.*}}, [%r{{[0-9]+}}]39; G64: ld.global.b16 %{{.*}}, [%rd{{[0-9]+}}]40; ALL: ret41 %a = load i16, ptr addrspace(1) %ptr42 ret i16 %a43}44define i16 @ld_shared_i16(ptr addrspace(3) %ptr) {45; ALL-LABEL: ld_shared_i1646; LS32: ld.shared.b16 %{{.*}}, [%r{{[0-9]+}}]47; LS64: ld.shared.b16 %{{.*}}, [%rd{{[0-9]+}}]48; ALL: ret49 %a = load i16, ptr addrspace(3) %ptr50 ret i16 %a51}52define i16 @ld_local_i16(ptr addrspace(5) %ptr) {53; ALL-LABEL: ld_local_i1654; LS32: ld.local.b16 %{{.*}}, [%r{{[0-9]+}}]55; LS64: ld.local.b16 %{{.*}}, [%rd{{[0-9]+}}]56; ALL: ret57 %a = load i16, ptr addrspace(5) %ptr58 ret i16 %a59}60 61;; i3262define i32 @ld_global_i32(ptr addrspace(1) %ptr) {63; ALL-LABEL: ld_global_i3264; G32: ld.global.b32 %{{.*}}, [%r{{[0-9]+}}]65; G64: ld.global.b32 %{{.*}}, [%rd{{[0-9]+}}]66; ALL: ret67 %a = load i32, ptr addrspace(1) %ptr68 ret i32 %a69}70define i32 @ld_shared_i32(ptr addrspace(3) %ptr) {71; ALL-LABEL: ld_shared_i3272; LS32: ld.shared.b32 %{{.*}}, [%r{{[0-9]+}}]73; LS64: ld.shared.b32 %{{.*}}, [%rd{{[0-9]+}}]74; PTX64: ret75 %a = load i32, ptr addrspace(3) %ptr76 ret i32 %a77}78define i32 @ld_local_i32(ptr addrspace(5) %ptr) {79; ALL-LABEL: ld_local_i3280; LS32: ld.local.b32 %{{.*}}, [%r{{[0-9]+}}]81; LS64: ld.local.b32 %{{.*}}, [%rd{{[0-9]+}}]82; ALL: ret83 %a = load i32, ptr addrspace(5) %ptr84 ret i32 %a85}86 87;; i6488define i64 @ld_global_i64(ptr addrspace(1) %ptr) {89; ALL-LABEL: ld_global_i6490; G32: ld.global.b64 %{{.*}}, [%r{{[0-9]+}}]91; G64: ld.global.b64 %{{.*}}, [%rd{{[0-9]+}}]92; ALL: ret93 %a = load i64, ptr addrspace(1) %ptr94 ret i64 %a95}96define i64 @ld_shared_i64(ptr addrspace(3) %ptr) {97; ALL-LABEL: ld_shared_i6498; LS32: ld.shared.b64 %{{.*}}, [%r{{[0-9]+}}]99; LS64: ld.shared.b64 %{{.*}}, [%rd{{[0-9]+}}]100; ALL: ret101 %a = load i64, ptr addrspace(3) %ptr102 ret i64 %a103}104define i64 @ld_local_i64(ptr addrspace(5) %ptr) {105; ALL-LABEL: ld_local_i64106; LS32: ld.local.b64 %{{.*}}, [%r{{[0-9]+}}]107; LS64: ld.local.b64 %{{.*}}, [%rd{{[0-9]+}}]108; ALL: ret109 %a = load i64, ptr addrspace(5) %ptr110 ret i64 %a111}112 113;; f32114define float @ld_global_f32(ptr addrspace(1) %ptr) {115; ALL-LABEL: ld_global_f32116; G32: ld.global.b32 %{{.*}}, [%r{{[0-9]+}}]117; G64: ld.global.b32 %{{.*}}, [%rd{{[0-9]+}}]118; ALL: ret119 %a = load float, ptr addrspace(1) %ptr120 ret float %a121}122define float @ld_shared_f32(ptr addrspace(3) %ptr) {123; ALL-LABEL: ld_shared_f32124; LS32: ld.shared.b32 %{{.*}}, [%r{{[0-9]+}}]125; LS64: ld.shared.b32 %{{.*}}, [%rd{{[0-9]+}}]126; ALL: ret127 %a = load float, ptr addrspace(3) %ptr128 ret float %a129}130define float @ld_local_f32(ptr addrspace(5) %ptr) {131; ALL-LABEL: ld_local_f32132; LS32: ld.local.b32 %{{.*}}, [%r{{[0-9]+}}]133; LS64: ld.local.b32 %{{.*}}, [%rd{{[0-9]+}}]134; ALL: ret135 %a = load float, ptr addrspace(5) %ptr136 ret float %a137}138 139;; f64140define double @ld_global_f64(ptr addrspace(1) %ptr) {141; ALL-LABEL: ld_global_f64142; G32: ld.global.b64 %{{.*}}, [%r{{[0-9]+}}]143; G64: ld.global.b64 %{{.*}}, [%rd{{[0-9]+}}]144; ALL: ret145 %a = load double, ptr addrspace(1) %ptr146 ret double %a147}148define double @ld_shared_f64(ptr addrspace(3) %ptr) {149; ALL-LABEL: ld_shared_f64150; LS32: ld.shared.b64 %{{.*}}, [%r{{[0-9]+}}]151; LS64: ld.shared.b64 %{{.*}}, [%rd{{[0-9]+}}]152; ALL: ret153 %a = load double, ptr addrspace(3) %ptr154 ret double %a155}156define double @ld_local_f64(ptr addrspace(5) %ptr) {157; ALL-LABEL: ld_local_f64158; LS32: ld.local.b64 %{{.*}}, [%r{{[0-9]+}}]159; LS64: ld.local.b64 %{{.*}}, [%rd{{[0-9]+}}]160; ALL: ret161 %a = load double, ptr addrspace(5) %ptr162 ret double %a163}164