[HIP] Fix regressions due to fp contract change

[yaxunl]

Summary

Recently HIP toolchain made a change to use clang instead of opt/llc to do compilation
(https://reviews.llvm.org/D81861). The intention is to make HIP toolchain canonical like
other toolchains.

However, this change introduced an unintentional change regarding backend fp fuse
option, which caused regressions in some HIP applications.

Basically before the change, HIP toolchain used clang to generate bitcode, then use
opt/llc to optimize bitcode and generate ISA. As such, the amdgpu backend takes
the default fp fuse mode which is 'Standard'. This mode respect contract flag of
fmul/fadd instructions and do not fuse fmul/fadd instructions without contract flag.

However, after the change, HIP toolchain now use clang to generate IR, do optimization,
and generate ISA as one process. Now amdgpu backend fp fuse option is determined
by -ffp-contract option, which is 'fast' by default. And this -ffp-contract=fast language option
is translated to 'Fast' fp fuse option in backend. Suddenly backend starts to fuse fmul/fadd
instructions without contract flag.

This causes wrong result for some device library functions, e.g. tan(-1e20), which should
return 0.8446, now returns -0.933. What is worse is that since backend with 'Fast' fp fuse
option does not respect contract flag, there is no way to use #pragma clang fp contract
directive to enforce fp contract requirements.

This patch fixes the regression by introducing a new value 'fast-honor-pragmas' for -ffp-contract
and use it for HIP by default. 'fast-honor-pragmas' is equivalent to 'fast' in frontend but
let the backend to use 'Standard' fp fuse option. 'fast-honor-pragmas' is useful since 'Fast'
fp fuse option in backend does not honor contract flag, it is of little use to HIP
applications since all code with #pragma STDC FP_CONTRACT or any IR from a
source compiled with -ffp-contract=on is broken.

[rjmccall]

I have objections to the code change here. I'll leave the conceptual question to other people interested in the HIP toolchain.

[yaxunl]

In D90174#2354249, @rjmccall wrote:

I have objections to the code change here. I'll leave the conceptual question to other people interested in the HIP toolchain.

Is it OK to introduce a clang codegen option e.g. -fp-contract-backend=x to control the backend fp fuse option? By default it matches -ffp-contract language option, but allows being overridden by explicit option. Then HIP toolchain can use it to override the backend fp fuse option.

[tra]

clang/lib/CodeGen/BackendUtil.cpp

486

I don't think it's a good idea to force this.

Perhaps a better way to address this would be to set HIP-specific default to Standard where CUDA does it:
https://github.com/llvm/llvm-project/blob/master/clang/lib/Frontend/CompilerInvocation.cpp#L2415

Currently HIP inherits this setting from CUDA.

[yaxunl]

clang/lib/CodeGen/BackendUtil.cpp

486

We want to keep -ffp-contract=fast for frontend so that we can continue emitting fmul/fadd insts with contract flag in IR for HIP programs. We only want to change the backend fp fuse option. Currently there is no separate clang option to set backend fp fuse option.

[rjmccall]

Argh, sorry! I meant to say "I have *no* objections".

[tra]

clang/lib/CodeGen/BackendUtil.cpp

486

I do not see any references to AllowFPOpFusion anywhere under clang/ other than in this function.
Perhaps I'm missing something. How/where does it make a difference in the front-end other than setting the option for the back-end?

[yaxunl]

clang/lib/CodeGen/BackendUtil.cpp

486

-ffp-contract not only sets backend fp fuse option, but also setFPContractMode

https://github.com/llvm/llvm-project/blob/2e204e23911b1f8bd1463535da40c6e48747a138/clang/include/clang/Basic/LangOptions.h#L413

which enables allowFPContractAcrossStatement

https://github.com/llvm/llvm-project/blob/2e204e23911b1f8bd1463535da40c6e48747a138/clang/include/clang/Basic/LangOptions.h#L439

which sets llvm::FastMathFlags

https://github.com/llvm/llvm-project/blob/d3205bbca3e0002d76282878986993e7e7994779/clang/lib/CodeGen/CodeGenFunction.cpp#L129

which causes fmul/fadd to have contract flag

[tra]

clang/lib/CodeGen/BackendUtil.cpp

486

Thank you. I see.
I'm still uncomfortable with growing a target-specific quirk in FP behavior that can't be overridden from command line, but considering that it's limted to HIP only, it may be OK short-term.

I think similar issue (fast-math does not allow changing contraction mode) was discussed last year:
http://lists.llvm.org/pipermail/llvm-dev/2019-July/133787.html

One thing I know about FP is that there are more nuances than I'm aware of. It may be worth asking with more experience with this.

@scanon -- do we have any better options to honor contraction pragmas with implicitly-enabled fast-math?

[rjmccall]

clang/lib/CodeGen/BackendUtil.cpp

486

In the abstract, it seems reasonable for pragmas to override a global fast-math setting.

[yaxunl]

clang/lib/CodeGen/BackendUtil.cpp

486

I think probably we need to introduce a new value for -ffp-contract= option faststd, which is fast for FE and Standard for BE. In this case, fp add/mult by default can be fused, unless disabled by pragma, and BE respect the restrictions imposed by pragmas in FE. I think this mode is more useful than the original fast mode.

[tra]

clang/lib/CodeGen/BackendUtil.cpp

486

This should work, I think.

[yaxunl]

introduce faststd as value for -ffp-contract and use it for HIP by default.

[tra]

LGTM, but I'll defer to @rjmccall for the approval.

[rjmccall]

I agree this is useful. However, you need to update the manual to cover faststd.

Could you also an IRGen test for this rather than only testing CUDA assembly output?

[yaxunl]

In D90174#2370336, @rjmccall wrote:

I agree this is useful. However, you need to update the manual to cover faststd.

will update the manual.

clang/test/CodeGenCUDA/fp-contract.cu

174

@rjmccall I have IRGen checks in this test. Are they sufficient? Thanks.

[yaxunl]

updated manual

[rjmccall]

Hmm. Do we actually want this behavior of fast overriding pragmas? What do other compilers do here? It might be reasonable to just treat this as a bug.

clang/docs/LanguageExtensions.rst

3214 ↗

(On Diff #302560)

Suggestion:

This can be useful when fast contraction is otherwise enabled for the translation unit
with the ``-ffp-contract=faststd`` flag. Note that ``-ffp-contract=fast`` will override
pragmas to fuse multiply and addition across statements regardless of any controlling
pragmas.

clang/docs/UsersManual.rst

1295 ↗

(On Diff #302560)

"...with the `FP_CONTRACT and clang fp contract` pragmas. Please refer..."

[scanon]

GCC doesn't respect the pragma, so "what other compilers do" is not a particularly useful metric.

[scanon]

(If you tell GCC to respect the pragma via -std=c17 or similar, then -ffp-contract=fast overrides it just like clang's current behavior: https://godbolt.org/z/5dxxGb)

[yaxunl]

revised manual by John's comments

[yaxunl]

In D90174#2371577, @rjmccall wrote:

Hmm. Do we actually want this behavior of fast overriding pragmas? What do other compilers do here? It might be reasonable to just treat this as a bug.

I think clang is just trying to follow gcc's behavior. However, this is undesirable in certain cases. Introducing 'faststd' gives us more choices to avoid the undesirable behavior.

[yaxunl]

In D90174#2373829, @yaxunl wrote:

In D90174#2371577, @rjmccall wrote:

Hmm. Do we actually want this behavior of fast overriding pragmas? What do other compilers do here? It might be reasonable to just treat this as a bug.

I think clang is just trying to follow gcc's behavior. However, this is undesirable in certain cases. Introducing 'faststd' gives us more choices to avoid the undesirable behavior.

For other compilers:

MSVC respects pragma with /fp:fast option:

https://godbolt.org/z/3rja55

Intel compiler also respects fp_contract pragma with -fp-model fast={1|2} option:

https://godbolt.org/z/fez86h

nvcc by default always fuse across statements and does not support pragmas to control fp contract

https://docs.nvidia.com/cuda/floating-point/index.html#controlling-fused-multiply-add

[yaxunl]

@rjmccall ping. Any further concerns for this patch? Thanks.

[rjmccall]

Okay. It sounds like strict compatibility with GCC implies ignoring pragmas in fast, and that's what we're most concerned with, since this is originally a GCC option. So the only question I have now is whether faststd is the best name for this. Does anyone want to suggest an alternative?

[scanon]

I do not much like faststd, as there's nothing "standard" about it. I do not, however, have a better suggestion off the top of my head. Let's pause and consider the name a little bit longer, please?

[yaxunl]

How about fast-constrained, fast-limited, fast-restricted, or fast-restrained?

[rjmccall]

fast-strict? Sounds sort of oxymoronic, but it's fast while also being strict about honoring pragmas. I don't have any great ideas here.

[scanon]

Strictly speaking, fp-contract=fast probably should have been a separate flag entirely (since there's no _expression_ being contracted in fast). Unfortunately, that ship has sailed, and it does constrain our ability to choose an accurate name somewhat.

What if we just spell it out? fast-respect-pragma? fast-when-unspecified? I don't think that we really need to try to be as brief as possible with this one.

[tra]

In D90174#2387518, @scanon wrote:

Strictly speaking, fp-contract=fast probably should have been a separate flag entirely (since there's no _expression_ being contracted in fast). Unfortunately, that ship has sailed, and it does constrain our ability to choose an accurate name somewhat.

What if we just spell it out? fast-respect-pragma? fast-when-unspecified? I don't think that we really need to try to be as brief as possible with this one.

This sounds reasonable. We already have -fhonor-nans and -fhonor-infinities. Should we make it fast-honor-pragma for consistency?

[yaxunl]

In D90174#2389269, @tra wrote:

In D90174#2387518, @scanon wrote:

Strictly speaking, fp-contract=fast probably should have been a separate flag entirely (since there's no _expression_ being contracted in fast). Unfortunately, that ship has sailed, and it does constrain our ability to choose an accurate name somewhat.

What if we just spell it out? fast-respect-pragma? fast-when-unspecified? I don't think that we really need to try to be as brief as possible with this one.

This sounds reasonable. We already have -fhonor-nans and -fhonor-infinities. Should we make it fast-honor-pragma for consistency?

+1 with fast-honor-pragma

[rjmccall]

Probably should be pluralized for consistency, fast-honor-pragmas, but yeah, that's fine with me.

[yaxunl]

rename faststd to fast-honor-pragmas

[yaxunl]

ping

[scanon]

I'm fine with this.

clang/lib/CodeGen/BackendUtil.cpp

static void initTargetOptions(DiagnosticsEngine &Diags,
   case LangOptions::FPM_On:
     Options.AllowFPOpFusion = llvm::FPOpFusion::Standard;
     break;
   case LangOptions::FPM_Fast:
     Options.AllowFPOpFusion = llvm::FPOpFusion::Fast;
     break;
   }
 
+  // HIP toolchain does not support 'Fast' FPOpFusion in backends since it fuses
+  // multiplication/addition instructions without contract flag from device
+  // library functions in LLVM bitcode, which causes accuracy loss in certain
+  // math functions, e.g. tan(-1e20) becomes -0.933 instead of 0.8446. For
+  // device library functions in bitcode to work, 'Strict' or 'Standard'
+  // FPOpFusion options in backends is needed.
+  //
+  // Clang may need a separate option to control FPOpFusion option in backend
+  // instead of using -ffp-contract to control both frontend and backend.
+  if (LangOpts.HIP)
+    Options.AllowFPOpFusion = llvm::FPOpFusion::Standard;
+
   Options.UseInitArray = CodeGenOpts.UseInitArray;
   Options.DisableIntegratedAS = CodeGenOpts.DisableIntegratedAS;
   Options.CompressDebugSections = CodeGenOpts.getCompressDebugSections();
   Options.RelaxELFRelocations = CodeGenOpts.RelaxELFRelocations;
 
   // Set EABI version.
   Options.EABIVersion = TargetOpts.EABIVersion;
 

clang/test/CodeGenCUDA/fp-contract.cu

 // REQUIRES: x86-registered-target
 // REQUIRES: nvptx-registered-target
+// REQUIRES: amdgpu-registered-target
 
-// By default we should fuse multiply/add into fma instruction.
+// By default CUDA/HIP uses -ffp-contract=fast.
+// we should fuse multiply/add into fma instruction.
+// In IR, fmul/fadd instructions with contract flag are emitted.
+// In backend
+//    nvptx -  assumes fast fp fuse option, which fuses
+//             mult/add insts disregarding contract flag and
+//             llvm.fmuladd intrinsics.
+//    amdgcn - assumes standard fp fuse option, which only
+//             fuses mult/add insts with contract flag and
+//             llvm.fmuladd intrinsics.
+
+// RUN: %clang_cc1 -fcuda-is-device -triple nvptx-nvidia-cuda -S \
+// RUN:   -disable-llvm-passes -o - %s \
+// RUN:   | FileCheck -check-prefixes=COMMON,NV-ON %s
+// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -S \
+// RUN:   -target-cpu gfx906 -disable-llvm-passes -o - -x hip %s \
+// RUN:   | FileCheck -check-prefixes=COMMON,AMD-ON %s
 // RUN: %clang_cc1 -fcuda-is-device -triple nvptx-nvidia-cuda -S \
-// RUN:   -disable-llvm-passes -o - %s | FileCheck -check-prefix ENABLED %s
+// RUN:   -O3 -o - %s \
+// RUN:   | FileCheck -check-prefixes=COMMON,NV-OPT-FAST %s
+// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -S \
+// RUN:   -O3 -target-cpu gfx906 -o - -x hip %s \
+// RUN:   | FileCheck -check-prefixes=COMMON,AMD-OPT-FAST %s
+
+// Check separate compile/backend steps corresponding to -save-temps.
+
+// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -emit-llvm \
+// RUN:   -O3 -disable-llvm-passes -target-cpu gfx906 -o %t.ll -x hip %s
+// RUN: cat %t.ll  | FileCheck -check-prefixes=COMMON,AMD-OPT-FAST-IR %s
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -S \
+// RUN:   -O3 -target-cpu gfx906 -o - -x ir %t.ll \
+// RUN:   | FileCheck -check-prefixes=COMMON,AMD-OPT-FAST %s
 
 // Explicit -ffp-contract=fast
+// In IR, fmul/fadd instructions with contract flag are emitted.
+// In backend
+//    nvptx -  assumes fast fp fuse option, which fuses
+//             mult/add insts disregarding contract flag and
+//             llvm.fmuladd intrinsics.
+//    amdgcn - assumes standard fp fuse option, which only
+//             fuses mult/add insts with contract flag and
+//             llvm.fmuladd intrinsics.
+
 // RUN: %clang_cc1 -fcuda-is-device -triple nvptx-nvidia-cuda -S \
 // RUN:   -ffp-contract=fast -disable-llvm-passes -o - %s \
-// RUN:   | FileCheck -check-prefix ENABLED %s
+// RUN:   | FileCheck -check-prefixes=COMMON,NV-ON %s
+// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -S \
+// RUN:   -target-cpu gfx906 -disable-llvm-passes -o - -x hip %s \
+// RUN:   -ffp-contract=fast \
+// RUN:   | FileCheck -check-prefixes=COMMON,AMD-ON %s
+// RUN: %clang_cc1 -fcuda-is-device -triple nvptx-nvidia-cuda -S \
+// RUN:   -O3 -o - %s \
+// RUN:   -ffp-contract=fast \
+// RUN:   | FileCheck -check-prefixes=COMMON,NV-OPT-FAST %s
+// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -S \
+// RUN:   -O3 -target-cpu gfx906 -o - -x hip %s \
+// RUN:   -ffp-contract=fast \
+// RUN:   | FileCheck -check-prefixes=COMMON,AMD-OPT-FAST %s
+
+// Check separate compile/backend steps corresponding to -save-temps.
+
+// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -emit-llvm \
+// RUN:   -ffp-contract=fast \
+// RUN:   -O3 -disable-llvm-passes -target-cpu gfx906 -o %t.ll -x hip %s
+// RUN: cat %t.ll  | FileCheck -check-prefixes=COMMON,AMD-OPT-FAST-IR %s
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -S \
+// RUN:   -ffp-contract=fast \
+// RUN:   -O3 -target-cpu gfx906 -o - -x ir %t.ll \
+// RUN:   | FileCheck -check-prefixes=COMMON,AMD-OPT-FAST %s
 
 // Explicit -ffp-contract=on -- fusing by front-end.
+// In IR,
+//    mult/add in the same statement - llvm.fmuladd instrinsic emitted
+//    mult/add in different statement -  fmul/fadd instructions without
+//                                       contract flag are emitted.
+// In backend
+//    nvptx/amdgcn - assumes standard fp fuse option, which only
+//                   fuses mult/add insts with contract flag or
+//                   llvm.fmuladd intrinsics.
+
 // RUN: %clang_cc1 -fcuda-is-device -triple nvptx-nvidia-cuda -S \
 // RUN:   -ffp-contract=on -disable-llvm-passes -o - %s \
-// RUN:   | FileCheck -check-prefix ENABLED %s
+// RUN:   | FileCheck -check-prefixes=COMMON,NV-ON %s
+// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -S \
+// RUN:   -target-cpu gfx906 -disable-llvm-passes -o - -x hip %s \
+// RUN:   -ffp-contract=on \
+// RUN:   | FileCheck -check-prefixes=COMMON,AMD-ON %s
+// RUN: %clang_cc1 -fcuda-is-device -triple nvptx-nvidia-cuda -S \
+// RUN:   -O3 -o - %s \
+// RUN:   -ffp-contract=on \
+// RUN:   | FileCheck -check-prefixes=COMMON,NV-OPT-ON %s
+// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -S \
+// RUN:   -O3 -target-cpu gfx906 -o - -x hip %s \
+// RUN:   -ffp-contract=on \
+// RUN:   | FileCheck -check-prefixes=COMMON,AMD-OPT-ON %s
+
+// Check separate compile/backend steps corresponding to -save-temps.
+
+// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -emit-llvm \
+// RUN:   -ffp-contract=on \
+// RUN:   -O3 -disable-llvm-passes -target-cpu gfx906 -o %t.ll -x hip %s
+// RUN: cat %t.ll  | FileCheck -check-prefixes=COMMON,AMD-OPT-ON-IR %s
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -S \
+// RUN:   -ffp-contract=on \
+// RUN:   -O3 -target-cpu gfx906 -o - -x ir %t.ll \
+// RUN:   | FileCheck -check-prefixes=COMMON,AMD-OPT-ON %s
 
 // Explicit -ffp-contract=off should disable instruction fusing.
+// In IR, fmul/fadd instructions without contract flag are emitted.
+// In backend
+//    nvptx/amdgcn - assumes standard fp fuse option, which only
+//                   fuses mult/add insts with contract flag or
+//                   llvm.fmuladd intrinsics.
+
 // RUN: %clang_cc1 -fcuda-is-device -triple nvptx-nvidia-cuda -S \
 // RUN:   -ffp-contract=off -disable-llvm-passes -o - %s \
-// RUN:   | FileCheck -check-prefix DISABLED %s
-
+// RUN:   | FileCheck -check-prefixes=COMMON,NV-OFF %s
+// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -S \
+// RUN:   -target-cpu gfx906 -disable-llvm-passes -o - -x hip %s \
+// RUN:   -ffp-contract=off \
+// RUN:   | FileCheck -check-prefixes=COMMON,AMD-OFF %s
+// RUN: %clang_cc1 -fcuda-is-device -triple nvptx-nvidia-cuda -S \
+// RUN:   -O3 -o - %s \
+// RUN:   -ffp-contract=off \
+// RUN:   | FileCheck -check-prefixes=COMMON,NV-OPT-ON %s
+// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -S \
+// RUN:   -O3 -target-cpu gfx906 -o - -x hip %s \
+// RUN:   -ffp-contract=off \
+// RUN:   | FileCheck -check-prefixes=COMMON,AMD-OPT-OFF %s
+
+// Check separate compile/backend steps corresponding to -save-temps.
+
+// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -emit-llvm \
+// RUN:   -ffp-contract=off \
+// RUN:   -O3 -disable-llvm-passes -target-cpu gfx906 -o %t.ll -x hip %s
+// RUN: cat %t.ll  | FileCheck -check-prefixes=COMMON,AMD-OPT-OFF-IR %s
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -S \
+// RUN:   -ffp-contract=off \
+// RUN:   -O3 -target-cpu gfx906 -o - -x ir %t.ll \
+// RUN:   | FileCheck -check-prefixes=COMMON,AMD-OPT-OFF %s
 
 #include "Inputs/cuda.h"
 
+// Test multiply/add in the same statement, which can be emitted as FMA when
+// fp-contract is on or fast.
 __host__ __device__ float func(float a, float b, float c) { return a + b * c; }
-// ENABLED:       fma.rn.f32
-// ENABLED-NEXT:  st.param.f32
-
-// DISABLED:      mul.rn.f32
-// DISABLED-NEXT: add.rn.f32
-// DISABLED-NEXT: st.param.f32
+// COMMON-LABEL: _Z4funcfff
+// NV-ON:       fma.rn.f32
+// NV-ON-NEXT:  st.param.f32
+// AMD-ON:       v_fmac_f32_e64
+// AMD-ON-NEXT:  s_setpc_b64
+
+// NV-OFF:      mul.rn.f32
+// NV-OFF-NEXT: add.rn.f32
+// NV-OFF-NEXT: st.param.f32
+// AMD-OFF:      v_mul_f32_e64
+// AMD-OFF-NEXT: v_add_f32_e64
+// AMD-OFF-NEXT: s_setpc_b64
+
+// AMD-OPT-FAST-IR: fmul contract float
+// AMD-OPT-FAST-IR: fadd contract float
+// AMD-OPT-FAST: v_fmac_f32_e32
+// AMD-OPT-FAST-NEXT: s_setpc_b64
+// AMD-OPT-ON-IR: @llvm.fmuladd.f32
+// AMD-OPT-ON: v_fmac_f32_e32
+// AMD-OPT-ON-NEXT: s_setpc_b64
+// AMD-OPT-OFF-IR: fmul float
+// AMD-OPT-OFF-IR: fadd float
+// AMD-OPT-OFF: v_mul_f32_e32
+// AMD-OPT-OFF-NEXT: v_add_f32_e32
+// AMD-OPT-OFF-NEXT: s_setpc_b64
+
+// Test multiply/add in the different statements, which can be emitted as
+// FMA when fp-contract is fast but not on.
+__host__ __device__ float func2(float a, float b, float c) {
+  float t = b * c;
+  return t + a;
+}
+// COMMON-LABEL: _Z5func2fff
+// NV-OPT-FAST:       fma.rn.f32
+// NV-OPT-FAST-NEXT:  st.param.f32
+
+// NV-OPT-ON:      mul.rn.f32
+// NV-OPT-ON:      add.rn.f32
+// NV-OPT-ON-NEXT: st.param.f32
+
+// AMD-OPT-FAST-IR: fmul contract float
+// AMD-OPT-FAST-IR: fadd contract float
+// AMD-OPT-FAST: v_fmac_f32_e32
+// AMD-OPT-FAST-NEXT: s_setpc_b64
+// AMD-OPT-ON-IR: fmul float
+// AMD-OPT-ON-IR: fadd float
+// AMD-OPT-ON: v_mul_f32_e32
+// AMD-OPT-ON-NEXT: v_add_f32_e32
+// AMD-OPT-ON-NEXT: s_setpc_b64
+// AMD-OPT-OFF-IR: fmul float
+// AMD-OPT-OFF-IR: fadd float
+// AMD-OPT-OFF: v_mul_f32_e32
+// AMD-OPT-OFF-NEXT: v_add_f32_e32
+// AMD-OPT-OFF-NEXT: s_setpc_b64
+
+// Test multiply/add in the different statements, which is forced
+// to be compiled with fp contract on. fmul/fadd without contract
+// flags are emitted in IR. In nvptx, they are emitted as FMA in
+// fp-contract is fast but not on, as nvptx backend uses the same
+// fp fuse option as front end, whereas fast fp fuse option in
+// backend fuses fadd/fmul disregarding contract flag. In amdgcn
+// they are not fused as amdgcn always use standard fp fusion
+// option which respects contract flag.
+  __host__ __device__ float func3(float a, float b, float c) {
+#pragma clang fp contract(on)
+  float t = b * c;
+  return t + a;
+}
+// COMMON-LABEL: _Z5func3fff
+// NV-OPT-FAST:       fma.rn.f32
+// NV-OPT-FAST-NEXT:  st.param.f32
+
+// NV-OPT-ON:      mul.rn.f32
+// NV-OPT-ON:      add.rn.f32
+// NV-OPT-ON-NEXT: st.param.f32
+
+// AMD-OPT-FAST-IR: fmul float
+// AMD-OPT-FAST-IR: fadd float
+// AMD-OPT-FAST: v_mul_f32_e32
+// AMD-OPT-FAST-NEXT: v_add_f32_e32
+// AMD-OPT-FAST-NEXT: s_setpc_b64
+// AMD-OPT-ON-IR: fmul float
+// AMD-OPT-ON-IR: fadd float
+// AMD-OPT-ON: v_mul_f32_e32
+// AMD-OPT-ON-NEXT: v_add_f32_e32
+// AMD-OPT-ON-NEXT: s_setpc_b64
+// AMD-OPT-OFF-IR: fmul float
+// AMD-OPT-OFF-IR: fadd float
+// AMD-OPT-OFF: v_mul_f32_e32
+// AMD-OPT-OFF-NEXT: v_add_f32_e32
+// AMD-OPT-OFF-NEXT: s_setpc_b64