[llvm-branch-commits] [llvm] release/23.x: [NVPTX] Set default value of nvptx-allow-ftz-atomics to true (#206154) (PR #229003)
via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Mon Oct 5 01:19:31 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-backend-nvptx
Author: llvmbot
<details>
<summary>Changes</summary>
Backport 2b6e9d261433f5ec8c49502f91d63dd7d07346c3
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
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]
``````````
</details>
https://github.com/llvm/llvm-project/pull/229003
More information about the llvm-branch-commits
mailing list