117 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_30 -fp-contract=fast | FileCheck %s --check-prefixes=CHECK,FAST3; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_30 | FileCheck %s --check-prefixes=CHECK,DEFAULT4; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_30 -fp-contract=fast | %ptxas-verify %}5; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_30 | %ptxas-verify %}6 7target triple = "nvptx64-unknown-cuda"8 9;; Make sure we are generating proper instruction sequences for fused ops10;; If fusion is allowed, we try to form fma.rn at the PTX level, and emit11;; add.f32 otherwise. Without an explicit rounding mode on add.f32, ptxas12;; is free to fuse with a multiply if it is able. If fusion is not allowed,13;; we do not form fma.rn at the PTX level and explicitly generate add.rn14;; for all adds to prevent ptxas from fusion the ops.15define float @t0(float %a, float %b, float %c) {16; FAST-LABEL: t0(17; FAST: {18; FAST-NEXT: .reg .b32 %r<5>;19; FAST-EMPTY:20; FAST-NEXT: // %bb.0:21; FAST-NEXT: ld.param.b32 %r1, [t0_param_0];22; FAST-NEXT: ld.param.b32 %r2, [t0_param_1];23; FAST-NEXT: ld.param.b32 %r3, [t0_param_2];24; FAST-NEXT: fma.rn.f32 %r4, %r1, %r2, %r3;25; FAST-NEXT: st.param.b32 [func_retval0], %r4;26; FAST-NEXT: ret;27;28; DEFAULT-LABEL: t0(29; DEFAULT: {30; DEFAULT-NEXT: .reg .b32 %r<6>;31; DEFAULT-EMPTY:32; DEFAULT-NEXT: // %bb.0:33; DEFAULT-NEXT: ld.param.b32 %r1, [t0_param_0];34; DEFAULT-NEXT: ld.param.b32 %r2, [t0_param_1];35; DEFAULT-NEXT: mul.rn.f32 %r3, %r1, %r2;36; DEFAULT-NEXT: ld.param.b32 %r4, [t0_param_2];37; DEFAULT-NEXT: add.rn.f32 %r5, %r3, %r4;38; DEFAULT-NEXT: st.param.b32 [func_retval0], %r5;39; DEFAULT-NEXT: ret;40 %v0 = fmul float %a, %b41 %v1 = fadd float %v0, %c42 ret float %v143}44 45;; We cannot form an fma here, but make sure we explicitly emit add.rn.f3246;; to prevent ptxas from fusing this with anything else.47define float @t1(float %a, float %b) {48; FAST-LABEL: t1(49; FAST: {50; FAST-NEXT: .reg .b32 %r<6>;51; FAST-EMPTY:52; FAST-NEXT: // %bb.0:53; FAST-NEXT: ld.param.b32 %r1, [t1_param_0];54; FAST-NEXT: ld.param.b32 %r2, [t1_param_1];55; FAST-NEXT: add.f32 %r3, %r1, %r2;56; FAST-NEXT: sub.f32 %r4, %r1, %r2;57; FAST-NEXT: mul.f32 %r5, %r3, %r4;58; FAST-NEXT: st.param.b32 [func_retval0], %r5;59; FAST-NEXT: ret;60;61; DEFAULT-LABEL: t1(62; DEFAULT: {63; DEFAULT-NEXT: .reg .b32 %r<6>;64; DEFAULT-EMPTY:65; DEFAULT-NEXT: // %bb.0:66; DEFAULT-NEXT: ld.param.b32 %r1, [t1_param_0];67; DEFAULT-NEXT: ld.param.b32 %r2, [t1_param_1];68; DEFAULT-NEXT: add.rn.f32 %r3, %r1, %r2;69; DEFAULT-NEXT: sub.rn.f32 %r4, %r1, %r2;70; DEFAULT-NEXT: mul.rn.f32 %r5, %r3, %r4;71; DEFAULT-NEXT: st.param.b32 [func_retval0], %r5;72; DEFAULT-NEXT: ret;73 %v1 = fadd float %a, %b74 %v2 = fsub float %a, %b75 %v3 = fmul float %v1, %v276 ret float %v377}78 79;; Make sure we generate the non ".rn" version when the "contract" flag is80;; present on the instructions81define float @t2(float %a, float %b) {82; CHECK-LABEL: t2(83; CHECK: {84; CHECK-NEXT: .reg .b32 %r<6>;85; CHECK-EMPTY:86; CHECK-NEXT: // %bb.0:87; CHECK-NEXT: ld.param.b32 %r1, [t2_param_0];88; CHECK-NEXT: ld.param.b32 %r2, [t2_param_1];89; CHECK-NEXT: add.f32 %r3, %r1, %r2;90; CHECK-NEXT: sub.f32 %r4, %r1, %r2;91; CHECK-NEXT: mul.f32 %r5, %r3, %r4;92; CHECK-NEXT: st.param.b32 [func_retval0], %r5;93; CHECK-NEXT: ret;94 %v1 = fadd contract float %a, %b95 %v2 = fsub contract float %a, %b96 %v3 = fmul contract float %v1, %v297 ret float %v398}99 100;; Make sure we always fold to fma when the "contract" flag is present101define float @t3(float %a, float %b, float %c) {102; CHECK-LABEL: t3(103; CHECK: {104; CHECK-NEXT: .reg .b32 %r<5>;105; CHECK-EMPTY:106; CHECK-NEXT: // %bb.0:107; CHECK-NEXT: ld.param.b32 %r1, [t3_param_0];108; CHECK-NEXT: ld.param.b32 %r2, [t3_param_1];109; CHECK-NEXT: ld.param.b32 %r3, [t3_param_2];110; CHECK-NEXT: fma.rn.f32 %r4, %r1, %r2, %r3;111; CHECK-NEXT: st.param.b32 [func_retval0], %r4;112; CHECK-NEXT: ret;113 %v0 = fmul contract float %a, %b114 %v1 = fadd contract float %v0, %c115 ret float %v1116}117