release/23.x: [NVPTX] Set default value of nvptx-allow-ftz-atomics to true (#206154) by llvmbot · Pull Request #229003 · llvm/llvm-project

llvmbot · 2026-10-05T08:18:46Z

Backport 2b6e9d2

Requested by: @ilya-biryukov

llvmbot · 2026-10-05T08:18:57Z

@rnk What do you think about merging this PR to the release branch?

llvmorg-github-actions · 2026-10-05T08:19:24Z

@llvm/pr-subscribers-backend-nvptx

Author: llvmbot

Changes

Backport 2b6e9d2

Requested by: @ilya-biryukov

Patch is 604.13 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/229003.diff

7 Files Affected:

(modified) llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp (+5-4)
(modified) llvm/test/CodeGen/NVPTX/atomicrmw-allow-ftz-atomics.ll (+53-49)
(modified) llvm/test/CodeGen/NVPTX/atomicrmw-sm60.ll (+2949-1487)
(modified) llvm/test/CodeGen/NVPTX/atomicrmw-sm70.ll (+2871-1461)
(modified) llvm/test/CodeGen/NVPTX/atomicrmw-sm90.ll (+1635-849)
(modified) llvm/test/CodeGen/NVPTX/atomicrmw.py (+10-2)
(modified) llvm/test/CodeGen/NVPTX/atomics.ll (+71-41)

diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index 59d04dde1eeb0f5..6cca7f4acd85896 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -110,14 +110,15 @@ 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 by default we fall back to a CAS loop
-// when they disagree. This flag is an escape hatch to use atom.add anyway,
-// trading correct denormal handling for the speed of the native instruction.
+// (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.
 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(false));
+    cl::init(true));
 
 /// 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 727c13bb4e47e70..c1d223082eb3020 100644
--- a/llvm/test/CodeGen/NVPTX/atomicrmw-allow-ftz-atomics.ll
+++ b/llvm/test/CodeGen/NVPTX/atomicrmw-allow-ftz-atomics.ll
@@ -1,36 +1,39 @@
 ; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
 ; -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 default lowering (DEFAULT) against
-; the flag (ALLOW).
-; 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 | FileCheck %s --check-prefixes=CHECK,ALLOW
+; we'd otherwise emit a CAS loop. Compare the lowering with the option disabled
+; (DISALLOW) against the option enabled or default (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: %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 %}
-; RUN: %if ptxas-sm_90 && ptxas-isa-8.7 %{ llc < %s -march=nvptx64 -mcpu=sm_90 -mattr=+ptx87 -nvptx-allow-ftz-atomics | %ptxas-verify -arch=sm_90 %}
 
