Repository navigation
[NVPTX] Set default value of nvptx-allow-ftz-atomics to true - #206154
Conversation
|
@llvm/pr-subscribers-backend-nvptx Author: Artem Belevich (Artem-B) ChangesThis 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. Patch is 615.98 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/206154.diff 7 Files Affected:
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index 1c5caf124c17e..315bf5b9f58b0 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -109,14 +109,14 @@ 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 727c13bb4e47e..c1d223082eb30 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 5eb508eb12e17..c85a1eaa9c982 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,322 +1747,608 @@ 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: Depth=1
+; SM60-FTZ...
[truncated]
|
|
✅ With the latest revision this PR passed the C/C++ code formatter. |
|
FWIW I'm in favor of this. The associativity of the float operations is already non-deterministic since it is dependent on race conditions anyway so it seems like users are probably also going to be fine with a little FTZ flushing here and there too. |
|
To add some clarity that I don't think is obvious from the PR description: The issue here is with code that does something like this: If atomic_add is issued in a convergent manner (i.e. all threads run it at the same time), then all threads will read the same value for *addr. If we change atomic_add into a CAS loop, then different threads may read different values for *addr. Any code which relies on getting the same value for *addr has a race condition. Moreover because atomics are not convergent operations in LLVM, the compiler is free to expose this race even in the case where the atomic is not lowered to a CAS loop. For example the compiler is free to do the following transformation This would expose the same problem in the user code. If we want not to expose this problem in user code, we'd need to mark atomic_add as convergent. If we don't mark atomic_add as convergent, and someone in LLVM checks in a change that makes us more likely to make a transformation like above, are we going to revert that change too? If not, why not? What is the principle here? Personally I don't think that we should disable a correctness fix (or any other fix) in order to preserve the behavior of programs that have race conditions. You have a clearly buggy program, the compiler change exposes the bug, that is not our problem. Disabling this correctness fix for performance reasons is a different story and I think that's kind of what #206154 (comment) is getting at. (The theory being, the correctness issue is not that big and in many cases not even noticeable. I'm not 100% sure that's true, but I think it could be correct.) Art says that he wants to flip it back once there's a PTX version that supports native ftz and non-ftz PTX atomic ops. Apparently this is coming soon. But I am skeptical that will actually solve his problem, for two reasons.
Having said all this, I don't object to this flag flip. It just doesn't make a lot of sense to me. |
|
I can't say that I'm happy to find myself in the position where I have to support preserving a somewhat broken status quo, but for me it boils down to logistical trade-off based on risk/reward balance, and how it affects the users' ability to reason about their code.
In the end, the correctness improvement is real, but its impact is negligible in practice, and the escape hatch is still there to enable the correct behavior in the infrequent cases where it matters for the user. On the other hand, the performance impact will be widespread, and will affect most of the users who likely want the native atomic ops. It's just whole lot of extra work for everyone (even if all they have to do is add Another argument for making it an opt-in is that the change exposed a number of higher level questions that we may need to figure out first, before making the new behavior the default. It's mostly about the users being able to reason about the FP atomics. E.g.
|
|
For what it's worth, we may not be the first ones discussing this. Clang appears to have llvm-project/clang/include/clang/Options/Options.td Lines 2538 to 2542 in 5447679 It may be interesting to dig up their motivation for the option and the rationale for choosing the default. Also, here's what Fable produced when I asked it's opinion on the choice of the default: https://artem-b.github.io/nvptx-ftz-atomics-default-assessment.html The key point it produced is that:
|
c363b47 to
689a0cb
Compare
|
Any objections to committing this patch? I will follow up with per-instruction |
This is a follow-up to llvm#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.
689a0cb to
8e13bdf
Compare
|
To check my understanding: there's a bug in the user code sample. Technically speaking, programmers should insert On the other hand, we have the issue of floating point accuracy and correctness of not flushing to zero when the programmer didn't explicitly ask for it, so we're weighing FP precision against breaking technically incorrect user code. I think the LLVM way is to worry less about breaking users relying on UB, but to offer migration paths for downstream toolchain vendors, which is what this flag does. It stands to reason that upstream would take the stance of being more technically correct, allow downstreams to complete this migration to align with the upstream default at some point in the future, and then we'd move on. I'm not an FP expert, but there's a bell going off in my head from all the graphics programmers I've met over the years who have insisted to me that denormals are a performance trap, so at least for real-time applications, you always want FTZ. Casual research shows that CUDA, Triton, and XLA all default to FTZ, so it seems like the blast radius here should be pretty small, and I'm surprised we're considering this workaround. It sounds like noftz is a build configuration you maintain for testing purposes to measure the impact of FTZ on the output. |
UB code is how we found the issue, but since then, I've got a few more users that were affected by performance regression on a valid code. We do have the escape hatch option, so it's not a showstopper, but I'm willing to bet that there are more affected users why didn't notice the problem yet. The more important issue is that we currently have no way to let users specify exactly what they want. I.e. if a function uses generic pointers (but the user knows the address space) LLVM will always use CAS. Or if the user mixes global and shared atomics in the same functions, only one of them will be fast, and the other will use CAS, regardless of FTZ setting.
I'd say that they are more keen on "you always want fast", and it's not the same as "you always want FTZ" (not on NVIDIA GPUs, at least), and we currently do not have sufficiently flexible controls to let them achieve that goal.
Whichever FTZ mode they pick, some of the atomics will end up as CAS. Most of the time the impact will be negligible, but there are cases where regression is noticeable-to-catastrophic. My proposal is to roll back default to the old behavior to avoid unnecessarily disrupting the users with valid code while we have no good solution for them, implement AMD-style per-instruction FTZ control, and then switch back to "correct by default". If/when somebody needs atomic to produce just-so variant, we can point to the docs and tell them how they can express that in IR or CUDA. In short, I agree with your points and and that LLVM should generate correct code by default, I'm trying to do it in the way that's less painful for the users. Considering that we lived with the old behavior for a long time with nobody complaining about it, and that the current correct code causes noticeable perf regressions, I think it makes sense to avoid disrupting users until we can give them a usable way to achieve what they need. An alternative is to tell every affected user to use the escape hatch option, but then we'll have a problem of removing that option later on, as we'll grow an unknown number of users that depend on it. |
rnk
left a comment
rnk
left a comment
There was a problem hiding this comment.
I talked to some folks here at NVIDIA, and I think I agree with Art. I can't make hard commitments about product release dates and which fixes will be in them, but let's revisit this in the fall, and see if there are new ptxas features that allow us to fix this in a way that gives users the best possible behavior with the least disruption.
I'll go look for this in the issue tracker and write that down as the next step.
NVPTX_LLVM_Backend_jll 23.1.2+0 backports llvm/llvm-project#206154, which lowers global Float32 atomic additions to native, subnormal-flushing atom.add.f32 regardless of the kernel's denormal mode. Fast math therefore no longer changes atomic lowering, and the tests that expected a CAS loop (and preserved subnormals) in IEEE mode fail. Probe the option with a plain addition instead, which selects add.ftz.f32 under fast math, drop the atomic lowering tests, and remove the atomics motivation from the performance tutorial.
|
It seems that #200732 made it into the Clang 23 release and this change did not. From the discussion here and the fact that the PR landed, should we patch it into the next minor Clang 23 release? |
Yes, we probably do want to avoid atomics regression in a new release. Cherry-picking this change is a low-hanging fruit. @chsigg would be the person to talk to about that. |
|
I agree, we should cherry pick it, I think we have some scripts to automate that. |
|
Following https://llvm.org/docs/GitHub.html#backporting-fixes-to-the-release-branches /cherry-pick 2b6e9d2 |
|
/pull-request #229003 |
|
#229003 needs a stamp and I don't feel bold enough to do this myself given limited experience with NVPTX. |
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. (cherry picked from commit 2b6e9d2)
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.
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.
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.
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.
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.
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, #217587, #217588, #232725, #232726), which is supported in Clang, Flang, 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.
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.