[HIP] Set default FP_CONTRACT to ON

[dfukalov]

Summary

With the FAST default FP_CONTRACT mode' setting for
func(float a, float b, float c) { return a + b * c; } FE generates pair
fmul contract + fadd contract that are fused to an fma operation in BE.

But OpenCL fuses these in FE. This approach seems more effective since avoids a
probabilty that these instructions are not fused in BE. Default setting can be
overridden with #pragma STDC FP_CONTRACT by a programmer.

[yaxunl]

I think the point of -ffp-contract=fast is that it allows emitting fma across statements if there are such opportunities whereas -ffp-contract=on only allows fma in one statement

https://llvm.org/docs/CompileCudaWithLLVM.html#flags-that-control-numerical-code

[dfukalov]

In D82650#2116849, @yaxunl wrote:

I think the point of -ffp-contract=fast is that it allows emitting fma across statements if there are such opportunities whereas -ffp-contract=on only allows fma in one statement

https://llvm.org/docs/CompileCudaWithLLVM.html#flags-that-control-numerical-code

Yes, you're right. I was confused by current FE behavior and a number of different commits, reverts and re-commits.

Currently FE checks allowFPContractWithinStatement() and it returns false for FAST mode. As I understand, FAST mode should be more aggressive and include this case.
I'll re-work the change.

[dfukalov]

I was completely wrong

clang/lib/Frontend/CompilerInvocation.cpp

#endif
   }
 
   Opts.HIP = IK.getLanguage() == Language::HIP;
   Opts.CUDA = IK.getLanguage() == Language::CUDA || Opts.HIP;
   if (Opts.CUDA)
     // Set default FP_CONTRACT to FAST.
     Opts.setDefaultFPContractMode(LangOptions::FPM_Fast);
 
+  // Set default FP_CONTRACT to ON like for OpenCL.
+  if (Opts.HIP)
+    Opts.setDefaultFPContractMode(LangOptions::FPM_On);
+
   Opts.RenderScript = IK.getLanguage() == Language::RenderScript;
   if (Opts.RenderScript) {
     Opts.NativeHalfType = 1;
     Opts.NativeHalfArgsAndReturns = 1;
   }
 
   // OpenCL and C++ both have bool, true, false keywords.
   Opts.Bool = Opts.OpenCL || Opts.CPlusPlus;

clang/test/CodeGenHIP/fp-contract.hip

+// By default we should fuse multiply/add into llvm.fmuladd instruction.
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -emit-llvm -fcuda-is-device -o - %s | FileCheck -check-prefixes ENABLED,ALL %s
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -emit-llvm -ffp-contract=on -fcuda-is-device -o - %s | FileCheck -check-prefixes ENABLED,ALL %s
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -emit-llvm -ffp-contract=fast -fcuda-is-device -o - %s | FileCheck -check-prefixes FAST,ALL %s
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -emit-llvm -ffp-contract=off -fcuda-is-device -o - %s | FileCheck -check-prefixes DISABLED,ALL %s
+
+#define __device__ __attribute__((device))
+
+// ALL-LABEL: func
+// ENABLED: call float @llvm.fmuladd.f32
+// FAST: fmul contract float
+// FAST-NEXT: fadd contract float
+// DISABLED: fmul float
+// DISABLED-NEXT: fadd float
+__device__ float func(float a, float b, float c) { return a + b * c; }
+
+// ALL-LABEL: func_on
+// ALL: call float @llvm.fmuladd.f32
+#pragma STDC FP_CONTRACT ON
+__device__ float func_on(float a, float b, float c) { return a + b * c; }
+
+// ALL-LABEL: func_off
+// ALL: fmul float
+// ALL-NEXT: fadd float
+#pragma STDC FP_CONTRACT OFF
+__device__ float func_off(float a, float b, float c) { return a + b * c; }
+
+// ALL-LABEL: func_def
+// ENABLED: call float @llvm.fmuladd.f32
+// FAST: fmul contract float
+// FAST-NEXT: fadd contract float
+// DISABLED: fmul float
+// DISABLED-NEXT: fadd float
+#pragma STDC FP_CONTRACT DEFAULT
+__device__ float func_def(float a, float b, float c) { return a + b * c; }

clang/test/CodeGenHIP/lit.local.cfg

+config.suffixes = ['.cpp', '.hip']