261 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 62; RUN: llc -march=nvptx64 -mcpu=sm_100f -o - %s | FileCheck %s3 4target triple = "nvptx64-nvidia-cuda"5target datalayout = "e-i64:64-i128:128-v16:16-v32:32-n16:32:64"6 7define void @test_fmaxnum3(float %a, float %b, float %c, ptr addrspace(1) %output) {8; CHECK-LABEL: test_fmaxnum3(9; CHECK: {10; CHECK-NEXT: .reg .b32 %r<5>;11; CHECK-NEXT: .reg .b64 %rd<2>;12; CHECK-EMPTY:13; CHECK-NEXT: // %bb.0: // %entry14; CHECK-NEXT: ld.param.b32 %r1, [test_fmaxnum3_param_0];15; CHECK-NEXT: ld.param.b32 %r2, [test_fmaxnum3_param_1];16; CHECK-NEXT: ld.param.b32 %r3, [test_fmaxnum3_param_2];17; CHECK-NEXT: max.f32 %r4, %r1, %r2, %r3;18; CHECK-NEXT: ld.param.b64 %rd1, [test_fmaxnum3_param_3];19; CHECK-NEXT: st.global.b32 [%rd1], %r4;20; CHECK-NEXT: ret;21entry:22 %max_ab = call float @llvm.maxnum.f32(float %a, float %b)23 %max_abc = call float @llvm.maxnum.f32(float %max_ab, float %c)24 store float %max_abc, ptr addrspace(1) %output, align 425 ret void26}27 28define void @test_fminnum3(float %a, float %b, float %c, ptr addrspace(1) %output) {29; CHECK-LABEL: test_fminnum3(30; CHECK: {31; CHECK-NEXT: .reg .b32 %r<5>;32; CHECK-NEXT: .reg .b64 %rd<2>;33; CHECK-EMPTY:34; CHECK-NEXT: // %bb.0: // %entry35; CHECK-NEXT: ld.param.b32 %r1, [test_fminnum3_param_0];36; CHECK-NEXT: ld.param.b32 %r2, [test_fminnum3_param_1];37; CHECK-NEXT: ld.param.b32 %r3, [test_fminnum3_param_2];38; CHECK-NEXT: min.f32 %r4, %r1, %r2, %r3;39; CHECK-NEXT: ld.param.b64 %rd1, [test_fminnum3_param_3];40; CHECK-NEXT: st.global.b32 [%rd1], %r4;41; CHECK-NEXT: ret;42entry:43 %min_ab = call float @llvm.minnum.f32(float %a, float %b)44 %min_abc = call float @llvm.minnum.f32(float %min_ab, float %c)45 store float %min_abc, ptr addrspace(1) %output, align 446 ret void47}48 49define void @test_fmaximum3(float %a, float %b, float %c, ptr addrspace(1) %output) {50; CHECK-LABEL: test_fmaximum3(51; CHECK: {52; CHECK-NEXT: .reg .b32 %r<5>;53; CHECK-NEXT: .reg .b64 %rd<2>;54; CHECK-EMPTY:55; CHECK-NEXT: // %bb.0: // %entry56; CHECK-NEXT: ld.param.b32 %r1, [test_fmaximum3_param_0];57; CHECK-NEXT: ld.param.b32 %r2, [test_fmaximum3_param_1];58; CHECK-NEXT: ld.param.b32 %r3, [test_fmaximum3_param_2];59; CHECK-NEXT: max.NaN.f32 %r4, %r1, %r2, %r3;60; CHECK-NEXT: ld.param.b64 %rd1, [test_fmaximum3_param_3];61; CHECK-NEXT: st.global.b32 [%rd1], %r4;62; CHECK-NEXT: ret;63entry:64 %max_ab = call float @llvm.maximum.f32(float %a, float %b)65 %max_abc = call float @llvm.maximum.f32(float %max_ab, float %c)66 store float %max_abc, ptr addrspace(1) %output, align 467 ret void68}69 70define void @test_fminimum3(float %a, float %b, float %c, ptr addrspace(1) %output) {71; CHECK-LABEL: test_fminimum3(72; CHECK: {73; CHECK-NEXT: .reg .b32 %r<5>;74; CHECK-NEXT: .reg .b64 %rd<2>;75; CHECK-EMPTY:76; CHECK-NEXT: // %bb.0: // %entry77; CHECK-NEXT: ld.param.b32 %r1, [test_fminimum3_param_0];78; CHECK-NEXT: ld.param.b32 %r2, [test_fminimum3_param_1];79; CHECK-NEXT: ld.param.b32 %r3, [test_fminimum3_param_2];80; CHECK-NEXT: min.NaN.f32 %r4, %r1, %r2, %r3;81; CHECK-NEXT: ld.param.b64 %rd1, [test_fminimum3_param_3];82; CHECK-NEXT: st.global.b32 [%rd1], %r4;83; CHECK-NEXT: ret;84entry:85 %min_ab = call float @llvm.minimum.f32(float %a, float %b)86 %min_abc = call float @llvm.minimum.f32(float %min_ab, float %c)87 store float %min_abc, ptr addrspace(1) %output, align 488 ret void89}90 91define void @test_fmaximumnum3(float %a, float %b, float %c, ptr addrspace(1) %output) {92; CHECK-LABEL: test_fmaximumnum3(93; CHECK: {94; CHECK-NEXT: .reg .b32 %r<5>;95; CHECK-NEXT: .reg .b64 %rd<2>;96; CHECK-EMPTY:97; CHECK-NEXT: // %bb.0: // %entry98; CHECK-NEXT: ld.param.b32 %r1, [test_fmaximumnum3_param_0];99; CHECK-NEXT: ld.param.b32 %r2, [test_fmaximumnum3_param_1];100; CHECK-NEXT: ld.param.b32 %r3, [test_fmaximumnum3_param_2];101; CHECK-NEXT: max.f32 %r4, %r1, %r2, %r3;102; CHECK-NEXT: ld.param.b64 %rd1, [test_fmaximumnum3_param_3];103; CHECK-NEXT: st.global.b32 [%rd1], %r4;104; CHECK-NEXT: ret;105entry:106 %max_ab = call float @llvm.maximumnum.f32(float %a, float %b)107 %max_abc = call float @llvm.maximumnum.f32(float %max_ab, float %c)108 store float %max_abc, ptr addrspace(1) %output, align 4109 ret void110}111 112define void @test_fminimumnum3(float %a, float %b, float %c, ptr addrspace(1) %output) {113; CHECK-LABEL: test_fminimumnum3(114; CHECK: {115; CHECK-NEXT: .reg .b32 %r<5>;116; CHECK-NEXT: .reg .b64 %rd<2>;117; CHECK-EMPTY:118; CHECK-NEXT: // %bb.0: // %entry119; CHECK-NEXT: ld.param.b32 %r1, [test_fminimumnum3_param_0];120; CHECK-NEXT: ld.param.b32 %r2, [test_fminimumnum3_param_1];121; CHECK-NEXT: ld.param.b32 %r3, [test_fminimumnum3_param_2];122; CHECK-NEXT: min.f32 %r4, %r1, %r2, %r3;123; CHECK-NEXT: ld.param.b64 %rd1, [test_fminimumnum3_param_3];124; CHECK-NEXT: st.global.b32 [%rd1], %r4;125; CHECK-NEXT: ret;126entry:127 %min_ab = call float @llvm.minimumnum.f32(float %a, float %b)128 %min_abc = call float @llvm.minimumnum.f32(float %min_ab, float %c)129 store float %min_abc, ptr addrspace(1) %output, align 4130 ret void131}132 133; Test commuted operands (second operand is the nested operation)134define void @test_fmaxnum3_commuted(float %a, float %b, float %c, ptr addrspace(1) %output) {135; CHECK-LABEL: test_fmaxnum3_commuted(136; CHECK: {137; CHECK-NEXT: .reg .b32 %r<5>;138; CHECK-NEXT: .reg .b64 %rd<2>;139; CHECK-EMPTY:140; CHECK-NEXT: // %bb.0: // %entry141; CHECK-NEXT: ld.param.b32 %r1, [test_fmaxnum3_commuted_param_0];142; CHECK-NEXT: ld.param.b32 %r2, [test_fmaxnum3_commuted_param_1];143; CHECK-NEXT: ld.param.b32 %r3, [test_fmaxnum3_commuted_param_2];144; CHECK-NEXT: max.f32 %r4, %r1, %r2, %r3;145; CHECK-NEXT: ld.param.b64 %rd1, [test_fmaxnum3_commuted_param_3];146; CHECK-NEXT: st.global.b32 [%rd1], %r4;147; CHECK-NEXT: ret;148entry:149 %max_bc = call float @llvm.maxnum.f32(float %b, float %c)150 %max_abc = call float @llvm.maxnum.f32(float %a, float %max_bc)151 store float %max_abc, ptr addrspace(1) %output, align 4152 ret void153}154 155; NEGATIVE TEST: Mixed min/max operations should not combine156define void @test_mixed_minmax_no_combine(float %a, float %b, float %c, ptr addrspace(1) %output) {157; CHECK-LABEL: test_mixed_minmax_no_combine(158; CHECK: {159; CHECK-NEXT: .reg .b32 %r<6>;160; CHECK-NEXT: .reg .b64 %rd<2>;161; CHECK-EMPTY:162; CHECK-NEXT: // %bb.0: // %entry163; CHECK-NEXT: ld.param.b32 %r1, [test_mixed_minmax_no_combine_param_0];164; CHECK-NEXT: ld.param.b32 %r2, [test_mixed_minmax_no_combine_param_1];165; CHECK-NEXT: min.f32 %r3, %r1, %r2;166; CHECK-NEXT: ld.param.b32 %r4, [test_mixed_minmax_no_combine_param_2];167; CHECK-NEXT: max.f32 %r5, %r3, %r4;168; CHECK-NEXT: ld.param.b64 %rd1, [test_mixed_minmax_no_combine_param_3];169; CHECK-NEXT: st.global.b32 [%rd1], %r5;170; CHECK-NEXT: ret;171entry:172 %min_ab = call float @llvm.minnum.f32(float %a, float %b)173 %max_result = call float @llvm.maxnum.f32(float %min_ab, float %c)174 store float %max_result, ptr addrspace(1) %output, align 4175 ret void176}177 178; NEGATIVE TEST: Mixed maxnum/maximum operations should not combine179define void @test_mixed_maxnum_maximum_no_combine(float %a, float %b, float %c, ptr addrspace(1) %output) {180; CHECK-LABEL: test_mixed_maxnum_maximum_no_combine(181; CHECK: {182; CHECK-NEXT: .reg .b32 %r<6>;183; CHECK-NEXT: .reg .b64 %rd<2>;184; CHECK-EMPTY:185; CHECK-NEXT: // %bb.0: // %entry186; CHECK-NEXT: ld.param.b32 %r1, [test_mixed_maxnum_maximum_no_combine_param_0];187; CHECK-NEXT: ld.param.b32 %r2, [test_mixed_maxnum_maximum_no_combine_param_1];188; CHECK-NEXT: max.f32 %r3, %r1, %r2;189; CHECK-NEXT: ld.param.b32 %r4, [test_mixed_maxnum_maximum_no_combine_param_2];190; CHECK-NEXT: max.NaN.f32 %r5, %r3, %r4;191; CHECK-NEXT: ld.param.b64 %rd1, [test_mixed_maxnum_maximum_no_combine_param_3];192; CHECK-NEXT: st.global.b32 [%rd1], %r5;193; CHECK-NEXT: ret;194entry:195 %maxnum_ab = call float @llvm.maxnum.f32(float %a, float %b)196 %maximum_result = call float @llvm.maximum.f32(float %maxnum_ab, float %c)197 store float %maximum_result, ptr addrspace(1) %output, align 4198 ret void199}200 201; NEGATIVE TEST: f16 should not be combined (only f32 supported)202define void @test_f16_no_combine(half %a, half %b, half %c, ptr addrspace(1) %output) {203; CHECK-LABEL: test_f16_no_combine(204; CHECK: {205; CHECK-NEXT: .reg .b16 %rs<6>;206; CHECK-NEXT: .reg .b64 %rd<2>;207; CHECK-EMPTY:208; CHECK-NEXT: // %bb.0: // %entry209; CHECK-NEXT: ld.param.b16 %rs1, [test_f16_no_combine_param_0];210; CHECK-NEXT: ld.param.b16 %rs2, [test_f16_no_combine_param_1];211; CHECK-NEXT: max.f16 %rs3, %rs1, %rs2;212; CHECK-NEXT: ld.param.b16 %rs4, [test_f16_no_combine_param_2];213; CHECK-NEXT: max.f16 %rs5, %rs3, %rs4;214; CHECK-NEXT: ld.param.b64 %rd1, [test_f16_no_combine_param_3];215; CHECK-NEXT: st.global.b16 [%rd1], %rs5;216; CHECK-NEXT: ret;217entry:218 %max_ab = call half @llvm.maxnum.f16(half %a, half %b)219 %max_abc = call half @llvm.maxnum.f16(half %max_ab, half %c)220 store half %max_abc, ptr addrspace(1) %output, align 2221 ret void222}223 224; NEGATIVE TEST: Multiple uses of intermediate result should not combine225define void @test_multiple_uses_no_combine(float %a, float %b, float %c, ptr addrspace(1) %output1, ptr addrspace(1) %output2) {226; CHECK-LABEL: test_multiple_uses_no_combine(227; CHECK: {228; CHECK-NEXT: .reg .b32 %r<6>;229; CHECK-NEXT: .reg .b64 %rd<3>;230; CHECK-EMPTY:231; CHECK-NEXT: // %bb.0: // %entry232; CHECK-NEXT: ld.param.b32 %r1, [test_multiple_uses_no_combine_param_0];233; CHECK-NEXT: ld.param.b32 %r2, [test_multiple_uses_no_combine_param_1];234; CHECK-NEXT: max.f32 %r3, %r1, %r2;235; CHECK-NEXT: ld.param.b32 %r4, [test_multiple_uses_no_combine_param_2];236; CHECK-NEXT: max.f32 %r5, %r3, %r4;237; CHECK-NEXT: ld.param.b64 %rd1, [test_multiple_uses_no_combine_param_3];238; CHECK-NEXT: st.global.b32 [%rd1], %r3;239; CHECK-NEXT: ld.param.b64 %rd2, [test_multiple_uses_no_combine_param_4];240; CHECK-NEXT: st.global.b32 [%rd2], %r5;241; CHECK-NEXT: ret;242entry:243 %max_ab = call float @llvm.maxnum.f32(float %a, float %b)244 %max_abc = call float @llvm.maxnum.f32(float %max_ab, float %c)245 ; Multiple uses of %max_ab should prevent combining246 store float %max_ab, ptr addrspace(1) %output1, align 4247 store float %max_abc, ptr addrspace(1) %output2, align 4248 ret void249}250 251; Declare all the intrinsics we need252declare float @llvm.maxnum.f32(float, float) #0253declare float @llvm.minnum.f32(float, float) #0254declare float @llvm.maximum.f32(float, float) #0255declare float @llvm.minimum.f32(float, float) #0256declare float @llvm.maximumnum.f32(float, float) #0257declare float @llvm.minimumnum.f32(float, float) #0258declare half @llvm.maxnum.f16(half, half) #0259 260attributes #0 = { nounwind readnone speculatable willreturn }261