IR: Promote "denormal-fp-math" to a first class attribute #174293
+2,369
−938
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Contributor
Author
This stack of pull requests is managed by Graphite. Learn more about stacking. |
arsenm
requested review from
andykaylor,
davemgreen,
dtcxzyw,
jcranmer-intel,
nikic and
spavloff
Member
|
@llvm/pr-subscribers-llvm-support @llvm/pr-subscribers-llvm-ir Author: Matt Arsenault (arsenm) ChangesConvert "denormal-fp-math" and "denormal-fp-math-f32" into a first The syntax in the common cases looks like this: I wasn't sure about reusing the float type name instead of adding a This has a behavior change when using the command flag debug This also does not perform all the code cleanups this enables. I also guessed at how to support this in MLIR. I followed Patch is 439.46 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/174293.diff 219 Files Affected:
diff --git a/clang/lib/CodeGen/CGCall.cpp b/clang/lib/CodeGen/CGCall.cpp
index 4a9025b6e0b0f..c4e2e334134e6 100644
--- a/clang/lib/CodeGen/CGCall.cpp
+++ b/clang/lib/CodeGen/CGCall.cpp
@@ -1941,11 +1941,9 @@ static bool HasStrictReturn(const CodeGenModule &Module, QualType RetTy,
static void addDenormalModeAttrs(llvm::DenormalMode FPDenormalMode,
llvm::DenormalMode FP32DenormalMode,
llvm::AttrBuilder &FuncAttrs) {
- if (FPDenormalMode != llvm::DenormalMode::getDefault())
- FuncAttrs.addAttribute("denormal-fp-math", FPDenormalMode.str());
-
- if (FP32DenormalMode != FPDenormalMode && FP32DenormalMode.isValid())
- FuncAttrs.addAttribute("denormal-fp-math-f32", FP32DenormalMode.str());
+ llvm::DenormalFPEnv FPEnv(FPDenormalMode, FP32DenormalMode);
+ if (FPEnv != llvm::DenormalFPEnv::getDefault())
+ FuncAttrs.addDenormalFPEnvAttr(FPEnv);
}
/// Add default attributes to a function, which have merge semantics under
@@ -2178,24 +2176,16 @@ void CodeGen::mergeDefaultFunctionDefinitionAttributes(
CodeGenOpts.FP32DenormalMode.mergeCalleeMode(DenormModeToMergeF32);
}
- if (Merged == llvm::DenormalMode::getDefault()) {
- AttrsToRemove.addAttribute("denormal-fp-math");
- } else if (Merged != DenormModeToMerge) {
- // Overwrite existing attribute
- FuncAttrs.addAttribute("denormal-fp-math",
- CodeGenOpts.FPDenormalMode.str());
- }
+ llvm::DenormalFPEnv MergedFPEnv(Merged, MergedF32);
- if (MergedF32 == llvm::DenormalMode::getDefault()) {
- AttrsToRemove.addAttribute("denormal-fp-math-f32");
- } else if (MergedF32 != DenormModeToMergeF32) {
+ if (MergedFPEnv == llvm::DenormalFPEnv::getDefault()) {
+ AttrsToRemove.addAttribute(llvm::Attribute::DenormalFPMath);
+ } else {
// Overwrite existing attribute
- FuncAttrs.addAttribute("denormal-fp-math-f32",
- CodeGenOpts.FP32DenormalMode.str());
+ FuncAttrs.addDenormalFPEnvAttr(MergedFPEnv);
}
F.removeFnAttrs(AttrsToRemove);
- addDenormalModeAttrs(Merged, MergedF32, FuncAttrs);
overrideFunctionFeaturesWithTargetFeatures(FuncAttrs, F, TargetOpts);
diff --git a/clang/lib/CodeGen/CGCall.h b/clang/lib/CodeGen/CGCall.h
index 4a86d58895dd9..145992652934f 100644
--- a/clang/lib/CodeGen/CGCall.h
+++ b/clang/lib/CodeGen/CGCall.h
@@ -410,7 +410,7 @@ class ReturnValueSlot {
/// This is useful for adding attrs to bitcode modules that you want to link
/// with but don't control, such as CUDA's libdevice. When linking with such
/// a bitcode library, you might want to set e.g. its functions'
-/// "denormal-fp-math" attribute to match the attr of the functions you're
+/// denormal_fp_math attribute to match the attr of the functions you're
/// codegen'ing. Otherwise, LLVM will interpret the bitcode module's lack of
/// denormal-fp-math attrs as tantamount to denormal-fp-math=ieee, and then LLVM
/// will propagate denormal-fp-math=ieee up to every transitive caller of a
diff --git a/clang/test/CodeGen/denormalfpmode-f32.c b/clang/test/CodeGen/denormalfpmode-f32.c
index 312d1c9277722..9f0340a0f55a8 100644
--- a/clang/test/CodeGen/denormalfpmode-f32.c
+++ b/clang/test/CodeGen/denormalfpmode-f32.c
@@ -1,48 +1,54 @@
-// RUN: %clang_cc1 %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-NONE,CHECK-F32-NONE
-// RUN: %clang_cc1 -fdenormal-fp-math=ieee %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-NONE,CHECK-F32-NONE
-// RUN: %clang_cc1 -fdenormal-fp-math=preserve-sign %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-PS,CHECK-F32-NONE
-// RUN: %clang_cc1 -fdenormal-fp-math=positive-zero %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-PZ,CHECK-F32-NONE
-// RUN: %clang_cc1 -fdenormal-fp-math=dynamic %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-DYNAMIC,CHECK-F32-NONE
+// RUN: %clang_cc1 %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-NONE
+// RUN: %clang_cc1 -fdenormal-fp-math=ieee %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-NONE
+// RUN: %clang_cc1 -fdenormal-fp-math=preserve-sign %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-PS
+// RUN: %clang_cc1 -fdenormal-fp-math=positive-zero %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-PZ
+// RUN: %clang_cc1 -fdenormal-fp-math=dynamic %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-DYNAMIC
-// RUN: %clang_cc1 -fdenormal-fp-math-f32=ieee %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-NONE,CHECK-F32-NONE
-// RUN: %clang_cc1 -fdenormal-fp-math=ieee -fdenormal-fp-math-f32=ieee %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-NONE,CHECK-F32-NONE
-// RUN: %clang_cc1 -fdenormal-fp-math=preserve-sign -fdenormal-fp-math-f32=ieee %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-PS,CHECK-F32-IEEE
-// RUN: %clang_cc1 -fdenormal-fp-math=positive-zero -fdenormal-fp-math-f32=ieee %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-PZ,CHECK-F32-IEEE
-// RUN: %clang_cc1 -fdenormal-fp-math=positive-zero -fdenormal-fp-math-f32=dynamic %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-PZ,CHECK-F32-DYNAMIC
+// RUN: %clang_cc1 -fdenormal-fp-math-f32=ieee %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-NONE
+// RUN: %clang_cc1 -fdenormal-fp-math=ieee -fdenormal-fp-math-f32=ieee %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-NONE
+// RUN: %clang_cc1 -fdenormal-fp-math=preserve-sign -fdenormal-fp-math-f32=ieee %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-PS-F32-IEEE
+// RUN: %clang_cc1 -fdenormal-fp-math=positive-zero -fdenormal-fp-math-f32=ieee %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-PZ-F32-IEEE
+// RUN: %clang_cc1 -fdenormal-fp-math=positive-zero -fdenormal-fp-math-f32=dynamic %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-PZ-F32-DYNAMIC
// RUN: %clang_cc1 -fdenormal-fp-math-f32=preserve-sign %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-NONE,CHECK-F32-PS
// RUN: %clang_cc1 -fdenormal-fp-math=ieee -fdenormal-fp-math-f32=preserve-sign %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-NONE,CHECK-F32-PS
-// RUN: %clang_cc1 -fdenormal-fp-math=preserve-sign -fdenormal-fp-math-f32=preserve-sign %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-PS,CHECK-F32-NONE
-// RUN: %clang_cc1 -fdenormal-fp-math=positive-zero -fdenormal-fp-math-f32=preserve-sign %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-PZ,CHECK-F32-PS
+// RUN: %clang_cc1 -fdenormal-fp-math=preserve-sign -fdenormal-fp-math-f32=preserve-sign %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-PS
+// RUN: %clang_cc1 -fdenormal-fp-math=positive-zero -fdenormal-fp-math-f32=preserve-sign %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-PZ-F32-PS
// RUN: %clang_cc1 -fdenormal-fp-math=ieee -fdenormal-fp-math-f32=dynamic %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-NONE,CHECK-F32-DYNAMIC
// RUN: %clang_cc1 -fdenormal-fp-math-f32=positive-zero %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-NONE,CHECK-F32-PZ
// RUN: %clang_cc1 -fdenormal-fp-math-f32=dynamic %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-NONE,CHECK-F32-DYNAMIC
-// RUN: %clang_cc1 -fdenormal-fp-math=ieee -fdenormal-fp-math-f32=positive-zero %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-NONE,CHECK-F32-PZ
-// RUN: %clang_cc1 -fdenormal-fp-math=dynamic -fdenormal-fp-math-f32=positive-zero %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-DYNAMIC,CHECK-F32-PZ
-// RUN: %clang_cc1 -fdenormal-fp-math=preserve-sign -fdenormal-fp-math-f32=positive-zero %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-PS,CHECK-F32-PZ
-// RUN: %clang_cc1 -fdenormal-fp-math=positive-zero -fdenormal-fp-math-f32=positive-zero %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-PZ,CHECK-F32-NONE
-// RUN: %clang_cc1 -fdenormal-fp-math=dynamic -fdenormal-fp-math-f32=dynamic %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-DYNAMIC,CHECK-F32-NONE
+// RUN: %clang_cc1 -fdenormal-fp-math=ieee -fdenormal-fp-math-f32=positive-zero %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-F32-PZ
+// RUN: %clang_cc1 -fdenormal-fp-math=dynamic -fdenormal-fp-math-f32=positive-zero %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-DYNAMIC-F32-PZ
+// RUN: %clang_cc1 -fdenormal-fp-math=preserve-sign -fdenormal-fp-math-f32=positive-zero %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-PS-F32-PZ
+// RUN: %clang_cc1 -fdenormal-fp-math=positive-zero -fdenormal-fp-math-f32=positive-zero %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-PZ
+// RUN: %clang_cc1 -fdenormal-fp-math=dynamic -fdenormal-fp-math-f32=dynamic %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-DYNAMIC
// CHECK-LABEL: main
// CHECK-ATTR: attributes #0 =
-// CHECK-NONE-NOT:"denormal-fp-math"
-// CHECK-IEEE: "denormal-fp-math"="ieee,ieee"
-// CHECK-PS: "denormal-fp-math"="preserve-sign,preserve-sign"
-// CHECK-PZ: "denormal-fp-math"="positive-zero,positive-zero"
-// CHECK-DYNAMIC: "denormal-fp-math"="dynamic,dynamic"
+// CHECK-NONE-NOT: denormal_fpenv
-// CHECK-F32-NONE-NOT:"denormal-fp-math-f32"
-// CHECK-F32-IEEE: "denormal-fp-math-f32"="ieee,ieee"
-// CHECK-F32-PS: "denormal-fp-math-f32"="preserve-sign,preserve-sign"
-// CHECK-F32-PZ: "denormal-fp-math-f32"="positive-zero,positive-zero"
+// CHECK-IEEE: denormal_fpenv(ieee,ieee)
+// CHECK-PS: denormal_fpenv(preservesign,preservesign)
+// CHECK-PZ: denormal_fpenv(positivezero,positivezero)
+// CHECK-DYNAMIC: denormal_fpenv(dynamic,dynamic)
-// CHECK-F32-DYNAMIC: "denormal-fp-math-f32"="dynamic,dynamic"
+// CHECK-PS-F32-IEEE: denormal_fpenv(preservesign,preservesign float: ieee,ieee)
+// CHECK-PZ-F32-IEEE: denormal_fpenv(positivezero,positivezero float: ieee,ieee)
+// CHECK-PZ-F32-DYNAMIC: denormal_fpenv(positivezero,positivezero float: dynamic,dynamic)
+// CHECK-PZ-F32-PS: denormal_fpenv(positivezero,positivezero float: preservesign,preservesign)
+// CHECK-DYNAMIC-F32-PZ: denormal_fpenv(dynamic,dynamic float: positivezero,positivezero)
+// CHECK: CHECK-PS-F32-PZ: denormal_fpenv(preservesign,preservesign float: positivezero,positivezero)
+
+// CHECK-F32-IEEE: denormal_fpenv(float: ieee,ieee)
+// CHECK-F32-PS: denormal_fpenv(float: preservesign,preservesign)
+// CHECK-F32-PZ: denormal_fpenv(float: positivezero,positivezero)
+// CHECK-F32-DYNAMIC: denormal_fpenv(float: dynamic,dynamic)
int main(void) {
return 0;
diff --git a/clang/test/CodeGen/denormalfpmode.c b/clang/test/CodeGen/denormalfpmode.c
index cffff90d6fbe7..8d0ce644c4da1 100644
--- a/clang/test/CodeGen/denormalfpmode.c
+++ b/clang/test/CodeGen/denormalfpmode.c
@@ -6,10 +6,10 @@
// CHECK-LABEL: main
// The ieee,ieee is the default, so omit the attribute
-// CHECK-IEEE-NOT:"denormal-fp-math"
-// CHECK-PS: attributes #0 = {{.*}}"denormal-fp-math"="preserve-sign,preserve-sign"{{.*}}
-// CHECK-PZ: attributes #0 = {{.*}}"denormal-fp-math"="positive-zero,positive-zero"{{.*}}
-// CHECK-DYNAMIC: attributes #0 = {{.*}}"denormal-fp-math"="dynamic,dynamic"{{.*}}
+// CHECK-IEEE-NOT:denormal_fpenv
+// CHECK-PS: attributes #0 = {{.*}}denormal_fpenv(preservesign,preservesign){{.*}}
+// CHECK-PZ: attributes #0 = {{.*}}denormal_fpenv(positivezero,positivezero){{.*}}
+// CHECK-DYNAMIC: attributes #0 = {{.*}}denormal_fpenv(dynamic,dynamic){{.*}}
int main(void) {
return 0;
diff --git a/clang/test/CodeGenCUDA/flush-denormals.cu b/clang/test/CodeGenCUDA/flush-denormals.cu
index b5abc29dea14b..7e1700ac64331 100644
--- a/clang/test/CodeGenCUDA/flush-denormals.cu
+++ b/clang/test/CodeGenCUDA/flush-denormals.cu
@@ -24,7 +24,7 @@
#include "Inputs/cuda.h"
-// Checks that device function calls get emitted with the "denormal-fp-math-f32"
+// Checks that device function calls get emitted with the denormal_fpenv
// attribute set when we compile CUDA device code with
// -fdenormal-fp-math-f32. Further, check that we reflect the presence or
// absence of -fgpu-flush-denormals-to-zero in a module flag.
@@ -41,8 +41,8 @@
// 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"
+// FTZ: attributes #0 = {{.*}} denormal_fpenv(float: preservesign,preservesign)
+// NOFTZ-NOT: denormal_fpenv
// PTXFTZ:!llvm.module.flags = !{{{.*}}[[MODFLAG:![0-9]+]]}
// PTXFTZ:[[MODFLAG]] = !{i32 4, !"nvvm-reflect-ftz", i32 1}
diff --git a/clang/test/CodeGenCUDA/link-builtin-bitcode-denormal-fp-mode.cu b/clang/test/CodeGenCUDA/link-builtin-bitcode-denormal-fp-mode.cu
index ef02668c3697b..39460ad92d2b2 100644
--- a/clang/test/CodeGenCUDA/link-builtin-bitcode-denormal-fp-mode.cu
+++ b/clang/test/CodeGenCUDA/link-builtin-bitcode-denormal-fp-mode.cu
@@ -127,39 +127,39 @@ __global__ void kernel_f64(double* out, double* a, double* b, double* c) {
// We should not be littering call sites with the attribute
// Everything should use the default ieee with no explicit attribute
-// FIXME: Should check-not "denormal-fp-math" within the denormal-fp-math-f32
+// FIXME: Should check-not denormal_fpenv within the denormal-fp-math-f32
// lines.
// Default mode relies on the implicit check-not for the denormal-fp-math.
-// PSZ: #[[$KERNELATTR]] = { {{.*}} "denormal-fp-math"="preserve-sign,preserve-sign"
+// PSZ: #[[$KERNELATTR]] = { {{.*}} denormal_fpenv(preservesign,preservesign)
// PSZ-SAME: "target-cpu"="gfx803"
-// PSZ: #[[$FUNCATTR]] = { {{.*}} "denormal-fp-math-f32"="preserve-sign,preserve-sign"
+// PSZ: #[[$FUNCATTR]] = { {{.*}} denormal_fpenv(float: preservesign,preservesign)
// PSZ-SAME: "target-cpu"="gfx803"
-// PSZ: #[[$WEAK_FUNCATTR]] = { {{.*}} "denormal-fp-math-f32"="preserve-sign,preserve-sign"
+// PSZ: #[[$WEAK_FUNCATTR]] = { {{.*}} denormal_fpenv(float: preservesign,preservesign)
// PSZ-SAME: "target-cpu"="gfx803"
-// FIXME: Should check-not "denormal-fp-math" within the line
-// IEEEF64-PSZF32: #[[$KERNELATTR]] = { {{.*}} "denormal-fp-math-f32"="preserve-sign,preserve-sign"
+// FIXME: Should check-not denormal_fpenv within the line
+// IEEEF64-PSZF32: #[[$KERNELATTR]] = { {{.*}} denormal_fpenv(float: preservesign,preservesign)
// IEEEF64-PSZF32-SAME: "target-cpu"="gfx803"
-// IEEEF64-PSZF32: #[[$FUNCATTR]] = { {{.*}} "denormal-fp-math-f32"="preserve-sign,preserve-sign"
+// IEEEF64-PSZF32: #[[$FUNCATTR]] = { {{.*}} denormal_fpenv(float: preservesign,preservesign)
// IEEEF64-PSZF32-SAME: "target-cpu"="gfx803"
-// IEEEF64-PSZF32: #[[$WEAK_FUNCATTR]] = { {{.*}} "denormal-fp-math-f32"="preserve-sign,preserve-sign"
+// IEEEF64-PSZF32: #[[$WEAK_FUNCATTR]] = { {{.*}} denormal_fpenv(float: preservesign,preservesign)
// IEEEF64-PSZF32-SAME: "target-cpu"="gfx803"
-// IEEEF32-PSZF64-DYNF32: #[[$KERNELATTR]] = { {{.*}} "denormal-fp-math"="preserve-sign,preserve-sign" "denormal-fp-math-f32"="ieee,ieee" {{.*}} "target-cpu"="gfx803" {{.*}} }
+// IEEEF32-PSZF64-DYNF32: #[[$KERNELATTR]] = { {{.*}} denormal_fpenv(preservesign,preservesign float: ieee,ieee) {{.*}} "target-cpu"="gfx803" {{.*}} }
// implicit check-not
// implicit check-not
-// IEEEF32-PSZF64-DYNFULL: #[[$KERNELATTR]] = { {{.*}} "denormal-fp-math"="preserve-sign,preserve-sign" "denormal-fp-math-f32"="ieee,ieee"
+// IEEEF32-PSZF64-DYNFULL: #[[$KERNELATTR]] = { {{.*}} denormal_fpenv(preservesign,preservesign float: ieee,ieee)
// IEEEF32-PSZF64-DYNFULL-SAME: "target-cpu"="gfx803"
-// IEEEF32-PSZF64-DYNFULL: #[[$FUNCATTR]] = { {{.*}} "denormal-fp-math"="preserve-sign,preserve-sign" "denormal-fp-math-f32"="ieee,ieee"
+// IEEEF32-PSZF64-DYNFULL: #[[$FUNCATTR]] = { {{.*}} denormal_fpenv(preservesign,preservesign float: ieee,ieee)
// IEEEF32-PSZF64-DYNFULL-SAME: "target-cpu"="gfx803"
-// IEEEF32-PSZF64-DYNFULL: #[[$WEAK_FUNCATTR]] = { {{.*}} "denormal-fp-math"="preserve-sign,preserve-sign" "denormal-fp-math-f32"="ieee,ieee"
+// IEEEF32-PSZF64-DYNFULL: #[[$WEAK_FUNCATTR]] = { {{.*}} denormal_fpenv(preservesign,preservesign float: ieee,ieee)
// IEEEF32-PSZF64-DYNFULL-SAME: "target-cpu"="gfx803"
// -mlink-bitcode-file doesn't internalize or propagate attributes.
-// NOINTERNALIZE-IEEEF32-PSZF64-DYNFULL: #[[$KERNELATTR]] = { {{.*}} "denormal-fp-math"="preserve-sign,preserve-sign" "denormal-fp-math-f32"="ieee,ieee" {{.*}} "target-cpu"="gfx803" {{.*}} }
-// NOINTERNALIZE-IEEEF32-PSZF64-DYNFULL: #[[$FUNCATTR]] = { {{.*}} "denormal-fp-math"="dynamic,dynamic" {{.*}} }
-// NOINTERNALIZE-IEEEF32-PSZF64-DYNFULL: #[[$WEAK_FUNCATTR]] = { {{.*}} "denormal-fp-math"="dynamic,dynamic" {{.*}} }
+// NOINTERNALIZE-IEEEF32-PSZF64-DYNFULL: #[[$KERNELATTR]] = { {{.*}} denormal_fpenv(preservesign,preservesign float: ieee,ieee) {{.*}} "target-cpu"="gfx803" {{.*}} }
+// NOINTERNALIZE-IEEEF32-PSZF64-DYNFULL: #[[$FUNCATTR]] = { {{.*}} denormal_fpenv(dynamic,dynamic) {{.*}} }
+// NOINTERNALIZE-IEEEF32-PSZF64-DYNFULL: #[[$WEAK_FUNCATTR]] = { {{.*}} denormal_fpenv(dynamic,dynamic) {{.*}} }
diff --git a/clang/test/CodeGenCUDA/propagate-attributes.cu b/clang/test/CodeGenCUDA/propagate-attributes.cu
index a7e6b09ff97db..40a8c32e53d21 100644
--- a/clang/test/CodeGenCUDA/propagate-attributes.cu
+++ b/clang/test/CodeGenCUDA/propagate-attributes.cu
@@ -53,14 +53,14 @@ __global__ void kernel() { lib_fn(); }
// line.
// Check the attribute list for kernel.
+// NOFTZ-NOT: denormal_fpenv
+
// CHECK: attributes [[kattr]] = {
// CHECK-SAME: convergent
// CHECK-SAME: norecurse
-// FTZ-NOT: "denormal-fp-math"
-// FTZ-SAME: "denormal-fp-math-f32"="preserve-sign,preserve-sign"
-// NOFTZ-NOT: "denormal-fp-math-f32"
+// FTZ-SAME: denormal_fpenv(float: preservesign,preservesign)
// CHECK-SAME: "no-trapping-math"="true"
@@ -70,10 +70,4 @@ __global__ void kernel() { lib_fn(); }
// CHECK-SAME: convergent
// CHECK-NOT: norecurse
-// FTZ-NOT: "denormal-fp-math"
-// NOFTZ-NOT: "denormal-fp-math"
-
-// FTZ-SAME: "denormal-fp-math-f32"="preserve-sign,preserve-sign"
-// NOFTZ-NOT: "denormal-fp-math-f32"
-
// CHECK-SAME: "no-trapping-math"="true"
diff --git a/clang/test/CodeGenOpenCL/amdgpu-enqueue-kernel.cl b/clang/test/CodeGenOpenCL/amdgpu-enqueue-kernel.cl
index 2cbc9787a04b0..8fe3404fed366 100644
--- a/clang/test/CodeGenOpenCL/amdgpu-enqueue-kernel.cl
+++ b/clang/test/CodeGenOpenCL/amdgpu-enqueue-kernel.cl
@@ -69,7 +69,7 @@ kernel void test_target_features_kernel(global int *i) {
// CHECK: @__test_target_features_kernel_block_invoke_kernel.runtime.handle = internal addrspace(1) externally_initialized constant %block.runtime.handle.t.3 zeroinitializer, section ".amdgpu.kernel.runtime.handle"
// CHECK: @llvm.used = appending addrspace(1) global [10 x ptr] [ptr @__test_block_invoke_kernel, ptr addrspacecast (ptr addrspace(1) @__test_block_invoke_kernel.runtime.handle to ptr), ptr @__test_block_invoke_2_kernel, ptr addrspacecast (ptr addrspace(1) @__test_block_invoke_2_kernel.runtime.handle to ptr), ptr @__test_block_invoke_3_kernel, ptr addrspacecast (ptr addrspace(1) @__test_block_invoke_3_kernel.runtime.handle to ptr), ptr @__test_block_invoke_4_kernel, ptr addrspacecast (ptr addrspace(1) @__test_block_invoke_4_kernel.runtime.handle to ptr), ptr @__test_target_features_kernel_block_invoke_kernel, ptr addrspacecast (ptr addrspace(1) @__test_target_features_kernel_block_invoke_kernel.runtime.handle to ptr)], section "llvm.metadata"
//.
-// NOCPU: Function Attrs: convergent noinline norecurse nounwind optnone
+// NOCPU: Function Attrs: convergent noinline norecurse nounwind optnone denormal_fpenv(float: preservesign,preservesign)
// NOCPU-LABEL: define dso_local void @callee(
// NOCPU-SAME: i64 noundef [[ID:%.*]], ptr addrspace(1) noundef [[OUT:%.*]]) #[[ATTR1:[0-9]+]] {
// NOCPU-NEXT: [[ENTRY:.*:]]
@@ -87,7 +87,7 @@ kernel void test_target_features_kernel(global int *i) {
// NOCPU-NEXT: ret void
//
//
-// NOCPU: Function Attrs: convergent noinline norecurse nounwind optnone
+// NOCPU: Function Attrs: convergent noinline norecurse nounwind optnone denormal_fpenv(float: preservesign,prese...
[truncated]
|
llvmbot
added
backend:ARM
backend:AArch64
backend:AMDGPU
backend:PowerPC
backend:X86
clang:codegen
🐧 Linux x64 Test Results
✅ The build succeeded and all tests passed. |
🪟 Windows x64 Test Results
✅ The build succeeded and all tests passed. |
Base automatically changed from
users/arsenm/codegen/remove-targetoptions-denormal-mode
to
main
January 4, 2026 15:54
arsenm
force-pushed
the
users/arsenm/ir/promote-denormal-fp-math-attribute
branch
3 times, most recently
from
b71bf1f to
30bb38b
Compare
Convert "denormal-fp-math" and "denormal-fp-math-f32" into a first class denormal_fpenv attribute. Previously the query for the effective deormal mode involved two string attribute queries with parsing. I'm introducing more uses of this, so it makes sense to convert this to a more efficient encoding. The old representation was also awkward since it was split across two separate attributes. The new encoding just stores the default and float modes as bitfields, largely avoiding the need to consider if the other mode is set. The syntax in the common cases looks like this: `denormal_fpenv(preservesign,preservesign)` `denormal_fpenv(float: preservesign,preservesign)` `denormal_fpenv(dynamic,dynamic float: preservesign,preservesign)` I wasn't sure about reusing the float type name instead of adding a new keyword. It's parsed as a type but only accepts float. I'm also debating switching the name to subnormal to match the current preferred IEEE terminology (also used by nofpclass and other contexts). This has a behavior change when using the command flag debug options to set the denormal mode. The behavior of the flag ignored functions with an explicit attribute set, per the default and f32 version. Now that these are one attribute, the flag logic can't distinguish which of the two components were explicitly set on the function. Only one test appeared to rely on this behavior, so I just avoided using the flags in it. This also does not perform all the code cleanups this enables. In particular the attributor handling could be cleaned up. I also guessed at how to support this in MLIR. I followed MemoryEffects as a reference; it appears bitfields are expanded into arguments to attributes, so the representation there is a bit uglier with the 2 2-element fields flattened into 4 arguments.
arsenm
force-pushed
the
users/arsenm/ir/promote-denormal-fp-math-attribute
branch
from
30bb38b to
e8fe4c1
Compare
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
Convert "denormal-fp-math" and "denormal-fp-math-f32" into a first
class denormal_fpenv attribute. Previously the query for the effective
deormal mode involved two string attribute queries with parsing. I'm
introducing more uses of this, so it makes sense to convert this
to a more efficient encoding. The old representation was also awkward
since it was split across two separate attributes. The new encoding
just stores the default and float modes as bitfields, largely avoiding
the need to consider if the other mode is set.
The syntax in the common cases looks like this:
denormal_fpenv(preservesign,preservesign)denormal_fpenv(float: preservesign,preservesign)denormal_fpenv(dynamic,dynamic float: preservesign,preservesign)I wasn't sure about reusing the float type name instead of adding a
new keyword. It's parsed as a type but only accepts float. I'm also
debating switching the name to subnormal to match the current
preferred IEEE terminology (also used by nofpclass and other
contexts).
This has a behavior change when using the command flag debug
options to set the denormal mode. The behavior of the flag
ignored functions with an explicit attribute set, per
the default and f32 version. Now that these are one attribute,
the flag logic can't distinguish which of the two components
were explicitly set on the function. Only one test appeared to
rely on this behavior, so I just avoided using the flags in it.
This also does not perform all the code cleanups this enables.
In particular the attributor handling could be cleaned up.
I also guessed at how to support this in MLIR. I followed
MemoryEffects as a reference; it appears bitfields are expanded
into arguments to attributes, so the representation there is
a bit uglier with the 2 2-element fields flattened into 4 arguments.