Skip to content

IR: Promote "denormal-fp-math" to a first class attribute #174293

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
arsenm wants to merge 1 commit into
base: main
Choose a base branch
from
Open

IR: Promote "denormal-fp-math" to a first class attribute #174293

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
40 changes: 11 additions & 29 deletions clang/lib/CodeGen/CGCall.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down Expand Up @@ -2167,35 +2165,19 @@ void CodeGen::mergeDefaultFunctionDefinitionAttributes(

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);
}
llvm::DenormalFPEnv OptsFPEnv(CodeGenOpts.FPDenormalMode,
CodeGenOpts.FP32DenormalMode);
llvm::DenormalFPEnv MergedFPEnv =
OptsFPEnv.mergeCalleeMode(F.getDenormalFPEnv());

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) {
if (MergedFPEnv == llvm::DenormalFPEnv::getDefault()) {
AttrsToRemove.addAttribute(llvm::Attribute::DenormalFPEnv);
} 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);

Expand Down
2 changes: 1 addition & 1 deletion clang/lib/CodeGen/CGCall.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down
60 changes: 33 additions & 27 deletions clang/test/CodeGen/denormalfpmode-f32.c
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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;
Expand Down
8 changes: 4 additions & 4 deletions clang/test/CodeGen/denormalfpmode.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down
6 changes: 3 additions & 3 deletions clang/test/CodeGenCUDA/flush-denormals.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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.
Expand All @@ -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}
Expand Down
30 changes: 15 additions & 15 deletions clang/test/CodeGenCUDA/link-builtin-bitcode-denormal-fp-mode.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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) {{.*}} }
12 changes: 3 additions & 9 deletions clang/test/CodeGenCUDA/propagate-attributes.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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"

Expand All @@ -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"
Loading