[llvm-branch-commits] [llvm] [NVPTX] Honor !atomic.ignore.denormal.mode on atomicrmw fadd (PR #217586)
via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Fri Aug 21 01:28:45 PDT 2026
github-actions[bot] wrote:
<!--PREMERGE ADVISOR COMMENT: Linux-->
# :penguin: Linux x64 Test Results
* 181014 tests passed
* 3709 tests skipped
* 1 test failed
## Failed Tests
(click on a test name to see its output)
### LLVM
<details>
<summary>LLVM.CodeGen/NVPTX/atomicrmw-ignore-denormal-mode.ll</summary>
```
Exit Code: 1
Command Output (stdout):
--
# RUN: at line 2
/home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/llc < /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/atomicrmw-ignore-denormal-mode.ll -march=nvptx64 -mcpu=sm_90 -mattr=+ptx87 -nvptx-allow-ftz-atomics=false | /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/atomicrmw-ignore-denormal-mode.ll --check-prefixes=CHECK,STRICT
# executed command: /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/llc -march=nvptx64 -mcpu=sm_90 -mattr=+ptx87 -nvptx-allow-ftz-atomics=false
# note: command had no output on stdout or stderr
# executed command: /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/atomicrmw-ignore-denormal-mode.ll --check-prefixes=CHECK,STRICT
# .---command stderr------------
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/atomicrmw-ignore-denormal-mode.ll:28:16: error: STRICT-NEXT: expected string not found in input
# | ; STRICT-NEXT: ld.param.b32 %r2, [fadd_f32_global_ieee_no_md_param_1];
# | ^
# | <stdin>:20:10: note: scanning from here
# | // %bb.0:
# | ^
# | <stdin>:21:8: note: possible intended match here
# | ld.param::func.b32 %r2, [fadd_f32_global_ieee_no_md_param_1];
# | ^
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/atomicrmw-ignore-denormal-mode.ll:64:15: error: CHECK-NEXT: expected string not found in input
# | ; CHECK-NEXT: ld.param.b64 %rd1, [fadd_f32_global_ieee_md_param_0];
# | ^
# | <stdin>:45:10: note: scanning from here
# | // %bb.0:
# | ^
# | <stdin>:46:8: note: possible intended match here
# | ld.param::func.b64 %rd1, [fadd_f32_global_ieee_md_param_0];
# | ^
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/atomicrmw-ignore-denormal-mode.ll:82:15: error: CHECK-NEXT: expected string not found in input
# | ; CHECK-NEXT: ld.param.b64 %rd1, [fadd_f32_global_ftz_no_md_param_0];
# | ^
# | <stdin>:62:10: note: scanning from here
# | // %bb.0:
# | ^
# | <stdin>:63:8: note: possible intended match here
# | ld.param::func.b64 %rd1, [fadd_f32_global_ftz_no_md_param_0];
# | ^
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/atomicrmw-ignore-denormal-mode.ll:102:16: error: STRICT-NEXT: expected string not found in input
# | ; STRICT-NEXT: ld.param.b32 %r2, [fadd_f32_generic_ieee_no_md_param_1];
# | ^
# | <stdin>:80:10: note: scanning from here
# | // %bb.0:
# | ^
# | <stdin>:81:8: note: possible intended match here
# | ld.param::func.b32 %r2, [fadd_f32_generic_ieee_no_md_param_1];
# | ^
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/atomicrmw-ignore-denormal-mode.ll:138:15: error: CHECK-NEXT: expected string not found in input
# | ; CHECK-NEXT: ld.param.b64 %rd1, [fadd_f32_generic_ieee_md_param_0];
# | ^
# | <stdin>:105:10: note: scanning from here
# | // %bb.0:
# | ^
# | <stdin>:106:8: note: possible intended match here
# | ld.param::func.b64 %rd1, [fadd_f32_generic_ieee_md_param_0];
# | ^
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/atomicrmw-ignore-denormal-mode.ll:156:15: error: CHECK-NEXT: expected string not found in input
# | ; CHECK-NEXT: ld.param.b64 %rd1, [fadd_f32_shared_ieee_no_md_param_0];
# | ^
# | <stdin>:122:10: note: scanning from here
# | // %bb.0:
# | ^
# | <stdin>:123:8: note: possible intended match here
# | ld.param::func.b64 %rd1, [fadd_f32_shared_ieee_no_md_param_0];
# | ^
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/atomicrmw-ignore-denormal-mode.ll:176:16: error: STRICT-NEXT: expected string not found in input
# | ; STRICT-NEXT: ld.param.b16 %rs1, [fadd_f16_global_ftz_no_md_param_1];
# | ^
# | <stdin>:141:10: note: scanning from here
# | // %bb.0:
# | ^
# | <stdin>:142:8: note: possible intended match here
# | ld.param::func.b16 %rs1, [fadd_f16_global_ftz_no_md_param_1];
# | ^
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/atomicrmw-ignore-denormal-mode.ll:226:15: error: CHECK-NEXT: expected string not found in input
# | ; CHECK-NEXT: ld.param.b64 %rd1, [fadd_f16_global_ftz_md_param_0];
# | ^
# | <stdin>:180:10: note: scanning from here
# | // %bb.0:
# | ^
# | <stdin>:181:8: note: possible intended match here
# | ld.param::func.b64 %rd1, [fadd_f16_global_ftz_md_param_0];
# | ^
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/atomicrmw-ignore-denormal-mode.ll:246:15: error: CHECK-NEXT: expected string not found in input
# | ; CHECK-NEXT: ld.param.b64 %rd1, [fadd_f16_global_f32_ftz_only_no_md_param_0];
# | ^
# | <stdin>:197:10: note: scanning from here
# | // %bb.0:
# | ^
# | <stdin>:198:8: note: possible intended match here
# | ld.param::func.b64 %rd1, [fadd_f16_global_f32_ftz_only_no_md_param_0];
# | ^
# |
# | Input file: <stdin>
# | Check file: /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/atomicrmw-ignore-denormal-mode.ll
# |
# | -dump-input=help explains the following input dump.
# |
# | Input was:
# | <<<<<<
# | .
# | .
# | .
# | 15: {
# | 16: .reg .pred %p<2>;
# | 17: .reg .b32 %r<5>;
# | 18: .reg .b64 %rd<2>;
# | 19:
# | 20: // %bb.0:
# | next:28'0 { search range start (exclusive)
# | next:28'1 error: no match found in search range
# | 21: ld.param::func.b32 %r2, [fadd_f32_global_ieee_no_md_param_1];
# | next:28'2 ? possible intended match
# | 22: ld.param::func.b64 %rd1, [fadd_f32_global_ieee_no_md_param_0];
# | 23: ld.relaxed.sys.global.b32 %r4, [%rd1];
# | 24: $L__BB0_1: // %atomicrmw.start
# | 25: // =>This Inner Loop Header: Depth=1
# | 26: add.rn.f32 %r3, %r4, %r2;
# | .
# | .
# | .
# | 32: st.param::func.b32 [func_retval0], %r1;
# | 33: ret;
# | 34: // -- End function
# | 35: }
# | 36: // .globl fadd_f32_global_ieee_md // -- Begin function fadd_f32_global_ieee_md
# | 37: .visible .func (.param .b32 func_retval0) fadd_f32_global_ieee_md(
# | next:28'3 } search range end (exclusive)
# | 38: .param .b64 fadd_f32_global_ieee_md_param_0,
# | 39: .param .b32 fadd_f32_global_ieee_md_param_1
# | 40: ) // @fadd_f32_global_ieee_md
# | 41: {
# | 42: .reg .b32 %r<3>;
# | 43: .reg .b64 %rd<2>;
# | 44:
# | 45: // %bb.0:
# | next:64'0 { search range start (exclusive)
# | next:64'1 error: no match found in search range
# | 46: ld.param::func.b64 %rd1, [fadd_f32_global_ieee_md_param_0];
# | next:64'2 ? possible intended match
# | 47: ld.param::func.b32 %r1, [fadd_f32_global_ieee_md_param_1];
# | 48: atom.relaxed.sys.global.add.f32 %r2, [%rd1], %r1;
# | 49: st.param::func.b32 [func_retval0], %r2;
# | 50: ret;
# | 51: // -- End function
# | 52: }
# | 53: // .globl fadd_f32_global_ftz_no_md // -- Begin function fadd_f32_global_ftz_no_md
# | 54: .visible .func (.param .b32 func_retval0) fadd_f32_global_ftz_no_md(
# | next:64'3 } search range end (exclusive)
# | 55: .param .b64 fadd_f32_global_ftz_no_md_param_0,
# | 56: .param .b32 fadd_f32_global_ftz_no_md_param_1
# | 57: ) // @fadd_f32_global_ftz_no_md
# | 58: {
# | 59: .reg .b32 %r<3>;
# | 60: .reg .b64 %rd<2>;
# | 61:
# | 62: // %bb.0:
# | next:82'0 { search range start (exclusive)
# | next:82'1 error: no match found in search range
# | 63: ld.param::func.b64 %rd1, [fadd_f32_global_ftz_no_md_param_0];
# | next:82'2 ? possible intended match
# | 64: ld.param::func.b32 %r1, [fadd_f32_global_ftz_no_md_param_1];
# | 65: atom.relaxed.sys.global.add.f32 %r2, [%rd1], %r1;
# | 66: st.param::func.b32 [func_retval0], %r2;
# | 67: ret;
# | 68: // -- End function
# | 69: }
# | 70: // .globl fadd_f32_generic_ieee_no_md // -- Begin function fadd_f32_generic_ieee_no_md
# | 71: .visible .func (.param .b32 func_retval0) fadd_f32_generic_ieee_no_md(
# | next:82'3 } search range end (exclusive)
# | 72: .param .b64 fadd_f32_generic_ieee_no_md_param_0,
# | 73: .param .b32 fadd_f32_generic_ieee_no_md_param_1
# | 74: ) // @fadd_f32_generic_ieee_no_md
# | 75: {
# | 76: .reg .pred %p<2>;
# | 77: .reg .b32 %r<5>;
# | 78: .reg .b64 %rd<2>;
# | 79:
# | 80: // %bb.0:
# | next:102'0 { search range start (exclusive)
# | next:102'1 error: no match found in search range
# | 81: ld.param::func.b32 %r2, [fadd_f32_generic_ieee_no_md_param_1];
# | next:102'2 ? possible intended match
# | 82: ld.param::func.b64 %rd1, [fadd_f32_generic_ieee_no_md_param_0];
# | 83: ld.relaxed.sys.b32 %r4, [%rd1];
# | 84: $L__BB3_1: // %atomicrmw.start
# | 85: // =>This Inner Loop Header: Depth=1
# | 86: add.rn.f32 %r3, %r4, %r2;
# | .
# | .
# | .
# | 92: st.param::func.b32 [func_retval0], %r1;
# | 93: ret;
# | 94: // -- End function
# | 95: }
# | 96: // .globl fadd_f32_generic_ieee_md // -- Begin function fadd_f32_generic_ieee_md
# | 97: .visible .func (.param .b32 func_retval0) fadd_f32_generic_ieee_md(
# | next:102'3 } search range end (exclusive)
# | 98: .param .b64 fadd_f32_generic_ieee_md_param_0,
# | 99: .param .b32 fadd_f32_generic_ieee_md_param_1
# | 100: ) // @fadd_f32_generic_ieee_md
# | 101: {
# | 102: .reg .b32 %r<3>;
# | 103: .reg .b64 %rd<2>;
# | 104:
# | 105: // %bb.0:
# | next:138'0 { search range start (exclusive)
# | next:138'1 error: no match found in search range
# | 106: ld.param::func.b64 %rd1, [fadd_f32_generic_ieee_md_param_0];
# | next:138'2 ? possible intended match
# | 107: ld.param::func.b32 %r1, [fadd_f32_generic_ieee_md_param_1];
# | 108: atom.relaxed.sys.add.f32 %r2, [%rd1], %r1;
# | 109: st.param::func.b32 [func_retval0], %r2;
# | 110: ret;
# | 111: // -- End function
# | 112: }
# | 113: // .globl fadd_f32_shared_ieee_no_md // -- Begin function fadd_f
# | ...
# `---data was truncated (10240/15367) (change limit with -D output_limit=N)
# error: command failed with exit status: 1
--
```
</details>
If these failures are unrelated to your changes (for example tests are broken or flaky at HEAD), please open an issue at https://github.com/llvm/llvm-project/issues and add the `infrastructure` label.
https://github.com/llvm/llvm-project/pull/217586
More information about the llvm-branch-commits
mailing list