[CUDA] Add --cuda-flush-denormals-to-zero.

[jlebar]

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.

[tra]

include/clang/Driver/Options.td

385

We need a provide a way to both enable and disable this. We either need a "-no" variant or make it an option with value.

Also, can we shorten it to --cuda-ftz? I would probably mistype current name more often than not.

386

Is there an equivalent for ftz fo host-side FP operations? It would be good to keep identical host and device side calculations as close as we can.

lib/Driver/ToolChains.cpp

4212

Perhaps we don't need different flags at driver and CC1 levels. Top-level "-f*" options in OPT_f_group are passed to CC1 automatically.

[jlebar]

Address tra's review comments.

I also decided no longer to turn this on when -menable-unsafe-fp-math is on:
That's a cc1 flag that's implied by various clang flags, but now ftz is a clang
flag, so turning it on implicitly didn't make a ton of sense to me.

[jlebar]

Thank you for the review, Art!

include/clang/Driver/Options.td

385

We need a provide a way to both enable and disable this. We either need a "-no" variant or make it an option with value.

Many (most) of the -f flags don't have -fno variants -- how do we decide which ones get an -fno and which don't?

Also, can we shorten it to --cuda-ftz? I would probably mistype current name more often than not.

Well, you and I both were calling it "ctz" about 50% of the time, so I'm not sure --cuda-ftz would solve the problem! :) (In all seriousness, that was one of the reasons I chose not to abbreviate it.)

Maybe "ftz" is a well-known acronym. Doesn't quite look like it from googling, though.

I looked through the flags and concluded that "ftz" was more abbreviated than most of them. Although "flush-denormals-to-zero" is at the verbose end of the spectrum. I considered "flush-denormals", thought that was a big ambiguous -- flush them how?

386

Is there an equivalent for ftz fo host-side FP operations?

Not that I can tell. The only other one I saw was opencl's equivalent flag, which does nothing at the moment.

lib/Driver/ToolChains.cpp

4212

Aha, much better, thank you!

[rnk]

include/clang/Driver/Options.td

385

Let's stick with the long descriptive name for now.

[rnk]

include/clang/Driver/Options.td

385

Personally I would add an -fno variant, and add code to the driver that conditionally forwards the flag to cc1.

[jlebar]

Add -fno variant.

[jlebar]

include/clang/Driver/Options.td

385

OK, added an -fno variant. But it seems simpler to pass both the -f flag and -fno flag to cc1 and let it figure things out? lmk what you think.

[rnk]

include/clang/Driver/Options.td

385

I agree it would be simpler, but it's not what we do elsewhere. There's this idea that the -cc1 line is a canonicalized command line, where no-op options are removed, like -fno-exceptions when EH is off by default.

lib/Frontend/CompilerInvocation.cpp

1567

The right idiom for this is:

Opts.Thing = Args.hasFlag(OPT_fthing, OPT_fno_thing, false);

[jlebar]

Update flags so we only pass -fcuda-flush-denormals-to-zero to cc1 if appropriate.

[jlebar]

lib/Frontend/CompilerInvocation.cpp

1567

Aha, I knew there had to be a better way to do that. Thanks.

[jlebar]

Thank you for explaining that, Reid!

[jlebar]

Reid, are we good here?

[rnk]

Yep, lgtm

include/clang/Basic/LangOptions.def

 LANGOPT(CUDA              , 1, 0, "CUDA")
 LANGOPT(OpenMP            , 1, 0, "OpenMP support")
 LANGOPT(OpenMPUseTLS      , 1, 0, "Use TLS for threadprivates or runtime calls")
 LANGOPT(OpenMPIsDevice    , 1, 0, "Generate code only for OpenMP target device")
 
 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")
 LANGOPT(ConceptsTS , 1, 0, "enable C++ Extensions for Concepts")
 BENIGN_LANGOPT(ElideConstructors , 1, 1, "C++ copy constructor elision")
 BENIGN_LANGOPT(DumpRecordLayouts , 1, 0, "dumping the layout of IRgen'd records")
 BENIGN_LANGOPT(DumpRecordLayoutsSimple , 1, 0, "dumping the layout of IRgen'd records in a simple form")
 BENIGN_LANGOPT(DumpVTableLayouts , 1, 0, "dumping the layouts of emitted vtables")

include/clang/Driver/CC1Options.td

 def fcuda_is_device : Flag<["-"], "fcuda-is-device">,
   HelpText<"Generate code for CUDA device">;
 def fcuda_include_gpubinary : Separate<["-"], "fcuda-include-gpubinary">,
   HelpText<"Incorporate CUDA device-side binary into host object file.">;
 def fcuda_allow_variadic_functions : Flag<["-"], "fcuda-allow-variadic-functions">,
   HelpText<"Allow variadic functions in CUDA device code.">;
 def fno_cuda_host_device_constexpr : Flag<["-"], "fno-cuda-host-device-constexpr">,
   HelpText<"Don't treat unattributed constexpr functions as __host__ __device__.">;
