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

This is a follow-up to #200732.

Setting the default value of nvptx-allow-ftz-atomics to true preserves
the status quo. It allows the NVPTX backend to emit native
floating-point atomic add instructions by default, rather than falling
back to the CAS loop required for strict FTZ compliance.

Correct FTZ handling forces a lot of existing code to emit CAS loops.
This introduces thread divergence, breaking code that relies on atomic
operations being non-divergent (for example, warp-wide atomic additions
followed by a load of the reduced value without explicit
synchronization).

While FTZ-correct atomic behavior is desirable as a long-term default,
introducing thread divergence into a common operation that previously
"happened to work" without synchronization is highly disruptive.
Defaulting to true makes the FTZ-correct behavior opt-in for now,
allowing users time to audit and adapt their code.

Test cases are updated to verify both enabled and disabled states of the
option, alongside the new default behavior.

Author: Artem-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.

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:       {