[clang] bc37be1 - LangRef: Add "dynamic" option to "denormal-fp-math"
Matt Arsenault via cfe-commits
cfe-commits at lists.llvm.org
Sat Apr 29 06:30:12 PDT 2023
Author: Matt Arsenault
Date: 2023-04-29T08:44:59-04:00
New Revision: bc37be1855773c1dcf8c6bf577a096a81fd58652
URL: https://github.com/llvm/llvm-project/commit/bc37be1855773c1dcf8c6bf577a096a81fd58652
DIFF: https://github.com/llvm/llvm-project/commit/bc37be1855773c1dcf8c6bf577a096a81fd58652.diff
LOG: LangRef: Add "dynamic" option to "denormal-fp-math"
This is stricter than the default "ieee", and should probably be the
default. This patch leaves the default alone. I can change this in a
future patch.
There are non-reversible transforms I would like to perform which are
legal under IEEE denormal handling, but illegal with flushing zero
behavior. Namely, conversions between llvm.is.fpclass and fcmp with
zeroes.
Under "ieee" handling, it is legal to translate between
llvm.is.fpclass(x, fcZero) and fcmp x, 0.
Under "preserve-sign" handling, it is legal to translate between
llvm.is.fpclass(x, fcSubnormal|fcZero) and fcmp x, 0.
I would like to compile and distribute some math library functions in
a mode where it's callable from code with and without denormals
enabled, which requires not changing the compares with denormals or
zeroes.
If an IEEE function transforms an llvm.is.fpclass call into an fcmp 0,
it is no longer possible to call the function from code with denormals
enabled, or write an optimization to move the function into a denormal
flushing mode. For the original function, if x was a denormal, the
class would evaluate to false. If the function compiled with denormal
handling was converted to or called from a preserve-sign function, the
fcmp now evaluates to true.
This could also be of use for strictfp handling, where code may be
changing the denormal mode.
Alternative name could be "unknown".
Replaces the old AMDGPU custom inlining logic with more conservative
logic which tries to permit inlining for callees with dynamic handling
and avoids inlining other mismatched modes.
Added:
clang/test/CodeGenCUDA/Inputs/ocml-sample.cl
clang/test/CodeGenCUDA/link-builtin-bitcode-denormal-fp-mode.cu
llvm/test/CodeGen/Generic/denormal-fp-math-cl-opt.ll
Modified:
clang/lib/CodeGen/CGCall.cpp
clang/lib/CodeGen/CodeGenAction.cpp
clang/lib/CodeGen/CodeGenModule.h
clang/test/CodeGen/denormalfpmode-f32.c
clang/test/CodeGen/denormalfpmode.c
clang/test/Driver/denormal-fp-math.c
llvm/docs/LangRef.rst
llvm/include/llvm/ADT/FloatingPointMode.h
llvm/include/llvm/Analysis/ConstantFolding.h
llvm/include/llvm/IR/Attributes.td
llvm/include/llvm/IR/Function.h
llvm/lib/Analysis/ConstantFolding.cpp
llvm/lib/CodeGen/CommandFlags.cpp
llvm/lib/CodeGen/SelectionDAG/TargetLowering.cpp
llvm/lib/IR/Attributes.cpp
llvm/lib/IR/Function.cpp
llvm/lib/Target/AMDGPU/SIModeRegisterDefaults.h
llvm/test/CodeGen/X86/sqrt-fastmath.ll
llvm/test/Transforms/Inline/AMDGPU/inline-denormal-fp-math.ll
llvm/test/Transforms/InstSimplify/canonicalize.ll
llvm/test/Transforms/InstSimplify/constant-fold-fp-denormal.ll
llvm/unittests/ADT/FloatingPointMode.cpp
llvm/utils/TableGen/Attributes.cpp
Removed:
################################################################################
diff --git a/clang/lib/CodeGen/CGCall.cpp b/clang/lib/CodeGen/CGCall.cpp
index 63296aeec5b9f..9cdf689b74440 100644
--- a/clang/lib/CodeGen/CGCall.cpp
+++ b/clang/lib/CodeGen/CGCall.cpp
@@ -1829,10 +1829,32 @@ static bool HasStrictReturn(const CodeGenModule &Module, QualType RetTy,
Module.getLangOpts().Sanitize.has(SanitizerKind::Return);
}
-void CodeGenModule::getDefaultFunctionAttributes(StringRef Name,
- bool HasOptnone,
- bool AttrOnCallSite,
- llvm::AttrBuilder &FuncAttrs) {
+/// Add denormal-fp-math and denormal-fp-math-f32 as appropriate for the
+/// requested denormal behavior, accounting for the overriding behavior of the
+/// -f32 case.
+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());
+}
+
+/// Add default attributes to a function, which have merge semantics under
+/// -mlink-builtin-bitcode and should not simply overwrite any existing
+/// attributes in the linked library.
+static void
+addMergableDefaultFunctionAttributes(const CodeGenOptions &CodeGenOpts,
+ llvm::AttrBuilder &FuncAttrs) {
+ addDenormalModeAttrs(CodeGenOpts.FPDenormalMode, CodeGenOpts.FP32DenormalMode,
+ FuncAttrs);
+}
+
+void CodeGenModule::getTrivialDefaultFunctionAttributes(
+ StringRef Name, bool HasOptnone, bool AttrOnCallSite,
+ llvm::AttrBuilder &FuncAttrs) {
// OptimizeNoneAttr takes precedence over -Os or -Oz. No warning needed.
if (!HasOptnone) {
if (CodeGenOpts.OptimizeSize)
@@ -1874,15 +1896,6 @@ void CodeGenModule::getDefaultFunctionAttributes(StringRef Name,
if (CodeGenOpts.NullPointerIsValid)
FuncAttrs.addAttribute(llvm::Attribute::NullPointerIsValid);
- if (CodeGenOpts.FPDenormalMode != llvm::DenormalMode::getIEEE())
- FuncAttrs.addAttribute("denormal-fp-math",
- CodeGenOpts.FPDenormalMode.str());
- if (CodeGenOpts.FP32DenormalMode != CodeGenOpts.FPDenormalMode) {
- FuncAttrs.addAttribute(
- "denormal-fp-math-f32",
- CodeGenOpts.FP32DenormalMode.str());
- }
-
if (LangOpts.getDefaultExceptionMode() == LangOptions::FPE_Ignore)
FuncAttrs.addAttribute("no-trapping-math", "true");
@@ -1984,6 +1997,19 @@ void CodeGenModule::getDefaultFunctionAttributes(StringRef Name,
}
}
+void CodeGenModule::getDefaultFunctionAttributes(StringRef Name,
+ bool HasOptnone,
+ bool AttrOnCallSite,
+ llvm::AttrBuilder &FuncAttrs) {
+ getTrivialDefaultFunctionAttributes(Name, HasOptnone, AttrOnCallSite,
+ FuncAttrs);
+ if (!AttrOnCallSite) {
+ // If we're just getting the default, get the default values for mergeable
+ // attributes.
+ addMergableDefaultFunctionAttributes(CodeGenOpts, FuncAttrs);
+ }
+}
+
void CodeGenModule::addDefaultFunctionDefinitionAttributes(llvm::Function &F) {
llvm::AttrBuilder FuncAttrs(F.getContext());
getDefaultFunctionAttributes(F.getName(), F.hasOptNone(),
@@ -1992,8 +2018,60 @@ void CodeGenModule::addDefaultFunctionDefinitionAttributes(llvm::Function &F) {
F.addFnAttrs(FuncAttrs);
}
+/// Apply default attributes to \p F, accounting for merge semantics of
+/// attributes that should not overwrite existing attributes.
+void CodeGenModule::mergeDefaultFunctionDefinitionAttributes(
+ llvm::Function &F, bool WillInternalize) {
+ llvm::AttrBuilder FuncAttrs(F.getContext());
+ getTrivialDefaultFunctionAttributes(F.getName(), F.hasOptNone(),
+ /*AttrOnCallSite=*/false, FuncAttrs);
+ GetCPUAndFeaturesAttributes(GlobalDecl(), FuncAttrs);
+
+ if (!WillInternalize && F.isInterposable()) {
+ // Do not promote "dynamic" denormal-fp-math to this translation unit's
+ // setting for weak functions that won't be internalized. The user has no
+ // real control for how builtin bitcode is linked, so we shouldn't assume
+ // later copies will use a consistent mode.
+ F.addFnAttrs(FuncAttrs);
+ return;
+ }
+
+ llvm::AttributeMask AttrsToRemove;
+
+ llvm::DenormalMode DenormModeToMerge = F.getDenormalModeRaw();
+ llvm::DenormalMode DenormModeToMergeF32 = F.getDenormalModeF32Raw();
+ llvm::DenormalMode Merged =
+ CodeGenOpts.FPDenormalMode.mergeCalleeMode(DenormModeToMerge);
+ llvm::DenormalMode MergedF32 = CodeGenOpts.FP32DenormalMode;
+
+ if (DenormModeToMergeF32.isValid()) {
+ MergedF32 =
+ 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());
+ }
+
+ if (MergedF32 == llvm::DenormalMode::getDefault()) {
+ AttrsToRemove.addAttribute("denormal-fp-math-f32");
+ } else if (MergedF32 != DenormModeToMergeF32) {
+ // Overwrite existing attribute
+ FuncAttrs.addAttribute("denormal-fp-math-f32",
+ CodeGenOpts.FP32DenormalMode.str());
+ }
+
+ F.removeFnAttrs(AttrsToRemove);
+ addDenormalModeAttrs(Merged, MergedF32, FuncAttrs);
+ F.addFnAttrs(FuncAttrs);
+}
+
void CodeGenModule::addDefaultFunctionDefinitionAttributes(
- llvm::AttrBuilder &attrs) {
+ llvm::AttrBuilder &attrs) {
getDefaultFunctionAttributes(/*function name*/ "", /*optnone*/ false,
/*for call*/ false, attrs);
GetCPUAndFeaturesAttributes(GlobalDecl(), attrs);
diff --git a/clang/lib/CodeGen/CodeGenAction.cpp b/clang/lib/CodeGen/CodeGenAction.cpp
index 17e781181ee26..29adf88acd704 100644
--- a/clang/lib/CodeGen/CodeGenAction.cpp
+++ b/clang/lib/CodeGen/CodeGenAction.cpp
@@ -270,7 +270,8 @@ namespace clang {
// in LLVM IR.
if (F.isIntrinsic())
continue;
- Gen->CGM().addDefaultFunctionDefinitionAttributes(F);
+ Gen->CGM().mergeDefaultFunctionDefinitionAttributes(F,
+ LM.Internalize);
}
CurLinkModule = LM.Module.get();
diff --git a/clang/lib/CodeGen/CodeGenModule.h b/clang/lib/CodeGen/CodeGenModule.h
index 0a0fdc89202f1..e4ef209195275 100644
--- a/clang/lib/CodeGen/CodeGenModule.h
+++ b/clang/lib/CodeGen/CodeGenModule.h
@@ -1272,6 +1272,8 @@ class CodeGenModule : public CodeGenTypeCache {
/// function which relies on particular fast-math attributes for correctness.
/// It's up to you to ensure that this is safe.
void addDefaultFunctionDefinitionAttributes(llvm::Function &F);
+ void mergeDefaultFunctionDefinitionAttributes(llvm::Function &F,
+ bool WillInternalize);
/// Like the overload taking a `Function &`, but intended specifically
/// for frontends that want to build on Clang's target-configuration logic.
@@ -1734,6 +1736,12 @@ class CodeGenModule : public CodeGenTypeCache {
/// function.
void SimplifyPersonality();
+ /// Helper function for getDefaultFunctionAttributes. Builds a set of function
+ /// attributes which can be simply added to a function.
+ void getTrivialDefaultFunctionAttributes(StringRef Name, bool HasOptnone,
+ bool AttrOnCallSite,
+ llvm::AttrBuilder &FuncAttrs);
+
/// Helper function for ConstructAttributeList and
/// addDefaultFunctionDefinitionAttributes. Builds a set of function
/// attributes to add to a function with the given properties.
diff --git a/clang/test/CodeGen/denormalfpmode-f32.c b/clang/test/CodeGen/denormalfpmode-f32.c
index f89202bcc9ed8..2e2306f5a3478 100644
--- a/clang/test/CodeGen/denormalfpmode-f32.c
+++ b/clang/test/CodeGen/denormalfpmode-f32.c
@@ -2,21 +2,30 @@
// RUN: %clang_cc1 -S -fdenormal-fp-math=ieee %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-NONE,CHECK-F32-NONE
// RUN: %clang_cc1 -S -fdenormal-fp-math=preserve-sign %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-PS,CHECK-F32-NONE
// RUN: %clang_cc1 -S -fdenormal-fp-math=positive-zero %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-PZ,CHECK-F32-NONE
+// RUN: %clang_cc1 -S -fdenormal-fp-math=dynamic %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-DYNAMIC,CHECK-F32-NONE
// RUN: %clang_cc1 -S -fdenormal-fp-math-f32=ieee %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-NONE,CHECK-F32-NONE
// RUN: %clang_cc1 -S -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 -S -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 -S -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 -S -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 -S -fdenormal-fp-math-f32=preserve-sign %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-NONE,CHECK-F32-PS
// RUN: %clang_cc1 -S -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 -S -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 -S -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 -S -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 -S -fdenormal-fp-math-f32=positive-zero %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-NONE,CHECK-F32-PZ
+// RUN: %clang_cc1 -S -fdenormal-fp-math-f32=dynamic %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-NONE,CHECK-F32-DYNAMIC
// RUN: %clang_cc1 -S -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 -S -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 -S -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 -S -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 -S -fdenormal-fp-math=dynamic -fdenormal-fp-math-f32=dynamic %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK-ATTR,CHECK-DYNAMIC,CHECK-F32-NONE
+
// CHECK-LABEL: main
@@ -25,11 +34,16 @@
// 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-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-F32-DYNAMIC: "denormal-fp-math-f32"="dynamic,dynamic"
+
int main(void) {
return 0;
}
diff --git a/clang/test/CodeGen/denormalfpmode.c b/clang/test/CodeGen/denormalfpmode.c
index 6213454a675dc..36f25038ce2bc 100644
--- a/clang/test/CodeGen/denormalfpmode.c
+++ b/clang/test/CodeGen/denormalfpmode.c
@@ -1,6 +1,7 @@
// RUN: %clang_cc1 -S -fdenormal-fp-math=ieee %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK-IEEE
// RUN: %clang_cc1 -S -fdenormal-fp-math=preserve-sign %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK-PS
// RUN: %clang_cc1 -S -fdenormal-fp-math=positive-zero %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK-PZ
+// RUN: %clang_cc1 -S -fdenormal-fp-math=dynamic %s -emit-llvm -o - | FileCheck %s --check-prefix=CHECK-DYNAMIC
// CHECK-LABEL: main
@@ -8,6 +9,7 @@
// 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"{{.*}}
int main(void) {
return 0;
diff --git a/clang/test/CodeGenCUDA/Inputs/ocml-sample.cl b/clang/test/CodeGenCUDA/Inputs/ocml-sample.cl
new file mode 100644
index 0000000000000..46bd33a1c5947
--- /dev/null
+++ b/clang/test/CodeGenCUDA/Inputs/ocml-sample.cl
@@ -0,0 +1,18 @@
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+
+half do_f16_stuff(half a, half b, half c) {
+ return __builtin_fmaf16(a, b, c) + 4.0h;
+}
+
+float do_f32_stuff(float a, float b, float c) {
+ return __builtin_fmaf(a, b, c) + 4.0f;
+}
+
+double do_f64_stuff(double a, double b, double c) {
+ return __builtin_fma(a, b, c) + 4.0;
+}
+
+__attribute__((weak))
+float weak_do_f32_stuff(float a, float b, float c) {
+ return c * (a / b);
+}
diff --git a/clang/test/CodeGenCUDA/link-builtin-bitcode-denormal-fp-mode.cu b/clang/test/CodeGenCUDA/link-builtin-bitcode-denormal-fp-mode.cu
new file mode 100644
index 0000000000000..910d2470315f1
--- /dev/null
+++ b/clang/test/CodeGenCUDA/link-builtin-bitcode-denormal-fp-mode.cu
@@ -0,0 +1,157 @@
+// Verify the behavior of the denormal-fp-mode attributes in the way that
+// rocm-device-libs should be built with. The bitcode should be compiled with
+// denormal-fp-math-f32=dynamic, and should be replaced with the denormal mode
+// of the final TU.
+
+// Build the fake device library in the way rocm-device-libs should be built.
+//
+// RUN: %clang_cc1 -x cl -triple amdgcn-amd-amdhsa -fdenormal-fp-math-f32=dynamic \
+// RUN: -mcode-object-version=none -emit-llvm-bc \
+// RUN: %S/Inputs/ocml-sample.cl -o %t.dynamic.f32.bc
+//
+// RUN: %clang_cc1 -x cl -triple amdgcn-amd-amdhsa -fdenormal-fp-math=dynamic \
+// RUN: -mcode-object-version=none -emit-llvm-bc \
+// RUN: %S/Inputs/ocml-sample.cl -o %t.dynamic.full.bc
+
+
+
+// Check the default behavior with no denormal-fp-math arguments.
+// RUN: %clang_cc1 -x hip -triple amdgcn-amd-amdhsa -target-cpu gfx803 -fcuda-is-device \
+// RUN: -mlink-builtin-bitcode %t.dynamic.f32.bc \
+// RUN: -emit-llvm %s -o - | FileCheck -implicit-check-not=denormal-fp-math %s --check-prefixes=CHECK,INTERNALIZE
+
+
+// Check an explicit full ieee request
+// RUN: %clang_cc1 -x hip -triple amdgcn-amd-amdhsa -target-cpu gfx803 -fcuda-is-device \
+// RUN: -fdenormal-fp-math=ieee \
+// RUN: -mlink-builtin-bitcode %t.dynamic.f32.bc \
+// RUN: -emit-llvm %s -o - | FileCheck -implicit-check-not=denormal-fp-math %s --check-prefixes=CHECK,INTERNALIZE
+
+
+// Check explicit f32-only flushing request
+// RUN: %clang_cc1 -x hip -triple amdgcn-amd-amdhsa -target-cpu gfx803 \
+// RUN: -fcuda-is-device -fdenormal-fp-math-f32=preserve-sign \
+// RUN: -mlink-builtin-bitcode %t.dynamic.f32.bc -emit-llvm %s -o - \
+// RUN: | FileCheck -implicit-check-not=denormal-fp-math --enable-var-scope %s --check-prefixes=CHECK,INTERNALIZE,IEEEF64-PSZF32
+
+
+// Check explicit flush all request. Only the f32 component of the library is
+// dynamic, so the linked functions should use IEEE as the base mode and the new
+// functions preserve-sign.
+// RUN: %clang_cc1 -x hip -triple amdgcn-amd-amdhsa -target-cpu gfx803 \
+// RUN: -fcuda-is-device -fdenormal-fp-math=preserve-sign \
+// RUN: -mlink-builtin-bitcode %t.dynamic.f32.bc -emit-llvm %s -o - \
+// RUN: | FileCheck -implicit-check-not=denormal-fp-math --enable-var-scope %s --check-prefixes=CHECK,INTERNALIZE,PSZ
+
+
+// Check explicit f32-only, ieee-other flushing request
+// RUN: %clang_cc1 -x hip -triple amdgcn-amd-amdhsa -target-cpu gfx803 \
+// RUN: -fcuda-is-device -fdenormal-fp-math=ieee -fdenormal-fp-math-f32=preserve-sign \
+// RUN: -mlink-builtin-bitcode %t.dynamic.f32.bc -emit-llvm %s -o - \
+// RUN: | FileCheck -implicit-check-not=denormal-fp-math --enable-var-scope %s --check-prefixes=CHECK,INTERNALIZE,IEEEF64-PSZF32
+
+
+// Check inverse of normal usage. Requesting IEEE f32, with flushed f16/f64
+// RUN: %clang_cc1 -x hip -triple amdgcn-amd-amdhsa -target-cpu gfx803 \
+// RUN: -fcuda-is-device -fdenormal-fp-math=preserve-sign -fdenormal-fp-math-f32=ieee \
+// RUN: -mlink-builtin-bitcode %t.dynamic.f32.bc -emit-llvm %s -o - \
+// RUN: | FileCheck -implicit-check-not=denormal-fp-math --enable-var-scope %s --check-prefixes=CHECK,INTERNALIZE,IEEEF32-PSZF64-DYNF32
+
+
+// Check backwards from the normal usage where both library components can be
+// overridden.
+// RUN: %clang_cc1 -x hip -triple amdgcn-amd-amdhsa -target-cpu gfx803 \
+// RUN: -fcuda-is-device -fdenormal-fp-math=preserve-sign -fdenormal-fp-math-f32=ieee \
+// RUN: -mlink-builtin-bitcode %t.dynamic.full.bc -emit-llvm %s -o - \
+// RUN: | FileCheck -implicit-check-not=denormal-fp-math --enable-var-scope %s --check-prefixes=CHECK,INTERNALIZE,IEEEF32-PSZF64-DYNFULL
+
+
+
+// Check the case where no internalization is performed
+// RUN: %clang_cc1 -x hip -triple amdgcn-amd-amdhsa -target-cpu gfx803 \
+// RUN: -fcuda-is-device -fdenormal-fp-math=preserve-sign -fdenormal-fp-math-f32=ieee \
+// RUN: -mlink-bitcode-file %t.dynamic.full.bc -emit-llvm %s -o - \
+// RUN: | FileCheck -implicit-check-not=denormal-fp-math --enable-var-scope %s --check-prefixes=CHECK,NOINTERNALIZE,NOINTERNALIZE-IEEEF32-PSZF64-DYNFULL
+
+
+
+#define __device__ __attribute__((device))
+#define __global__ __attribute__((global))
+
+typedef _Float16 half;
+
+extern "C" {
+__device__ half do_f16_stuff(half a, half b, half c);
+__device__ float do_f32_stuff(float a, float b, float c);
+
+// Currently all library functions are internalized. Check a weak function in
+// case we ever choose to not internalize these. In that case, the safest thing
+// to do would likely be to preserve the dynamic denormal-fp-math.
+__attribute__((weak)) __device__ float weak_do_f32_stuff(float a, float b, float c);
+__device__ double do_f64_stuff(double a, double b, double c);
+
+
+ // CHECK: kernel_f16({{.*}}) #[[$KERNELATTR:[0-9]+]]
+__global__ void kernel_f16(float* out, float* a, float* b, float* c) {
+ int id = 0;
+ out[id] = do_f16_stuff(a[id], b[id], c[id]);
+}
+
+// CHECK: kernel_f32({{.*}}) #[[$KERNELATTR]]
+__global__ void kernel_f32(float* out, float* a, float* b, float* c) {
+ int id = 0;
+ out[id] = do_f32_stuff(a[id], b[id], c[id]);
+ out[id] += weak_do_f32_stuff(a[id], b[id], c[id]);
+}
+
+// CHECK: kernel_f64({{.*}}) #[[$KERNELATTR]]
+__global__ void kernel_f64(double* out, double* a, double* b, double* c) {
+ int id = 0;
+ out[id] = do_f64_stuff(a[id], b[id], c[id]);
+}
+}
+
+// INTERNALIZE: define internal half @do_f16_stuff({{.*}}) #[[$FUNCATTR:[0-9]+]]
+// INTERNALIZE: define internal float @do_f32_stuff({{.*}}) #[[$FUNCATTR]]
+// INTERNALIZE: define internal double @do_f64_stuff({{.*}}) #[[$FUNCATTR]]
+// INTERNALIZE: define internal float @weak_do_f32_stuff({{.*}}) #[[$WEAK_FUNCATTR:[0-9]+]]
+
+
+// NOINTERNALIZE: define dso_local half @do_f16_stuff({{.*}}) #[[$FUNCATTR:[0-9]+]]
+// NOINTERNALIZE: define dso_local float @do_f32_stuff({{.*}}) #[[$FUNCATTR]]
+// NOINTERNALIZE: define dso_local double @do_f64_stuff({{.*}}) #[[$FUNCATTR]]
+// NOINTERNALIZE: define weak float @weak_do_f32_stuff({{.*}}) #[[$WEAK_FUNCATTR:[0-9]+]]
+
+
+
+// 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
+// lines.
+
+// Default mode relies on the implicit check-not for the denormal-fp-math.
+
+// PSZ: #[[$KERNELATTR]] = { {{.*}} "denormal-fp-math"="preserve-sign,preserve-sign" {{.*}} "target-cpu"="gfx803" {{.*}} }
+// PSZ: #[[$FUNCATTR]] = { {{.*}} "denormal-fp-math-f32"="preserve-sign,preserve-sign" {{.*}} "target-cpu"="gfx803" {{.*}} }
+// PSZ: #[[$WEAK_FUNCATTR]] = { {{.*}} "denormal-fp-math-f32"="preserve-sign,preserve-sign" {{.*}} "target-cpu"="gfx803" {{.*}} }
+
+// FIXME: Should check-not "denormal-fp-math" within the line
+// IEEEF64-PSZF32: #[[$KERNELATTR]] = { {{.*}} "denormal-fp-math-f32"="preserve-sign,preserve-sign" {{.*}} "target-cpu"="gfx803" {{.*}} }
+// IEEEF64-PSZF32: #[[$FUNCATTR]] = { {{.*}} "denormal-fp-math-f32"="preserve-sign,preserve-sign" {{.*}} "target-cpu"="gfx803" {{.*}} }
+// IEEEF64-PSZF32: #[[$WEAK_FUNCATTR]] = { {{.*}} "denormal-fp-math-f32"="preserve-sign,preserve-sign" {{.*}} "target-cpu"="gfx803" {{.*}} }
+
+// IEEEF32-PSZF64-DYNF32: #[[$KERNELATTR]] = { {{.*}} "denormal-fp-math"="preserve-sign,preserve-sign" "denormal-fp-math-f32"="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" {{.*}} "target-cpu"="gfx803" {{.*}} }
+// IEEEF32-PSZF64-DYNFULL: #[[$FUNCATTR]] = { {{.*}} "denormal-fp-math"="preserve-sign,preserve-sign" "denormal-fp-math-f32"="ieee,ieee" {{.*}} "target-cpu"="gfx803" {{.*}} }
+// IEEEF32-PSZF64-DYNFULL: #[[$WEAK_FUNCATTR]] = { {{.*}} "denormal-fp-math"="preserve-sign,preserve-sign" "denormal-fp-math-f32"="ieee,ieee" {{.*}} "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" {{.*}} }
diff --git a/clang/test/Driver/denormal-fp-math.c b/clang/test/Driver/denormal-fp-math.c
index ea4dc8699ecc8..243de3b203fc4 100644
--- a/clang/test/Driver/denormal-fp-math.c
+++ b/clang/test/Driver/denormal-fp-math.c
@@ -1,6 +1,9 @@
// RUN: %clang -### -target arm-unknown-linux-gnu -c %s -fdenormal-fp-math=ieee -v 2>&1 | FileCheck -check-prefix=CHECK-IEEE %s
// RUN: %clang -### -target arm-unknown-linux-gnu -c %s -fdenormal-fp-math=preserve-sign -v 2>&1 | FileCheck -check-prefix=CHECK-PS %s
// RUN: %clang -### -target arm-unknown-linux-gnu -c %s -fdenormal-fp-math=positive-zero -v 2>&1 | FileCheck -check-prefix=CHECK-PZ %s
+
+// RUN: %clang -### -target arm-unknown-linux-gnu -c %s -fdenormal-fp-math=dynamic -v 2>&1 | FileCheck -check-prefix=CHECK-DYNAMIC %s
+
// RUN: %clang -### -target arm-unknown-linux-gnu -c %s -fdenormal-fp-math=ieee -fno-fast-math -v 2>&1 | FileCheck -check-prefix=CHECK-NO-UNSAFE %s
// RUN: %clang -### -target arm-unknown-linux-gnu -c %s -fdenormal-fp-math=ieee -fno-unsafe-math-optimizations -v 2>&1 | FileCheck -check-prefix=CHECK-NO-UNSAFE %s
// RUN: not %clang -target arm-unknown-linux-gnu -c %s -fdenormal-fp-math=foo -v 2>&1 | FileCheck -check-prefix=CHECK-INVALID0 %s
@@ -12,6 +15,9 @@
// CHECK-IEEE-NOT: -fdenormal-fp-math=
// CHECK-PS: "-fdenormal-fp-math=preserve-sign,preserve-sign"
// CHECK-PZ: "-fdenormal-fp-math=positive-zero,positive-zero"
+// CHECK-DYNAMIC: "-fdenormal-fp-math=dynamic,dynamic"
+
+
// CHECK-NO-UNSAFE-NOT: "-fdenormal-fp-math=ieee"
// CHECK-INVALID0: error: invalid value 'foo' in '-fdenormal-fp-math=foo'
// CHECK-INVALID1: error: invalid value 'ieee,foo' in '-fdenormal-fp-math=ieee,foo'
diff --git a/llvm/docs/LangRef.rst b/llvm/docs/LangRef.rst
index 48e658d08e7b3..0d51f7eb6e320 100644
--- a/llvm/docs/LangRef.rst
+++ b/llvm/docs/LangRef.rst
@@ -2233,31 +2233,36 @@ example:
This indicates the denormal (subnormal) handling that may be
assumed for the default floating-point environment. This is a
comma separated pair. The elements may be one of ``"ieee"``,
- ``"preserve-sign"``, or ``"positive-zero"``. The first entry
- indicates the flushing mode for the result of floating point
- operations. The second indicates the handling of denormal inputs
+ ``"preserve-sign"``, ``"positive-zero"``, or ``"dynamic"``. The
+ first entry indicates the flushing mode for the result of floating
+ point operations. The second indicates the handling of denormal inputs
to floating point instructions. For compatibility with older
bitcode, if the second value is omitted, both input and output
modes will assume the same mode.
- If this is attribute is not specified, the default is
- ``"ieee,ieee"``.
+ If this is attribute is not specified, the default is ``"ieee,ieee"``.
If the output mode is ``"preserve-sign"``, or ``"positive-zero"``,
denormal outputs may be flushed to zero by standard floating-point
operations. It is not mandated that flushing to zero occurs, but if
a denormal output is flushed to zero, it must respect the sign
- mode. Not all targets support all modes. While this indicates the
- expected floating point mode the function will be executed with,
- this does not make any attempt to ensure the mode is
- consistent. User or platform code is expected to set the floating
- point mode appropriately before function entry.
-
- If the input mode is ``"preserve-sign"``, or ``"positive-zero"``, a
- floating-point operation must treat any input denormal value as
- zero. In some situations, if an instruction does not respect this
- mode, the input may need to be converted to 0 as if by
- ``@llvm.canonicalize`` during lowering for correctness.
+ mode. Not all targets support all modes.
+
+ If the mode is ``"dynamic"``, the behavior is derived from the
+ dynamic state of the floating-point environment. Transformations
+ which depend on the behavior of denormal values should not be
+ performed.
+
+ While this indicates the expected floating point mode the function
+ will be executed with, this does not make any attempt to ensure
+ the mode is consistent. User or platform code is expected to set
+ the floating point mode appropriately before function entry.
+
+ If the input mode is ``"preserve-sign"``, or ``"positive-zero"``,
+ a floating-point operation must treat any input denormal value as
+ zero. In some situations, if an instruction does not respect this
+ mode, the input may need to be converted to 0 as if by
+ ``@llvm.canonicalize`` during lowering for correctness.
``"denormal-fp-math-f32"``
Same as ``"denormal-fp-math"``, but only controls the behavior of
diff --git a/llvm/include/llvm/ADT/FloatingPointMode.h b/llvm/include/llvm/ADT/FloatingPointMode.h
index a7198babc69bf..61e57094fdbb9 100644
--- a/llvm/include/llvm/ADT/FloatingPointMode.h
+++ b/llvm/include/llvm/ADT/FloatingPointMode.h
@@ -80,7 +80,10 @@ struct DenormalMode {
PreserveSign,
/// Denormals are flushed to positive zero.
- PositiveZero
+ PositiveZero,
+
+ /// Denormals have unknown treatment.
+ Dynamic
};
/// Denormal flushing mode for floating point instruction results in the
@@ -101,6 +104,11 @@ struct DenormalMode {
return DenormalMode(DenormalModeKind::Invalid, DenormalModeKind::Invalid);
}
+ /// Return the assumed default mode for a function without denormal-fp-math.
+ static constexpr DenormalMode getDefault() {
+ return getIEEE();
+ }
+
static constexpr DenormalMode getIEEE() {
return DenormalMode(DenormalModeKind::IEEE, DenormalModeKind::IEEE);
}
@@ -115,6 +123,10 @@ struct DenormalMode {
DenormalModeKind::PositiveZero);
}
+ static constexpr DenormalMode getDynamic() {
+ return DenormalMode(DenormalModeKind::Dynamic, DenormalModeKind::Dynamic);
+ }
+
bool operator==(DenormalMode Other) const {
return Output == Other.Output && Input == Other.Input;
}
@@ -144,6 +156,18 @@ struct DenormalMode {
Output == DenormalModeKind::PositiveZero;
}
+ /// Get the effective denormal mode if the mode if this caller calls into a
+ /// function with \p Callee. This promotes dynamic modes to the mode of the
+ /// caller.
+ DenormalMode mergeCalleeMode(DenormalMode Callee) const {
+ DenormalMode MergedMode = Callee;
+ if (Callee.Input == DenormalMode::Dynamic)
+ MergedMode.Input = Input;
+ if (Callee.Output == DenormalMode::Dynamic)
+ MergedMode.Output = Output;
+ return MergedMode;
+ }
+
inline void print(raw_ostream &OS) const;
inline std::string str() const {
@@ -164,10 +188,11 @@ inline DenormalMode::DenormalModeKind
parseDenormalFPAttributeComponent(StringRef Str) {
// Assume ieee on unspecified attribute.
return StringSwitch<DenormalMode::DenormalModeKind>(Str)
- .Cases("", "ieee", DenormalMode::IEEE)
- .Case("preserve-sign", DenormalMode::PreserveSign)
- .Case("positive-zero", DenormalMode::PositiveZero)
- .Default(DenormalMode::Invalid);
+ .Cases("", "ieee", DenormalMode::IEEE)
+ .Case("preserve-sign", DenormalMode::PreserveSign)
+ .Case("positive-zero", DenormalMode::PositiveZero)
+ .Case("dynamic", DenormalMode::Dynamic)
+ .Default(DenormalMode::Invalid);
}
/// Return the name used for the denormal handling mode used by the the
@@ -180,6 +205,8 @@ inline StringRef denormalModeKindName(DenormalMode::DenormalModeKind Mode) {
return "preserve-sign";
case DenormalMode::PositiveZero:
return "positive-zero";
+ case DenormalMode::Dynamic:
+ return "dynamic";
default:
return "";
}
diff --git a/llvm/include/llvm/Analysis/ConstantFolding.h b/llvm/include/llvm/Analysis/ConstantFolding.h
index 23ec7d6b70ec4..b11981c609ce1 100644
--- a/llvm/include/llvm/Analysis/ConstantFolding.h
+++ b/llvm/include/llvm/Analysis/ConstantFolding.h
@@ -99,6 +99,9 @@ Constant *ConstantFoldFPInstOperands(unsigned Opcode, Constant *LHS,
/// correct sign, otherwise return the original constant. Inputs and outputs to
/// floating point instructions can have their mode set separately, so the
/// direction is also needed.
+///
+/// If the calling function's "denormal-fp-math" input mode is "dynamic" for the
+/// floating-point type, returns nullptr for denormal inputs.
Constant *FlushFPConstant(Constant *Operand, const Instruction *I,
bool IsOutput);
diff --git a/llvm/include/llvm/IR/Attributes.td b/llvm/include/llvm/IR/Attributes.td
index e72054cfd475e..aba1d718f7f72 100644
--- a/llvm/include/llvm/IR/Attributes.td
+++ b/llvm/include/llvm/IR/Attributes.td
@@ -41,6 +41,9 @@ class TypeAttr<string S, list<AttrProperty> P> : Attr<S, P>;
/// StringBool attribute.
class StrBoolAttr<string S> : Attr<S, []>;
+/// Arbitrary string attribute.
+class ComplexStrAttr<string S, list<AttrProperty> P> : Attr<S, P>;
+
/// Target-independent enum attributes.
/// Alignment of parameter (5 bits) stored as log2 of alignment with +1 bias.
@@ -321,6 +324,9 @@ def NoInlineLineTables : StrBoolAttr<"no-inline-line-tables">;
def ProfileSampleAccurate : StrBoolAttr<"profile-sample-accurate">;
def UseSampleProfile : StrBoolAttr<"use-sample-profile">;
+def DenormalFPMath : ComplexStrAttr<"denormal-fp-math", [FnAttr]>;
+def DenormalFPMathF32 : ComplexStrAttr<"denormal-fp-math-f32", [FnAttr]>;
+
class CompatRule<string F> {
// The name of the function called to check the attribute of the caller and
// callee and decide whether inlining should be allowed. The function's
@@ -340,6 +346,8 @@ def : CompatRule<"isEqual<SafeStackAttr>">;
def : CompatRule<"isEqual<ShadowCallStackAttr>">;
def : CompatRule<"isEqual<UseSampleProfileAttr>">;
def : CompatRule<"isEqual<NoProfileAttr>">;
+def : CompatRule<"checkDenormMode">;
+
class MergeRule<string F> {
// The name of the function called to merge the attributes of the caller and
diff --git a/llvm/include/llvm/IR/Function.h b/llvm/include/llvm/IR/Function.h
index 86bdc241a5ffe..8b514528daf38 100644
--- a/llvm/include/llvm/IR/Function.h
+++ b/llvm/include/llvm/IR/Function.h
@@ -654,6 +654,15 @@ class LLVM_EXTERNAL_VISIBILITY Function : public GlobalObject,
/// function.
DenormalMode getDenormalMode(const fltSemantics &FPType) const;
+ /// Return the representational value of "denormal-fp-math". Code interested
+ /// in the semantics of the function should use getDenormalMode instead.
+ DenormalMode getDenormalModeRaw() const;
+
+ /// Return the representational value of "denormal-fp-math-f32". Code
+ /// interested in the semantics of the function should use getDenormalMode
+ /// instead.
+ DenormalMode getDenormalModeF32Raw() const;
+
/// copyAttributesFrom - copy all additional attributes (those not needed to
/// create a Function) from the Function Src to this one.
void copyAttributesFrom(const Function *Src);
diff --git a/llvm/lib/Analysis/ConstantFolding.cpp b/llvm/lib/Analysis/ConstantFolding.cpp
index f085b16c2622c..ce6f334cdc173 100644
--- a/llvm/lib/Analysis/ConstantFolding.cpp
+++ b/llvm/lib/Analysis/ConstantFolding.cpp
@@ -1326,7 +1326,11 @@ Constant *llvm::ConstantFoldCompareInstOperands(
// Flush any denormal constant float input according to denormal handling
// mode.
Ops0 = FlushFPConstant(Ops0, I, /* IsOutput */ false);
+ if (!Ops0)
+ return nullptr;
Ops1 = FlushFPConstant(Ops1, I, /* IsOutput */ false);
+ if (!Ops1)
+ return nullptr;
return ConstantExpr::getCompare(Predicate, Ops0, Ops1);
}
@@ -1361,6 +1365,10 @@ Constant *llvm::FlushFPConstant(Constant *Operand, const Instruction *I,
return Operand;
const APFloat &APF = CFP->getValueAPF();
+ // TODO: Should this canonicalize nans?
+ if (!APF.isDenormal())
+ return Operand;
+
Type *Ty = CFP->getType();
DenormalMode DenormMode =
I->getFunction()->getDenormalMode(Ty->getFltSemantics());
@@ -1369,7 +1377,8 @@ Constant *llvm::FlushFPConstant(Constant *Operand, const Instruction *I,
switch (Mode) {
default:
llvm_unreachable("unknown denormal mode");
- return Operand;
+ case DenormalMode::Dynamic:
+ return nullptr;
case DenormalMode::IEEE:
return Operand;
case DenormalMode::PreserveSign:
@@ -1395,7 +1404,11 @@ Constant *llvm::ConstantFoldFPInstOperands(unsigned Opcode, Constant *LHS,
if (Instruction::isBinaryOp(Opcode)) {
// Flush denormal inputs if needed.
Constant *Op0 = FlushFPConstant(LHS, I, /* IsOutput */ false);
+ if (!Op0)
+ return nullptr;
Constant *Op1 = FlushFPConstant(RHS, I, /* IsOutput */ false);
+ if (!Op1)
+ return nullptr;
// Calculate constant result.
Constant *C = ConstantFoldBinaryOpOperands(Opcode, Op0, Op1, DL);
@@ -1969,13 +1982,26 @@ static Constant *constantFoldCanonicalize(const Type *Ty, const CallBase *CI,
if (Src.isDenormal() && CI->getParent() && CI->getFunction()) {
DenormalMode DenormMode =
CI->getFunction()->getDenormalMode(Src.getSemantics());
+
+ // TODO: Should allow folding for pure IEEE.
if (DenormMode == DenormalMode::getIEEE())
return nullptr;
+ if (DenormMode == DenormalMode::getDynamic())
+ return nullptr;
+
+ // If we know if either input or output is flushed, we can fold.
+ if ((DenormMode.Input == DenormalMode::Dynamic &&
+ DenormMode.Output == DenormalMode::IEEE) ||
+ (DenormMode.Input == DenormalMode::IEEE &&
+ DenormMode.Output == DenormalMode::Dynamic))
+ return nullptr;
+
bool IsPositive =
(!Src.isNegative() || DenormMode.Input == DenormalMode::PositiveZero ||
(DenormMode.Output == DenormalMode::PositiveZero &&
DenormMode.Input == DenormalMode::IEEE));
+
return ConstantFP::get(CI->getContext(),
APFloat::getZero(Src.getSemantics(), !IsPositive));
}
diff --git a/llvm/lib/CodeGen/CommandFlags.cpp b/llvm/lib/CodeGen/CommandFlags.cpp
index 6dafe23955c62..35263eabfc4f9 100644
--- a/llvm/lib/CodeGen/CommandFlags.cpp
+++ b/llvm/lib/CodeGen/CommandFlags.cpp
@@ -241,14 +241,15 @@ codegen::RegisterCodeGenFlags::RegisterCodeGenFlags() {
cl::init(false));
CGBINDOPT(EnableNoTrappingFPMath);
- static const auto DenormFlagEnumOptions =
- cl::values(clEnumValN(DenormalMode::IEEE, "ieee",
- "IEEE 754 denormal numbers"),
- clEnumValN(DenormalMode::PreserveSign, "preserve-sign",
- "the sign of a flushed-to-zero number is preserved "
- "in the sign of 0"),
- clEnumValN(DenormalMode::PositiveZero, "positive-zero",
- "denormals are flushed to positive zero"));
+ static const auto DenormFlagEnumOptions = cl::values(
+ clEnumValN(DenormalMode::IEEE, "ieee", "IEEE 754 denormal numbers"),
+ clEnumValN(DenormalMode::PreserveSign, "preserve-sign",
+ "the sign of a flushed-to-zero number is preserved "
+ "in the sign of 0"),
+ clEnumValN(DenormalMode::PositiveZero, "positive-zero",
+ "denormals are flushed to positive zero"),
+ clEnumValN(DenormalMode::Dynamic, "dynamic",
+ "denormals have unknown treatment"));
// FIXME: Doesn't have way to specify separate input and output modes.
static cl::opt<DenormalMode::DenormalModeKind> DenormalFPMath(
diff --git a/llvm/lib/CodeGen/SelectionDAG/TargetLowering.cpp b/llvm/lib/CodeGen/SelectionDAG/TargetLowering.cpp
index 5a295a7957dae..3580e131d7840 100644
--- a/llvm/lib/CodeGen/SelectionDAG/TargetLowering.cpp
+++ b/llvm/lib/CodeGen/SelectionDAG/TargetLowering.cpp
@@ -6755,20 +6755,23 @@ SDValue TargetLowering::getSqrtInputTest(SDValue Op, SelectionDAG &DAG,
EVT VT = Op.getValueType();
EVT CCVT = getSetCCResultType(DAG.getDataLayout(), *DAG.getContext(), VT);
SDValue FPZero = DAG.getConstantFP(0.0, DL, VT);
+
+ // This is specifically a check for the handling of denormal inputs, not the
+ // result.
+ if (Mode.Input == DenormalMode::PreserveSign ||
+ Mode.Input == DenormalMode::PositiveZero) {
+ // Test = X == 0.0
+ return DAG.getSetCC(DL, CCVT, Op, FPZero, ISD::SETEQ);
+ }
+
// Testing it with denormal inputs to avoid wrong estimate.
- if (Mode.Input == DenormalMode::IEEE) {
- // This is specifically a check for the handling of denormal inputs,
- // not the result.
-
- // Test = fabs(X) < SmallestNormal
- const fltSemantics &FltSem = DAG.EVTToAPFloatSemantics(VT);
- APFloat SmallestNorm = APFloat::getSmallestNormalized(FltSem);
- SDValue NormC = DAG.getConstantFP(SmallestNorm, DL, VT);
- SDValue Fabs = DAG.getNode(ISD::FABS, DL, VT, Op);
- return DAG.getSetCC(DL, CCVT, Fabs, NormC, ISD::SETLT);
- }
- // Test = X == 0.0
- return DAG.getSetCC(DL, CCVT, Op, FPZero, ISD::SETEQ);
+ //
+ // Test = fabs(X) < SmallestNormal
+ const fltSemantics &FltSem = DAG.EVTToAPFloatSemantics(VT);
+ APFloat SmallestNorm = APFloat::getSmallestNormalized(FltSem);
+ SDValue NormC = DAG.getConstantFP(SmallestNorm, DL, VT);
+ SDValue Fabs = DAG.getNode(ISD::FABS, DL, VT, Op);
+ return DAG.getSetCC(DL, CCVT, Fabs, NormC, ISD::SETLT);
}
SDValue TargetLowering::getNegatedExpression(SDValue Op, SelectionDAG &DAG,
diff --git a/llvm/lib/IR/Attributes.cpp b/llvm/lib/IR/Attributes.cpp
index 57944c35836a2..c1c4c0fc81a52 100644
--- a/llvm/lib/IR/Attributes.cpp
+++ b/llvm/lib/IR/Attributes.cpp
@@ -2001,6 +2001,37 @@ AttributeMask AttributeFuncs::getUBImplyingAttributes() {
return AM;
}
+/// Callees with dynamic denormal modes are compatible with any caller mode.
+static bool denormModeCompatible(DenormalMode CallerMode,
+ DenormalMode CalleeMode) {
+ if (CallerMode == CalleeMode || CalleeMode == DenormalMode::getDynamic())
+ return true;
+
+ // If they don't exactly match, it's OK if the mismatched component is
+ // dynamic.
+ if (CalleeMode.Input == CallerMode.Input &&
+ CalleeMode.Output == DenormalMode::Dynamic)
+ return true;
+
+ if (CalleeMode.Output == CallerMode.Output &&
+ CalleeMode.Input == DenormalMode::Dynamic)
+ return true;
+ return false;
+}
+
+static bool checkDenormMode(const Function &Caller, const Function &Callee) {
+ DenormalMode CallerMode = Caller.getDenormalModeRaw();
+ DenormalMode CalleeMode = Callee.getDenormalModeRaw();
+
+ if (denormModeCompatible(CallerMode, CalleeMode)) {
+ DenormalMode CallerModeF32 = Caller.getDenormalModeF32Raw();
+ DenormalMode CalleeModeF32 = Callee.getDenormalModeF32Raw();
+ return denormModeCompatible(CallerModeF32, CalleeModeF32);
+ }
+
+ return false;
+}
+
template<typename AttrClass>
static bool isEqual(const Function &Caller, const Function &Callee) {
return Caller.getFnAttribute(AttrClass::getKind()) ==
diff --git a/llvm/lib/IR/Function.cpp b/llvm/lib/IR/Function.cpp
index bb097e8bec68d..300e0227b66f8 100644
--- a/llvm/lib/IR/Function.cpp
+++ b/llvm/lib/IR/Function.cpp
@@ -702,17 +702,30 @@ void Function::addDereferenceableOrNullParamAttr(unsigned ArgNo,
DenormalMode Function::getDenormalMode(const fltSemantics &FPType) const {
if (&FPType == &APFloat::IEEEsingle()) {
- Attribute Attr = getFnAttribute("denormal-fp-math-f32");
- StringRef Val = Attr.getValueAsString();
- if (!Val.empty())
- return parseDenormalFPAttribute(Val);
-
+ DenormalMode Mode = getDenormalModeF32Raw();
// If the f32 variant of the attribute isn't specified, try to use the
// generic one.
+ if (Mode.isValid())
+ return Mode;
}
+ return getDenormalModeRaw();
+}
+
+DenormalMode Function::getDenormalModeRaw() const {
Attribute Attr = getFnAttribute("denormal-fp-math");
- return parseDenormalFPAttribute(Attr.getValueAsString());
+ StringRef Val = Attr.getValueAsString();
+ return parseDenormalFPAttribute(Val);
+}
+
+DenormalMode Function::getDenormalModeF32Raw() const {
+ Attribute Attr = getFnAttribute("denormal-fp-math-f32");
+ if (Attr.isValid()) {
+ StringRef Val = Attr.getValueAsString();
+ return parseDenormalFPAttribute(Val);
+ }
+
+ return DenormalMode::getInvalid();
}
const std::string &Function::getGC() const {
diff --git a/llvm/lib/Target/AMDGPU/SIModeRegisterDefaults.h b/llvm/lib/Target/AMDGPU/SIModeRegisterDefaults.h
index c834a83047fc0..6693bd5150600 100644
--- a/llvm/lib/Target/AMDGPU/SIModeRegisterDefaults.h
+++ b/llvm/lib/Target/AMDGPU/SIModeRegisterDefaults.h
@@ -34,9 +34,11 @@ struct SIModeRegisterDefaults {
/// and f16/v2f16 instructions.
DenormalMode FP64FP16Denormals;
- SIModeRegisterDefaults()
- : IEEE(true), DX10Clamp(true), FP32Denormals(DenormalMode::getIEEE()),
- FP64FP16Denormals(DenormalMode::getIEEE()) {}
+ SIModeRegisterDefaults() :
+ IEEE(true),
+ DX10Clamp(true),
+ FP32Denormals(DenormalMode::getIEEE()),
+ FP64FP16Denormals(DenormalMode::getIEEE()) {}
SIModeRegisterDefaults(const Function &F);
@@ -84,35 +86,10 @@ struct SIModeRegisterDefaults {
return FP_DENORM_FLUSH_NONE;
}
- /// Returns true if a flag is compatible if it's enabled in the callee, but
- /// disabled in the caller.
- static bool oneWayCompatible(bool CallerMode, bool CalleeMode) {
- return CallerMode == CalleeMode || (!CallerMode && CalleeMode);
- }
-
// FIXME: Inlining should be OK for dx10-clamp, since the caller's mode should
// be able to override.
bool isInlineCompatible(SIModeRegisterDefaults CalleeMode) const {
- if (DX10Clamp != CalleeMode.DX10Clamp)
- return false;
- if (IEEE != CalleeMode.IEEE)
- return false;
-
- // Allow inlining denormals enabled into denormals flushed functions.
- return oneWayCompatible(FP64FP16Denormals.Input !=
- DenormalMode::PreserveSign,
- CalleeMode.FP64FP16Denormals.Input !=
- DenormalMode::PreserveSign) &&
- oneWayCompatible(FP64FP16Denormals.Output !=
- DenormalMode::PreserveSign,
- CalleeMode.FP64FP16Denormals.Output !=
- DenormalMode::PreserveSign) &&
- oneWayCompatible(FP32Denormals.Input != DenormalMode::PreserveSign,
- CalleeMode.FP32Denormals.Input !=
- DenormalMode::PreserveSign) &&
- oneWayCompatible(FP32Denormals.Output != DenormalMode::PreserveSign,
- CalleeMode.FP32Denormals.Output !=
- DenormalMode::PreserveSign);
+ return DX10Clamp == CalleeMode.DX10Clamp && IEEE == CalleeMode.IEEE;
}
};
diff --git a/llvm/test/CodeGen/Generic/denormal-fp-math-cl-opt.ll b/llvm/test/CodeGen/Generic/denormal-fp-math-cl-opt.ll
new file mode 100644
index 0000000000000..7e0922c7d0810
--- /dev/null
+++ b/llvm/test/CodeGen/Generic/denormal-fp-math-cl-opt.ll
@@ -0,0 +1,9 @@
+; RUN: llc -denormal-fp-math=dynamic --denormal-fp-math-f32=preserve-sign -stop-after=finalize-isel < %s | FileCheck %s
+
+; Check that the command line flag annotates the IR with the
+; appropriate attributes.
+
+; CHECK: attributes #0 = { "denormal-fp-math"="dynamic,dynamic" "denormal-fp-math-f32"="preserve-sign,preserve-sign" }
+define float @foo(float %var) {
+ ret float %var
+}
diff --git a/llvm/test/CodeGen/X86/sqrt-fastmath.ll b/llvm/test/CodeGen/X86/sqrt-fastmath.ll
index 333aebc503f58..54ea207ac5dc6 100644
--- a/llvm/test/CodeGen/X86/sqrt-fastmath.ll
+++ b/llvm/test/CodeGen/X86/sqrt-fastmath.ll
@@ -183,8 +183,8 @@ define <4 x float> @sqrt_v4f32_check_denorms(<4 x float> %x) #3 {
ret <4 x float> %call
}
-define <4 x float> @sqrt_v4f32_check_denorms_ninf(<4 x float> %x) #3 {
-; SSE-LABEL: sqrt_v4f32_check_denorms_ninf:
+define <4 x float> @sqrt_v4f32_check_denorms_ieee_ninf(<4 x float> %x) #3 {
+; SSE-LABEL: sqrt_v4f32_check_denorms_ieee_ninf:
; SSE: # %bb.0:
; SSE-NEXT: rsqrtps %xmm0, %xmm1
; SSE-NEXT: movaps %xmm0, %xmm2
@@ -201,7 +201,7 @@ define <4 x float> @sqrt_v4f32_check_denorms_ninf(<4 x float> %x) #3 {
; SSE-NEXT: movaps %xmm1, %xmm0
; SSE-NEXT: retq
;
-; AVX1-LABEL: sqrt_v4f32_check_denorms_ninf:
+; AVX1-LABEL: sqrt_v4f32_check_denorms_ieee_ninf:
; AVX1: # %bb.0:
; AVX1-NEXT: vrsqrtps %xmm0, %xmm1
; AVX1-NEXT: vmulps %xmm1, %xmm0, %xmm2
@@ -215,7 +215,58 @@ define <4 x float> @sqrt_v4f32_check_denorms_ninf(<4 x float> %x) #3 {
; AVX1-NEXT: vandps %xmm1, %xmm0, %xmm0
; AVX1-NEXT: retq
;
-; AVX512-LABEL: sqrt_v4f32_check_denorms_ninf:
+; AVX512-LABEL: sqrt_v4f32_check_denorms_ieee_ninf:
+; AVX512: # %bb.0:
+; AVX512-NEXT: vrsqrtps %xmm0, %xmm1
+; AVX512-NEXT: vmulps %xmm1, %xmm0, %xmm2
+; AVX512-NEXT: vbroadcastss {{.*#+}} xmm3 = [-3.0E+0,-3.0E+0,-3.0E+0,-3.0E+0]
+; AVX512-NEXT: vfmadd231ps {{.*#+}} xmm3 = (xmm2 * xmm1) + xmm3
+; AVX512-NEXT: vbroadcastss {{.*#+}} xmm1 = [-5.0E-1,-5.0E-1,-5.0E-1,-5.0E-1]
+; AVX512-NEXT: vmulps %xmm1, %xmm2, %xmm1
+; AVX512-NEXT: vmulps %xmm3, %xmm1, %xmm1
+; AVX512-NEXT: vbroadcastss {{.*#+}} xmm2 = [NaN,NaN,NaN,NaN]
+; AVX512-NEXT: vandps %xmm2, %xmm0, %xmm0
+; AVX512-NEXT: vbroadcastss {{.*#+}} xmm2 = [1.17549435E-38,1.17549435E-38,1.17549435E-38,1.17549435E-38]
+; AVX512-NEXT: vcmpleps %xmm0, %xmm2, %xmm0
+; AVX512-NEXT: vandps %xmm1, %xmm0, %xmm0
+; AVX512-NEXT: retq
+ %call = tail call ninf afn <4 x float> @llvm.sqrt.v4f32(<4 x float> %x) #2
+ ret <4 x float> %call
+}
+
+define <4 x float> @sqrt_v4f32_check_denorms_dynamic_ninf(<4 x float> %x) #6 {
+; SSE-LABEL: sqrt_v4f32_check_denorms_dynamic_ninf:
+; SSE: # %bb.0:
+; SSE-NEXT: rsqrtps %xmm0, %xmm1
+; SSE-NEXT: movaps %xmm0, %xmm2
+; SSE-NEXT: mulps %xmm1, %xmm2
+; SSE-NEXT: movaps {{.*#+}} xmm3 = [-5.0E-1,-5.0E-1,-5.0E-1,-5.0E-1]
+; SSE-NEXT: mulps %xmm2, %xmm3
+; SSE-NEXT: mulps %xmm1, %xmm2
+; SSE-NEXT: addps {{\.?LCPI[0-9]+_[0-9]+}}(%rip), %xmm2
+; SSE-NEXT: mulps %xmm3, %xmm2
+; SSE-NEXT: andps {{\.?LCPI[0-9]+_[0-9]+}}(%rip), %xmm0
+; SSE-NEXT: movaps {{.*#+}} xmm1 = [1.17549435E-38,1.17549435E-38,1.17549435E-38,1.17549435E-38]
+; SSE-NEXT: cmpleps %xmm0, %xmm1
+; SSE-NEXT: andps %xmm2, %xmm1
+; SSE-NEXT: movaps %xmm1, %xmm0
+; SSE-NEXT: retq
+;
+; AVX1-LABEL: sqrt_v4f32_check_denorms_dynamic_ninf:
+; AVX1: # %bb.0:
+; AVX1-NEXT: vrsqrtps %xmm0, %xmm1
+; AVX1-NEXT: vmulps %xmm1, %xmm0, %xmm2
+; AVX1-NEXT: vmulps {{\.?LCPI[0-9]+_[0-9]+}}(%rip), %xmm2, %xmm3
+; AVX1-NEXT: vmulps %xmm1, %xmm2, %xmm1
+; AVX1-NEXT: vaddps {{\.?LCPI[0-9]+_[0-9]+}}(%rip), %xmm1, %xmm1
+; AVX1-NEXT: vmulps %xmm1, %xmm3, %xmm1
+; AVX1-NEXT: vandps {{\.?LCPI[0-9]+_[0-9]+}}(%rip), %xmm0, %xmm0
+; AVX1-NEXT: vmovaps {{.*#+}} xmm2 = [1.17549435E-38,1.17549435E-38,1.17549435E-38,1.17549435E-38]
+; AVX1-NEXT: vcmpleps %xmm0, %xmm2, %xmm0
+; AVX1-NEXT: vandps %xmm1, %xmm0, %xmm0
+; AVX1-NEXT: retq
+;
+; AVX512-LABEL: sqrt_v4f32_check_denorms_dynamic_ninf:
; AVX512: # %bb.0:
; AVX512-NEXT: vrsqrtps %xmm0, %xmm1
; AVX512-NEXT: vmulps %xmm1, %xmm0, %xmm2
@@ -971,6 +1022,7 @@ define double @sqrt_simplify_before_recip_order(double %x, ptr %p) nounwind {
attributes #0 = { "unsafe-fp-math"="true" "reciprocal-estimates"="!sqrtf,!vec-sqrtf,!divf,!vec-divf" }
attributes #1 = { "unsafe-fp-math"="true" "reciprocal-estimates"="sqrt,vec-sqrt" }
attributes #2 = { nounwind readnone }
-attributes #3 = { "unsafe-fp-math"="true" "reciprocal-estimates"="sqrt,vec-sqrt" "denormal-fp-math"="ieee" }
+attributes #3 = { "unsafe-fp-math"="true" "reciprocal-estimates"="sqrt,vec-sqrt" "denormal-fp-math"="preserve-sign,ieee" }
attributes #4 = { "unsafe-fp-math"="true" "reciprocal-estimates"="sqrt,vec-sqrt" "denormal-fp-math"="ieee,preserve-sign" }
attributes #5 = { "unsafe-fp-math"="true" "reciprocal-estimates"="all:0" }
+attributes #6 = { "unsafe-fp-math"="true" "reciprocal-estimates"="sqrt,vec-sqrt" "denormal-fp-math"="preserve-sign,dynamic" }
diff --git a/llvm/test/Transforms/Inline/AMDGPU/inline-denormal-fp-math.ll b/llvm/test/Transforms/Inline/AMDGPU/inline-denormal-fp-math.ll
index ed63486edb01f..4c802ef57267e 100644
--- a/llvm/test/Transforms/Inline/AMDGPU/inline-denormal-fp-math.ll
+++ b/llvm/test/Transforms/Inline/AMDGPU/inline-denormal-fp-math.ll
@@ -37,9 +37,45 @@ define i32 @func_ieee_psz() #3 {
ret i32 4
}
+define i32 @func_dynamic_dynamic() #4 {
+; CHECK-LABEL: @func_dynamic_dynamic(
+; CHECK-NEXT: ret i32 5
+;
+ ret i32 5
+}
+
+define i32 @func_dynamic_ieee() #5 {
+; CHECK-LABEL: @func_dynamic_ieee(
+; CHECK-NEXT: ret i32 6
+;
+ ret i32 6
+}
+
+define i32 @func_ieee_dynamic() #6 {
+; CHECK-LABEL: @func_ieee_dynamic(
+; CHECK-NEXT: ret i32 7
+;
+ ret i32 7
+}
+
+define i32 @func_psz_dynamic() #7 {
+; CHECK-LABEL: @func_psz_dynamic(
+; CHECK-NEXT: ret i32 8
+;
+ ret i32 8
+}
+
+define i32 @func_dynamic_psz() #8 {
+; CHECK-LABEL: @func_dynamic_psz(
+; CHECK-NEXT: ret i32 9
+;
+ ret i32 9
+}
+
define i32 @call_default_from_psz_psz() #1 {
; CHECK-LABEL: @call_default_from_psz_psz(
-; CHECK-NEXT: ret i32 0
+; CHECK-NEXT: [[CALL:%.*]] = call i32 @func_default()
+; CHECK-NEXT: ret i32 [[CALL]]
;
%call = call i32 @func_default()
ret i32 %call
@@ -55,7 +91,8 @@ define i32 @call_ieee_ieee_from_ieee_ieee() #0 {
define i32 @call_ieee_ieee_from_psz_psz() #1 {
; CHECK-LABEL: @call_ieee_ieee_from_psz_psz(
-; CHECK-NEXT: ret i32 1
+; CHECK-NEXT: [[CALL:%.*]] = call i32 @func_ieee_ieee()
+; CHECK-NEXT: ret i32 [[CALL]]
;
%call = call i32 @func_ieee_ieee()
ret i32 %call
@@ -80,7 +117,8 @@ define i32 @call_psz_psz_from_psz_psz() #1 {
define i32 @call_psz_ieee_from_psz_psz() #1 {
; CHECK-LABEL: @call_psz_ieee_from_psz_psz(
-; CHECK-NEXT: ret i32 3
+; CHECK-NEXT: [[CALL:%.*]] = call i32 @func_psz_ieee()
+; CHECK-NEXT: ret i32 [[CALL]]
;
%call = call i32 @func_psz_ieee()
ret i32 %call
@@ -88,7 +126,8 @@ define i32 @call_psz_ieee_from_psz_psz() #1 {
define i32 @call_ieee_psz_from_psz_psz() #1 {
; CHECK-LABEL: @call_ieee_psz_from_psz_psz(
-; CHECK-NEXT: ret i32 4
+; CHECK-NEXT: [[CALL:%.*]] = call i32 @func_ieee_psz()
+; CHECK-NEXT: ret i32 [[CALL]]
;
%call = call i32 @func_ieee_psz()
ret i32 %call
@@ -114,7 +153,8 @@ define i32 @call_ieee_psz_from_ieee_ieee() #0 {
define i32 @call_ieee_ieee_from_psz_ieee() #2 {
; CHECK-LABEL: @call_ieee_ieee_from_psz_ieee(
-; CHECK-NEXT: ret i32 1
+; CHECK-NEXT: [[CALL:%.*]] = call i32 @func_ieee_ieee()
+; CHECK-NEXT: ret i32 [[CALL]]
;
%call = call i32 @func_ieee_ieee()
ret i32 %call
@@ -148,7 +188,8 @@ define i32 @call_psz_psz_from_psz_ieee() #2 {
define i32 @call_ieee_ieee_from_ieee_psz() #3 {
; CHECK-LABEL: @call_ieee_ieee_from_ieee_psz(
-; CHECK-NEXT: ret i32 1
+; CHECK-NEXT: [[CALL:%.*]] = call i32 @func_ieee_ieee()
+; CHECK-NEXT: ret i32 [[CALL]]
;
%call = call i32 @func_ieee_ieee()
ret i32 %call
@@ -180,7 +221,614 @@ define i32 @call_psz_psz_from_ieee_psz() #3 {
ret i32 %call
}
+define i32 @call_dynamic_dynamic_from_ieee_ieee() #0 {
+; CHECK-LABEL: @call_dynamic_dynamic_from_ieee_ieee(
+; CHECK-NEXT: ret i32 5
+;
+ %call = call i32 @func_dynamic_dynamic()
+ ret i32 %call
+}
+
+define i32 @call_dynamic_ieee_from_ieee_ieee() #0 {
+; CHECK-LABEL: @call_dynamic_ieee_from_ieee_ieee(
+; CHECK-NEXT: ret i32 6
+;
+ %call = call i32 @func_dynamic_ieee()
+ ret i32 %call
+}
+
+define i32 @call_ieee_dynamic_from_ieee_ieee() #0 {
+; CHECK-LABEL: @call_ieee_dynamic_from_ieee_ieee(
+; CHECK-NEXT: ret i32 7
+;
+ %call = call i32 @func_ieee_dynamic()
+ ret i32 %call
+}
+
+define i32 @call_dynamic_psz_from_ieee_ieee() #0 {
+; CHECK-LABEL: @call_dynamic_psz_from_ieee_ieee(
+; CHECK-NEXT: [[CALL:%.*]] = call i32 @func_dynamic_psz()
+; CHECK-NEXT: ret i32 [[CALL]]
+;
+ %call = call i32 @func_dynamic_psz()
+ ret i32 %call
+}
+
+define i32 @call_psz_dynamic_from_ieee_ieee() #0 {
+; CHECK-LABEL: @call_psz_dynamic_from_ieee_ieee(
+; CHECK-NEXT: [[CALL:%.*]] = call i32 @func_psz_dynamic()
+; CHECK-NEXT: ret i32 [[CALL]]
+;
+ %call = call i32 @func_psz_dynamic()
+ ret i32 %call
+}
+
+define i32 @call_dynamic_dynamic_from_psz_psz() #1 {
+; CHECK-LABEL: @call_dynamic_dynamic_from_psz_psz(
+; CHECK-NEXT: ret i32 5
+;
+ %call = call i32 @func_dynamic_dynamic()
+ ret i32 %call
+}
+
+define i32 @call_ieee_dynamic_from_psz_psz() #1 {
+; CHECK-LABEL: @call_ieee_dynamic_from_psz_psz(
+; CHECK-NEXT: [[CALL:%.*]] = call i32 @func_ieee_dynamic()
+; CHECK-NEXT: ret i32 [[CALL]]
+;
+ %call = call i32 @func_ieee_dynamic()
+ ret i32 %call
+}
+
+define i32 @call_dynamic_ieee_from_psz_psz() #1 {
+; CHECK-LABEL: @call_dynamic_ieee_from_psz_psz(
+; CHECK-NEXT: [[CALL:%.*]] = call i32 @func_dynamic_ieee()
+; CHECK-NEXT: ret i32 [[CALL]]
+;
+ %call = call i32 @func_dynamic_ieee()
+ ret i32 %call
+}
+
+define i32 @call_psz_dynamic_from_psz_psz() #1 {
+; CHECK-LABEL: @call_psz_dynamic_from_psz_psz(
+; CHECK-NEXT: ret i32 8
+;
+ %call = call i32 @func_psz_dynamic()
+ ret i32 %call
+}
+
+define i32 @call_dynamic_psz_from_psz_psz() #1 {
+; CHECK-LABEL: @call_dynamic_psz_from_psz_psz(
+; CHECK-NEXT: ret i32 9
+;
+ %call = call i32 @func_dynamic_psz()
+ ret i32 %call
+}
+
+define i32 @call_dynamic_dynamic_from_psz_ieee() #2 {
+; CHECK-LABEL: @call_dynamic_dynamic_from_psz_ieee(
+; CHECK-NEXT: ret i32 5
+;
+ %call = call i32 @func_dynamic_dynamic()
+ ret i32 %call
+}
+
+define i32 @call_dynamic_ieee_from_psz_ieee() #2 {
+; CHECK-LABEL: @call_dynamic_ieee_from_psz_ieee(
+; CHECK-NEXT: ret i32 6
+;
+ %call = call i32 @func_dynamic_ieee()
+ ret i32 %call
+}
+
+define i32 @call_ieee_dynamic_from_psz_ieee() #2 {
+; CHECK-LABEL: @call_ieee_dynamic_from_psz_ieee(
+; CHECK-NEXT: [[CALL:%.*]] = call i32 @func_ieee_dynamic()
+; CHECK-NEXT: ret i32 [[CALL]]
+;
+ %call = call i32 @func_ieee_dynamic()
+ ret i32 %call
+}
+
+define i32 @call_dynamic_psz_from_psz_ieee() #2 {
+; CHECK-LABEL: @call_dynamic_psz_from_psz_ieee(
+; CHECK-NEXT: [[CALL:%.*]] = call i32 @func_dynamic_psz()
+; CHECK-NEXT: ret i32 [[CALL]]
+;
+ %call = call i32 @func_dynamic_psz()
+ ret i32 %call
+}
+
+define i32 @call_psz_dynamic_from_psz_ieee() #2 {
+; CHECK-LABEL: @call_psz_dynamic_from_psz_ieee(
+; CHECK-NEXT: ret i32 8
+;
+ %call = call i32 @func_psz_dynamic()
+ ret i32 %call
+}
+
+define i32 @call_dynamic_dynamic_from_ieee_psz() #3 {
+; CHECK-LABEL: @call_dynamic_dynamic_from_ieee_psz(
+; CHECK-NEXT: ret i32 5
+;
+ %call = call i32 @func_dynamic_dynamic()
+ ret i32 %call
+}
+
+define i32 @call_ieee_dynamic_from_ieee_psz() #3 {
+; CHECK-LABEL: @call_ieee_dynamic_from_ieee_psz(
+; CHECK-NEXT: ret i32 7
+;
+ %call = call i32 @func_ieee_dynamic()
+ ret i32 %call
+}
+
+define i32 @call_dynamic_ieee_from_ieee_psz() #3 {
+; CHECK-LABEL: @call_dynamic_ieee_from_ieee_psz(
+; CHECK-NEXT: [[CALL:%.*]] = call i32 @func_dynamic_ieee()
+; CHECK-NEXT: ret i32 [[CALL]]
+;
+ %call = call i32 @func_dynamic_ieee()
+ ret i32 %call
+}
+
+define i32 @call_dynamic_psz_from_ieee_psz() #3 {
+; CHECK-LABEL: @call_dynamic_psz_from_ieee_psz(
+; CHECK-NEXT: ret i32 9
+;
+ %call = call i32 @func_dynamic_psz()
+ ret i32 %call
+}
+
+define i32 @call_psz_dynamic_from_ieee_psz() #3 {
+; CHECK-LABEL: @call_psz_dynamic_from_ieee_psz(
+; CHECK-NEXT: [[CALL:%.*]] = call i32 @func_psz_dynamic()
+; CHECK-NEXT: ret i32 [[CALL]]
+;
+ %call = call i32 @func_psz_dynamic()
+ ret i32 %call
+}
+
+define i32 @call_dynamic_dynamic_from_dynamic_dynamic() #4 {
+; CHECK-LABEL: @call_dynamic_dynamic_from_dynamic_dynamic(
+; CHECK-NEXT: ret i32 5
+;
+ %call = call i32 @func_dynamic_dynamic()
+ ret i32 %call
+}
+
+define i32 @call_ieee_ieee_from_dynamic_dynamic() #4 {
+; CHECK-LABEL: @call_ieee_ieee_from_dynamic_dynamic(
+; CHECK-NEXT: [[CALL:%.*]] = call i32 @func_ieee_ieee()
+; CHECK-NEXT: ret i32 [[CALL]]
+;
+ %call = call i32 @func_ieee_ieee()
+ ret i32 %call
+}
+
+define i32 @call_ieee_dynamic_from_dynamic_dynamic() #4 {
+; CHECK-LABEL: @call_ieee_dynamic_from_dynamic_dynamic(
+; CHECK-NEXT: [[CALL:%.*]] = call i32 @func_ieee_dynamic()
+; CHECK-NEXT: ret i32 [[CALL]]
+;
+ %call = call i32 @func_ieee_dynamic()
+ ret i32 %call
+}
+
+define i32 @call_dynamic_ieee_from_dynamic_dynamic() #4 {
+; CHECK-LABEL: @call_dynamic_ieee_from_dynamic_dynamic(
+; CHECK-NEXT: [[CALL:%.*]] = call i32 @func_dynamic_ieee()
+; CHECK-NEXT: ret i32 [[CALL]]
+;
+ %call = call i32 @func_dynamic_ieee()
+ ret i32 %call
+}
+
+define i32 @call_psz_dynamic_from_dynamic_dynamic() #4 {
+; CHECK-LABEL: @call_psz_dynamic_from_dynamic_dynamic(
+; CHECK-NEXT: [[CALL:%.*]] = call i32 @func_psz_dynamic()
+; CHECK-NEXT: ret i32 [[CALL]]
+;
+ %call = call i32 @func_psz_dynamic()
+ ret i32 %call
+}
+
+define i32 @call_dynamic_psz_from_dynamic_dynamic() #4 {
+; CHECK-LABEL: @call_dynamic_psz_from_dynamic_dynamic(
+; CHECK-NEXT: [[CALL:%.*]] = call i32 @func_dynamic_psz()
+; CHECK-NEXT: ret i32 [[CALL]]
+;
+ %call = call i32 @func_dynamic_psz()
+ ret i32 %call
+}
+
+define i32 @call_psz_psz_from_dynamic_dynamic() #4 {
+; CHECK-LABEL: @call_psz_psz_from_dynamic_dynamic(
+; CHECK-NEXT: [[CALL:%.*]] = call i32 @func_psz_psz()
+; CHECK-NEXT: ret i32 [[CALL]]
+;
+ %call = call i32 @func_psz_psz()
+ ret i32 %call
+}
+
+define i32 @call_psz_ieee_from_dynamic_dynamic() #4 {
+; CHECK-LABEL: @call_psz_ieee_from_dynamic_dynamic(
+; CHECK-NEXT: [[CALL:%.*]] = call i32 @func_psz_ieee()
+; CHECK-NEXT: ret i32 [[CALL]]
+;
+ %call = call i32 @func_psz_ieee()
+ ret i32 %call
+}
+
+define i32 @call_ieee_psz_from_dynamic_dynamic() #4 {
+; CHECK-LABEL: @call_ieee_psz_from_dynamic_dynamic(
+; CHECK-NEXT: [[CALL:%.*]] = call i32 @func_ieee_psz()
+; CHECK-NEXT: ret i32 [[CALL]]
+;
+ %call = call i32 @func_ieee_psz()
+ ret i32 %call
+}
+
+define i32 @call_ieee_ieee_from_dynamic_ieee() #5 {
+; CHECK-LABEL: @call_ieee_ieee_from_dynamic_ieee(
+; CHECK-NEXT: [[CALL:%.*]] = call i32 @func_ieee_ieee()
+; CHECK-NEXT: ret i32 [[CALL]]
+;
+ %call = call i32 @func_ieee_ieee()
+ ret i32 %call
+}
+
+define i32 @call_ieee_dynamic_from_dynamic_ieee() #5 {
+; CHECK-LABEL: @call_ieee_dynamic_from_dynamic_ieee(
+; CHECK-NEXT: [[CALL:%.*]] = call i32 @func_ieee_dynamic()
+; CHECK-NEXT: ret i32 [[CALL]]
+;
+ %call = call i32 @func_ieee_dynamic()
+ ret i32 %call
+}
+
+define i32 @call_dynamic_ieee_from_dynamic_ieee() #5 {
+; CHECK-LABEL: @call_dynamic_ieee_from_dynamic_ieee(
+; CHECK-NEXT: ret i32 6
+;
+ %call = call i32 @func_dynamic_ieee()
+ ret i32 %call
+}
+
+define i32 @call_psz_dynamic_from_dynamic_ieee() #5 {
+; CHECK-LABEL: @call_psz_dynamic_from_dynamic_ieee(
+; CHECK-NEXT: [[CALL:%.*]] = call i32 @func_psz_dynamic()
+; CHECK-NEXT: ret i32 [[CALL]]
+;
+ %call = call i32 @func_psz_dynamic()
+ ret i32 %call
+}
+
+define i32 @call_dynamic_psz_from_dynamic_ieee() #5 {
+; CHECK-LABEL: @call_dynamic_psz_from_dynamic_ieee(
+; CHECK-NEXT: [[CALL:%.*]] = call i32 @func_dynamic_psz()
+; CHECK-NEXT: ret i32 [[CALL]]
+;
+ %call = call i32 @func_dynamic_psz()
+ ret i32 %call
+}
+
+define i32 @call_psz_psz_from_dynamic_ieee() #5 {
+; CHECK-LABEL: @call_psz_psz_from_dynamic_ieee(
+; CHECK-NEXT: [[CALL:%.*]] = call i32 @func_psz_psz()
+; CHECK-NEXT: ret i32 [[CALL]]
+;
+ %call = call i32 @func_psz_psz()
+ ret i32 %call
+}
+
+define i32 @call_psz_ieee_from_dynamic_ieee() #5 {
+; CHECK-LABEL: @call_psz_ieee_from_dynamic_ieee(
+; CHECK-NEXT: [[CALL:%.*]] = call i32 @func_psz_ieee()
+; CHECK-NEXT: ret i32 [[CALL]]
+;
+ %call = call i32 @func_psz_ieee()
+ ret i32 %call
+}
+
+define i32 @call_ieee_psz_from_dynamic_ieee() #5 {
+; CHECK-LABEL: @call_ieee_psz_from_dynamic_ieee(
+; CHECK-NEXT: [[CALL:%.*]] = call i32 @func_ieee_psz()
+; CHECK-NEXT: ret i32 [[CALL]]
+;
+ %call = call i32 @func_ieee_psz()
+ ret i32 %call
+}
+
+define i32 @call_ieee_ieee_from_ieee_dynamic() #6 {
+; CHECK-LABEL: @call_ieee_ieee_from_ieee_dynamic(
+; CHECK-NEXT: [[CALL:%.*]] = call i32 @func_ieee_ieee()
+; CHECK-NEXT: ret i32 [[CALL]]
+;
+ %call = call i32 @func_ieee_ieee()
+ ret i32 %call
+}
+
+define i32 @call_ieee_dynamic_from_ieee_dynamic() #6 {
+; CHECK-LABEL: @call_ieee_dynamic_from_ieee_dynamic(
+; CHECK-NEXT: ret i32 7
+;
+ %call = call i32 @func_ieee_dynamic()
+ ret i32 %call
+}
+
+define i32 @call_dynamic_ieee_from_ieee_dynamic() #6 {
+; CHECK-LABEL: @call_dynamic_ieee_from_ieee_dynamic(
+; CHECK-NEXT: [[CALL:%.*]] = call i32 @func_dynamic_ieee()
+; CHECK-NEXT: ret i32 [[CALL]]
+;
+ %call = call i32 @func_dynamic_ieee()
+ ret i32 %call
+}
+
+define i32 @call_dynamic_dynamic_from_ieee_dynamic() #6 {
+; CHECK-LABEL: @call_dynamic_dynamic_from_ieee_dynamic(
+; CHECK-NEXT: ret i32 5
+;
+ %call = call i32 @func_dynamic_dynamic()
+ ret i32 %call
+}
+
+define i32 @call_dynamic_psz_from_ieee_dynamic() #6 {
+; CHECK-LABEL: @call_dynamic_psz_from_ieee_dynamic(
+; CHECK-NEXT: [[CALL:%.*]] = call i32 @func_dynamic_psz()
+; CHECK-NEXT: ret i32 [[CALL]]
+;
+ %call = call i32 @func_dynamic_psz()
+ ret i32 %call
+}
+
+define i32 @call_psz_dynamic_from_ieee_dynamic() #6 {
+; CHECK-LABEL: @call_psz_dynamic_from_ieee_dynamic(
+; CHECK-NEXT: [[CALL:%.*]] = call i32 @func_psz_dynamic()
+; CHECK-NEXT: ret i32 [[CALL]]
+;
+ %call = call i32 @func_psz_dynamic()
+ ret i32 %call
+}
+
+define i32 @call_psz_psz_from_ieee_dynamic() #6 {
+; CHECK-LABEL: @call_psz_psz_from_ieee_dynamic(
+; CHECK-NEXT: [[CALL:%.*]] = call i32 @func_psz_psz()
+; CHECK-NEXT: ret i32 [[CALL]]
+;
+ %call = call i32 @func_psz_psz()
+ ret i32 %call
+}
+
+define i32 @call_psz_ieee_from_ieee_dynamic() #6 {
+; CHECK-LABEL: @call_psz_ieee_from_ieee_dynamic(
+; CHECK-NEXT: [[CALL:%.*]] = call i32 @func_psz_ieee()
+; CHECK-NEXT: ret i32 [[CALL]]
+;
+ %call = call i32 @func_psz_ieee()
+ ret i32 %call
+}
+
+define i32 @call_ieee_psz_from_ieee_dynamic() #6 {
+; CHECK-LABEL: @call_ieee_psz_from_ieee_dynamic(
+; CHECK-NEXT: [[CALL:%.*]] = call i32 @func_ieee_psz()
+; CHECK-NEXT: ret i32 [[CALL]]
+;
+ %call = call i32 @func_ieee_psz()
+ ret i32 %call
+}
+
+define i32 @call_ieee_ieee_from_psz_dynamic() #7 {
+; CHECK-LABEL: @call_ieee_ieee_from_psz_dynamic(
+; CHECK-NEXT: [[CALL:%.*]] = call i32 @func_ieee_ieee()
+; CHECK-NEXT: ret i32 [[CALL]]
+;
+ %call = call i32 @func_ieee_ieee()
+ ret i32 %call
+}
+
+define i32 @call_ieee_dynamic_from_psz_dynamic() #7 {
+; CHECK-LABEL: @call_ieee_dynamic_from_psz_dynamic(
+; CHECK-NEXT: [[CALL:%.*]] = call i32 @func_ieee_dynamic()
+; CHECK-NEXT: ret i32 [[CALL]]
+;
+ %call = call i32 @func_ieee_dynamic()
+ ret i32 %call
+}
+
+define i32 @call_dynamic_ieee_from_psz_dynamic() #7 {
+; CHECK-LABEL: @call_dynamic_ieee_from_psz_dynamic(
+; CHECK-NEXT: [[CALL:%.*]] = call i32 @func_dynamic_ieee()
+; CHECK-NEXT: ret i32 [[CALL]]
+;
+ %call = call i32 @func_dynamic_ieee()
+ ret i32 %call
+}
+
+define i32 @call_dynamic_psz_from_psz_dynamic() #7 {
+; CHECK-LABEL: @call_dynamic_psz_from_psz_dynamic(
+; CHECK-NEXT: [[CALL:%.*]] = call i32 @func_dynamic_psz()
+; CHECK-NEXT: ret i32 [[CALL]]
+;
+ %call = call i32 @func_dynamic_psz()
+ ret i32 %call
+}
+
+define i32 @call_psz_dynamic_from_psz_dynamic() #7 {
+; CHECK-LABEL: @call_psz_dynamic_from_psz_dynamic(
+; CHECK-NEXT: ret i32 8
+;
+ %call = call i32 @func_psz_dynamic()
+ ret i32 %call
+}
+
+define i32 @call_dynamic_dynamic_from_psz_dynamic() #7 {
+; CHECK-LABEL: @call_dynamic_dynamic_from_psz_dynamic(
+; CHECK-NEXT: ret i32 5
+;
+ %call = call i32 @func_dynamic_dynamic()
+ ret i32 %call
+}
+
+define i32 @call_psz_ieee_from_psz_dynamic() #7 {
+; CHECK-LABEL: @call_psz_ieee_from_psz_dynamic(
+; CHECK-NEXT: [[CALL:%.*]] = call i32 @func_psz_ieee()
+; CHECK-NEXT: ret i32 [[CALL]]
+;
+ %call = call i32 @func_psz_ieee()
+ ret i32 %call
+}
+
+define i32 @call_ieee_psz_from_psz_dynamic() #7 {
+; CHECK-LABEL: @call_ieee_psz_from_psz_dynamic(
+; CHECK-NEXT: [[CALL:%.*]] = call i32 @func_ieee_psz()
+; CHECK-NEXT: ret i32 [[CALL]]
+;
+ %call = call i32 @func_ieee_psz()
+ ret i32 %call
+}
+
+define i32 @call_psz_psz_from_psz_dynamic() #7 {
+; CHECK-LABEL: @call_psz_psz_from_psz_dynamic(
+; CHECK-NEXT: [[CALL:%.*]] = call i32 @func_psz_psz()
+; CHECK-NEXT: ret i32 [[CALL]]
+;
+ %call = call i32 @func_psz_psz()
+ ret i32 %call
+}
+
+define i32 @call_ieee_ieee_from_dynamic_psz() #8 {
+; CHECK-LABEL: @call_ieee_ieee_from_dynamic_psz(
+; CHECK-NEXT: [[CALL:%.*]] = call i32 @func_ieee_ieee()
+; CHECK-NEXT: ret i32 [[CALL]]
+;
+ %call = call i32 @func_ieee_ieee()
+ ret i32 %call
+}
+
+define i32 @call_ieee_psz_from_dynamic_psz() #8 {
+; CHECK-LABEL: @call_ieee_psz_from_dynamic_psz(
+; CHECK-NEXT: [[CALL:%.*]] = call i32 @func_ieee_psz()
+; CHECK-NEXT: ret i32 [[CALL]]
+;
+ %call = call i32 @func_ieee_psz()
+ ret i32 %call
+}
+
+define i32 @call_psz_ieee_from_dynamic_psz() #8 {
+; CHECK-LABEL: @call_psz_ieee_from_dynamic_psz(
+; CHECK-NEXT: [[CALL:%.*]] = call i32 @func_psz_ieee()
+; CHECK-NEXT: ret i32 [[CALL]]
+;
+ %call = call i32 @func_psz_ieee()
+ ret i32 %call
+}
+
+define i32 @call_dynamic_dynamic_from_dynamic_psz() #8 {
+; CHECK-LABEL: @call_dynamic_dynamic_from_dynamic_psz(
+; CHECK-NEXT: ret i32 5
+;
+ %call = call i32 @func_dynamic_dynamic()
+ ret i32 %call
+}
+
+define i32 @call_dynamic_psz_from_dynamic_psz() #8 {
+; CHECK-LABEL: @call_dynamic_psz_from_dynamic_psz(
+; CHECK-NEXT: ret i32 9
+;
+ %call = call i32 @func_dynamic_psz()
+ ret i32 %call
+}
+
+define i32 @call_psz_dynamic_from_dynamic_psz() #8 {
+; CHECK-LABEL: @call_psz_dynamic_from_dynamic_psz(
+; CHECK-NEXT: [[CALL:%.*]] = call i32 @func_psz_dynamic()
+; CHECK-NEXT: ret i32 [[CALL]]
+;
+ %call = call i32 @func_psz_dynamic()
+ ret i32 %call
+}
+
+define i32 @call_dynamic_ieee_from_dynamic_psz() #8 {
+; CHECK-LABEL: @call_dynamic_ieee_from_dynamic_psz(
+; CHECK-NEXT: [[CALL:%.*]] = call i32 @func_dynamic_ieee()
+; CHECK-NEXT: ret i32 [[CALL]]
+;
+ %call = call i32 @func_dynamic_ieee()
+ ret i32 %call
+}
+
+define i32 @call_ieee_dynamic_from_dynamic_psz() #8 {
+; CHECK-LABEL: @call_ieee_dynamic_from_dynamic_psz(
+; CHECK-NEXT: [[CALL:%.*]] = call i32 @func_ieee_dynamic()
+; CHECK-NEXT: ret i32 [[CALL]]
+;
+ %call = call i32 @func_ieee_dynamic()
+ ret i32 %call
+}
+
+define i32 @call_psz_psz_from_dynamic_psz() #8 {
+; CHECK-LABEL: @call_psz_psz_from_dynamic_psz(
+; CHECK-NEXT: [[CALL:%.*]] = call i32 @func_psz_psz()
+; CHECK-NEXT: ret i32 [[CALL]]
+;
+ %call = call i32 @func_psz_psz()
+ ret i32 %call
+}
+
+; --------------------------------------------------------------------
+; denormal-fp-math-f32
+; --------------------------------------------------------------------
+
+define i32 @func_dynamic_dynamic_f32() #9 {
+; CHECK-LABEL: @func_dynamic_dynamic_f32(
+; CHECK-NEXT: ret i32 10
+;
+ ret i32 10
+}
+
+define i32 @func_psz_psz_f32() #10 {
+; CHECK-LABEL: @func_psz_psz_f32(
+; CHECK-NEXT: ret i32 11
+;
+ ret i32 11
+}
+
+define i32 @call_dynamic_dynamic_from_psz_psz_f32() #10 {
+; CHECK-LABEL: @call_dynamic_dynamic_from_psz_psz_f32(
+; CHECK-NEXT: ret i32 10
+;
+ %result = call i32 @func_dynamic_dynamic_f32()
+ ret i32 %result
+}
+
+define i32 @call_psz_psz_from_psz_psz_f32() #10 {
+; CHECK-LABEL: @call_psz_psz_from_psz_psz_f32(
+; CHECK-NEXT: ret i32 10
+;
+ %result = call i32 @func_dynamic_dynamic_f32()
+ ret i32 %result
+}
+
+define i32 @call_psz_psz_from_ieee_ieee_f32() #11 {
+; CHECK-LABEL: @call_psz_psz_from_ieee_ieee_f32(
+; CHECK-NEXT: [[RESULT:%.*]] = call i32 @func_psz_psz_f32()
+; CHECK-NEXT: ret i32 [[RESULT]]
+;
+ %result = call i32 @func_psz_psz_f32()
+ ret i32 %result
+}
+
attributes #0 = { "denormal-fp-math"="ieee,ieee" }
attributes #1 = { "denormal-fp-math"="preserve-sign,preserve-sign" }
attributes #2 = { "denormal-fp-math"="preserve-sign,ieee" }
attributes #3 = { "denormal-fp-math"="ieee,preserve-sign" }
+attributes #4 = { "denormal-fp-math"="dynamic,dynamic" }
+attributes #5 = { "denormal-fp-math"="dynamic,ieee" }
+attributes #6 = { "denormal-fp-math"="ieee,dynamic" }
+attributes #7 = { "denormal-fp-math"="preserve-sign,dynamic" }
+attributes #8 = { "denormal-fp-math"="dynamic,preserve-sign" }
+attributes #9 = { "denormal-fp-math-f32"="dynamic,dynamic" }
+attributes #10 = { "denormal-fp-math-f32"="preserve-sign,preserve-sign" }
+attributes #11 = { "denormal-fp-math-f32"="ieee,ieee" "denormal-fp-math"="preserve-sign,preserve-sign" }
diff --git a/llvm/test/Transforms/InstSimplify/canonicalize.ll b/llvm/test/Transforms/InstSimplify/canonicalize.ll
index f3d3ce9c8ad0c..96b23acdf1c73 100644
--- a/llvm/test/Transforms/InstSimplify/canonicalize.ll
+++ b/llvm/test/Transforms/InstSimplify/canonicalize.ll
@@ -146,6 +146,100 @@ define float @canonicalize_neg_denorm_positive_zero_input() "denormal-fp-math"="
ret float %ret
}
+define float @canonicalize_pos_denorm_dynamic_dynamic() "denormal-fp-math"="dynamic,dynamic" {
+; CHECK-LABEL: @canonicalize_pos_denorm_dynamic_dynamic(
+; CHECK-NEXT: [[RET:%.*]] = call float @llvm.canonicalize.f32(float 0x380FFFFFC0000000)
+; CHECK-NEXT: ret float [[RET]]
+;
+ %ret = call float @llvm.canonicalize.f32(float bitcast (i32 8388607 to float))
+ ret float %ret
+}
+
+define float @canonicalize_neg_denorm_dynamic_dynamic() "denormal-fp-math"="dynamic,dynamic" {
+; CHECK-LABEL: @canonicalize_neg_denorm_dynamic_dynamic(
+; CHECK-NEXT: [[RET:%.*]] = call float @llvm.canonicalize.f32(float 0xB80FFFFFC0000000)
+; CHECK-NEXT: ret float [[RET]]
+;
+ %ret = call float @llvm.canonicalize.f32(float bitcast (i32 -2139095041 to float))
+ ret float %ret
+}
+
+; Dynamic output - cannot flush
+define float @canonicalize_pos_denorm_dynamic_output() "denormal-fp-math"="dynamic,ieee" {
+; CHECK-LABEL: @canonicalize_pos_denorm_dynamic_output(
+; CHECK-NEXT: [[RET:%.*]] = call float @llvm.canonicalize.f32(float 0x380FFFFFC0000000)
+; CHECK-NEXT: ret float [[RET]]
+;
+ %ret = call float @llvm.canonicalize.f32(float bitcast (i32 8388607 to float))
+ ret float %ret
+}
+
+; Dynamic output - cannot flush
+define float @canonicalize_neg_denorm_dynamic_output() "denormal-fp-math"="dynamic,ieee" {
+; CHECK-LABEL: @canonicalize_neg_denorm_dynamic_output(
+; CHECK-NEXT: [[RET:%.*]] = call float @llvm.canonicalize.f32(float 0xB80FFFFFC0000000)
+; CHECK-NEXT: ret float [[RET]]
+;
+ %ret = call float @llvm.canonicalize.f32(float bitcast (i32 -2139095041 to float))
+ ret float %ret
+}
+
+; Dynamic input - cannot flush
+define float @canonicalize_pos_denorm_dynamic_input() "denormal-fp-math"="ieee,dynamic" {
+; CHECK-LABEL: @canonicalize_pos_denorm_dynamic_input(
+; CHECK-NEXT: [[RET:%.*]] = call float @llvm.canonicalize.f32(float 0x380FFFFFC0000000)
+; CHECK-NEXT: ret float [[RET]]
+;
+ %ret = call float @llvm.canonicalize.f32(float bitcast (i32 8388607 to float))
+ ret float %ret
+}
+
+; Dynamic input - cannot flush
+define float @canonicalize_neg_denorm_dynamic_input() "denormal-fp-math"="ieee,dynamic" {
+; CHECK-LABEL: @canonicalize_neg_denorm_dynamic_input(
+; CHECK-NEXT: [[RET:%.*]] = call float @llvm.canonicalize.f32(float 0xB80FFFFFC0000000)
+; CHECK-NEXT: ret float [[RET]]
+;
+ %ret = call float @llvm.canonicalize.f32(float bitcast (i32 -2139095041 to float))
+ ret float %ret
+}
+
+; Input is flushed, can fold
+define float @canonicalize_pos_denorm_dynamic_output_preserve_sign_input() "denormal-fp-math"="dynamic,preserve-sign" {
+; CHECK-LABEL: @canonicalize_pos_denorm_dynamic_output_preserve_sign_input(
+; CHECK-NEXT: ret float 0.000000e+00
+;
+ %ret = call float @llvm.canonicalize.f32(float bitcast (i32 8388607 to float))
+ ret float %ret
+}
+
+; Input is flushed, can fold
+define float @canonicalize_neg_denorm_dynamic_output_preserve_sign_input() "denormal-fp-math"="dynamic,preserve-sign" {
+; CHECK-LABEL: @canonicalize_neg_denorm_dynamic_output_preserve_sign_input(
+; CHECK-NEXT: ret float -0.000000e+00
+;
+ %ret = call float @llvm.canonicalize.f32(float bitcast (i32 -2139095041 to float))
+ ret float %ret
+}
+
+; Output is known flushed, can fold
+define float @canonicalize_pos_preserve_sign_output_denorm_dynamic_input() "denormal-fp-math"="preserve-sign,dynamic" {
+; CHECK-LABEL: @canonicalize_pos_preserve_sign_output_denorm_dynamic_input(
+; CHECK-NEXT: ret float 0.000000e+00
+;
+ %ret = call float @llvm.canonicalize.f32(float bitcast (i32 8388607 to float))
+ ret float %ret
+}
+
+; Output is known flushed, can fold
+define float @canonicalize_neg_denorm_preserve_sign_output_dynamic_input() "denormal-fp-math"="preserve-sign,dynamic" {
+; CHECK-LABEL: @canonicalize_neg_denorm_preserve_sign_output_dynamic_input(
+; CHECK-NEXT: ret float -0.000000e+00
+;
+ %ret = call float @llvm.canonicalize.f32(float bitcast (i32 -2139095041 to float))
+ ret float %ret
+}
+
define float @canonicalize_inf() {
; CHECK-LABEL: @canonicalize_inf(
; CHECK-NEXT: ret float 0x7FF0000000000000
diff --git a/llvm/test/Transforms/InstSimplify/constant-fold-fp-denormal.ll b/llvm/test/Transforms/InstSimplify/constant-fold-fp-denormal.ll
index 09159a14c6f25..e3ff9416abdd1 100644
--- a/llvm/test/Transforms/InstSimplify/constant-fold-fp-denormal.ll
+++ b/llvm/test/Transforms/InstSimplify/constant-fold-fp-denormal.ll
@@ -1100,6 +1100,109 @@ entry:
ret i1 %cmp
}
+; ============================================================================ ;
+; dynamic mode tests
+; ============================================================================ ;
+
+define float @test_float_fadd_dynamic_ieee() #9 {
+; CHECK-LABEL: @test_float_fadd_dynamic_ieee(
+; CHECK-NEXT: [[RESULT:%.*]] = fadd float 0xB810000000000000, 0x3800000000000000
+; CHECK-NEXT: ret float [[RESULT]]
+;
+ %result = fadd float 0xB810000000000000, 0x3800000000000000
+ ret float %result
+}
+
+define float @test_float_fadd_ieee_dynamic() #10 {
+; CHECK-LABEL: @test_float_fadd_ieee_dynamic(
+; CHECK-NEXT: [[RESULT:%.*]] = fadd float 0xB810000000000000, 0x3800000000000000
+; CHECK-NEXT: ret float [[RESULT]]
+;
+ %result = fadd float 0xB810000000000000, 0x3800000000000000
+ ret float %result
+}
+
+define float @test_float_fadd_dynamic_dynamic() #11 {
+; CHECK-LABEL: @test_float_fadd_dynamic_dynamic(
+; CHECK-NEXT: [[RESULT:%.*]] = fadd float 0xB810000000000000, 0x3800000000000000
+; CHECK-NEXT: ret float [[RESULT]]
+;
+ %result = fadd float 0xB810000000000000, 0x3800000000000000
+ ret float %result
+}
+
+; Check for failed to fold on each operand
+define float @test_float_fadd_dynamic_dynamic_commute() #11 {
+; CHECK-LABEL: @test_float_fadd_dynamic_dynamic_commute(
+; CHECK-NEXT: [[RESULT:%.*]] = fadd float 0x3800000000000000, 0xB810000000000000
+; CHECK-NEXT: ret float [[RESULT]]
+;
+ %result = fadd float 0x3800000000000000, 0xB810000000000000
+ ret float %result
+}
+
+define i1 @fcmp_double_dynamic_ieee() #9 {
+; CHECK-LABEL: @fcmp_double_dynamic_ieee(
+; CHECK-NEXT: ret i1 true
+;
+ %cmp = fcmp une double 0x0008000000000000, 0x0
+ ret i1 %cmp
+}
+
+define i1 @fcmp_double_ieee_dynamic() #10 {
+; CHECK-LABEL: @fcmp_double_ieee_dynamic(
+; CHECK-NEXT: [[CMP:%.*]] = fcmp une double 0x8000000000000, 0.000000e+00
+; CHECK-NEXT: ret i1 [[CMP]]
+;
+ %cmp = fcmp une double 0x0008000000000000, 0x0
+ ret i1 %cmp
+}
+
+define i1 @fcmp_double_dynamic_dynamic() #11 {
+; CHECK-LABEL: @fcmp_double_dynamic_dynamic(
+; CHECK-NEXT: [[CMP:%.*]] = fcmp une double 0x8000000000000, 0.000000e+00
+; CHECK-NEXT: ret i1 [[CMP]]
+;
+ %cmp = fcmp une double 0x0008000000000000, 0x0
+ ret i1 %cmp
+}
+
+define i1 @fcmp_double_dynamic_dynamic_commute() #11 {
+; CHECK-LABEL: @fcmp_double_dynamic_dynamic_commute(
+; CHECK-NEXT: [[CMP:%.*]] = fcmp une double 0.000000e+00, 0x8000000000000
+; CHECK-NEXT: ret i1 [[CMP]]
+;
+ %cmp = fcmp une double 0x0, 0x0008000000000000
+ ret i1 %cmp
+}
+
+; Output doesn't matter.
+define i1 @fcmp_double_dynamic_psz() #12 {
+; CHECK-LABEL: @fcmp_double_dynamic_psz(
+; CHECK-NEXT: ret i1 false
+;
+ %cmp = fcmp une double 0x0008000000000000, 0x0
+ ret i1 %cmp
+}
+
+; Non-denormal values should fold
+define float @test_float_fadd_dynamic_dynamic_normals() #11 {
+; CHECK-LABEL: @test_float_fadd_dynamic_dynamic_normals(
+; CHECK-NEXT: ret float 3.000000e+00
+;
+ %result = fadd float 1.0, 2.0
+ ret float %result
+}
+
+; Non-denormal values should fold
+define i1 @fcmp_double_dynamic_dynamic_normals() #11 {
+; CHECK-LABEL: @fcmp_double_dynamic_dynamic_normals(
+; CHECK-NEXT: ret i1 true
+;
+ %cmp = fcmp une double 1.0, 2.0
+ ret i1 %cmp
+}
+
attributes #0 = { nounwind "denormal-fp-math"="ieee,ieee" }
attributes #1 = { nounwind "denormal-fp-math"="positive-zero,ieee" }
attributes #2 = { nounwind "denormal-fp-math"="preserve-sign,ieee" }
@@ -1109,3 +1212,7 @@ attributes #5 = { nounwind "denormal-fp-math"="ieee,ieee" "denormal-fp-math-f32"
attributes #6 = { nounwind "denormal-fp-math"="positive-zero,positive-zero" }
attributes #7 = { nounwind "denormal-fp-math"="preserve-sign,preserve-sign" }
attributes #8 = { nounwind "denormal-fp-math"="ieee,ieee" "denormal-fp-math-f32"="positive-zero,positive-zero" }
+attributes #9 = { nounwind "denormal-fp-math"="dynamic,ieee" }
+attributes #10 = { nounwind "denormal-fp-math"="ieee,dynamic" }
+attributes #11 = { nounwind "denormal-fp-math"="dynamic,dynamic" }
+attributes #12 = { nounwind "denormal-fp-math"="dynamic,preserve-sign" }
diff --git a/llvm/unittests/ADT/FloatingPointMode.cpp b/llvm/unittests/ADT/FloatingPointMode.cpp
index 5b819f3495afc..5eff9de92ee84 100644
--- a/llvm/unittests/ADT/FloatingPointMode.cpp
+++ b/llvm/unittests/ADT/FloatingPointMode.cpp
@@ -20,6 +20,8 @@ TEST(FloatingPointModeTest, ParseDenormalFPAttributeComponent) {
parseDenormalFPAttributeComponent("preserve-sign"));
EXPECT_EQ(DenormalMode::PositiveZero,
parseDenormalFPAttributeComponent("positive-zero"));
+ EXPECT_EQ(DenormalMode::Dynamic,
+ parseDenormalFPAttributeComponent("dynamic"));
EXPECT_EQ(DenormalMode::Invalid, parseDenormalFPAttributeComponent("foo"));
}
@@ -27,6 +29,7 @@ TEST(FloatingPointModeTest, DenormalAttributeName) {
EXPECT_EQ("ieee", denormalModeKindName(DenormalMode::IEEE));
EXPECT_EQ("preserve-sign", denormalModeKindName(DenormalMode::PreserveSign));
EXPECT_EQ("positive-zero", denormalModeKindName(DenormalMode::PositiveZero));
+ EXPECT_EQ("dynamic", denormalModeKindName(DenormalMode::Dynamic));
EXPECT_EQ("", denormalModeKindName(DenormalMode::Invalid));
}
@@ -54,6 +57,10 @@ TEST(FloatingPointModeTest, ParseDenormalFPAttribute) {
EXPECT_EQ(DenormalMode(DenormalMode::PositiveZero, DenormalMode::PositiveZero),
parseDenormalFPAttribute("positive-zero,positive-zero"));
+ EXPECT_EQ(DenormalMode(DenormalMode::Dynamic, DenormalMode::Dynamic),
+ parseDenormalFPAttribute("dynamic"));
+ EXPECT_EQ(DenormalMode(DenormalMode::Dynamic, DenormalMode::Dynamic),
+ parseDenormalFPAttribute("dynamic,dynamic"));
EXPECT_EQ(DenormalMode(DenormalMode::IEEE, DenormalMode::PositiveZero),
parseDenormalFPAttribute("ieee,positive-zero"));
@@ -65,6 +72,10 @@ TEST(FloatingPointModeTest, ParseDenormalFPAttribute) {
EXPECT_EQ(DenormalMode(DenormalMode::IEEE, DenormalMode::PreserveSign),
parseDenormalFPAttribute("ieee,preserve-sign"));
+ EXPECT_EQ(DenormalMode(DenormalMode::Dynamic, DenormalMode::PreserveSign),
+ parseDenormalFPAttribute("dynamic,preserve-sign"));
+ EXPECT_EQ(DenormalMode(DenormalMode::PreserveSign, DenormalMode::Dynamic),
+ parseDenormalFPAttribute("preserve-sign,dynamic"));
EXPECT_EQ(DenormalMode(DenormalMode::Invalid, DenormalMode::Invalid),
parseDenormalFPAttribute("foo"));
@@ -102,6 +113,13 @@ TEST(FloatingPointModeTest, RenderDenormalFPAttribute) {
EXPECT_EQ(
"preserve-sign,positive-zero",
DenormalMode(DenormalMode::PreserveSign, DenormalMode::PositiveZero).str());
+
+ EXPECT_EQ("dynamic,dynamic",
+ DenormalMode(DenormalMode::Dynamic, DenormalMode::Dynamic).str());
+ EXPECT_EQ("ieee,dynamic",
+ DenormalMode(DenormalMode::IEEE, DenormalMode::Dynamic).str());
+ EXPECT_EQ("dynamic,ieee",
+ DenormalMode(DenormalMode::Dynamic, DenormalMode::IEEE).str());
}
TEST(FloatingPointModeTest, DenormalModeIsSimple) {
@@ -110,6 +128,10 @@ TEST(FloatingPointModeTest, DenormalModeIsSimple) {
DenormalMode::Invalid).isSimple());
EXPECT_FALSE(DenormalMode(DenormalMode::PreserveSign,
DenormalMode::PositiveZero).isSimple());
+ EXPECT_FALSE(DenormalMode(DenormalMode::PreserveSign, DenormalMode::Dynamic)
+ .isSimple());
+ EXPECT_FALSE(DenormalMode(DenormalMode::Dynamic, DenormalMode::PreserveSign)
+ .isSimple());
}
TEST(FloatingPointModeTest, DenormalModeIsValid) {
@@ -125,10 +147,76 @@ TEST(FloatingPointModeTest, DenormalModeConstructor) {
DenormalMode::getInvalid());
EXPECT_EQ(DenormalMode(DenormalMode::IEEE, DenormalMode::IEEE),
DenormalMode::getIEEE());
+ EXPECT_EQ(DenormalMode::getIEEE(), DenormalMode::getDefault());
+ EXPECT_EQ(DenormalMode(DenormalMode::Dynamic, DenormalMode::Dynamic),
+ DenormalMode::getDynamic());
EXPECT_EQ(DenormalMode(DenormalMode::PreserveSign, DenormalMode::PreserveSign),
DenormalMode::getPreserveSign());
EXPECT_EQ(DenormalMode(DenormalMode::PositiveZero, DenormalMode::PositiveZero),
DenormalMode::getPositiveZero());
}
+TEST(FloatingPointModeTest, DenormalModeMerge) {
+ EXPECT_EQ(
+ DenormalMode::getInvalid(),
+ DenormalMode::getInvalid().mergeCalleeMode(DenormalMode::getInvalid()));
+ EXPECT_EQ(DenormalMode::getIEEE(), DenormalMode::getInvalid().mergeCalleeMode(
+ DenormalMode::getIEEE()));
+ EXPECT_EQ(DenormalMode::getInvalid(), DenormalMode::getIEEE().mergeCalleeMode(
+ DenormalMode::getInvalid()));
+
+ EXPECT_EQ(DenormalMode::getIEEE(), DenormalMode::getIEEE().mergeCalleeMode(
+ DenormalMode::getDynamic()));
+ EXPECT_EQ(DenormalMode::getPreserveSign(),
+ DenormalMode::getPreserveSign().mergeCalleeMode(
+ DenormalMode::getDynamic()));
+ EXPECT_EQ(DenormalMode::getPositiveZero(),
+ DenormalMode::getPositiveZero().mergeCalleeMode(
+ DenormalMode::getDynamic()));
+ EXPECT_EQ(
+ DenormalMode::getDynamic(),
+ DenormalMode::getDynamic().mergeCalleeMode(DenormalMode::getDynamic()));
+
+ EXPECT_EQ(DenormalMode(DenormalMode::IEEE, DenormalMode::PreserveSign),
+ DenormalMode(DenormalMode::IEEE, DenormalMode::PreserveSign)
+ .mergeCalleeMode(
+ DenormalMode(DenormalMode::IEEE, DenormalMode::Dynamic)));
+
+ EXPECT_EQ(DenormalMode(DenormalMode::PreserveSign, DenormalMode::IEEE),
+ DenormalMode(DenormalMode::PreserveSign, DenormalMode::IEEE)
+ .mergeCalleeMode(
+ DenormalMode(DenormalMode::Dynamic, DenormalMode::IEEE)));
+
+ EXPECT_EQ(
+ DenormalMode(DenormalMode::PositiveZero, DenormalMode::PreserveSign),
+ DenormalMode(DenormalMode::PositiveZero, DenormalMode::PreserveSign)
+ .mergeCalleeMode(
+ DenormalMode(DenormalMode::Dynamic, DenormalMode::Dynamic)));
+
+ EXPECT_EQ(
+ DenormalMode(DenormalMode::PositiveZero, DenormalMode::PreserveSign),
+ DenormalMode(DenormalMode::PositiveZero, DenormalMode::PreserveSign)
+ .mergeCalleeMode(
+ DenormalMode(DenormalMode::PositiveZero, DenormalMode::Dynamic)));
+
+ EXPECT_EQ(
+ DenormalMode(DenormalMode::PositiveZero, DenormalMode::PreserveSign),
+ DenormalMode(DenormalMode::PositiveZero, DenormalMode::PreserveSign)
+ .mergeCalleeMode(
+ DenormalMode(DenormalMode::Dynamic, DenormalMode::PreserveSign)));
+
+ // Test some invalid / undefined behavior cases
+ EXPECT_EQ(
+ DenormalMode::getPreserveSign(),
+ DenormalMode::getIEEE().mergeCalleeMode(DenormalMode::getPreserveSign()));
+ EXPECT_EQ(
+ DenormalMode::getPreserveSign(),
+ DenormalMode::getIEEE().mergeCalleeMode(DenormalMode::getPreserveSign()));
+ EXPECT_EQ(
+ DenormalMode::getIEEE(),
+ DenormalMode::getPreserveSign().mergeCalleeMode(DenormalMode::getIEEE()));
+ EXPECT_EQ(
+ DenormalMode::getIEEE(),
+ DenormalMode::getPreserveSign().mergeCalleeMode(DenormalMode::getIEEE()));
+}
}
diff --git a/llvm/utils/TableGen/Attributes.cpp b/llvm/utils/TableGen/Attributes.cpp
index 5f8dd1594a1f2..474042a3e9a33 100644
--- a/llvm/utils/TableGen/Attributes.cpp
+++ b/llvm/utils/TableGen/Attributes.cpp
@@ -55,6 +55,7 @@ void Attributes::emitTargetIndependentNames(raw_ostream &OS) {
// Emit attribute enums in the same order llvm::Attribute::operator< expects.
Emit({"EnumAttr", "TypeAttr", "IntAttr"}, "ATTRIBUTE_ENUM");
Emit({"StrBoolAttr"}, "ATTRIBUTE_STRBOOL");
+ Emit({"ComplexStrAttr"}, "ATTRIBUTE_COMPLEXSTR");
OS << "#undef ATTRIBUTE_ALL\n";
OS << "#endif\n\n";
More information about the cfe-commits
mailing list