[llvm-branch-commits] [llvm] [NVPTX] Set default value of nvptx-allow-ftz-atomics to false (PR #231429)
Christian Sigg via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Sat Oct 17 08:55:08 PDT 2026
https://github.com/chsigg updated https://github.com/llvm/llvm-project/pull/231429
>From f782788f6d306cf3495ef915b6ad3d88a7aa3df3 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] [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-elementwise-sm60.ll | 4 +-
.../NVPTX/atomicrmw-elementwise-sm70.ll | 4 +-
.../NVPTX/atomicrmw-elementwise-sm90.ll | 4 +-
.../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 +-
10 files changed, 60 insertions(+), 60 deletions(-)
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index 58340d62bb6fea..537850b883b9eb 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-elementwise-sm60.ll b/llvm/test/CodeGen/NVPTX/atomicrmw-elementwise-sm60.ll
index 5b631504b96257..a1d20e8e821f03 100644
--- a/llvm/test/CodeGen/NVPTX/atomicrmw-elementwise-sm60.ll
+++ b/llvm/test/CodeGen/NVPTX/atomicrmw-elementwise-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-elementwise-sm70.ll b/llvm/test/CodeGen/NVPTX/atomicrmw-elementwise-sm70.ll
index c5cebdf5946657..f1c34e825077d5 100644
--- a/llvm/test/CodeGen/NVPTX/atomicrmw-elementwise-sm70.ll
+++ b/llvm/test/CodeGen/NVPTX/atomicrmw-elementwise-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-elementwise-sm90.ll b/llvm/test/CodeGen/NVPTX/atomicrmw-elementwise-sm90.ll
index 19bd7eab664888..613169d23e0099 100644
--- a/llvm/test/CodeGen/NVPTX/atomicrmw-elementwise-sm90.ll
+++ b/llvm/test/CodeGen/NVPTX/atomicrmw-elementwise-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-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 llvm-branch-commits
mailing list