+def fcuda_flush_denormals_to_zero : Flag<["-"], "fcuda-flush-denormals-to-zero">,
+  HelpText<"Flush denormal floating point values to zero in CUDA device code. "
+           "-menable-unsafe-fp-math implies this option.">;
 
 //===----------------------------------------------------------------------===//
 // OpenMP Options
 //===----------------------------------------------------------------------===//
 
 def fopenmp_is_device : Flag<["-"], "fopenmp-is-device">,
   HelpText<"Generate code only for an OpenMP target device.">;
 def fomp_host_ir_file_path : Separate<["-"], "fomp-host-ir-file-path">,

include/clang/Driver/Options.td

 def cuda_gpu_arch_EQ : Joined<["--"], "cuda-gpu-arch=">,
   Flags<[DriverOption, HelpHidden]>, HelpText<"CUDA GPU architecture">;
 def cuda_host_only : Flag<["--"], "cuda-host-only">,
   HelpText<"Do host-side CUDA compilation only">;
 def cuda_noopt_device_debug : Flag<["--"], "cuda-noopt-device-debug">,
   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 cuda_flush_denormals_to_zero : Flag<["--"], "cuda-flush-denormals-to-zero">,
+  HelpText<"Flush denormal floating point values to zero in CUDA device mode.">;
 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">;
 def dM : Flag<["-"], "dM">, Group<d_Group>, Flags<[CC1Option]>,
   HelpText<"Print macro definitions in -E mode instead of normal output">;
 def dead__strip : Flag<["-"], "dead_strip">;
 def dependency_file : Separate<["-"], "dependency-file">, Flags<[CC1Option]>,
   HelpText<"Filename (or -) to write dependency output to">;

lib/CodeGen/CGCall.cpp

void CodeGenModule::ConstructAttributeList(
   }
 
   if (getLangOpts().CUDA && getLangOpts().CUDAIsDevice) {
     // Conservatively, mark all functions and calls in CUDA as convergent
     // (meaning, they may call an intrinsically convergent op, such as
     // __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);
 
   QualType RetTy = FI.getReturnType();
   const ABIArgInfo &RetAI = FI.getReturnInfo();
   switch (RetAI.getKind()) {
   case ABIArgInfo::Extend:

lib/CodeGen/CodeGenModule.cpp

if (   Arch == llvm::Triple::arm
     getModule().addModuleFlag(llvm::Module::Error, "min_enum_size", EnumWidth);
   }
 
   if (CodeGenOpts.SanitizeCfiCrossDso) {
     // Indicate that we want cross-DSO control flow integrity checks.
     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) {
     case 0: break;
     case 1: PL = llvm::PICLevel::Small; break;
     case 2: PL = llvm::PICLevel::Large; break;
     default: llvm_unreachable("Invalid PIC Level");
     }

lib/Driver/ToolChains.cpp

 }
 
 void
 CudaToolChain::addClangTargetOptions(const llvm::opt::ArgList &DriverArgs,
                                      llvm::opt::ArgStringList &CC1Args) const {
   Linux::addClangTargetOptions(DriverArgs, CC1Args);
   CC1Args.push_back("-fcuda-is-device");
 
+  if (DriverArgs.hasArg(options::OPT_cuda_flush_denormals_to_zero))
+    CC1Args.push_back("-fcuda-flush-denormals-to-zero");
+
   if (DriverArgs.hasArg(options::OPT_nocudalib))
     return;
 
   std::string LibDeviceFile = CudaInstallation.getLibDeviceFile(
       DriverArgs.getLastArgValue(options::OPT_march_EQ));
   if (!LibDeviceFile.empty()) {
     CC1Args.push_back("-mlink-cuda-bitcode");
     CC1Args.push_back(DriverArgs.MakeArgString(LibDeviceFile));

lib/Frontend/CompilerInvocation.cpp

if (Args.hasArg(OPT_fcuda_is_device))
     Opts.CUDAIsDevice = 1;
 
   if (Args.hasArg(OPT_fcuda_allow_variadic_functions))
     Opts.CUDAAllowVariadicFunctions = 1;
 
   if (Args.hasArg(OPT_fno_cuda_host_device_constexpr))
     Opts.CUDAHostDeviceConstexpr = 0;
 
+  if (Opts.CUDAIsDevice &&
+      (Args.hasArg(OPT_fcuda_flush_denormals_to_zero) || Opts.UnsafeFPMath))
+    Opts.CUDADeviceFlushDenormalsToZero = 1;
+
   if (Opts.ObjC1) {
     if (Arg *arg = Args.getLastArg(OPT_fobjc_runtime_EQ)) {
       StringRef value = arg->getValue();
       if (Opts.ObjCRuntime.tryParse(value))
         Diags.Report(diag::err_drv_unknown_objc_runtime) << value;
     }
 
     if (Args.hasArg(OPT_fobjc_gc_only))

test/CodeGenCUDA/flush-denormals.cu

+// 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}