r265435 - [CUDA] Add -fcuda-flush-denormals-to-zero.
Justin Lebar via cfe-commits
cfe-commits at lists.llvm.org
Tue Apr 5 11:26:22 PDT 2016
Author: jlebar
Date: Tue Apr 5 13:26:20 2016
New Revision: 265435
URL: http://llvm.org/viewvc/llvm-project?rev=265435&view=rev
Log:
[CUDA] Add -fcuda-flush-denormals-to-zero.
Summary:
Setting this flag causes all functions are annotated with the
"nvvm-f32ftz" = "true" attribute.
In addition, we annotate the module with "nvvm-reflect-ftz" set
to 0 or 1, depending on whether -cuda-flush-denormals-to-zero is set.
This is read by the NVVMReflect pass.
Reviewers: tra, rnk
Subscribers: cfe-commits
Differential Revision: http://reviews.llvm.org/D18671
Added:
cfe/trunk/test/CodeGenCUDA/flush-denormals.cu
Modified:
cfe/trunk/include/clang/Basic/LangOptions.def
cfe/trunk/include/clang/Driver/Options.td
cfe/trunk/lib/CodeGen/CGCall.cpp
cfe/trunk/lib/CodeGen/CodeGenModule.cpp
cfe/trunk/lib/Driver/ToolChains.cpp
cfe/trunk/lib/Frontend/CompilerInvocation.cpp
Modified: cfe/trunk/include/clang/Basic/LangOptions.def
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/LangOptions.def?rev=265435&r1=265434&r2=265435&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/LangOptions.def (original)
+++ cfe/trunk/include/clang/Basic/LangOptions.def Tue Apr 5 13:26:20 2016
@@ -173,6 +173,7 @@ LANGOPT(OpenMPIsDevice , 1, 0, "Gener
LANGOPT(CUDAIsDevice , 1, 0, "compiling for CUDA device")
LANGOPT(CUDAAllowVariadicFunctions, 1, 0, "allowing variadic functions in CUDA device code")
LANGOPT(CUDAHostDeviceConstexpr, 1, 1, "treating unattributed constexpr functions as __host__ __device__")
+LANGOPT(CUDADeviceFlushDenormalsToZero, 1, 0, "flushing denormals to zero")
LANGOPT(AssumeSaneOperatorNew , 1, 1, "implicit __attribute__((malloc)) for C++'s new operators")
LANGOPT(SizedDeallocation , 1, 0, "enable sized deallocation functions")
Modified: cfe/trunk/include/clang/Driver/Options.td
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Driver/Options.td?rev=265435&r1=265434&r2=265435&view=diff
==============================================================================
--- cfe/trunk/include/clang/Driver/Options.td (original)
+++ cfe/trunk/include/clang/Driver/Options.td Tue Apr 5 13:26:20 2016
@@ -382,6 +382,9 @@ def cuda_noopt_device_debug : Flag<["--"
HelpText<"Enable device-side debug info generation. Disables ptxas optimizations.">;
def cuda_path_EQ : Joined<["--"], "cuda-path=">, Group<i_Group>,
HelpText<"CUDA installation path">;
+def fcuda_flush_denormals_to_zero : Flag<["-"], "fcuda-flush-denormals-to-zero">,
+ Flags<[CC1Option]>, HelpText<"Flush denormal floating point values to zero in CUDA device mode.">;
+def fno_cuda_flush_denormals_to_zero : Flag<["-"], "fno-cuda-flush-denormals-to-zero">;
def dA : Flag<["-"], "dA">, Group<d_Group>;
def dD : Flag<["-"], "dD">, Group<d_Group>, Flags<[CC1Option]>,
HelpText<"Print macro definitions in -E mode in addition to normal output">;
Modified: cfe/trunk/lib/CodeGen/CGCall.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGCall.cpp?rev=265435&r1=265434&r2=265435&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGCall.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGCall.cpp Tue Apr 5 13:26:20 2016
@@ -1768,6 +1768,10 @@ void CodeGenModule::ConstructAttributeLi
// __syncthreads(), and so can't have certain optimizations applied around
// them). LLVM will remove this attribute where it safely can.
FuncAttrs.addAttribute(llvm::Attribute::Convergent);
+
+ // Respect -fcuda-flush-denormals-to-zero.
+ if (getLangOpts().CUDADeviceFlushDenormalsToZero)
+ FuncAttrs.addAttribute("nvptx-f32ftz", "true");
}
ClangToLLVMArgMapping IRFunctionArgs(getContext(), FI);
Modified: cfe/trunk/lib/CodeGen/CodeGenModule.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CodeGenModule.cpp?rev=265435&r1=265434&r2=265435&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CodeGenModule.cpp (original)
+++ cfe/trunk/lib/CodeGen/CodeGenModule.cpp Tue Apr 5 13:26:20 2016
@@ -472,6 +472,14 @@ void CodeGenModule::Release() {
getModule().addModuleFlag(llvm::Module::Override, "Cross-DSO CFI", 1);
}
+ if (LangOpts.CUDAIsDevice && getTarget().getTriple().isNVPTX()) {
+ // Indicate whether __nvvm_reflect should be configured to flush denormal
+ // floating point values to 0. (This corresponds to its "__CUDA_FTZ"
+ // property.)
+ getModule().addModuleFlag(llvm::Module::Override, "nvvm-reflect-ftz",
+ LangOpts.CUDADeviceFlushDenormalsToZero ? 1 : 0);
+ }
+
if (uint32_t PLevel = Context.getLangOpts().PICLevel) {
llvm::PICLevel::Level PL = llvm::PICLevel::Default;
switch (PLevel) {
Modified: cfe/trunk/lib/Driver/ToolChains.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Driver/ToolChains.cpp?rev=265435&r1=265434&r2=265435&view=diff
==============================================================================
--- cfe/trunk/lib/Driver/ToolChains.cpp (original)
+++ cfe/trunk/lib/Driver/ToolChains.cpp Tue Apr 5 13:26:20 2016
@@ -4208,6 +4208,10 @@ CudaToolChain::addClangTargetOptions(con
Linux::addClangTargetOptions(DriverArgs, CC1Args);
CC1Args.push_back("-fcuda-is-device");
+ if (DriverArgs.hasFlag(options::OPT_fcuda_flush_denormals_to_zero,
+ options::OPT_fno_cuda_flush_denormals_to_zero, false))
+ CC1Args.push_back("-fcuda-flush-denormals-to-zero");
+
if (DriverArgs.hasArg(options::OPT_nocudalib))
return;
Modified: cfe/trunk/lib/Frontend/CompilerInvocation.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Frontend/CompilerInvocation.cpp?rev=265435&r1=265434&r2=265435&view=diff
==============================================================================
--- cfe/trunk/lib/Frontend/CompilerInvocation.cpp (original)
+++ cfe/trunk/lib/Frontend/CompilerInvocation.cpp Tue Apr 5 13:26:20 2016
@@ -1571,6 +1571,9 @@ static void ParseLangArgs(LangOptions &O
if (Args.hasArg(OPT_fno_cuda_host_device_constexpr))
Opts.CUDAHostDeviceConstexpr = 0;
+ if (Opts.CUDAIsDevice && Args.hasArg(OPT_fcuda_flush_denormals_to_zero))
+ Opts.CUDADeviceFlushDenormalsToZero = 1;
+
if (Opts.ObjC1) {
if (Arg *arg = Args.getLastArg(OPT_fobjc_runtime_EQ)) {
StringRef value = arg->getValue();
Added: cfe/trunk/test/CodeGenCUDA/flush-denormals.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenCUDA/flush-denormals.cu?rev=265435&view=auto
==============================================================================
--- cfe/trunk/test/CodeGenCUDA/flush-denormals.cu (added)
+++ cfe/trunk/test/CodeGenCUDA/flush-denormals.cu Tue Apr 5 13:26:20 2016
@@ -0,0 +1,23 @@
+// RUN: %clang_cc1 -fcuda-is-device \
+// RUN: -triple nvptx-nvidia-cuda -emit-llvm -o - %s | FileCheck %s -check-prefix NOFTZ
+// RUN: %clang_cc1 -fcuda-is-device -fcuda-flush-denormals-to-zero \
+// RUN: -triple nvptx-nvidia-cuda -emit-llvm -o - %s | FileCheck %s -check-prefix FTZ
+
+#include "Inputs/cuda.h"
+
+// Checks that device function calls get emitted with the "ntpvx-f32ftz"
+// attribute set to "true" when we compile CUDA device code with
+// -fcuda-flush-denormals-to-zero. Further, check that we reflect the presence
+// or absence of -fcuda-flush-denormals-to-zero in a module flag.
+
+// CHECK: define void @foo() #0
+extern "C" __device__ void foo() {}
+
+// FTZ: attributes #0 = {{.*}} "nvptx-f32ftz"="true"
+// NOFTZ-NOT: attributes #0 = {{.*}} "nvptx-f32ftz"
+
+// FTZ:!llvm.module.flags = !{[[MODFLAG:![0-9]+]]}
+// FTZ:[[MODFLAG]] = !{i32 4, !"nvvm-reflect-ftz", i32 1}
+
+// NOFTZ:!llvm.module.flags = !{[[MODFLAG:![0-9]+]]}
+// NOFTZ:[[MODFLAG]] = !{i32 4, !"nvvm-reflect-ftz", i32 0}
More information about the cfe-commits
mailing list