maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

[CUDA] Enable fusing FP ops (-ffp-contract=fast) for CUDA by default. 

This matches default nvcc behavior and gives substantial
performance boost on GPU where fmad is much cheaper compared to add+mul.

Differential Revision: http://reviews.llvm.org/D20341

llvm-svn: 270094
This commit is contained in:
Artem Belevich 2016-05-19 18:44:45 +00:00
parent cb2d266360
commit 31c3bad499
2 changed files with 41 additions and 4 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

13
clang/lib/Frontend/CompilerInvocation.cpp
View File

@ -2255,10 +2255,15 @@ bool CompilerInvocation::CreateFromArgs(CompilerInvocation &Res,
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

32
clang/test/CodeGenCUDA/fp-contract.cu Normal file
View File

@ -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