diff --git a/clang/lib/Frontend/CompilerInvocation.cpp b/clang/lib/Frontend/CompilerInvocation.cpp
--- a/clang/lib/Frontend/CompilerInvocation.cpp
+++ b/clang/lib/Frontend/CompilerInvocation.cpp
@@ -2374,6 +2374,10 @@
     // 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;
diff --git a/clang/test/CodeGenHIP/fp-contract.hip b/clang/test/CodeGenHIP/fp-contract.hip
new file mode 100644
--- /dev/null
+++ b/clang/test/CodeGenHIP/fp-contract.hip
@@ -0,0 +1,36 @@
+// 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; }
diff --git a/clang/test/CodeGenHIP/lit.local.cfg b/clang/test/CodeGenHIP/lit.local.cfg
new file mode 100644
--- /dev/null
+++ b/clang/test/CodeGenHIP/lit.local.cfg
@@ -0,0 +1 @@
+config.suffixes = ['.cpp', '.hip']