clang/AMDGPU: Stop setting old denormal subtarget features

[yaxunl]

LGTM. Thanks!

[arsenm]

ce2258c1cd5dc9cf20040d1b1e540d80250c1435

clang/lib/Basic/Targets/AMDGPU.h

std::string convertConstraint(const char *&Constraint) const override {
     return std::string(1, *Constraint);
   }
 
   bool
   initFeatureMap(llvm::StringMap<bool> &Features, DiagnosticsEngine &Diags,
                  StringRef CPU,
                  const std::vector<std::string> &FeatureVec) const override;
 
-  void adjustTargetOptions(const CodeGenOptions &CGOpts,
-                           TargetOptions &TargetOpts) const override;
-
   ArrayRef<Builtin::Info> getTargetBuiltins() const override;
 
   void getTargetDefines(const LangOptions &Opts,
                         MacroBuilder &Builder) const override;
 
   BuiltinVaListKind getBuiltinVaListKind() const override {
     return TargetInfo::CharPtrBuiltinVaList;
   }

clang/lib/Basic/Targets/AMDGPU.cpp

if (isAMDGCN(getTriple())) {
     default:
       llvm_unreachable("Unhandled GPU!");
     }
   }
 
   return TargetInfo::initFeatureMap(Features, Diags, CPU, FeatureVec);
 }
 
-void AMDGPUTargetInfo::adjustTargetOptions(const CodeGenOptions &CGOpts,
-                                           TargetOptions &TargetOpts) const {
-  bool hasFP32Denormals = false;
-  bool hasFP64Denormals = false;
-
-  for (auto &I : TargetOpts.FeaturesAsWritten) {
-    if (I == "+fp32-denormals" || I == "-fp32-denormals")
-      hasFP32Denormals = true;
-    if (I == "+fp64-fp16-denormals" || I == "-fp64-fp16-denormals")
-      hasFP64Denormals = true;
-  }
-  if (!hasFP32Denormals)
-    TargetOpts.Features.push_back(
-      (Twine(hasFastFMAF() && hasFullRateDenormalsF32() &&
-             CGOpts.FP32DenormalMode.Output == llvm::DenormalMode::IEEE
-             ? '+' : '-') + Twine("fp32-denormals"))
-            .str());
-  // Always do not flush fp64 or fp16 denorms.
-  if (!hasFP64Denormals && hasFP64())
-    TargetOpts.Features.push_back("+fp64-fp16-denormals");
-}
-
 void AMDGPUTargetInfo::fillValidCPUList(
     SmallVectorImpl<StringRef> &Values) const {
   if (isAMDGCN(getTriple()))
     llvm::AMDGPU::fillValidArchListAMDGCN(Values);
   else
     llvm::AMDGPU::fillValidArchListR600(Values);
 }
 

clang/test/CodeGenCUDA/flush-denormals.cu

 // RUN: %clang_cc1 -fcuda-is-device \
 // RUN:   -triple nvptx-nvidia-cuda -emit-llvm -o - %s | \
-// RUN:   FileCheck -check-prefix=NOFTZ %s
+// RUN:   FileCheck -check-prefixes=NOFTZ,PTXNOFTZ %s
 
 // RUN: %clang_cc1 -fcuda-is-device -fdenormal-fp-math-f32=ieee \
 // RUN:   -triple nvptx-nvidia-cuda -emit-llvm -o - %s | \
-// RUN:   FileCheck -check-prefix=NOFTZ %s
+// RUN:   FileCheck -check-prefixes=NOFTZ,PTXNOFTZ %s
 
 // RUN: %clang_cc1 -fcuda-is-device -fdenormal-fp-math-f32=preserve-sign \
 // RUN:   -triple nvptx-nvidia-cuda -emit-llvm -o - %s | \
-// RUN:   FileCheck -check-prefix=FTZ %s
+// RUN:   FileCheck -check-prefixes=FTZ,PTXFTZ %s
 
 // RUN: %clang_cc1 -fcuda-is-device -x hip \
 // RUN:   -triple amdgcn-amd-amdhsa -target-cpu gfx900 -emit-llvm -o - %s | \
-// RUN:   FileCheck -check-prefix=AMDNOFTZ %s
+// RUN:   FileCheck -check-prefix=NOFTZ %s
 
 // RUN: %clang_cc1 -fcuda-is-device -x hip \
 // RUN:   -triple amdgcn-amd-amdhsa -target-cpu gfx900 -fdenormal-fp-math-f32=ieee -emit-llvm -o - %s | \
-// RUN:   FileCheck -check-prefix=AMDNOFTZ %s
+// RUN:   FileCheck -check-prefix=NOFTZ %s
 
 // RUN: %clang_cc1 -fcuda-is-device -x hip -fdenormal-fp-math-f32=preserve-sign \
 // RUN:   -triple amdgcn-amd-amdhsa -target-cpu gfx900 -emit-llvm -o - %s | \
-// RUN:   FileCheck -check-prefix=AMDFTZ %s
+// RUN:   FileCheck -check-prefix=FTZ %s
 
 #include "Inputs/cuda.h"
 
 // Checks that device function calls get emitted with the "denormal-fp-math-f32"
 // attribute set when we compile CUDA device code with
 // -fdenormal-fp-math-f32. 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
