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 _______________________________________________ llvm-branch-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
