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} _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
