321 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_32 | FileCheck %s3; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_32 | %ptxas-verify %}4 5 6declare i8 @llvm.nvvm.ldu.global.i.i8.p1(ptr addrspace(1) %ptr, i32 %align)7declare i16 @llvm.nvvm.ldu.global.i.i16.p1(ptr addrspace(1) %ptr, i32 %align)8declare i32 @llvm.nvvm.ldu.global.i.i32.p1(ptr addrspace(1) %ptr, i32 %align)9declare i64 @llvm.nvvm.ldu.global.i.i64.p1(ptr addrspace(1) %ptr, i32 %align)10declare ptr @llvm.nvvm.ldu.global.p.p1(ptr addrspace(1) %ptr, i32 %align)11declare float @llvm.nvvm.ldu.global.f.f32.p1(ptr addrspace(1) %ptr, i32 %align)12declare double @llvm.nvvm.ldu.global.f.f64.p1(ptr addrspace(1) %ptr, i32 %align)13declare half @llvm.nvvm.ldu.global.f.f16.p1(ptr addrspace(1) %ptr, i32 %align)14declare <2 x half> @llvm.nvvm.ldu.global.f.v2f16.p1(ptr addrspace(1) %ptr, i32 %align)15 16declare i8 @llvm.nvvm.ldg.global.i.i8.p1(ptr addrspace(1) %ptr, i32 %align)17declare i16 @llvm.nvvm.ldg.global.i.i16.p1(ptr addrspace(1) %ptr, i32 %align)18declare i32 @llvm.nvvm.ldg.global.i.i32.p1(ptr addrspace(1) %ptr, i32 %align)19declare i64 @llvm.nvvm.ldg.global.i.i64.p1(ptr addrspace(1) %ptr, i32 %align)20declare ptr @llvm.nvvm.ldg.global.p.p1(ptr addrspace(1) %ptr, i32 %align)21declare float @llvm.nvvm.ldg.global.f.f32.p1(ptr addrspace(1) %ptr, i32 %align)22declare double @llvm.nvvm.ldg.global.f.f64.p1(ptr addrspace(1) %ptr, i32 %align)23declare half @llvm.nvvm.ldg.global.f.f16.p1(ptr addrspace(1) %ptr, i32 %align)24declare <2 x half> @llvm.nvvm.ldg.global.f.v2f16.p1(ptr addrspace(1) %ptr, i32 %align)25 26define i8 @test_ldu_i8(ptr addrspace(1) %ptr) {27; CHECK-LABEL: test_ldu_i8(28; CHECK: {29; CHECK-NEXT: .reg .b16 %rs<2>;30; CHECK-NEXT: .reg .b32 %r<2>;31; CHECK-NEXT: .reg .b64 %rd<2>;32; CHECK-EMPTY:33; CHECK-NEXT: // %bb.0:34; CHECK-NEXT: ld.param.b64 %rd1, [test_ldu_i8_param_0];35; CHECK-NEXT: ldu.global.b8 %rs1, [%rd1];36; CHECK-NEXT: cvt.u32.u16 %r1, %rs1;37; CHECK-NEXT: st.param.b32 [func_retval0], %r1;38; CHECK-NEXT: ret;39 %val = tail call i8 @llvm.nvvm.ldu.global.i.i8.p1(ptr addrspace(1) %ptr, i32 4)40 ret i8 %val41}42 43define i16 @test_ldu_i16(ptr addrspace(1) %ptr) {44; CHECK-LABEL: test_ldu_i16(45; CHECK: {46; CHECK-NEXT: .reg .b16 %rs<2>;47; CHECK-NEXT: .reg .b32 %r<2>;48; CHECK-NEXT: .reg .b64 %rd<2>;49; CHECK-EMPTY:50; CHECK-NEXT: // %bb.0:51; CHECK-NEXT: ld.param.b64 %rd1, [test_ldu_i16_param_0];52; CHECK-NEXT: ldu.global.b16 %rs1, [%rd1];53; CHECK-NEXT: cvt.u32.u16 %r1, %rs1;54; CHECK-NEXT: st.param.b32 [func_retval0], %r1;55; CHECK-NEXT: ret;56 %val = tail call i16 @llvm.nvvm.ldu.global.i.i16.p1(ptr addrspace(1) %ptr, i32 2)57 ret i16 %val58}59 60define i32 @test_ldu_i32(ptr addrspace(1) %ptr) {61; CHECK-LABEL: test_ldu_i32(62; CHECK: {63; CHECK-NEXT: .reg .b32 %r<2>;64; CHECK-NEXT: .reg .b64 %rd<2>;65; CHECK-EMPTY:66; CHECK-NEXT: // %bb.0:67; CHECK-NEXT: ld.param.b64 %rd1, [test_ldu_i32_param_0];68; CHECK-NEXT: ldu.global.b32 %r1, [%rd1];69; CHECK-NEXT: st.param.b32 [func_retval0], %r1;70; CHECK-NEXT: ret;71 %val = tail call i32 @llvm.nvvm.ldu.global.i.i32.p1(ptr addrspace(1) %ptr, i32 4)72 ret i32 %val73}74 75define i64 @test_ldu_i64(ptr addrspace(1) %ptr) {76; CHECK-LABEL: test_ldu_i64(77; CHECK: {78; CHECK-NEXT: .reg .b64 %rd<3>;79; CHECK-EMPTY:80; CHECK-NEXT: // %bb.0:81; CHECK-NEXT: ld.param.b64 %rd1, [test_ldu_i64_param_0];82; CHECK-NEXT: ldu.global.b64 %rd2, [%rd1];83; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;84; CHECK-NEXT: ret;85 %val = tail call i64 @llvm.nvvm.ldu.global.i.i64.p1(ptr addrspace(1) %ptr, i32 8)86 ret i64 %val87}88 89define ptr @test_ldu_p(ptr addrspace(1) %ptr) {90; CHECK-LABEL: test_ldu_p(91; CHECK: {92; CHECK-NEXT: .reg .b64 %rd<3>;93; CHECK-EMPTY:94; CHECK-NEXT: // %bb.0:95; CHECK-NEXT: ld.param.b64 %rd1, [test_ldu_p_param_0];96; CHECK-NEXT: ldu.global.b64 %rd2, [%rd1];97; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;98; CHECK-NEXT: ret;99 %val = tail call ptr @llvm.nvvm.ldu.global.p.p1(ptr addrspace(1) %ptr, i32 8)100 ret ptr %val101}102 103define float @test_ldu_f32(ptr addrspace(1) %ptr) {104; CHECK-LABEL: test_ldu_f32(105; CHECK: {106; CHECK-NEXT: .reg .b32 %r<2>;107; CHECK-NEXT: .reg .b64 %rd<2>;108; CHECK-EMPTY:109; CHECK-NEXT: // %bb.0:110; CHECK-NEXT: ld.param.b64 %rd1, [test_ldu_f32_param_0];111; CHECK-NEXT: ldu.global.b32 %r1, [%rd1];112; CHECK-NEXT: st.param.b32 [func_retval0], %r1;113; CHECK-NEXT: ret;114 %val = tail call float @llvm.nvvm.ldu.global.f.f32.p1(ptr addrspace(1) %ptr, i32 4)115 ret float %val116}117 118define double @test_ldu_f64(ptr addrspace(1) %ptr) {119; CHECK-LABEL: test_ldu_f64(120; CHECK: {121; CHECK-NEXT: .reg .b64 %rd<3>;122; CHECK-EMPTY:123; CHECK-NEXT: // %bb.0:124; CHECK-NEXT: ld.param.b64 %rd1, [test_ldu_f64_param_0];125; CHECK-NEXT: ldu.global.b64 %rd2, [%rd1];126; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;127; CHECK-NEXT: ret;128 %val = tail call double @llvm.nvvm.ldu.global.f.f64.p1(ptr addrspace(1) %ptr, i32 8)129 ret double %val130}131 132define half @test_ldu_f16(ptr addrspace(1) %ptr) {133; CHECK-LABEL: test_ldu_f16(134; CHECK: {135; CHECK-NEXT: .reg .b16 %rs<2>;136; CHECK-NEXT: .reg .b64 %rd<2>;137; CHECK-EMPTY:138; CHECK-NEXT: // %bb.0:139; CHECK-NEXT: ld.param.b64 %rd1, [test_ldu_f16_param_0];140; CHECK-NEXT: ldu.global.b16 %rs1, [%rd1];141; CHECK-NEXT: st.param.b16 [func_retval0], %rs1;142; CHECK-NEXT: ret;143 %val = tail call half @llvm.nvvm.ldu.global.f.f16.p1(ptr addrspace(1) %ptr, i32 2)144 ret half %val145}146 147define <2 x half> @test_ldu_v2f16(ptr addrspace(1) %ptr) {148; CHECK-LABEL: test_ldu_v2f16(149; CHECK: {150; CHECK-NEXT: .reg .b32 %r<2>;151; CHECK-NEXT: .reg .b64 %rd<2>;152; CHECK-EMPTY:153; CHECK-NEXT: // %bb.0:154; CHECK-NEXT: ld.param.b64 %rd1, [test_ldu_v2f16_param_0];155; CHECK-NEXT: ldu.global.b32 %r1, [%rd1];156; CHECK-NEXT: st.param.b32 [func_retval0], %r1;157; CHECK-NEXT: ret;158 %val = tail call <2 x half> @llvm.nvvm.ldu.global.f.v2f16.p1(ptr addrspace(1) %ptr, i32 4)159 ret <2 x half> %val160}161 162define i8 @test_ldg_i8(ptr addrspace(1) %ptr) {163; CHECK-LABEL: test_ldg_i8(164; CHECK: {165; CHECK-NEXT: .reg .b32 %r<2>;166; CHECK-NEXT: .reg .b64 %rd<2>;167; CHECK-EMPTY:168; CHECK-NEXT: // %bb.0:169; CHECK-NEXT: ld.param.b64 %rd1, [test_ldg_i8_param_0];170; CHECK-NEXT: ld.global.nc.b8 %r1, [%rd1];171; CHECK-NEXT: st.param.b32 [func_retval0], %r1;172; CHECK-NEXT: ret;173 %val = tail call i8 @llvm.nvvm.ldg.global.i.i8.p1(ptr addrspace(1) %ptr, i32 4)174 ret i8 %val175}176 177define i16 @test_ldg_i16(ptr addrspace(1) %ptr) {178; CHECK-LABEL: test_ldg_i16(179; CHECK: {180; CHECK-NEXT: .reg .b32 %r<2>;181; CHECK-NEXT: .reg .b64 %rd<2>;182; CHECK-EMPTY:183; CHECK-NEXT: // %bb.0:184; CHECK-NEXT: ld.param.b64 %rd1, [test_ldg_i16_param_0];185; CHECK-NEXT: ld.global.nc.b16 %r1, [%rd1];186; CHECK-NEXT: st.param.b32 [func_retval0], %r1;187; CHECK-NEXT: ret;188 %val = tail call i16 @llvm.nvvm.ldg.global.i.i16.p1(ptr addrspace(1) %ptr, i32 2)189 ret i16 %val190}191 192define i32 @test_ldg_i32(ptr addrspace(1) %ptr) {193; CHECK-LABEL: test_ldg_i32(194; CHECK: {195; CHECK-NEXT: .reg .b32 %r<2>;196; CHECK-NEXT: .reg .b64 %rd<2>;197; CHECK-EMPTY:198; CHECK-NEXT: // %bb.0:199; CHECK-NEXT: ld.param.b64 %rd1, [test_ldg_i32_param_0];200; CHECK-NEXT: ld.global.nc.b32 %r1, [%rd1];201; CHECK-NEXT: st.param.b32 [func_retval0], %r1;202; CHECK-NEXT: ret;203 %val = tail call i32 @llvm.nvvm.ldg.global.i.i32.p1(ptr addrspace(1) %ptr, i32 4)204 ret i32 %val205}206 207define i64 @test_ldg_i64(ptr addrspace(1) %ptr) {208; CHECK-LABEL: test_ldg_i64(209; CHECK: {210; CHECK-NEXT: .reg .b64 %rd<3>;211; CHECK-EMPTY:212; CHECK-NEXT: // %bb.0:213; CHECK-NEXT: ld.param.b64 %rd1, [test_ldg_i64_param_0];214; CHECK-NEXT: ld.global.nc.b64 %rd2, [%rd1];215; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;216; CHECK-NEXT: ret;217 %val = tail call i64 @llvm.nvvm.ldg.global.i.i64.p1(ptr addrspace(1) %ptr, i32 8)218 ret i64 %val219}220 221define ptr @test_ldg_p(ptr addrspace(1) %ptr) {222; CHECK-LABEL: test_ldg_p(223; CHECK: {224; CHECK-NEXT: .reg .b64 %rd<3>;225; CHECK-EMPTY:226; CHECK-NEXT: // %bb.0:227; CHECK-NEXT: ld.param.b64 %rd1, [test_ldg_p_param_0];228; CHECK-NEXT: ld.global.nc.b64 %rd2, [%rd1];229; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;230; CHECK-NEXT: ret;231 %val = tail call ptr @llvm.nvvm.ldg.global.p.p1(ptr addrspace(1) %ptr, i32 8)232 ret ptr %val233}234 235define float @test_ldg_f32(ptr addrspace(1) %ptr) {236; CHECK-LABEL: test_ldg_f32(237; CHECK: {238; CHECK-NEXT: .reg .b32 %r<2>;239; CHECK-NEXT: .reg .b64 %rd<2>;240; CHECK-EMPTY:241; CHECK-NEXT: // %bb.0:242; CHECK-NEXT: ld.param.b64 %rd1, [test_ldg_f32_param_0];243; CHECK-NEXT: ld.global.nc.b32 %r1, [%rd1];244; CHECK-NEXT: st.param.b32 [func_retval0], %r1;245; CHECK-NEXT: ret;246 %val = tail call float @llvm.nvvm.ldg.global.f.f32.p1(ptr addrspace(1) %ptr, i32 4)247 ret float %val248}249 250define double @test_ldg_f64(ptr addrspace(1) %ptr) {251; CHECK-LABEL: test_ldg_f64(252; CHECK: {253; CHECK-NEXT: .reg .b64 %rd<3>;254; CHECK-EMPTY:255; CHECK-NEXT: // %bb.0:256; CHECK-NEXT: ld.param.b64 %rd1, [test_ldg_f64_param_0];257; CHECK-NEXT: ld.global.nc.b64 %rd2, [%rd1];258; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;259; CHECK-NEXT: ret;260 %val = tail call double @llvm.nvvm.ldg.global.f.f64.p1(ptr addrspace(1) %ptr, i32 8)261 ret double %val262}263 264define half @test_ldg_f16(ptr addrspace(1) %ptr) {265; CHECK-LABEL: test_ldg_f16(266; CHECK: {267; CHECK-NEXT: .reg .b16 %rs<2>;268; CHECK-NEXT: .reg .b64 %rd<2>;269; CHECK-EMPTY:270; CHECK-NEXT: // %bb.0:271; CHECK-NEXT: ld.param.b64 %rd1, [test_ldg_f16_param_0];272; CHECK-NEXT: ld.global.nc.b16 %rs1, [%rd1];273; CHECK-NEXT: st.param.b16 [func_retval0], %rs1;274; CHECK-NEXT: ret;275 %val = tail call half @llvm.nvvm.ldg.global.f.f16.p1(ptr addrspace(1) %ptr, i32 2)276 ret half %val277}278 279define <2 x half> @test_ldg_v2f16(ptr addrspace(1) %ptr) {280; CHECK-LABEL: test_ldg_v2f16(281; CHECK: {282; CHECK-NEXT: .reg .b32 %r<2>;283; CHECK-NEXT: .reg .b64 %rd<2>;284; CHECK-EMPTY:285; CHECK-NEXT: // %bb.0:286; CHECK-NEXT: ld.param.b64 %rd1, [test_ldg_v2f16_param_0];287; CHECK-NEXT: ld.global.nc.b32 %r1, [%rd1];288; CHECK-NEXT: st.param.b32 [func_retval0], %r1;289; CHECK-NEXT: ret;290 %val = tail call <2 x half> @llvm.nvvm.ldg.global.f.v2f16.p1(ptr addrspace(1) %ptr, i32 4)291 ret <2 x half> %val292}293 294@g = addrspace(1) global i32 0295 296define i32 @test_ldg_asi() {297; CHECK-LABEL: test_ldg_asi(298; CHECK: {299; CHECK-NEXT: .reg .b32 %r<2>;300; CHECK-EMPTY:301; CHECK-NEXT: // %bb.0:302; CHECK-NEXT: ld.global.nc.b32 %r1, [g+4];303; CHECK-NEXT: st.param.b32 [func_retval0], %r1;304; CHECK-NEXT: ret;305 %val = tail call i32 @llvm.nvvm.ldg.global.i.i32.p1(ptr addrspace(1) getelementptr (i8, ptr addrspace(1) @g, i32 4), i32 4)306 ret i32 %val307}308 309define i32 @test_lug_asi() {310; CHECK-LABEL: test_lug_asi(311; CHECK: {312; CHECK-NEXT: .reg .b32 %r<2>;313; CHECK-EMPTY:314; CHECK-NEXT: // %bb.0:315; CHECK-NEXT: ldu.global.b32 %r1, [g+4];316; CHECK-NEXT: st.param.b32 [func_retval0], %r1;317; CHECK-NEXT: ret;318 %val = tail call i32 @llvm.nvvm.ldu.global.i.i32.p1(ptr addrspace(1) getelementptr (i8, ptr addrspace(1) @g, i32 4), i32 4)319 ret i32 %val320}321