+// AMDGCN targets always have f64/f16 denormals enabled.
+//
+// AMDGCN targets without fast FMAF (e.g. gfx803) always have f32 denormal
+// flushing by default.
+//
+// For AMDGCN target with fast FMAF (e.g. gfx900), it has ieee denormals by
+// default and preserve-sign when there with the option
 // -fcuda-flush-denormals-to-zero.
 
 // CHECK-LABEL: define void @foo() #0
 extern "C" __device__ void foo() {}
 
 // FTZ: attributes #0 = {{.*}} "denormal-fp-math-f32"="preserve-sign,preserve-sign"
 // NOFTZ-NOT: "denormal-fp-math-f32"
 
-// 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}
+// PTXFTZ:!llvm.module.flags = !{{{.*}}[[MODFLAG:![0-9]+]]}
+// PTXFTZ:[[MODFLAG]] = !{i32 4, !"nvvm-reflect-ftz", i32 1}
 
-// NOFTZ:!llvm.module.flags = !{{{.*}}[[MODFLAG:![0-9]+]]}
-// NOFTZ:[[MODFLAG]] = !{i32 4, !"nvvm-reflect-ftz", i32 0}
+// PTXNOFTZ:!llvm.module.flags = !{{{.*}}[[MODFLAG:![0-9]+]]}
+// PTXNOFTZ:[[MODFLAG]] = !{i32 4, !"nvvm-reflect-ftz", i32 0}

clang/test/CodeGenOpenCL/amdgpu-features.cl

 // REQUIRES: amdgpu-registered-target
 
 // Check that appropriate features are defined for every supported AMDGPU
 // "-target" and "-mcpu" options.
 
 // RUN: %clang_cc1 -triple amdgcn -target-cpu gfx904 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX904 %s
 // RUN: %clang_cc1 -triple amdgcn -target-cpu gfx906 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX906 %s
 // RUN: %clang_cc1 -triple amdgcn -target-cpu gfx908 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX908 %s
 // RUN: %clang_cc1 -triple amdgcn -target-cpu gfx1010 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX1010 %s
 // RUN: %clang_cc1 -triple amdgcn -target-cpu gfx1011 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX1011 %s
 // RUN: %clang_cc1 -triple amdgcn -target-cpu gfx1012 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX1012 %s
 // RUN: %clang_cc1 -triple amdgcn -target-cpu gfx801 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX801 %s
 // RUN: %clang_cc1 -triple amdgcn -target-cpu gfx700 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX700 %s
 // RUN: %clang_cc1 -triple amdgcn -target-cpu gfx600 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX600 %s
 // RUN: %clang_cc1 -triple amdgcn -target-cpu gfx601 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX601 %s
 
-// GFX904: "target-features"="+16-bit-insts,+ci-insts,+dpp,+flat-address-space,+fp32-denormals,+fp64-fp16-denormals,+gfx8-insts,+gfx9-insts,+s-memrealtime"
-// GFX906: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dpp,+flat-address-space,+fp32-denormals,+fp64-fp16-denormals,+gfx8-insts,+gfx9-insts,+s-memrealtime"
-// GFX908: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dpp,+flat-address-space,+fp32-denormals,+fp64-fp16-denormals,+gfx8-insts,+gfx9-insts,+mai-insts,+s-memrealtime"
-// GFX1010: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dpp,+flat-address-space,+fp32-denormals,+fp64-fp16-denormals,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime"
-// GFX1011: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dpp,+flat-address-space,+fp32-denormals,+fp64-fp16-denormals,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime"
-// GFX1012: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dpp,+flat-address-space,+fp32-denormals,+fp64-fp16-denormals,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime"
-// GFX801: "target-features"="+16-bit-insts,+ci-insts,+dpp,+flat-address-space,+fp32-denormals,+fp64-fp16-denormals,+gfx8-insts,+s-memrealtime"
-// GFX700: "target-features"="+ci-insts,+flat-address-space,+fp64-fp16-denormals,-fp32-denormals"
-// GFX600: "target-features"="+fp64-fp16-denormals,-fp32-denormals"
-// GFX601: "target-features"="+fp64-fp16-denormals,-fp32-denormals"
+// GFX904: "target-features"="+16-bit-insts,+ci-insts,+dpp,+flat-address-space,+gfx8-insts,+gfx9-insts,+s-memrealtime"
+// GFX906: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dpp,+flat-address-space,+gfx8-insts,+gfx9-insts,+s-memrealtime"
+// GFX908: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dpp,+flat-address-space,+gfx8-insts,+gfx9-insts,+mai-insts,+s-memrealtime"
+// GFX1010: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dpp,+flat-address-space,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime"
+// GFX1011: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dpp,+flat-address-space,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime"
+// GFX1012: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dpp,+flat-address-space,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime"
+// GFX801: "target-features"="+16-bit-insts,+ci-insts,+dpp,+flat-address-space,+gfx8-insts,+s-memrealtime"
+// GFX700: "target-features"="+ci-insts,+flat-address-space"
+// GFX600-NOT: "target-features"
+// GFX601-NOT: "target-features"
 
 kernel void test() {}