-; global f32: atom.add flushes, so default uses a CAS loop (function is not FTZ);
-; the flag forces atom.add.
+; 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.
 define float @fadd_f32_global(ptr addrspace(1) %addr, float %val) {
-; DEFAULT-LABEL: fadd_f32_global(
-; DEFAULT:       {
-; DEFAULT-NEXT:    .reg .pred %p<2>;
-; DEFAULT-NEXT:    .reg .b32 %r<5>;
-; DEFAULT-NEXT:    .reg .b64 %rd<2>;
-; DEFAULT-EMPTY:
-; DEFAULT-NEXT:  // %bb.0:
-; DEFAULT-NEXT:    ld.param.b32 %r2, [fadd_f32_global_param_1];
-; DEFAULT-NEXT:    ld.param.b64 %rd1, [fadd_f32_global_param_0];
-; DEFAULT-NEXT:    ld.relaxed.sys.global.b32 %r4, [%rd1];
-; DEFAULT-NEXT:  $L__BB0_1: // %atomicrmw.start
-; DEFAULT-NEXT:    // =>This Inner Loop Header: Depth=1
-; DEFAULT-NEXT:    add.rn.f32 %r3, %r4, %r2;
-; DEFAULT-NEXT:    atom.relaxed.sys.global.cas.b32 %r1, [%rd1], %r4, %r3;
-; DEFAULT-NEXT:    setp.ne.b32 %p1, %r1, %r4;
-; DEFAULT-NEXT:    mov.b32 %r4, %r1;
-; DEFAULT-NEXT:    @%p1 bra $L__BB0_1;
-; DEFAULT-NEXT:  // %bb.2: // %atomicrmw.end
-; DEFAULT-NEXT:    st.param.b32 [func_retval0], %r1;
-; DEFAULT-NEXT:    ret;
+;
+; DISALLOW-LABEL: fadd_f32_global(
+; 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, [fadd_f32_global_param_1];
+; DISALLOW-NEXT:    ld.param.b64 %rd1, [fadd_f32_global_param_0];
+; DISALLOW-NEXT:    ld.relaxed.sys.global.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.relaxed.sys.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__BB0_1;
+; DISALLOW-NEXT:  // %bb.2: // %atomicrmw.end
+; DISALLOW-NEXT:    st.param.b32 [func_retval0], %r1;
+; DISALLOW-NEXT:    ret;
 ;
 ; ALLOW-LABEL: fadd_f32_global(
 ; ALLOW:       {
@@ -47,29 +50,30 @@ define float @fadd_f32_global(ptr addrspace(1) %addr, float %val) {
   ret float %r
 }
 
-; generic f32: we can't know the memory space, so default always uses a CAS loop;
-; the flag forces atom.add.
+; 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.
 define float @fadd_f32_generic(ptr %addr, float %val) {
-; DEFAULT-LABEL: fadd_f32_generic(
-; DEFAULT:       {
-; DEFAULT-NEXT:    .reg .pred %p<2>;
-; DEFAULT-NEXT:    .reg .b32 %r<5>;
-; DEFAULT-NEXT:    .reg .b64 %rd<2>;
-; DEFAULT-EMPTY:
-; DEFAULT-NEXT:  // %bb.0:
-; DEFAULT-NEXT:    ld.param.b32 %r2, [fadd_f32_generic_param_1];
-; DEFAULT-NEXT:    ld.param.b64 %rd1, [fadd_f32_generic_param_0];
-; DEFAULT-NEXT:    ld.relaxed.sys.b32 %r4, [%rd1];
-; DEFAULT-NEXT:  $L__BB1_1: // %atomicrmw.start
-; DEFAULT-NEXT:    // =>This Inner Loop Header: Depth=1
-; DEFAULT-NEXT:    add.rn.f32 %r3, %r4, %r2;
-; DEFAULT-NEXT:    atom.relaxed.sys.cas.b32 %r1, [%rd1], %r4, %r3;
-; DEFAULT-NEXT:    setp.ne.b32 %p1, %r1, %r4;
-; DEFAULT-NEXT:    mov.b32 %r4, %r1;
-; DEFAULT-NEXT:    @%p1 bra $L__BB1_1;
-; DEFAULT-NEXT:  // %bb.2: // %atomicrmw.end
-; DEFAULT-NEXT:    st.param.b32 [func_retval0], %r1;
-; DEFAULT-NEXT:    ret;
+;
+; DISALLOW-LABEL: fadd_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, [fadd_f32_generic_param_1];
+; DISALLOW-NEXT:    ld.param.b64 %rd1, [fadd_f32_generic_param_0];
+; DISALLOW-NEXT:    ld.relaxed.sys.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.relaxed.sys.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: fadd_f32_generic(
 ; ALLOW:       {
diff --git a/llvm/test/CodeGen/NVPTX/atomicrmw-sm60.ll b/llvm/test/CodeGen/NVPTX/atomicrmw-sm60.ll
index 40927ea1f68b485..5a43d78501ae0ea 100644
--- a/llvm/test/CodeGen/NVPTX/atomicrmw-sm60.ll
+++ b/llvm/test/CodeGen/NVPTX/atomicrmw-sm60.ll
@@ -1,7 +1,15 @@
 ; 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 | FileCheck %s --check-prefixes=SM60,SM60-NOFTZ
-; RUN: llc < %s -march=nvptx64 -mcpu=sm_60 -mattr=+ptx50 -denormal-fp-math-f32=preserve-sign | FileCheck %s --check-prefixes=SM60,SM60-FTZ
+; 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 -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: %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 %}
+; RUN: %if ptxas-sm_60 && ptxas-isa-5.0 %{ llc < %s -march=nvptx64 -mcpu=sm_60 -mattr=+ptx50 -denormal-fp-math-f32=preserve-sign -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 -denormal-fp-math-f32=preserve-sign -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 -denormal-fp-math-f32=preserve-sign | %ptxas-verify -arch=sm_60 %}
 
 define i8 @xchg_acq_rel_i8_global_cta(ptr addrspace(1) %addr, i8 %val) {
@@ -1739,314 +1747,592 @@ define i64 @usub_sat_acq_rel_i64_global_cta(ptr addrspace(1) %addr, i64 %val) {
 }
 
 define float @fadd_acq_rel_float_global_cta(ptr addrspace(1) %addr, float %val) {
-; SM60-NOFTZ-LABEL: fadd_acq_rel_float_global_cta(
-; SM60-NOFTZ:       {
-; SM60-NOFTZ-NEXT:    .reg .pred %p<2>;
-; SM60-NOFTZ-NEXT:    .reg .b32 %r<5>;
-; SM60-NOFTZ-NEXT:    .reg .b64 %rd<2>;
-; SM60-NOFTZ-EMPTY:
-; SM60-NOFTZ-NEXT:  // %bb.0:
-; SM60-NOFTZ-NEXT:    ld.param.b32 %r2, [fadd_acq_rel_float_global_cta_param_1];
-; SM60-NOFTZ-NEXT:    ld.param.b64 %rd1, [fadd_acq_rel_float_global_cta_param_0];
-; SM60-NOFTZ-NEXT:    membar.cta;
-; SM60-NOFTZ-NEXT:    ld.volatile.global.b32 %r4, [%rd1];
-; SM60-NOFTZ-NEXT:  $L__BB60_1: // %atomicrmw.start
-; SM60-NOFTZ-NEXT:    // =>This Inner Loop Header: Depth=1
-; SM60-NOFTZ-NEXT:    add.rn.f32 %r3, %r4, %r2;
-; SM60-NOFTZ-NEXT:    atom.cta.global.cas.b32 %r1, [%rd1], %r4, %r3;
-; SM60-NOFTZ-NEXT:    setp.ne.b32 %p1, %r1, %r4;
-; SM60-NOFTZ-NEXT:    mov.b32 %r4, %r1;
-; SM60-NOFTZ-NEXT:    @%p1 bra $L__BB60_1;
-; SM60-NOFTZ-NEXT:  // %bb.2: // %atomicrmw.end
-; SM60-NOFTZ-NEXT:    membar.cta;
-; SM60-NOFTZ-NEXT:    st.param.b32 [func_retval0], %r1;
-; SM60-NOFTZ-NEXT:    ret;
+; SM60-NOFTZ-DISALLOW-LABEL: fadd_acq_rel_float_global_cta(
+; SM60-NOFTZ-DISALLOW:       {
+; SM60-NOFTZ-DISALLOW-NEXT:    .reg .pred %p<2>;
+; SM60-NOFTZ-DISALLOW-NEXT:    .reg .b32 %r<5>;
+; SM60-NOFTZ-DISALLOW-NEXT:    .reg .b64 %rd<2>;
+; SM60-NOFTZ-DISALLOW-EMPTY:
+; SM60-NOFTZ-DISALLOW-NEXT:  // %bb.0:
+; SM60-NOFTZ-DISALLOW-NEXT:    ld.param.b32 %r2, [fadd_acq_rel_float_global_cta_param_1];
+; SM60-NOFTZ-DISALLOW-NEXT:    ld.param.b64 %rd1, [fadd_acq_rel_float_global_cta_param_0];
+; SM60-NOFTZ-DISALLOW-NEXT:    membar.cta;
+; SM60-NOFTZ-DISALLOW-NEXT:    ld.volatile.global.b32 %r4, [%rd1];
+; SM60-NOFTZ-DISALLOW-NEXT:  $L__BB60_1: // %atomicrmw.start
+; SM60-NOFTZ-DISALLOW-NEXT:    // =>This Inner Loop Header: Depth=1
+; SM60-NOFTZ-DISALLOW-NEXT:    add.rn.f32 %r3, %r4, %r2;
+; SM60-NOFTZ-DISALLOW-NEXT:    atom.cta.global.cas.b32 %r1, [%rd1], %r4, %r3;
+; SM60-NOFTZ-DISALLOW-NEXT:    setp.ne.b32 %p1, %r1, %r4;
+; SM60-NOFTZ-DISALLOW-NEXT:    mov.b32 %r4, %r1;
+; SM60-NOFTZ-DISALLOW-NEXT:    @%p1 bra $L__BB60_1;
+; SM60-NOFTZ-DISALLOW-NEXT:  // %bb.2: // %atomicrmw.end
+; SM60-NOFTZ-DISALLOW-NEXT:    membar.cta;
+; SM60-NOFTZ-DISALLOW-NEXT:    st.param.b32 [func_retval0], %r1;
+; SM60-NOFTZ-DISALLOW-NEXT:    ret;
 ;
-; SM60-FTZ-LABEL: fadd_acq_rel_float_global_cta(
-; SM60-FTZ:       {
-; SM60-FTZ-NEXT:    .reg .b32 %r<3>;
-; SM60-FTZ-NEXT:    .reg .b64 %rd<2>;
-; SM60-FTZ-EMPTY:
-; SM60-FTZ-NEXT:  // %bb.0:
-; SM60-FTZ-NEXT:    ld.param.b64 %rd1, [fadd_acq_rel_float_global_cta_param_0];
-; SM60-FTZ-NEXT:    ld.param.b32 %r1, [fadd_acq_rel_float_global_cta_param_1];
-; SM60-FTZ-NEXT:    atom.cta.global.add.f32 %r2, [%rd1], %r1;
-; SM60-FTZ-NEXT:    st.param.b32 [func_retval0], %r2;
-; SM60-FTZ-NEXT:    ret;
+; SM60-NOFTZ-ALLOW-LABEL: fadd_acq_rel_float_global_cta(
+; SM60-NOFTZ-ALLOW:       {
+; SM60-NOFTZ-ALLOW-NEXT:    .reg .b32 %r<3>;
+; SM60-NOFTZ-ALLOW-NEXT:    .reg .b64 %rd<2>;
+; SM60-NOFTZ-ALLOW-EMPTY:
+; SM60-NOFTZ-ALLOW-NEXT:  // %bb.0:
+; SM60-NOFTZ-ALLOW-NEXT:    ld.param.b64 %rd1, [fadd_acq_rel_float_global_cta_param_0];
+; SM60-NOFTZ-ALLOW-NEXT:    ld.param.b32 %r1, [fadd_acq_rel_float_global_cta_param_1];
+; SM60-NOFTZ-ALLOW-NEXT:    atom.cta.global.add.f32 %r2, [%rd1], %r1;
+; SM60-NOFTZ-ALLOW-NEXT:    st.param.b32 [func_retval0], %r2;
+; SM60-NOFTZ-ALLOW-NEXT:    ret;
+;
+; SM60-FTZ-DISALLOW-LABEL: fadd_acq_rel_float_global_cta(
+; SM60-FTZ-DISALLOW:       {
+; SM60-FTZ-DISALLOW-NEXT:    .reg .b32 %r<3>;
+; SM60-FTZ-DISALLOW-NEXT:    .reg .b64 %rd<2>;
+; SM60-FTZ-DISALLOW-EMPTY:
+; SM60-FTZ-DISALLOW-NEXT:  // %bb.0:
+; SM60-FTZ-DISALLOW-NEXT:    ld.param.b64 %rd1, [fadd_acq_rel_float_global_cta_param_0];
+; SM60-FTZ-DISALLOW-NEXT:    ld.param.b32 %r1, [fadd_acq_rel_float_global_cta_param_1];
+; SM60-FTZ-DISALLOW-NEXT:    atom.cta.global.add.f32 %r2, [%rd1], %r1;
+; SM60-FTZ-DISALLOW-NEXT:    st.param.b32 [func_retval0], %r2;
+; SM60-FTZ-DISALLOW-NEXT:    ret;
+;
+; SM60-FTZ-ALLOW-LABEL: fadd_acq_rel_float_global_cta(
+; SM60-FTZ-ALLOW:       {
+; SM60-FTZ-ALLOW-NEXT:    .reg .b32 %r<3>;
+; SM60-FTZ-ALLOW-NEXT:    .reg .b64 %rd<2>;
+; SM60-FTZ-ALLOW-EMPTY:
+; SM60-FTZ-ALLOW-NEXT:  // %bb.0:
+; SM60-FTZ-ALLOW-NEXT:    ld.param.b64 %rd1, [fadd_acq_rel_float_global_cta_param_0];
+; SM60-FTZ-ALLOW-NEXT:    ld.param.b32 %r1, [fadd_acq_rel_float_global_cta_param_1];
+; SM60-FTZ-ALLOW-NEXT:    atom.cta.global.add.f32 %r2, [%rd1], %r1;
+; SM60-FTZ-ALLOW-NEXT:    st.param.b32 [func_retval0], %r2;
+; SM60-FTZ-ALLOW-NEXT:    ret;
         %retval = atomicrmw fadd ptr  addrspace(1) %addr, float %val syncscope("block") acq_rel
         ret float %retval
 }
 
 define float @fsub_acq_rel_float_global_cta(ptr addrspace(1) %addr, float %val) {
-; SM60-NOFTZ-LABEL: fsub_acq_rel_float_global_cta(
-; SM60-NOFTZ:       {
-; SM60-NOFTZ-NEXT:    .reg .pred %p<2>;
-; SM60-NOFTZ-NEXT:    .reg .b32 %r<5>;
-; SM60-NOFTZ-NEXT:    .reg .b64 %rd<2>;
-; SM60-NOFTZ-EMPTY:
-; SM60-NOFTZ-NEXT:  // %bb.0:
-; SM60-NOFTZ-NEXT:    ld.param.b32 %r2, [fsub_acq_rel_float_global_cta_param_1];
-; SM60-NOFTZ-NEXT:    ld.param.b64 %rd1, [fsub_acq_rel_float_global_cta_param_0];
-; SM60-NOFTZ-NEXT:    membar.cta;
-; SM60-NOFTZ-NEXT:    ld.volatile.global.b32 %r4, [%rd1];
-; SM60-NOFTZ-NEXT:  $L__BB61_1: // %atomicrmw.start
-; SM60-NOFTZ-NEXT:    // =>This Inner Loop Header: Depth=1
-; SM60-NOFTZ-NEXT:    sub.rn.f32 %r3, %r4, %r2;
-; SM60-NOFTZ-NEXT:    atom.cta.global.cas.b32 %r1, [%rd1], %r4, %r3;
-; SM60-NOFTZ-NEXT:    setp.ne.b32 %p1, %r1, %r4;
-; SM60-NOFTZ-NEXT:    mov.b32 %r4, %r1;
-; SM60-NOFTZ-NEXT:    @%p1 bra $L__BB61_1;
-; SM60-NOFTZ-NEXT:  // %bb.2: // %atomicrmw.end
-; SM60-NOFTZ-NEXT:    membar.cta;
-; SM60-NOFTZ-NEXT:    st.param.b32 [func_retval0], %r1;
-; SM60-NOFTZ-NEXT:    ret;
+; SM60-NOFTZ-DISALLOW-LABEL: fsub_acq_rel_float_global_cta(
+; SM60-NOFTZ-DISALLOW:       {
+; SM60-NOFTZ-DISALLOW-NEXT:    .reg .pred %p<2>;
+; SM60-NOFTZ-DISALLOW-NEXT:    .reg .b32 %r<5>;
+; SM60-NOFTZ-DISALLOW-NEXT:    .reg .b64 %rd<2>;
+; SM60-NOFTZ-DISALLOW-EMPTY:
+; SM60-NOFTZ-DISALLOW-NEXT:  // %bb.0:
+; SM60-NOFTZ-DISALLOW-NEXT:    ld.param.b32 %r2, [fsub_acq_rel_float_global_cta_param_1];
+; SM60-NOFTZ-DISALLOW-NEXT:    ld.param.b64 %rd1, [fsub_acq_rel_float_global_cta_param_0];
+; SM60-NOFTZ-DISALLOW-NEXT:    membar.cta;
+; SM60-NOFTZ-DISALLOW-NEXT:    ld.volatile.global.b32 %r4, [%rd1];
+; SM60-NOFTZ-DISALLOW-NEXT:  $L__BB61_1: // %atomicrmw.start
+; SM60-NOFTZ-DISALLOW-NEXT:    // =>This Inner Loop Header: Depth=1
+; SM60-NOFTZ-DISALLOW-NEXT:    sub.rn.f32 %r3, %r4, %r2;
+; SM60-NOFTZ-DISALLOW-NEXT:    atom.cta.global.cas.b32 %r1, [%rd1], %r4, %r3;
+; SM60-NOFTZ-DISALLOW-NEXT:    setp.ne.b32 %p1, %r1, %r4;
+; SM60-NOFTZ-DISALLOW-NEXT:    mov.b32 %r4, %r1;
+; SM60-NOFTZ-DISALLOW-NEXT:    @%p1 bra $L__BB61_1;
+; SM60-NOFTZ-DISALLOW-NEXT:  // %bb.2: // %atomicrmw.end
+; SM60-NOFTZ-DISALLOW-NEXT:    membar.cta;
+; SM60-NOFTZ-DISALLOW-NEXT:    st.param.b32 [func_retval0], %r1;
+; SM60-NOFTZ-DISALLOW-NEXT:    ret;
+;
+; SM60-NOFTZ-ALLOW-LABEL: fsub_acq_rel_float_global_cta(
+; SM60-NOFTZ-ALLOW:       {
+; SM60-NOFTZ-ALLOW-NEXT:    .reg .pred %p<2>;
+; SM60-NOFTZ-ALLOW-NEXT:    .reg .b32 %r<5>;
+; SM60-NOFTZ-ALLOW-NEXT:    .reg .b64 %rd<2>;
+; SM60-NOFTZ-ALLOW-EMPTY:
+; SM60-NOFTZ-ALLOW-NEXT:  // %bb.0:
+; SM60-NOFTZ-ALLOW-NEXT:    ld.param.b32 %r2, [fsub_acq_rel_float_global_cta_param_1];
+; SM60-NOFTZ-ALLOW-NEXT:    ld.param.b64 %rd1, [fsub_acq_rel_float_global_cta_param_0];
+; SM60-NOFTZ-ALLOW-NEXT:    membar.cta;
+; SM60-NOFTZ-ALLOW-NEXT:    ld.volatile.global.b32 %r4, [%rd1];
+; SM60-NOFTZ-ALLOW-NEXT:  $L__BB61_1: // %atomicrmw.start
+; SM60-NOFTZ-ALLOW-NEXT:    // =>This Inner Loop Header: Depth=1
+; SM60-NOFTZ-ALLOW-NEXT:    sub.rn.f32 %r3, %r4, %r2;
+; SM60-NOFTZ-ALLOW-NEXT:    atom.cta.global.cas.b32 %r1, [%rd1], %r4, %r3;
+; SM60-NOFTZ-ALLOW-NEXT:    setp.ne.b32 %p1, %r1, %r4;
+; SM60-NOFTZ-ALLOW-NEXT:    mov.b32 %r4, %r1;
+; SM60-NOFTZ-ALLOW-NEXT:    @%p1 bra $L__BB61_1;
+; SM60-NOFTZ-ALLOW-NEXT:  // %bb.2: // %atomicrmw.end
+; SM60-NOFTZ-ALLOW-NEXT:    membar.cta;
+; SM60-NOFTZ-ALLOW-NEXT:    st.param.b32 [func_retval0], %r1;
+; SM60-NOFTZ-ALLOW-NEXT:    ret;
 ;
-; SM60-FTZ-LABEL: fsub_acq_rel_float_global_cta(
-; SM60-FTZ:       {
-; SM60-FTZ-NEXT:    .reg .pred %p<2>;
-; SM60-FTZ-NEXT:    .reg .b32 %r<5>;
-; SM60-FTZ-NEXT:    .reg .b64 %rd<2>;
-; SM60-FTZ-EMPTY:
-; SM60-FTZ-NEXT:  // %bb.0:
-; SM60-FTZ-NEXT:    ld.param.b32 %r2, [fsub_acq_rel_float_global_cta_param_1];
-; SM60-FTZ-NEXT:    ld.param.b64 %rd1, [fsub_acq_rel_float_global_cta_param_0];
-; SM60-FTZ-NEXT:    membar.cta;
-; SM60-FTZ-NEXT:    ld.volatile.global.b32 %r4, [%rd1];
-; SM60-FTZ-NEXT:  $L__BB61_1: // %atomicrmw.start
-; SM60-FTZ-NEXT:    // =>This Inner Loop Header: Depth=1
-; SM60-FTZ-NEXT:    sub.rn.ftz.f32 %r3, %r4, %r2;
-; SM60-FTZ-NEXT:    atom.cta.global.cas.b32 %r1, [%rd1], %r4, %r3;
-; SM60-FTZ-NEXT:    setp.ne.b32 %p1, %r1, %r4;
-; SM60-FTZ-NEXT:    mov.b32 %r4, %r1;
-; SM60-FTZ-NEXT:    @%p1 bra $L__BB61_1;
-; SM60-FTZ-NEXT:  // %bb.2: // %atomicrmw.end
-; SM60-FTZ-NEXT:    membar.cta;
-; SM60-FTZ-NEXT:    st.param.b32 [func_retval0], %r1;
-; SM60-FTZ-NEXT:    ret;
+; SM60-FTZ-DISALLOW-LABEL: fsub_acq_rel_float_global_cta(
+; SM60-FTZ-DISALLOW:       {
+; SM60-FTZ-DISALLOW-NEXT:    .reg .pred %p<2>;
+; SM60-FTZ-DISALLOW-NEXT:    .reg .b32 %r<5>;
+; SM60-FTZ-DISALLOW-NEXT:    .reg .b64 %rd<2>;
+; SM60-FTZ-DISALLOW-EMPTY:
+; SM60-FTZ-DISALLOW-NEXT:  // %bb.0:
+; SM60-FTZ-DISALLOW-NEXT:    ld.param.b32 %r2, [fsub_acq_rel_float_global_cta_param_1];
+; SM60-FTZ-DISALLOW-NEXT:    ld.param.b64 %rd1, [fsub_acq_rel_float_global_cta_param_0];
+; SM60-FTZ-DISALLOW-NEXT:    membar.cta;
+; SM60-FTZ-DISALLOW-NEXT:    ld.volatile.global.b32 %r4, [%rd1];
+; SM60-FTZ-DISALLOW-NEXT:  $L__BB61_1: // %atomicrmw.start
+; SM60-FTZ-DISALLOW-NEXT:    // =>This Inner Loop Header: Dep...
[truncated]

llvmbot added this to the LLVM 23.x Release milestone Oct 5, 2026

github-project-automation Bot added this to LLVM Release Status Oct 5, 2026

github-project-automation Bot moved this to Needs Triage in LLVM Release Status Oct 5, 2026

llvmbot requested a review from rnk October 5, 2026 08:18

llvmorg-github-actions Bot added the backend:NVPTX label Oct 5, 2026

llvmbot mentioned this pull request Oct 5, 2026

[NVPTX] Set default value of nvptx-allow-ftz-atomics to true #206154

Merged

Artem-B approved these changes Oct 10, 2026

View reviewed changes

rnk approved these changes Oct 10, 2026

View reviewed changes

tru closed this Oct 12, 2026

tru force-pushed the issue206154 branch from 8fecbd5 to 8e044a8 Compare October 12, 2026 14:55

tru moved this from Needs Merge to Done in LLVM Release Status Oct 13, 2026

Navigation Menu

Repository navigation

release/23.x: [NVPTX] Set default value of nvptx-allow-ftz-atomics to true (#206154) - #229003

release/23.x: [NVPTX] Set default value of nvptx-allow-ftz-atomics to true (#206154)#229003
llvmbot wants to merge 0 commit into
llvm:release/23.xfrom
llvmbot:issue206154

llvmbot commented Oct 5, 2026

Uh oh!

llvmbot commented Oct 5, 2026

Uh oh!

llvmorg-github-actions Bot commented Oct 5, 2026

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants

Conversation

llvmbot commented Oct 5, 2026

Uh oh!

llvmbot commented Oct 5, 2026

Uh oh!

llvmorg-github-actions Bot commented Oct 5, 2026

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants