[flang-commits] [flang] [llvm] [NVPTX] Set default value of nvptx-allow-ftz-atomics to false (PR #231429)
Christian Sigg via flang-commits
flang-commits at lists.llvm.org
Fri Oct 16 05:48:31 PDT 2026
https://github.com/chsigg updated https://github.com/llvm/llvm-project/pull/231429
>From 0c8a83b9116d7a8fe5374a87bc82c8ac090b241a Mon Sep 17 00:00:00 2001
From: Christian Sigg <csigg at google.com>
Date: Fri, 16 Oct 2026 07:56:21 +0200
Subject: [PATCH 1/3] [NVPTX] Attach !atomic.ignore.denormal.mode when
upgrading legacy NVVM atomic add intrinsics
---
llvm/lib/IR/AutoUpgrade.cpp | 8 ++
.../Assembler/auto_upgrade_nvvm_intrinsics.ll | 12 +--
llvm/test/CodeGen/NVPTX/atomics.ll | 89 ++++++-------------
3 files changed, 39 insertions(+), 70 deletions(-)
diff --git a/llvm/lib/IR/AutoUpgrade.cpp b/llvm/lib/IR/AutoUpgrade.cpp
index afb1cd1d48960e..4e754e34fcc643 100644
--- a/llvm/lib/IR/AutoUpgrade.cpp
+++ b/llvm/lib/IR/AutoUpgrade.cpp
@@ -3317,6 +3317,9 @@ static Value *upgradeNVVMIntrinsicCall(StringRef Name, CallBase *CI,
Rep = Builder.CreateAtomicRMW(
AtomicRMWInst::FAdd, Ptr, Val, MaybeAlign(), AtomicOrdering::Monotonic,
CI->getContext().getOrInsertSyncScopeID("device"));
+ cast<Instruction>(Rep)->setMetadata(
+ LLVMContext::MD_atomic_ignore_denormal_mode,
+ MDNode::get(CI->getContext(), {}));
// The default scope for atomic.load.* intrinsics is device
// (= gpu scope in ptx), but the default LLVM atomic scope is
// "system"
@@ -3363,6 +3366,11 @@ static Value *upgradeNVVMIntrinsicCall(StringRef Name, CallBase *CI,
"unexpected nvvm scoped atomic intrinsic");
Rep = Builder.CreateAtomicRMW(BinOp, Ptr, Val, MaybeAlign(),
AtomicOrdering::Monotonic, SSID);
+ if (BinOp == AtomicRMWInst::FAdd) {
+ cast<Instruction>(Rep)->setMetadata(
+ LLVMContext::MD_atomic_ignore_denormal_mode,
+ MDNode::get(CI->getContext(), {}));
+ }
}
} else if (Name == "clz.ll") {
// llvm.nvvm.clz.ll returns an i32, but llvm.ctlz.i64 returns an i64.
diff --git a/llvm/test/Assembler/auto_upgrade_nvvm_intrinsics.ll b/llvm/test/Assembler/auto_upgrade_nvvm_intrinsics.ll
index a2174877ac0914..1e201281cecfd4 100644
--- a/llvm/test/Assembler/auto_upgrade_nvvm_intrinsics.ll
+++ b/llvm/test/Assembler/auto_upgrade_nvvm_intrinsics.ll
@@ -399,8 +399,8 @@ define void @ldg(ptr %p0, ptr addrspace(1) %p1) {
define i32 @atomics(ptr %p0, i32 %a, float %b, double %c) {
; CHECK: %1 = atomicrmw uinc_wrap ptr %p0, i32 %a syncscope("device") monotonic
; CHECK: %2 = atomicrmw udec_wrap ptr %p0, i32 %a syncscope("device") monotonic
-; CHECK: %3 = atomicrmw fadd ptr %p0, float %b syncscope("device") monotonic
-; CHECK: %4 = atomicrmw fadd ptr %p0, double %c syncscope("device") monotonic
+; CHECK: %3 = atomicrmw fadd ptr %p0, float %b syncscope("device") monotonic, align 4, !atomic.ignore.denormal.mode !0
+; CHECK: %4 = atomicrmw fadd ptr %p0, double %c syncscope("device") monotonic, align 8, !atomic.ignore.denormal.mode !0
%r1 = call i32 @llvm.nvvm.atomic.load.inc.32.p0(ptr %p0, i32 %a)
%r2 = call i32 @llvm.nvvm.atomic.load.dec.32.p0(ptr %p0, i32 %a)
@@ -501,13 +501,13 @@ define void @atomics_scoped(ptr %p, i32 %a, i64 %b, float %c, double %d) {
; CHECK: [[CAS_sys_i64:%.*]] = cmpxchg ptr %p, i64 %b, i64 %b monotonic monotonic
; CHECK: extractvalue { i64, i1 } [[CAS_sys_i64]], 0
%cas.i.sys.i64 = call i64 @llvm.nvvm.atomic.cas.gen.i.sys.i64.p0(ptr %p, i64 %b, i64 %b)
-; CHECK: atomicrmw fadd ptr %p, float %c syncscope("block") monotonic
+; CHECK: atomicrmw fadd ptr %p, float %c syncscope("block") monotonic, align 4, !atomic.ignore.denormal.mode !0
%add.f.cta.f32 = call float @llvm.nvvm.atomic.add.gen.f.cta.f32.p0(ptr %p, float %c)
-; CHECK: atomicrmw fadd ptr %p, double %d syncscope("block") monotonic
+; CHECK: atomicrmw fadd ptr %p, double %d syncscope("block") monotonic, align 8, !atomic.ignore.denormal.mode !0
%add.f.cta.f64 = call double @llvm.nvvm.atomic.add.gen.f.cta.f64.p0(ptr %p, double %d)
-; CHECK: atomicrmw fadd ptr %p, float %c monotonic
+; CHECK: atomicrmw fadd ptr %p, float %c monotonic, align 4, !atomic.ignore.denormal.mode !0
%add.f.sys.f32 = call float @llvm.nvvm.atomic.add.gen.f.sys.f32.p0(ptr %p, float %c)
-; CHECK: atomicrmw fadd ptr %p, double %d monotonic
+; CHECK: atomicrmw fadd ptr %p, double %d monotonic, align 8, !atomic.ignore.denormal.mode !0
%add.f.sys.f64 = call double @llvm.nvvm.atomic.add.gen.f.sys.f64.p0(ptr %p, double %d)
ret void
}
diff --git a/llvm/test/CodeGen/NVPTX/atomics.ll b/llvm/test/CodeGen/NVPTX/atomics.ll
index e1853f1b00cf32..cb708fd855c77c 100644
--- a/llvm/test/CodeGen/NVPTX/atomics.ll
+++ b/llvm/test/CodeGen/NVPTX/atomics.ll
@@ -11,38 +11,17 @@ declare float @llvm.nvvm.atomic.load.add.f32.p0(ptr %addr, float %val)
; CHECK-LABEL: atomic_add_f32_generic
define float @atomic_add_f32_generic(ptr %addr, float %val) {
;
-; DISALLOW-LABEL: atomic_add_f32_generic(
-; DISALLOW: {
-; DISALLOW-NEXT: .reg .pred %p<2>;
-; DISALLOW-NEXT: .reg .b32 %r<5>;
-; DISALLOW-NEXT: .reg .b64 %rd<2>;
-; DISALLOW-EMPTY:
-; DISALLOW-NEXT: // %bb.0:
-; DISALLOW-NEXT: ld.param.b32 %r2, [atomic_add_f32_generic_param_1];
-; DISALLOW-NEXT: ld.param.b64 %rd1, [atomic_add_f32_generic_param_0];
-; DISALLOW-NEXT: ld.volatile.b32 %r4, [%rd1];
-; DISALLOW-NEXT: $L__BB0_1: // %atomicrmw.start
-; DISALLOW-NEXT: // =>This Inner Loop Header: Depth=1
-; DISALLOW-NEXT: add.rn.f32 %r3, %r4, %r2;
-; DISALLOW-NEXT: atom.cas.b32 %r1, [%rd1], %r4, %r3;
-; DISALLOW-NEXT: setp.ne.b32 %p1, %r1, %r4;
-; DISALLOW-NEXT: mov.b32 %r4, %r1;
-; DISALLOW-NEXT: @%p1 bra $L__BB0_1;
-; DISALLOW-NEXT: // %bb.2: // %atomicrmw.end
-; DISALLOW-NEXT: st.param.b32 [func_retval0], %r1;
-; DISALLOW-NEXT: ret;
-;
-; ALLOW-LABEL: atomic_add_f32_generic(
-; ALLOW: {
-; ALLOW-NEXT: .reg .b32 %r<3>;
-; ALLOW-NEXT: .reg .b64 %rd<2>;
-; ALLOW-EMPTY:
-; ALLOW-NEXT: // %bb.0:
-; ALLOW-NEXT: ld.param.b64 %rd1, [atomic_add_f32_generic_param_0];
-; ALLOW-NEXT: ld.param.b32 %r1, [atomic_add_f32_generic_param_1];
-; ALLOW-NEXT: atom.add.f32 %r2, [%rd1], %r1;
-; ALLOW-NEXT: st.param.b32 [func_retval0], %r2;
-; ALLOW-NEXT: ret;
+; CHECK-LABEL: atomic_add_f32_generic(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-NEXT: .reg .b64 %rd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b64 %rd1, [atomic_add_f32_generic_param_0];
+; CHECK-NEXT: ld.param.b32 %r1, [atomic_add_f32_generic_param_1];
+; CHECK-NEXT: atom.add.f32 %r2, [%rd1], %r1;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
%ret = call float @llvm.nvvm.atomic.load.add.f32.p0(ptr %addr, float %val)
ret float %ret
}
@@ -52,38 +31,17 @@ declare float @llvm.nvvm.atomic.load.add.f32.p1(ptr addrspace(1) %addr, float %v
; CHECK-LABEL: atomic_add_f32_addrspace1
define float @atomic_add_f32_addrspace1(ptr addrspace(1) %addr, float %val) {
;
-; DISALLOW-LABEL: atomic_add_f32_addrspace1(
-; DISALLOW: {
-; DISALLOW-NEXT: .reg .pred %p<2>;
-; DISALLOW-NEXT: .reg .b32 %r<5>;
-; DISALLOW-NEXT: .reg .b64 %rd<2>;
-; DISALLOW-EMPTY:
-; DISALLOW-NEXT: // %bb.0:
-; DISALLOW-NEXT: ld.param.b32 %r2, [atomic_add_f32_addrspace1_param_1];
-; DISALLOW-NEXT: ld.param.b64 %rd1, [atomic_add_f32_addrspace1_param_0];
-; DISALLOW-NEXT: ld.volatile.global.b32 %r4, [%rd1];
-; DISALLOW-NEXT: $L__BB1_1: // %atomicrmw.start
-; DISALLOW-NEXT: // =>This Inner Loop Header: Depth=1
-; DISALLOW-NEXT: add.rn.f32 %r3, %r4, %r2;
-; DISALLOW-NEXT: atom.global.cas.b32 %r1, [%rd1], %r4, %r3;
-; DISALLOW-NEXT: setp.ne.b32 %p1, %r1, %r4;
-; DISALLOW-NEXT: mov.b32 %r4, %r1;
-; DISALLOW-NEXT: @%p1 bra $L__BB1_1;
-; DISALLOW-NEXT: // %bb.2: // %atomicrmw.end
-; DISALLOW-NEXT: st.param.b32 [func_retval0], %r1;
-; DISALLOW-NEXT: ret;
-;
-; ALLOW-LABEL: atomic_add_f32_addrspace1(
-; ALLOW: {
-; ALLOW-NEXT: .reg .b32 %r<3>;
-; ALLOW-NEXT: .reg .b64 %rd<2>;
-; ALLOW-EMPTY:
-; ALLOW-NEXT: // %bb.0:
-; ALLOW-NEXT: ld.param.b64 %rd1, [atomic_add_f32_addrspace1_param_0];
-; ALLOW-NEXT: ld.param.b32 %r1, [atomic_add_f32_addrspace1_param_1];
-; ALLOW-NEXT: atom.global.add.f32 %r2, [%rd1], %r1;
-; ALLOW-NEXT: st.param.b32 [func_retval0], %r2;
-; ALLOW-NEXT: ret;
+; CHECK-LABEL: atomic_add_f32_addrspace1(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-NEXT: .reg .b64 %rd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b64 %rd1, [atomic_add_f32_addrspace1_param_0];
+; CHECK-NEXT: ld.param.b32 %r1, [atomic_add_f32_addrspace1_param_1];
+; CHECK-NEXT: atom.global.add.f32 %r2, [%rd1], %r1;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
%ret = call float @llvm.nvvm.atomic.load.add.f32.p1(ptr addrspace(1) %addr, float %val)
ret float %ret
}
@@ -106,3 +64,6 @@ define float @atomic_add_f32_addrspace3(ptr addrspace(3) %addr, float %val) {
%ret = call float @llvm.nvvm.atomic.load.add.f32.p3(ptr addrspace(3) %addr, float %val)
ret float %ret
}
+;; NOTE: These prefixes are unused and the list is autogenerated. Do not add tests below this line:
+; ALLOW: {{.*}}
+; DISALLOW: {{.*}}
>From f382e483219fef3194bd58633b1ed3567e194e16 Mon Sep 17 00:00:00 2001
From: Christian Sigg <csigg at google.com>
Date: Fri, 16 Oct 2026 08:16:41 +0200
Subject: [PATCH 2/3] [flang][CUDA] Emit ignore_denormal_mode for CUDA Fortran
atomicadd
---
.../Optimizer/Builder/CUDAIntrinsicCall.cpp | 7 +++--
flang/test/Lower/CUDA/cuda-atomicadd.cuf | 30 +++++++++++++++++--
flang/test/Lower/CUDA/cuda-device-proc.cuf | 6 ++--
3 files changed, 35 insertions(+), 8 deletions(-)
diff --git a/flang/lib/Optimizer/Builder/CUDAIntrinsicCall.cpp b/flang/lib/Optimizer/Builder/CUDAIntrinsicCall.cpp
index e979bbc44af98f..4e61335079711e 100644
--- a/flang/lib/Optimizer/Builder/CUDAIntrinsicCall.cpp
+++ b/flang/lib/Optimizer/Builder/CUDAIntrinsicCall.cpp
@@ -687,8 +687,11 @@ static mlir::Value genAtomBinOp(fir::FirOpBuilder &builder, mlir::Location &loc,
mlir::Value arg1) {
auto llvmPointerType = mlir::LLVM::LLVMPointerType::get(builder.getContext());
arg0 = builder.createConvert(loc, llvmPointerType, arg0);
- return mlir::LLVM::AtomicRMWOp::create(builder, loc, binOp, arg0, arg1,
- mlir::LLVM::AtomicOrdering::seq_cst);
+ bool ignoreDenormalMode = binOp == mlir::LLVM::AtomicBinOp::fadd;
+ return mlir::LLVM::AtomicRMWOp::create(
+ builder, loc, binOp, arg0, arg1, mlir::LLVM::AtomicOrdering::seq_cst,
+ /*syncscope=*/llvm::StringRef(), /*alignment=*/0, /*isVolatile=*/false,
+ /*ignoreDenormalMode=*/ignoreDenormalMode);
}
// ATOMICADD
diff --git a/flang/test/Lower/CUDA/cuda-atomicadd.cuf b/flang/test/Lower/CUDA/cuda-atomicadd.cuf
index dff029d33930d4..39e5a56808a707 100644
--- a/flang/test/Lower/CUDA/cuda-atomicadd.cuf
+++ b/flang/test/Lower/CUDA/cuda-atomicadd.cuf
@@ -8,7 +8,7 @@ attributes(global) subroutine test_atomicaddvector_r2()
end subroutine
! CHECK-LABEL: func.func @_QPtest_atomicaddvector_r2() attributes {cuf.proc_attr = #cuf.cuda_proc<global>{{.*}}}
-! CHECK: llvm.atomicrmw fadd %{{.*}}, %{{.*}} seq_cst : !llvm.ptr, vector<2xf16>
+! CHECK: llvm.atomicrmw ignore_denormal_mode fadd %{{.*}}, %{{.*}} seq_cst : !llvm.ptr, vector<2xf16>
attributes(global) subroutine test_atomicaddvector_r4()
real(4), device :: a(2), tmp1(2), tmp2(2)
@@ -16,7 +16,7 @@ attributes(global) subroutine test_atomicaddvector_r4()
end subroutine
! CHECK-LABEL: func.func @_QPtest_atomicaddvector_r4() attributes {cuf.proc_attr = #cuf.cuda_proc<global>{{.*}}}
-! CHECK: llvm.atomicrmw fadd %{{.*}}, %{{.*}} seq_cst : !llvm.ptr, vector<2xf32>
+! CHECK: llvm.atomicrmw ignore_denormal_mode fadd %{{.*}}, %{{.*}} seq_cst : !llvm.ptr, vector<2xf32>
attributes(global) subroutine test_atomicadd_r2x4()
real(4), device :: a(2), tmp1(2), tmp2(2)
@@ -24,7 +24,7 @@ attributes(global) subroutine test_atomicadd_r2x4()
end subroutine
! CHECK-LABEL: func.func @_QPtest_atomicadd_r2x4() attributes {cuf.proc_attr = #cuf.cuda_proc<global>{{.*}}}
-! CHECK: llvm.atomicrmw fadd %{{.*}}, %{{.*}} seq_cst : !llvm.ptr, vector<2xf32>
+! CHECK: llvm.atomicrmw ignore_denormal_mode fadd %{{.*}}, %{{.*}} seq_cst : !llvm.ptr, vector<2xf32>
attributes(global) subroutine test_atomicadd_r4x4()
real(4), device :: a(4), tmp1(4), tmp2(4)
@@ -33,3 +33,27 @@ end subroutine
! CHECK-LABEL: func.func @_QPtest_atomicadd_r4x4() attributes {cuf.proc_attr = #cuf.cuda_proc<global>{{.*}}}
! CHECK: atom.add.v4.f32
+
+attributes(global) subroutine test_atomicadd_r4()
+ real(4), device :: a, tmp1, tmp2
+ tmp1 = atomicadd(a, tmp2)
+end subroutine
+
+! CHECK-LABEL: func.func @_QPtest_atomicadd_r4() attributes {cuf.proc_attr = #cuf.cuda_proc<global>{{.*}}}
+! CHECK: llvm.atomicrmw ignore_denormal_mode fadd %{{.*}}, %{{.*}} seq_cst : !llvm.ptr, f32
+
+attributes(global) subroutine test_atomicadd_r8()
+ real(8), device :: a, tmp1, tmp2
+ tmp1 = atomicadd(a, tmp2)
+end subroutine
+
+! CHECK-LABEL: func.func @_QPtest_atomicadd_r8() attributes {cuf.proc_attr = #cuf.cuda_proc<global>{{.*}}}
+! CHECK: llvm.atomicrmw ignore_denormal_mode fadd %{{.*}}, %{{.*}} seq_cst : !llvm.ptr, f64
+
+attributes(global) subroutine test_atomicadd_i4()
+ integer(4), device :: a, tmp1, tmp2
+ tmp1 = atomicadd(a, tmp2)
+end subroutine
+
+! CHECK-LABEL: func.func @_QPtest_atomicadd_i4() attributes {cuf.proc_attr = #cuf.cuda_proc<global>{{.*}}}
+! CHECK: llvm.atomicrmw add %{{.*}}, %{{.*}} seq_cst : !llvm.ptr, i32
diff --git a/flang/test/Lower/CUDA/cuda-device-proc.cuf b/flang/test/Lower/CUDA/cuda-device-proc.cuf
index a1c3c33a76811d..b9081681569be9 100644
--- a/flang/test/Lower/CUDA/cuda-device-proc.cuf
+++ b/flang/test/Lower/CUDA/cuda-device-proc.cuf
@@ -123,9 +123,9 @@ end
! CHECK: %{{.*}} = nvvm.barrier.reduction #nvvm.reduction<or> %[[CONV]] -> i32
! CHECK: %{{.*}} = llvm.atomicrmw add %{{.*}}, %{{.*}} seq_cst : !llvm.ptr, i32
! CHECK: %{{.*}} = llvm.atomicrmw add %{{.*}}, %{{.*}} seq_cst : !llvm.ptr, i64
-! CHECK: %{{.*}} = llvm.atomicrmw fadd %{{.*}}, %{{.*}} seq_cst : !llvm.ptr, f32
-! CHECK: %{{.*}} = llvm.atomicrmw fadd %{{.*}}, %{{.*}} seq_cst : !llvm.ptr, f64
-! CHECK: %{{.*}} = llvm.atomicrmw fadd %{{.*}}, %{{.*}} seq_cst : !llvm.ptr, vector<2xf16>
+! CHECK: %{{.*}} = llvm.atomicrmw ignore_denormal_mode fadd %{{.*}}, %{{.*}} seq_cst : !llvm.ptr, f32
+! CHECK: %{{.*}} = llvm.atomicrmw ignore_denormal_mode fadd %{{.*}}, %{{.*}} seq_cst : !llvm.ptr, f64
+! CHECK: %{{.*}} = llvm.atomicrmw ignore_denormal_mode fadd %{{.*}}, %{{.*}} seq_cst : !llvm.ptr, vector<2xf16>
! CHECK: %{{.*}} = llvm.atomicrmw sub %{{.*}}, %{{.*}} seq_cst : !llvm.ptr, i32
! CHECK: %{{.*}} = llvm.atomicrmw sub %{{.*}}, %{{.*}} seq_cst : !llvm.ptr, i64
>From 9a618da6bddf2a736e81c2ebc9cf123bd266a623 Mon Sep 17 00:00:00 2001
From: Christian Sigg <csigg at google.com>
Date: Tue, 13 Oct 2026 11:42:50 +0200
Subject: [PATCH 3/3] [NVPTX] Set default value of nvptx-allow-ftz-atomics to
false
This restores the strict FTZ default originally introduced in #200732
and temporarily relaxed in #206154.
Frontends can now explicitly opt in to native FTZ-mismatched atomics on a
per-instruction basis using !atomic.ignore.denormal.mode metadata (#217585,
#217586), which is supported in Clang, MLIR, and XLA. Users who need the
previous behavior globally can still pass -nvptx-allow-ftz-atomics=true.
Update test cases to verify the new default disallow behavior.
---
llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp | 11 +--
.../NVPTX/atomicrmw-allow-ftz-atomics.ll | 8 +-
.../NVPTX/atomicrmw-ignore-denormal-mode.ll | 73 +++++++++----------
llvm/test/CodeGen/NVPTX/atomicrmw-sm60.ll | 4 +-
llvm/test/CodeGen/NVPTX/atomicrmw-sm70.ll | 4 +-
llvm/test/CodeGen/NVPTX/atomicrmw-sm90.ll | 4 +-
llvm/test/CodeGen/NVPTX/atomicrmw.py | 4 +-
7 files changed, 54 insertions(+), 54 deletions(-)
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index 7d55ba810e1065..13f1982d5ea9c6 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -112,15 +112,16 @@ static cl::opt<bool> UsePrecSqrtF32(
cl::init(true));
// PTX atom.add.f32 has fixed FTZ behavior that may not match the function's
-// (see shouldExpandAtomicRMWInIR), so we'd normally fall back to a CAS loop
-// when they disagree. This option (enabled by default) allows using atom.add
-// anyway, trading correct denormal handling for the speed of the native
-// instruction.
+// (see shouldExpandAtomicRMWInIR), so by default we fall back to a CAS loop
+// when they disagree (strict mode), unless declared insignificant via
+// !atomic.ignore.denormal.mode metadata or this option. This option allows
+// using atom.add anyway, trading correct denormal handling for the speed of the
+// native instruction.
static cl::opt<bool> AllowFTZAtomics(
"nvptx-allow-ftz-atomics", cl::Hidden,
cl::desc("NVPTX Specific: Lower atomicrmw fadd to atom.add even when its "
"FTZ behavior does not match the function's denormal mode."),
- cl::init(true));
+ cl::init(false));
/// Whereas CUDA's implementation (see libdevice) uses ex2.approx for exp2(), it
/// does NOT use lg2.approx for log2, so this is disabled by default.
diff --git a/llvm/test/CodeGen/NVPTX/atomicrmw-allow-ftz-atomics.ll b/llvm/test/CodeGen/NVPTX/atomicrmw-allow-ftz-atomics.ll
index 78c86bfb08e2fa..e78f2428bbde98 100644
--- a/llvm/test/CodeGen/NVPTX/atomicrmw-allow-ftz-atomics.ll
+++ b/llvm/test/CodeGen/NVPTX/atomicrmw-allow-ftz-atomics.ll
@@ -2,16 +2,16 @@
; -nvptx-allow-ftz-atomics forces atomicrmw fadd to lower to atom.add even when
; the hardware's FTZ behavior doesn't match the function's denormal mode, where
; we'd otherwise emit a CAS loop. Compare the lowering with the option disabled
-; (DISALLOW) against the option enabled or default (ALLOW).
+; or default (DISALLOW) against the option enabled (ALLOW).
; RUN: llc < %s -march=nvptx64 -mcpu=sm_90 -mattr=+ptx87 -nvptx-allow-ftz-atomics=false | FileCheck %s --check-prefixes=CHECK,DISALLOW
; RUN: llc < %s -march=nvptx64 -mcpu=sm_90 -mattr=+ptx87 -nvptx-allow-ftz-atomics=true | FileCheck %s --check-prefixes=CHECK,ALLOW
-; RUN: llc < %s -march=nvptx64 -mcpu=sm_90 -mattr=+ptx87 | FileCheck %s --check-prefixes=CHECK,ALLOW
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_90 -mattr=+ptx87 | FileCheck %s --check-prefixes=CHECK,DISALLOW
; RUN: %if ptxas-sm_90 && ptxas-isa-8.7 %{ llc < %s -march=nvptx64 -mcpu=sm_90 -mattr=+ptx87 -nvptx-allow-ftz-atomics=false | %ptxas-verify -arch=sm_90 %}
; RUN: %if ptxas-sm_90 && ptxas-isa-8.7 %{ llc < %s -march=nvptx64 -mcpu=sm_90 -mattr=+ptx87 -nvptx-allow-ftz-atomics=true | %ptxas-verify -arch=sm_90 %}
; RUN: %if ptxas-sm_90 && ptxas-isa-8.7 %{ llc < %s -march=nvptx64 -mcpu=sm_90 -mattr=+ptx87 | %ptxas-verify -arch=sm_90 %}
; global f32: atom.add flushes, so we'd normally use a CAS loop (function is not FTZ)
-; when the option is disabled; when enabled (default), it lowers to atom.add.
+; when the option is disabled (default); when enabled, it lowers to atom.add.
define float @fadd_f32_global(ptr addrspace(1) %addr, float %val) {
;
; DISALLOW-LABEL: fadd_f32_global(
@@ -51,7 +51,7 @@ define float @fadd_f32_global(ptr addrspace(1) %addr, float %val) {
}
; generic f32: we can't know the memory space, so we'd normally use a CAS loop
-; when the option is disabled; when enabled (default), it lowers to atom.add.
+; when the option is disabled (default); when enabled, it lowers to atom.add.
define float @fadd_f32_generic(ptr %addr, float %val) {
;
; DISALLOW-LABEL: fadd_f32_generic(
diff --git a/llvm/test/CodeGen/NVPTX/atomicrmw-ignore-denormal-mode.ll b/llvm/test/CodeGen/NVPTX/atomicrmw-ignore-denormal-mode.ll
index f3cf67199e8b91..212688e5d8ccf1 100644
--- a/llvm/test/CodeGen/NVPTX/atomicrmw-ignore-denormal-mode.ll
+++ b/llvm/test/CodeGen/NVPTX/atomicrmw-ignore-denormal-mode.ll
@@ -1,6 +1,7 @@
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
; RUN: llc < %s -march=nvptx64 -mcpu=sm_90 -mattr=+ptx87 -nvptx-allow-ftz-atomics=false | FileCheck %s --check-prefixes=CHECK,STRICT
-; RUN: llc < %s -march=nvptx64 -mcpu=sm_90 -mattr=+ptx87 | FileCheck %s --check-prefixes=CHECK,DEFAULT
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_90 -mattr=+ptx87 -nvptx-allow-ftz-atomics=true | FileCheck %s --check-prefixes=CHECK,ALLOW
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_90 -mattr=+ptx87 | FileCheck %s --check-prefixes=CHECK,STRICT
; RUN: %if ptxas-sm_90 && ptxas-isa-8.7 %{ llc < %s -march=nvptx64 -mcpu=sm_90 -mattr=+ptx87 -nvptx-allow-ftz-atomics=false | %ptxas-verify -arch=sm_90 %}
; PTX atom.add has a fixed FTZ behavior that need not match the function's
@@ -13,9 +14,7 @@
; atom.add.ftz.f16)
; When the two disagree the atomic is expanded into a CAS loop, unless the
; mismatch is explicitly declared insignificant -- either globally via
-; -nvptx-allow-ftz-atomics (on by default, hence the STRICT run lines above,
-; without which the metadata would never be the deciding factor) or per
-; instruction via !atomic.ignore.denormal.mode.
+; -nvptx-allow-ftz-atomics or per instruction via !atomic.ignore.denormal.mode.
; f32, global: the native instruction flushes, so an IEEE function needs the
; metadata to avoid the CAS loop.
@@ -42,17 +41,17 @@ define float @fadd_f32_global_ieee_no_md(ptr addrspace(1) %addr, float %val) {
; STRICT-NEXT: st.param::func.b32 [func_retval0], %r1;
; STRICT-NEXT: ret;
;
-; DEFAULT-LABEL: fadd_f32_global_ieee_no_md(
-; DEFAULT: {
-; DEFAULT-NEXT: .reg .b32 %r<3>;
-; DEFAULT-NEXT: .reg .b64 %rd<2>;
-; DEFAULT-EMPTY:
-; DEFAULT-NEXT: // %bb.0:
-; DEFAULT-NEXT: ld.param::func.b64 %rd1, [fadd_f32_global_ieee_no_md_param_0];
-; DEFAULT-NEXT: ld.param::func.b32 %r1, [fadd_f32_global_ieee_no_md_param_1];
-; DEFAULT-NEXT: atom.relaxed.sys.global.add.f32 %r2, [%rd1], %r1;
-; DEFAULT-NEXT: st.param::func.b32 [func_retval0], %r2;
-; DEFAULT-NEXT: ret;
+; ALLOW-LABEL: fadd_f32_global_ieee_no_md(
+; ALLOW: {
+; ALLOW-NEXT: .reg .b32 %r<3>;
+; ALLOW-NEXT: .reg .b64 %rd<2>;
+; ALLOW-EMPTY:
+; ALLOW-NEXT: // %bb.0:
+; ALLOW-NEXT: ld.param::func.b64 %rd1, [fadd_f32_global_ieee_no_md_param_0];
+; ALLOW-NEXT: ld.param::func.b32 %r1, [fadd_f32_global_ieee_no_md_param_1];
+; ALLOW-NEXT: atom.relaxed.sys.global.add.f32 %r2, [%rd1], %r1;
+; ALLOW-NEXT: st.param::func.b32 [func_retval0], %r2;
+; ALLOW-NEXT: ret;
%r = atomicrmw fadd ptr addrspace(1) %addr, float %val monotonic
ret float %r
}
@@ -116,17 +115,17 @@ define float @fadd_f32_generic_ieee_no_md(ptr %addr, float %val) {
; STRICT-NEXT: st.param::func.b32 [func_retval0], %r1;
; STRICT-NEXT: ret;
;
-; DEFAULT-LABEL: fadd_f32_generic_ieee_no_md(
-; DEFAULT: {
-; DEFAULT-NEXT: .reg .b32 %r<3>;
-; DEFAULT-NEXT: .reg .b64 %rd<2>;
-; DEFAULT-EMPTY:
-; DEFAULT-NEXT: // %bb.0:
-; DEFAULT-NEXT: ld.param::func.b64 %rd1, [fadd_f32_generic_ieee_no_md_param_0];
-; DEFAULT-NEXT: ld.param::func.b32 %r1, [fadd_f32_generic_ieee_no_md_param_1];
-; DEFAULT-NEXT: atom.relaxed.sys.add.f32 %r2, [%rd1], %r1;
-; DEFAULT-NEXT: st.param::func.b32 [func_retval0], %r2;
-; DEFAULT-NEXT: ret;
+; ALLOW-LABEL: fadd_f32_generic_ieee_no_md(
+; ALLOW: {
+; ALLOW-NEXT: .reg .b32 %r<3>;
+; ALLOW-NEXT: .reg .b64 %rd<2>;
+; ALLOW-EMPTY:
+; ALLOW-NEXT: // %bb.0:
+; ALLOW-NEXT: ld.param::func.b64 %rd1, [fadd_f32_generic_ieee_no_md_param_0];
+; ALLOW-NEXT: ld.param::func.b32 %r1, [fadd_f32_generic_ieee_no_md_param_1];
+; ALLOW-NEXT: atom.relaxed.sys.add.f32 %r2, [%rd1], %r1;
+; ALLOW-NEXT: st.param::func.b32 [func_retval0], %r2;
+; ALLOW-NEXT: ret;
%r = atomicrmw fadd ptr %addr, float %val monotonic
ret float %r
}
@@ -203,17 +202,17 @@ define half @fadd_f16_global_ftz_no_md(ptr addrspace(1) %addr, half %val) #1 {
; STRICT-NEXT: st.param::func.b16 [func_retval0], %r12;
; STRICT-NEXT: ret;
;
-; DEFAULT-LABEL: fadd_f16_global_ftz_no_md(
-; DEFAULT: {
-; DEFAULT-NEXT: .reg .b16 %rs<3>;
-; DEFAULT-NEXT: .reg .b64 %rd<2>;
-; DEFAULT-EMPTY:
-; DEFAULT-NEXT: // %bb.0:
-; DEFAULT-NEXT: ld.param::func.b64 %rd1, [fadd_f16_global_ftz_no_md_param_0];
-; DEFAULT-NEXT: ld.param::func.b16 %rs1, [fadd_f16_global_ftz_no_md_param_1];
-; DEFAULT-NEXT: atom.relaxed.sys.global.add.noftz.f16 %rs2, [%rd1], %rs1;
-; DEFAULT-NEXT: st.param::func.b16 [func_retval0], %rs2;
-; DEFAULT-NEXT: ret;
+; ALLOW-LABEL: fadd_f16_global_ftz_no_md(
+; ALLOW: {
+; ALLOW-NEXT: .reg .b16 %rs<3>;
+; ALLOW-NEXT: .reg .b64 %rd<2>;
+; ALLOW-EMPTY:
+; ALLOW-NEXT: // %bb.0:
+; ALLOW-NEXT: ld.param::func.b64 %rd1, [fadd_f16_global_ftz_no_md_param_0];
+; ALLOW-NEXT: ld.param::func.b16 %rs1, [fadd_f16_global_ftz_no_md_param_1];
+; ALLOW-NEXT: atom.relaxed.sys.global.add.noftz.f16 %rs2, [%rd1], %rs1;
+; ALLOW-NEXT: st.param::func.b16 [func_retval0], %rs2;
+; ALLOW-NEXT: ret;
%r = atomicrmw fadd ptr addrspace(1) %addr, half %val monotonic
ret half %r
}
diff --git a/llvm/test/CodeGen/NVPTX/atomicrmw-sm60.ll b/llvm/test/CodeGen/NVPTX/atomicrmw-sm60.ll
index fedb1ba54f8a8a..85d7f8f2f52433 100644
--- a/llvm/test/CodeGen/NVPTX/atomicrmw-sm60.ll
+++ b/llvm/test/CodeGen/NVPTX/atomicrmw-sm60.ll
@@ -1,10 +1,10 @@
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
; RUN: llc < %s -march=nvptx64 -mcpu=sm_60 -mattr=+ptx50 -nvptx-allow-ftz-atomics=false | FileCheck %s --check-prefixes=SM60,SM60-NOFTZ-DISALLOW
; RUN: llc < %s -march=nvptx64 -mcpu=sm_60 -mattr=+ptx50 -nvptx-allow-ftz-atomics=true | FileCheck %s --check-prefixes=SM60,SM60-NOFTZ-ALLOW
-; RUN: llc < %s -march=nvptx64 -mcpu=sm_60 -mattr=+ptx50 | FileCheck %s --check-prefixes=SM60,SM60-NOFTZ-ALLOW
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_60 -mattr=+ptx50 | FileCheck %s --check-prefixes=SM60,SM60-NOFTZ-DISALLOW
; RUN: llc < %s -march=nvptx64 -mcpu=sm_60 -mattr=+ptx50 -denormal-fp-math-f32=preserve-sign -nvptx-allow-ftz-atomics=false | FileCheck %s --check-prefixes=SM60,SM60-FTZ-DISALLOW
; RUN: llc < %s -march=nvptx64 -mcpu=sm_60 -mattr=+ptx50 -denormal-fp-math-f32=preserve-sign -nvptx-allow-ftz-atomics=true | FileCheck %s --check-prefixes=SM60,SM60-FTZ-ALLOW
-; RUN: llc < %s -march=nvptx64 -mcpu=sm_60 -mattr=+ptx50 -denormal-fp-math-f32=preserve-sign | FileCheck %s --check-prefixes=SM60,SM60-FTZ-ALLOW
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_60 -mattr=+ptx50 -denormal-fp-math-f32=preserve-sign | FileCheck %s --check-prefixes=SM60,SM60-FTZ-DISALLOW
; RUN: %if ptxas-sm_60 && ptxas-isa-5.0 %{ llc < %s -march=nvptx64 -mcpu=sm_60 -mattr=+ptx50 -nvptx-allow-ftz-atomics=false | %ptxas-verify -arch=sm_60 %}
; RUN: %if ptxas-sm_60 && ptxas-isa-5.0 %{ llc < %s -march=nvptx64 -mcpu=sm_60 -mattr=+ptx50 -nvptx-allow-ftz-atomics=true | %ptxas-verify -arch=sm_60 %}
; RUN: %if ptxas-sm_60 && ptxas-isa-5.0 %{ llc < %s -march=nvptx64 -mcpu=sm_60 -mattr=+ptx50 | %ptxas-verify -arch=sm_60 %}
diff --git a/llvm/test/CodeGen/NVPTX/atomicrmw-sm70.ll b/llvm/test/CodeGen/NVPTX/atomicrmw-sm70.ll
index 12ba9b5fb4f8ca..e40d78b4effaa9 100644
--- a/llvm/test/CodeGen/NVPTX/atomicrmw-sm70.ll
+++ b/llvm/test/CodeGen/NVPTX/atomicrmw-sm70.ll
@@ -1,10 +1,10 @@
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
; RUN: llc < %s -march=nvptx64 -mcpu=sm_70 -mattr=+ptx63 -nvptx-allow-ftz-atomics=false | FileCheck %s --check-prefixes=SM70,SM70-NOFTZ-DISALLOW
; RUN: llc < %s -march=nvptx64 -mcpu=sm_70 -mattr=+ptx63 -nvptx-allow-ftz-atomics=true | FileCheck %s --check-prefixes=SM70,SM70-NOFTZ-ALLOW
-; RUN: llc < %s -march=nvptx64 -mcpu=sm_70 -mattr=+ptx63 | FileCheck %s --check-prefixes=SM70,SM70-NOFTZ-ALLOW
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_70 -mattr=+ptx63 | FileCheck %s --check-prefixes=SM70,SM70-NOFTZ-DISALLOW
; RUN: llc < %s -march=nvptx64 -mcpu=sm_70 -mattr=+ptx63 -denormal-fp-math-f32=preserve-sign -nvptx-allow-ftz-atomics=false | FileCheck %s --check-prefixes=SM70,SM70-FTZ-DISALLOW
; RUN: llc < %s -march=nvptx64 -mcpu=sm_70 -mattr=+ptx63 -denormal-fp-math-f32=preserve-sign -nvptx-allow-ftz-atomics=true | FileCheck %s --check-prefixes=SM70,SM70-FTZ-ALLOW
-; RUN: llc < %s -march=nvptx64 -mcpu=sm_70 -mattr=+ptx63 -denormal-fp-math-f32=preserve-sign | FileCheck %s --check-prefixes=SM70,SM70-FTZ-ALLOW
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_70 -mattr=+ptx63 -denormal-fp-math-f32=preserve-sign | FileCheck %s --check-prefixes=SM70,SM70-FTZ-DISALLOW
; RUN: %if ptxas-sm_70 && ptxas-isa-6.3 %{ llc < %s -march=nvptx64 -mcpu=sm_70 -mattr=+ptx63 -nvptx-allow-ftz-atomics=false | %ptxas-verify -arch=sm_70 %}
; RUN: %if ptxas-sm_70 && ptxas-isa-6.3 %{ llc < %s -march=nvptx64 -mcpu=sm_70 -mattr=+ptx63 -nvptx-allow-ftz-atomics=true | %ptxas-verify -arch=sm_70 %}
; RUN: %if ptxas-sm_70 && ptxas-isa-6.3 %{ llc < %s -march=nvptx64 -mcpu=sm_70 -mattr=+ptx63 | %ptxas-verify -arch=sm_70 %}
diff --git a/llvm/test/CodeGen/NVPTX/atomicrmw-sm90.ll b/llvm/test/CodeGen/NVPTX/atomicrmw-sm90.ll
index ab31e4a297bfad..599d775d45da72 100644
--- a/llvm/test/CodeGen/NVPTX/atomicrmw-sm90.ll
+++ b/llvm/test/CodeGen/NVPTX/atomicrmw-sm90.ll
@@ -1,10 +1,10 @@
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
; RUN: llc < %s -march=nvptx64 -mcpu=sm_90 -mattr=+ptx87 -nvptx-allow-ftz-atomics=false | FileCheck %s --check-prefixes=SM90,SM90-NOFTZ-DISALLOW
; RUN: llc < %s -march=nvptx64 -mcpu=sm_90 -mattr=+ptx87 -nvptx-allow-ftz-atomics=true | FileCheck %s --check-prefixes=SM90,SM90-NOFTZ-ALLOW
-; RUN: llc < %s -march=nvptx64 -mcpu=sm_90 -mattr=+ptx87 | FileCheck %s --check-prefixes=SM90,SM90-NOFTZ-ALLOW
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_90 -mattr=+ptx87 | FileCheck %s --check-prefixes=SM90,SM90-NOFTZ-DISALLOW
; RUN: llc < %s -march=nvptx64 -mcpu=sm_90 -mattr=+ptx87 -denormal-fp-math-f32=preserve-sign -nvptx-allow-ftz-atomics=false | FileCheck %s --check-prefixes=SM90,SM90-FTZ-DISALLOW
; RUN: llc < %s -march=nvptx64 -mcpu=sm_90 -mattr=+ptx87 -denormal-fp-math-f32=preserve-sign -nvptx-allow-ftz-atomics=true | FileCheck %s --check-prefixes=SM90,SM90-FTZ-ALLOW
-; RUN: llc < %s -march=nvptx64 -mcpu=sm_90 -mattr=+ptx87 -denormal-fp-math-f32=preserve-sign | FileCheck %s --check-prefixes=SM90,SM90-FTZ-ALLOW
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_90 -mattr=+ptx87 -denormal-fp-math-f32=preserve-sign | FileCheck %s --check-prefixes=SM90,SM90-FTZ-DISALLOW
; RUN: %if ptxas-sm_90 && ptxas-isa-8.7 %{ llc < %s -march=nvptx64 -mcpu=sm_90 -mattr=+ptx87 -nvptx-allow-ftz-atomics=false | %ptxas-verify -arch=sm_90 %}
; RUN: %if ptxas-sm_90 && ptxas-isa-8.7 %{ llc < %s -march=nvptx64 -mcpu=sm_90 -mattr=+ptx87 -nvptx-allow-ftz-atomics=true | %ptxas-verify -arch=sm_90 %}
; RUN: %if ptxas-sm_90 && ptxas-isa-8.7 %{ llc < %s -march=nvptx64 -mcpu=sm_90 -mattr=+ptx87 | %ptxas-verify -arch=sm_90 %}
diff --git a/llvm/test/CodeGen/NVPTX/atomicrmw.py b/llvm/test/CodeGen/NVPTX/atomicrmw.py
index 05388afa53578e..832020d0e09795 100644
--- a/llvm/test/CodeGen/NVPTX/atomicrmw.py
+++ b/llvm/test/CodeGen/NVPTX/atomicrmw.py
@@ -80,10 +80,10 @@
run_statement = Template(
"""; RUN: llc < %s -march=nvptx64 -mcpu=sm_${sm} -mattr=+ptx${ptx} -nvptx-allow-ftz-atomics=false | FileCheck %s --check-prefixes=SM${sm},SM${sm}-NOFTZ-DISALLOW
; RUN: llc < %s -march=nvptx64 -mcpu=sm_${sm} -mattr=+ptx${ptx} -nvptx-allow-ftz-atomics=true | FileCheck %s --check-prefixes=SM${sm},SM${sm}-NOFTZ-ALLOW
-; RUN: llc < %s -march=nvptx64 -mcpu=sm_${sm} -mattr=+ptx${ptx} | FileCheck %s --check-prefixes=SM${sm},SM${sm}-NOFTZ-ALLOW
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_${sm} -mattr=+ptx${ptx} | FileCheck %s --check-prefixes=SM${sm},SM${sm}-NOFTZ-DISALLOW
; RUN: llc < %s -march=nvptx64 -mcpu=sm_${sm} -mattr=+ptx${ptx} -denormal-fp-math-f32=preserve-sign -nvptx-allow-ftz-atomics=false | FileCheck %s --check-prefixes=SM${sm},SM${sm}-FTZ-DISALLOW
; RUN: llc < %s -march=nvptx64 -mcpu=sm_${sm} -mattr=+ptx${ptx} -denormal-fp-math-f32=preserve-sign -nvptx-allow-ftz-atomics=true | FileCheck %s --check-prefixes=SM${sm},SM${sm}-FTZ-ALLOW
-; RUN: llc < %s -march=nvptx64 -mcpu=sm_${sm} -mattr=+ptx${ptx} -denormal-fp-math-f32=preserve-sign | FileCheck %s --check-prefixes=SM${sm},SM${sm}-FTZ-ALLOW
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_${sm} -mattr=+ptx${ptx} -denormal-fp-math-f32=preserve-sign | FileCheck %s --check-prefixes=SM${sm},SM${sm}-FTZ-DISALLOW
; RUN: %if ptxas-sm_${sm} && ptxas-isa-${ptxfp} %{ llc < %s -march=nvptx64 -mcpu=sm_${sm} -mattr=+ptx${ptx} -nvptx-allow-ftz-atomics=false | %ptxas-verify -arch=sm_${sm} %}
; RUN: %if ptxas-sm_${sm} && ptxas-isa-${ptxfp} %{ llc < %s -march=nvptx64 -mcpu=sm_${sm} -mattr=+ptx${ptx} -nvptx-allow-ftz-atomics=true | %ptxas-verify -arch=sm_${sm} %}
; RUN: %if ptxas-sm_${sm} && ptxas-isa-${ptxfp} %{ llc < %s -march=nvptx64 -mcpu=sm_${sm} -mattr=+ptx${ptx} | %ptxas-verify -arch=sm_${sm} %}
More information about the flang-commits
mailing list