[HIP] Support -fcuda-flush-denormals-to-zero for amdgcn

[yaxunl]

Correct comments in test.

[tra]

Using OpenCL's flag for the purpose adds a *third* way we handle denormals flushing in clang. Now it would be HIP (which is CUDA-like) using OpenCL's flag for denormals instead of CUDA's one.
You could change AMDGPUTargetInfo::adjustTargetOptions() to use CGOpts.getLangOpts().CUDADeviceFlushDenormalsToZero instead. That would at least make HIP and CUDA do the same thing.

I think it would work better if we could coalesce CUDADeviceFlushDenormalsToZero and CodeGenOpts.FlushDenorm and, maybe move the flag to LangOpts , so we could use LangOpts.CUDAIsDevice.

[yaxunl]

In D48287#1138262, @tra wrote:

Using OpenCL's flag for the purpose adds a *third* way we handle denormals flushing in clang. Now it would be HIP (which is CUDA-like) using OpenCL's flag for denormals instead of CUDA's one.
You could change AMDGPUTargetInfo::adjustTargetOptions() to use CGOpts.getLangOpts().CUDADeviceFlushDenormalsToZero instead. That would at least make HIP and CUDA do the same thing.

I think it would work better if we could coalesce CUDADeviceFlushDenormalsToZero and CodeGenOpts.FlushDenorm and, maybe move the flag to LangOpts , so we could use LangOpts.CUDAIsDevice.

Sorry for the delay.

CGOpts does not have member function getLangOpts(). It seems whereever we need to refer to LangOpts.CUDADeviceFlushDenormalsToZero, we can use Opts.FlushDenorm, but not true vice versa. Therefore if we want a unified option, Opts.FlushDenorm is a better choice.

[yaxunl]

Replace LangOpts.CUDADeviceFlushDenormalsToZero with CodeGenOpts.FlushDenorm.

[tra]

Thank you. That should work.

include/clang/Basic/LangOptions.def

 LANGOPT(OpenMPIsDevice    , 1, 0, "Generate code only for OpenMP target device")
 LANGOPT(OpenMPCUDAMode    , 1, 0, "Generate code for OpenMP pragmas in SIMT/SPMD mode")
 LANGOPT(OpenMPHostCXXExceptions    , 1, 0, "C++ exceptions handling in the host code.")
 LANGOPT(RenderScript      , 1, 0, "RenderScript")
 
 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(CUDADeviceApproxTranscendentals, 1, 0, "using approximate transcendental functions")
 LANGOPT(CUDARelocatableDeviceCode, 1, 0, "generate relocatable device code")
 
 LANGOPT(SizedDeallocation , 1, 0, "sized deallocation")
 LANGOPT(AlignedAllocation , 1, 0, "aligned allocation")
 LANGOPT(AlignedAllocationUnavailable, 1, 0, "aligned allocation functions are unavailable")
 LANGOPT(NewAlignOverride  , 32, 0, "maximum alignment guaranteed by '::operator new(size_t)'")
 LANGOPT(ConceptsTS , 1, 0, "enable C++ Extensions for Concepts")

