r337639 - [HIP] Support -fcuda-flush-denormals-to-zero for amdgcn
Yaxun Liu via cfe-commits
cfe-commits at lists.llvm.org
Fri Jul 20 19:02:22 PDT 2018
Author: yaxunl
Date: Fri Jul 20 19:02:22 2018
New Revision: 337639
URL: http://llvm.org/viewvc/llvm-project?rev=337639&view=rev
Log:
[HIP] Support -fcuda-flush-denormals-to-zero for amdgcn
Differential Revision: https://reviews.llvm.org/D48287
Modified:
cfe/trunk/include/clang/Basic/LangOptions.def
cfe/trunk/lib/CodeGen/CGCall.cpp
cfe/trunk/lib/CodeGen/CodeGenModule.cpp
cfe/trunk/lib/Frontend/CompilerInvocation.cpp
cfe/trunk/test/CodeGenCUDA/flush-denormals.cu
Modified: cfe/trunk/include/clang/Basic/LangOptions.def
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/LangOptions.def?rev=337639&r1=337638&r2=337639&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/LangOptions.def (original)
+++ cfe/trunk/include/clang/Basic/LangOptions.def Fri Jul 20 19:02:22 2018
@@ -209,7 +209,6 @@ LANGOPT(RenderScript , 1, 0, "Rende
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")
Modified: cfe/trunk/lib/CodeGen/CGCall.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGCall.cpp?rev=337639&r1=337638&r2=337639&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGCall.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGCall.cpp Fri Jul 20 19:02:22 2018
@@ -1800,7 +1800,7 @@ void CodeGenModule::ConstructDefaultFnAt
FuncAttrs.addAttribute(llvm::Attribute::NoUnwind);
// Respect -fcuda-flush-denormals-to-zero.
- if (getLangOpts().CUDADeviceFlushDenormalsToZero)
+ if (CodeGenOpts.FlushDenorm)
FuncAttrs.addAttribute("nvptx-f32ftz", "true");
}
}
Modified: cfe/trunk/lib/CodeGen/CodeGenModule.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CodeGenModule.cpp?rev=337639&r1=337638&r2=337639&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CodeGenModule.cpp (original)
+++ cfe/trunk/lib/CodeGen/CodeGenModule.cpp Fri Jul 20 19:02:22 2018
@@ -526,7 +526,7 @@ void CodeGenModule::Release() {
// 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.
Modified: cfe/trunk/lib/Frontend/CompilerInvocation.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Frontend/CompilerInvocation.cpp?rev=337639&r1=337638&r2=337639&view=diff
==============================================================================
--- cfe/trunk/lib/Frontend/CompilerInvocation.cpp (original)
+++ cfe/trunk/lib/Frontend/CompilerInvocation.cpp Fri Jul 20 19:02:22 2018
@@ -690,7 +690,9 @@ static bool ParseCodeGenArgs(CodeGenOpti
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 =
@@ -2191,9 +2193,6 @@ 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.CUDAIsDevice && Args.hasArg(OPT_fcuda_approx_transcendentals))
Opts.CUDADeviceApproxTranscendentals = 1;
Modified: cfe/trunk/test/CodeGenCUDA/flush-denormals.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenCUDA/flush-denormals.cu?rev=337639&r1=337638&r2=337639&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenCUDA/flush-denormals.cu (original)
+++ cfe/trunk/test/CodeGenCUDA/flush-denormals.cu Fri Jul 20 19:02:22 2018
@@ -5,6 +5,13 @@
// 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"
@@ -12,11 +19,19 @@
// -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}
More information about the cfe-commits
mailing list