[PATCH] D20341: [CUDA] Enable fusing FP ops for CUDA by default.

[Artem Belevich via cfe-commits]

This revision was automatically updated to reflect the committed changes.
Closed by commit rL270094: [CUDA] Enable fusing FP ops (-ffp-contract=fast) for CUDA by default. (authored by tra).

Changed prior to commit:
  http://reviews.llvm.org/D20341?vs=57541&id=57833#toc

Repository:
  rL LLVM

http://reviews.llvm.org/D20341

Files:
  cfe/trunk/lib/Frontend/CompilerInvocation.cpp
  cfe/trunk/test/CodeGenCUDA/fp-contract.cu

Index: cfe/trunk/test/CodeGenCUDA/fp-contract.cu
===================================================================
--- cfe/trunk/test/CodeGenCUDA/fp-contract.cu
+++ cfe/trunk/test/CodeGenCUDA/fp-contract.cu
@@ -0,0 +1,32 @@
+// REQUIRES: x86-registered-target
+// REQUIRES: nvptx-registered-target
+
+// By default we should fuse multiply/add into fma instruction.
+// RUN: %clang_cc1 -fcuda-is-device -triple nvptx-nvidia-cuda -S \
+// RUN:   -disable-llvm-passes -o - %s | FileCheck -check-prefix ENABLED %s
+
+// Explicit -ffp-contract=fast
+// 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
+
+// Explicit -ffp-contract=on -- fusing by front-end (disabled).
+// RUN: %clang_cc1 -fcuda-is-device -triple nvptx-nvidia-cuda -S \
+// RUN:   -ffp-contract=on -disable-llvm-passes -o - %s \
+// RUN:   | FileCheck -check-prefix DISABLED %s
+
+// Explicit -ffp-contract=off should disable instruction fusing.
+// 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
+
+
+#include "Inputs/cuda.h"
+
+__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
Index: cfe/trunk/lib/Frontend/CompilerInvocation.cpp
===================================================================
--- cfe/trunk/lib/Frontend/CompilerInvocation.cpp
+++ cfe/trunk/lib/Frontend/CompilerInvocation.cpp
@@ -2255,10 +2255,15 @@
       LangOpts.ObjCExceptions = 1;
   }
 
-  // During CUDA device-side compilation, the aux triple is the triple used for
-  // host compilation.
-  if (LangOpts.CUDA && LangOpts.CUDAIsDevice) {
-    Res.getTargetOpts().HostTriple = Res.getFrontendOpts().AuxTriple;
+  if (LangOpts.CUDA) {
+    // During CUDA device-side compilation, the aux triple is the
+    // triple used for host compilation.
+    if (LangOpts.CUDAIsDevice)
+      Res.getTargetOpts().HostTriple = Res.getFrontendOpts().AuxTriple;
+
+    // Set default FP_CONTRACT to FAST.
+    if (!Args.hasArg(OPT_ffp_contract))
+      Res.getCodeGenOpts().setFPContractMode(CodeGenOptions::FPC_Fast);
   }
 
   // FIXME: Override value name discarding when asan or msan is used because the


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D20341.57833.patch
Type: text/x-patch
Size: 2477 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20160519/e48e5051/attachment.bin>