lib/CodeGen/CGCall.cpp

     FuncAttrs.addAttribute(llvm::Attribute::Convergent);
   }
 
   if (getLangOpts().CUDA && getLangOpts().CUDAIsDevice) {
     // Exceptions aren't supported in CUDA device code.
     FuncAttrs.addAttribute(llvm::Attribute::NoUnwind);
 
     // Respect -fcuda-flush-denormals-to-zero.
-    if (getLangOpts().CUDADeviceFlushDenormalsToZero)
+    if (CodeGenOpts.FlushDenorm)
       FuncAttrs.addAttribute("nvptx-f32ftz", "true");
   }
 }
 
 void CodeGenModule::AddDefaultFnAttrs(llvm::Function &F) {
   llvm::AttrBuilder FuncAttrs;
   ConstructDefaultFnAttrList(F.getName(),
                              F.hasFnAttribute(llvm::Attribute::OptimizeNone),

lib/CodeGen/CodeGenModule.cpp

getModule().addModuleFlag(llvm::Module::Override, "cf-protection-branch",
                               1);
   }
 
   if (LangOpts.CUDAIsDevice && 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);
+                              CodeGenOpts.FlushDenorm ? 1 : 0);
   }
 
   // Emit OpenCL specific module metadata: OpenCL/SPIR version.
   if (LangOpts.OpenCL) {
     EmitOpenCLMetadata();
     // Emit SPIR version.
     if (getTriple().getArch() == llvm::Triple::spir ||
         getTriple().getArch() == llvm::Triple::spir64) {

lib/Frontend/CompilerInvocation.cpp

Opts.NoNaNsFPMath = (Args.hasArg(OPT_menable_no_nans) ||
                        Args.hasArg(OPT_cl_unsafe_math_optimizations) ||
                        Args.hasArg(OPT_cl_finite_math_only) ||
                        Args.hasArg(OPT_cl_fast_relaxed_math));
   Opts.NoSignedZeros = (Args.hasArg(OPT_fno_signed_zeros) ||
                         Args.hasArg(OPT_cl_no_signed_zeros) ||
                         Args.hasArg(OPT_cl_unsafe_math_optimizations) ||
                         Args.hasArg(OPT_cl_fast_relaxed_math));
   Opts.Reassociate = Args.hasArg(OPT_mreassociate);
-  Opts.FlushDenorm = Args.hasArg(OPT_cl_denorms_are_zero);
+  Opts.FlushDenorm = Args.hasArg(OPT_cl_denorms_are_zero) ||
+                     (Args.hasArg(OPT_fcuda_is_device) &&
+                      Args.hasArg(OPT_fcuda_flush_denormals_to_zero));
   Opts.CorrectlyRoundedDivSqrt =
       Args.hasArg(OPT_cl_fp32_correctly_rounded_divide_sqrt);
   Opts.UniformWGSize =
       Args.hasArg(OPT_cl_uniform_work_group_size);
   Opts.Reciprocals = Args.getAllArgValues(OPT_mrecip_EQ);
   Opts.ReciprocalMath = Args.hasArg(OPT_freciprocal_math);
   Opts.NoTrappingMath = Args.hasArg(OPT_fno_trapping_math);
   Opts.StrictFloatCastOverflow =
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.CUDADeviceFlushDenormalsToZero = 1;
-
   if (Opts.CUDAIsDevice && Args.hasArg(OPT_fcuda_approx_transcendentals))
     Opts.CUDADeviceApproxTranscendentals = 1;
 
   Opts.CUDARelocatableDeviceCode = Args.hasArg(OPT_fcuda_rdc);
 
   if (Opts.ObjC1) {
     if (Arg *arg = Args.getLastArg(OPT_fobjc_runtime_EQ)) {
       StringRef value = arg->getValue();

test/CodeGenCUDA/flush-denormals.cu

 // RUN: %clang_cc1 -fcuda-is-device \
 // RUN:   -triple nvptx-nvidia-cuda -emit-llvm -o - %s | \
 // RUN:   FileCheck %s -check-prefix CHECK -check-prefix NOFTZ
 // RUN: %clang_cc1 -fcuda-is-device -fcuda-flush-denormals-to-zero \
 // RUN:   -triple nvptx-nvidia-cuda -emit-llvm -o - %s | \
 // RUN:   FileCheck %s -check-prefix CHECK -check-prefix FTZ
 
+// RUN: %clang_cc1 -fcuda-is-device -x hip \
+// RUN:   -triple amdgcn-amd-amdhsa -target-cpu gfx900 -emit-llvm -o - %s | \
+// RUN:   FileCheck %s -check-prefix CHECK -check-prefix AMDNOFTZ
+// RUN: %clang_cc1 -fcuda-is-device -x hip -fcuda-flush-denormals-to-zero \
+// RUN:   -triple amdgcn-amd-amdhsa -target-cpu gfx900 -emit-llvm -o - %s | \
+// RUN:   FileCheck %s -check-prefix CHECK -check-prefix AMDFTZ
+
 #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.
 
+// AMDGCN targets always have +fp64-fp16-denormals.
+// AMDGCN targets without fast FMAF (e.g. gfx803) always have +fp32-denormals.
+// For AMDGCN target with fast FMAF (e.g. gfx900), it has +fp32-denormals
+// by default and -fp32-denormals when there is option
+// -fcuda-flush-denormals-to-zero.
+
 // CHECK-LABEL: define void @foo() #0
 extern "C" __device__ void foo() {}
 
 // FTZ: attributes #0 = {{.*}} "nvptx-f32ftz"="true"
 // NOFTZ-NOT: attributes #0 = {{.*}} "nvptx-f32ftz"
+// AMDNOFTZ: attributes #0 = {{.*}}+fp32-denormals{{.*}}+fp64-fp16-denormals
+// AMDFTZ: attributes #0 = {{.*}}+fp64-fp16-denormals{{.*}}-fp32-denormals
 
 // 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}