163 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -mtriple=nvptx -mcpu=sm_20 | FileCheck %s --check-prefix=PTX323; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 | FileCheck %s --check-prefix=PTX644; 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 7 8;; i89define i8 @ld_global_i8(ptr addrspace(0) %ptr) {10; PTX32-LABEL: ld_global_i8(11; PTX32: {12; PTX32-NEXT: .reg .b32 %r<3>;13; PTX32-EMPTY:14; PTX32-NEXT: // %bb.0:15; PTX32-NEXT: ld.param.b32 %r1, [ld_global_i8_param_0];16; PTX32-NEXT: ld.b8 %r2, [%r1];17; PTX32-NEXT: st.param.b32 [func_retval0], %r2;18; PTX32-NEXT: ret;19;20; PTX64-LABEL: ld_global_i8(21; PTX64: {22; PTX64-NEXT: .reg .b32 %r<2>;23; PTX64-NEXT: .reg .b64 %rd<2>;24; PTX64-EMPTY:25; PTX64-NEXT: // %bb.0:26; PTX64-NEXT: ld.param.b64 %rd1, [ld_global_i8_param_0];27; PTX64-NEXT: ld.b8 %r1, [%rd1];28; PTX64-NEXT: st.param.b32 [func_retval0], %r1;29; PTX64-NEXT: ret;30 %a = load i8, ptr addrspace(0) %ptr31 ret i8 %a32}33 34;; i1635define i16 @ld_global_i16(ptr addrspace(0) %ptr) {36; PTX32-LABEL: ld_global_i16(37; PTX32: {38; PTX32-NEXT: .reg .b32 %r<3>;39; PTX32-EMPTY:40; PTX32-NEXT: // %bb.0:41; PTX32-NEXT: ld.param.b32 %r1, [ld_global_i16_param_0];42; PTX32-NEXT: ld.b16 %r2, [%r1];43; PTX32-NEXT: st.param.b32 [func_retval0], %r2;44; PTX32-NEXT: ret;45;46; PTX64-LABEL: ld_global_i16(47; PTX64: {48; PTX64-NEXT: .reg .b32 %r<2>;49; PTX64-NEXT: .reg .b64 %rd<2>;50; PTX64-EMPTY:51; PTX64-NEXT: // %bb.0:52; PTX64-NEXT: ld.param.b64 %rd1, [ld_global_i16_param_0];53; PTX64-NEXT: ld.b16 %r1, [%rd1];54; PTX64-NEXT: st.param.b32 [func_retval0], %r1;55; PTX64-NEXT: ret;56 %a = load i16, ptr addrspace(0) %ptr57 ret i16 %a58}59 60;; i3261define i32 @ld_global_i32(ptr addrspace(0) %ptr) {62; PTX32-LABEL: ld_global_i32(63; PTX32: {64; PTX32-NEXT: .reg .b32 %r<3>;65; PTX32-EMPTY:66; PTX32-NEXT: // %bb.0:67; PTX32-NEXT: ld.param.b32 %r1, [ld_global_i32_param_0];68; PTX32-NEXT: ld.b32 %r2, [%r1];69; PTX32-NEXT: st.param.b32 [func_retval0], %r2;70; PTX32-NEXT: ret;71;72; PTX64-LABEL: ld_global_i32(73; PTX64: {74; PTX64-NEXT: .reg .b32 %r<2>;75; PTX64-NEXT: .reg .b64 %rd<2>;76; PTX64-EMPTY:77; PTX64-NEXT: // %bb.0:78; PTX64-NEXT: ld.param.b64 %rd1, [ld_global_i32_param_0];79; PTX64-NEXT: ld.b32 %r1, [%rd1];80; PTX64-NEXT: st.param.b32 [func_retval0], %r1;81; PTX64-NEXT: ret;82 %a = load i32, ptr addrspace(0) %ptr83 ret i32 %a84}85 86;; i6487define i64 @ld_global_i64(ptr addrspace(0) %ptr) {88; PTX32-LABEL: ld_global_i64(89; PTX32: {90; PTX32-NEXT: .reg .b32 %r<2>;91; PTX32-NEXT: .reg .b64 %rd<2>;92; PTX32-EMPTY:93; PTX32-NEXT: // %bb.0:94; PTX32-NEXT: ld.param.b32 %r1, [ld_global_i64_param_0];95; PTX32-NEXT: ld.b64 %rd1, [%r1];96; PTX32-NEXT: st.param.b64 [func_retval0], %rd1;97; PTX32-NEXT: ret;98;99; PTX64-LABEL: ld_global_i64(100; PTX64: {101; PTX64-NEXT: .reg .b64 %rd<3>;102; PTX64-EMPTY:103; PTX64-NEXT: // %bb.0:104; PTX64-NEXT: ld.param.b64 %rd1, [ld_global_i64_param_0];105; PTX64-NEXT: ld.b64 %rd2, [%rd1];106; PTX64-NEXT: st.param.b64 [func_retval0], %rd2;107; PTX64-NEXT: ret;108 %a = load i64, ptr addrspace(0) %ptr109 ret i64 %a110}111 112;; f32113define float @ld_global_f32(ptr addrspace(0) %ptr) {114; PTX32-LABEL: ld_global_f32(115; PTX32: {116; PTX32-NEXT: .reg .b32 %r<3>;117; PTX32-EMPTY:118; PTX32-NEXT: // %bb.0:119; PTX32-NEXT: ld.param.b32 %r1, [ld_global_f32_param_0];120; PTX32-NEXT: ld.b32 %r2, [%r1];121; PTX32-NEXT: st.param.b32 [func_retval0], %r2;122; PTX32-NEXT: ret;123;124; PTX64-LABEL: ld_global_f32(125; PTX64: {126; PTX64-NEXT: .reg .b32 %r<2>;127; PTX64-NEXT: .reg .b64 %rd<2>;128; PTX64-EMPTY:129; PTX64-NEXT: // %bb.0:130; PTX64-NEXT: ld.param.b64 %rd1, [ld_global_f32_param_0];131; PTX64-NEXT: ld.b32 %r1, [%rd1];132; PTX64-NEXT: st.param.b32 [func_retval0], %r1;133; PTX64-NEXT: ret;134 %a = load float, ptr addrspace(0) %ptr135 ret float %a136}137 138;; f64139define double @ld_global_f64(ptr addrspace(0) %ptr) {140; PTX32-LABEL: ld_global_f64(141; PTX32: {142; PTX32-NEXT: .reg .b32 %r<2>;143; PTX32-NEXT: .reg .b64 %rd<2>;144; PTX32-EMPTY:145; PTX32-NEXT: // %bb.0:146; PTX32-NEXT: ld.param.b32 %r1, [ld_global_f64_param_0];147; PTX32-NEXT: ld.b64 %rd1, [%r1];148; PTX32-NEXT: st.param.b64 [func_retval0], %rd1;149; PTX32-NEXT: ret;150;151; PTX64-LABEL: ld_global_f64(152; PTX64: {153; PTX64-NEXT: .reg .b64 %rd<3>;154; PTX64-EMPTY:155; PTX64-NEXT: // %bb.0:156; PTX64-NEXT: ld.param.b64 %rd1, [ld_global_f64_param_0];157; PTX64-NEXT: ld.b64 %rd2, [%rd1];158; PTX64-NEXT: st.param.b64 [func_retval0], %rd2;159; PTX64-NEXT: ret;160 %a = load double, ptr addrspace(0) %ptr161 ret double %a162}163