[llvm] [NVVM][NVPTX] Add tensor map override support in S2G and Reduce intrinsics (PR #215503)
via llvm-commits
llvm-commits at lists.llvm.org
Wed Aug 12 03:30:53 PDT 2026
github-actions[bot] wrote:
<!--PREMERGE ADVISOR COMMENT: Linux-->
# :penguin: Linux x64 Test Results
* 180398 tests passed
* 3665 tests skipped
* 2 tests failed
## Failed Tests
(click on a test name to see its output)
### LLVM
<details>
<summary>LLVM.CodeGen/NVPTX/cp-async-bulk-tensor-reduce-override.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/cp-async-bulk-tensor-reduce-override.ll -mtriple=nvptx64 -mcpu=sm_107f | /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/cp-async-bulk-tensor-reduce-override.ll
# executed command: /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/llc -mtriple=nvptx64 -mcpu=sm_107f
# 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/cp-async-bulk-tensor-reduce-override.ll
# .---command stderr------------
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/cp-async-bulk-tensor-reduce-override.ll:14:15: error: CHECK-NEXT: expected string not found in input
# | ; CHECK-NEXT: ld.param.b64 %rd1, [cp_async_bulk_tensor_reduce_tile_override_addr_param_0];
# | ^
# | <stdin>:26:10: note: scanning from here
# | // %bb.0:
# | ^
# | <stdin>:27:8: note: possible intended match here
# | ld.param::func.b64 %rd1, [cp_async_bulk_tensor_reduce_tile_override_addr_param_0];
# | ^
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/cp-async-bulk-tensor-reduce-override.ll:135:15: error: CHECK-NEXT: expected string not found in input
# | ; CHECK-NEXT: ld.param.b64 %rd1, [cp_async_bulk_tensor_reduce_im2col_override_addr_param_0];
# | ^
# | <stdin>:79:10: note: scanning from here
# | // %bb.0:
# | ^
# | <stdin>:80:8: note: possible intended match here
# | ld.param::func.b64 %rd1, [cp_async_bulk_tensor_reduce_im2col_override_addr_param_0];
# | ^
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/cp-async-bulk-tensor-reduce-override.ll:240:15: error: CHECK-NEXT: expected string not found in input
# | ; CHECK-NEXT: ld.param.b64 %rd1, [cp_async_bulk_tensor_reduce_im2col_w_override_addr_param_0];
# | ^
# | <stdin>:128:10: note: scanning from here
# | // %bb.0:
# | ^
# | <stdin>:129:8: note: possible intended match here
# | ld.param::func.b64 %rd1, [cp_async_bulk_tensor_reduce_im2col_w_override_addr_param_0];
# | ^
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/cp-async-bulk-tensor-reduce-override.ll:346:15: error: CHECK-NEXT: expected string not found in input
# | ; CHECK-NEXT: ld.param.b64 %rd1, [cp_async_bulk_tensor_reduce_tile_override_addr_dim_stride_param_0];
# | ^
# | <stdin>:188:10: note: scanning from here
# | // %bb.0:
# | ^
# | <stdin>:189:8: note: possible intended match here
# | ld.param::func.b64 %rd1, [cp_async_bulk_tensor_reduce_tile_override_addr_dim_stride_param_0];
# | ^
# |
# | Input file: <stdin>
# | Check file: /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/cp-async-bulk-tensor-reduce-override.ll
# |
# | -dump-input=help explains the following input dump.
# |
# | Input was:
# | <<<<<<
# | .
# | .
# | .
# | 21: )
# | 22: {
# | 23: .reg .b32 %r<6>;
# | 24: .reg .b64 %rd<5>;
# | 25:
# | 26: // %bb.0:
# | next:14'0 { search range start (exclusive)
# | next:14'1 error: no match found in search range
# | 27: ld.param::func.b64 %rd1, [cp_async_bulk_tensor_reduce_tile_override_addr_param_0];
# | next:14'2 ? possible intended match
# | 28: ld.param::func.b64 %rd2, [cp_async_bulk_tensor_reduce_tile_override_addr_param_1];
# | 29: ld.param::func.b64 %rd3, [cp_async_bulk_tensor_reduce_tile_override_addr_param_2];
# | 30: ld.param::func.b32 %r1, [cp_async_bulk_tensor_reduce_tile_override_addr_param_3];
# | 31: cp.reduce.async.bulk.tensor.1d.global.shared::cta.add.tile.bulk_group.override::global_address [%rd2, %rd3, {%r1}], [%rd1];
# | 32: ld.param::func.b32 %r2, [cp_async_bulk_tensor_reduce_tile_override_addr_param_4];
# | .
# | .
# | .
# | 59: cp.reduce.async.bulk.tensor.1d.global.shared::cta.xor.tile.bulk_group.override::global_address.L2::cache_hint [%rd2, %rd3, {%r1}], [%rd1], %rd4;
# | 60: ret;
# | 61: // -- End function
# | 62: }
# | 63: // .globl cp_async_bulk_tensor_reduce_im2col_override_addr // -- Begin function cp_async_bulk_tensor_reduce_im2col_override_addr
# | 64: .visible .func cp_async_bulk_tensor_reduce_im2col_override_addr(
# | next:14'3 } search range end (exclusive)
# | 65: .param .b64 cp_async_bulk_tensor_reduce_im2col_override_addr_param_0,
# | 66: .param .b64 cp_async_bulk_tensor_reduce_im2col_override_addr_param_1,
# | 67: .param .b64 cp_async_bulk_tensor_reduce_im2col_override_addr_param_2,
# | 68: .param .b32 cp_async_bulk_tensor_reduce_im2col_override_addr_param_3,
# | 69: .param .b32 cp_async_bulk_tensor_reduce_im2col_override_addr_param_4,
# | .
# | .
# | .
# | 74: ) // @cp_async_bulk_tensor_reduce_im2col_override_addr
# | 75: {
# | 76: .reg .b32 %r<6>;
# | 77: .reg .b64 %rd<5>;
# | 78:
# | 79: // %bb.0:
# | next:135'0 { search range start (exclusive)
# | next:135'1 error: no match found in search range
# | 80: ld.param::func.b64 %rd1, [cp_async_bulk_tensor_reduce_im2col_override_addr_param_0];
# | next:135'2 ? possible intended match
# | 81: ld.param::func.b64 %rd2, [cp_async_bulk_tensor_reduce_im2col_override_addr_param_1];
# | 82: ld.param::func.b64 %rd3, [cp_async_bulk_tensor_reduce_im2col_override_addr_param_2];
# | 83: ld.param::func.b32 %r1, [cp_async_bulk_tensor_reduce_im2col_override_addr_param_3];
# | 84: ld.param::func.b32 %r2, [cp_async_bulk_tensor_reduce_im2col_override_addr_param_4];
# | 85: ld.param::func.b32 %r3, [cp_async_bulk_tensor_reduce_im2col_override_addr_param_5];
# | .
# | .
# | .
# | 108: cp.reduce.async.bulk.tensor.3d.global.shared::cta.xor.im2col_no_offs.bulk_group.override::global_address.L2::cache_hint [%rd2, %rd3, {%r1, %r2, %r3}], [%rd1], %rd4;
# | 109: ret;
# | 110: // -- End function
# | 111: }
# | 112: // .globl cp_async_bulk_tensor_reduce_im2col_w_override_addr // -- Begin function cp_async_bulk_tensor_reduce_im2col_w_override_addr
# | 113: .visible .func cp_async_bulk_tensor_reduce_im2col_w_override_addr(
# | next:135'3 } search range end (exclusive)
# | 114: .param .b64 cp_async_bulk_tensor_reduce_im2col_w_override_addr_param_0,
# | 115: .param .b64 cp_async_bulk_tensor_reduce_im2col_w_override_addr_param_1,
# | 116: .param .b64 cp_async_bulk_tensor_reduce_im2col_w_override_addr_param_2,
# | 117: .param .b32 cp_async_bulk_tensor_reduce_im2col_w_override_addr_param_3,
# | 118: .param .b32 cp_async_bulk_tensor_reduce_im2col_w_override_addr_param_4,
# | .
# | .
# | .
# | 123: ) // @cp_async_bulk_tensor_reduce_im2col_w_override_addr
# | 124: {
# | 125: .reg .b32 %r<6>;
# | 126: .reg .b64 %rd<5>;
# | 127:
# | 128: // %bb.0:
# | next:240'0 { search range start (exclusive)
# | next:240'1 error: no match found in search range
# | 129: ld.param::func.b64 %rd1, [cp_async_bulk_tensor_reduce_im2col_w_override_addr_param_0];
# | next:240'2 ? possible intended match
# | 130: ld.param::func.b64 %rd2, [cp_async_bulk_tensor_reduce_im2col_w_override_addr_param_1];
# | 131: ld.param::func.b64 %rd3, [cp_async_bulk_tensor_reduce_im2col_w_override_addr_param_2];
# | 132: ld.param::func.b32 %r1, [cp_async_bulk_tensor_reduce_im2col_w_override_addr_param_3];
# | 133: ld.param::func.b32 %r2, [cp_async_bulk_tensor_reduce_im2col_w_override_addr_param_4];
# | 134: ld.param::func.b32 %r3, [cp_async_bulk_tensor_reduce_im2col_w_override_addr_param_5];
# | .
# | .
# | .
# | 157: cp.reduce.async.bulk.tensor.3d.global.shared::cta.xor.im2col_no_offs::w.bulk_group.override::global_address.L2::cache_hint [%rd2, %rd3, {%r1, %r2, %r3}], [%rd1], %rd4;
# | 158: ret;
# | 159: // -- End function
# | 160: }
# | 161: // .globl cp_async_bulk_tensor_reduce_tile_override_addr_dim_stride // -- Begin function cp_async_bulk_tensor_reduce_tile_override_addr_dim_stride
# | 162: .visible .func cp_async_bulk_tensor_reduce_tile_override_addr_dim_stride(
# | next:240'3 } search range end (exclusive)
# | 163: .param .b64 cp_async_bulk_tensor_reduce_tile_override_addr_dim_stride_param_0,
# | 164: .param .b64 cp_async_bulk_tensor_reduce_tile_override_addr_dim_stride_param_1,
# | 165: .param .b64 cp_async_bulk_tensor_reduce_tile_override_addr_dim_stride_param_2,
# | 166: .param .b32 cp_async_bulk_tensor_reduce_tile_override_addr_dim_stride_param_3,
# | 167: .param .b32 cp_async_bulk_tensor_reduce_tile_override_addr_dim_stride_param_4,
# | .
# | .
# | .
# | 183: {
# | 184: .reg .b16 %rs<7>;
# | 185: .reg .b32 %r<10>;
# | 186: .reg .b64 %rd<5>;
# | 187:
# | 188: // %bb.0:
# | next:346'0 { search range start (exclusive)
# | next:346'1 error: no match found in search range
# | 189: ld.param::func.b64 %rd1, [cp_async_bulk_tensor_reduce_tile_override_addr_dim_stride_param_0];
# | next:346'2 ? possible intended match
# | 190: ld.param::func.b64 %rd2, [cp_async_bulk_tensor_reduce_tile_override_addr_dim_stride_param_1];
# | 191: ld.param::func.b64 %rd3, [cp_async_bulk_tensor_reduce_tile_override_addr_dim_stride_param_2];
# | 192: ld.param::func.b16 %rs1, [cp_async_bulk_tensor_reduce_tile_override_addr_dim_stride_param_3];
# | 193: ld.param::func.b32 %r1, [cp_async_bulk_tensor_reduce_tile_override_addr_dim_stride_param_13];
# | 194: cp.reduce.async.bulk.tensor.1d.global.shared::cta.add.tile.bulk_group.override::global_address.override::global_dim [%rd2, %rd3, {%rs1}, {%r1}], [%rd1];
# | .
# | .
# | .
# | 243: cp.reduce.async.bulk.tensor.1d.global.shared::cta.xor.tile.bulk_group.override::global_address.override::global_dim.L2::cache_hint [%rd2, %rd3, {%rs1}, {%r1}
# | ...
# `---data was truncated (10240/10812) (change limit with -D output_limit=N)
# error: command failed with exit status: 1
--
```
</details>
<details>
<summary>LLVM.CodeGen/NVPTX/cp-async-bulk-tensor-s2g-override.ll</summary>
```
Exit Code: 1
Command Output (stdout):
--
# RUN: at line 3
/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/cp-async-bulk-tensor-s2g-override.ll -mtriple=nvptx64 -mcpu=sm_107f | /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/cp-async-bulk-tensor-s2g-override.ll
# executed command: /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/llc -mtriple=nvptx64 -mcpu=sm_107f
# 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/cp-async-bulk-tensor-s2g-override.ll
# .---command stderr------------
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/cp-async-bulk-tensor-s2g-override.ll:15:15: error: CHECK-NEXT: expected string not found in input
# | ; CHECK-NEXT: ld.param.b64 %rd1, [cp_async_bulk_tensor_s2g_tile_override_addr_param_0];
# | ^
# | <stdin>:26:10: note: scanning from here
# | // %bb.0:
# | ^
# | <stdin>:27:8: note: possible intended match here
# | ld.param::func.b64 %rd1, [cp_async_bulk_tensor_s2g_tile_override_addr_param_0];
# | ^
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/cp-async-bulk-tensor-s2g-override.ll:71:15: error: CHECK-NEXT: expected string not found in input
# | ; CHECK-NEXT: ld.param.b64 %rd1, [cp_async_bulk_tensor_s2g_im2col_override_addr_param_0];
# | ^
# | <stdin>:65:10: note: scanning from here
# | // %bb.0:
# | ^
# | <stdin>:66:8: note: possible intended match here
# | ld.param::func.b64 %rd1, [cp_async_bulk_tensor_s2g_im2col_override_addr_param_0];
# | ^
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/cp-async-bulk-tensor-s2g-override.ll:115:15: error: CHECK-NEXT: expected string not found in input
# | ; CHECK-NEXT: ld.param.b64 %rd1, [cp_async_bulk_tensor_s2g_im2col_w_override_addr_param_0];
# | ^
# | <stdin>:100:10: note: scanning from here
# | // %bb.0:
# | ^
# | <stdin>:101:8: note: possible intended match here
# | ld.param::func.b64 %rd1, [cp_async_bulk_tensor_s2g_im2col_w_override_addr_param_0];
# | ^
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/cp-async-bulk-tensor-s2g-override.ll:160:15: error: CHECK-NEXT: expected string not found in input
# | ; CHECK-NEXT: ld.param.b64 %rd1, [cp_async_bulk_tensor_s2g_tile_override_addr_dim_stride_param_0];
# | ^
# | <stdin>:146:10: note: scanning from here
# | // %bb.0:
# | ^
# | <stdin>:147:8: note: possible intended match here
# | ld.param::func.b64 %rd1, [cp_async_bulk_tensor_s2g_tile_override_addr_dim_stride_param_0];
# | ^
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/cp-async-bulk-tensor-s2g-override.ll:226:15: error: CHECK-NEXT: expected string not found in input
# | ; CHECK-NEXT: ld.param.b64 %rd1, [cp_async_bulk_tensor_s2g_tile_scatter4_2d_param_1];
# | ^
# | <stdin>:196:10: note: scanning from here
# | // %bb.0:
# | ^
# | <stdin>:197:8: note: possible intended match here
# | ld.param::func.b64 %rd1, [cp_async_bulk_tensor_s2g_tile_scatter4_2d_param_1];
# | ^
# |
# | Input file: <stdin>
# | Check file: /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/cp-async-bulk-tensor-s2g-override.ll
# |
# | -dump-input=help explains the following input dump.
# |
# | Input was:
# | <<<<<<
# | .
# | .
# | .
# | 21: )
# | 22: {
# | 23: .reg .b32 %r<6>;
# | 24: .reg .b64 %rd<5>;
# | 25:
# | 26: // %bb.0:
# | next:15'0 { search range start (exclusive)
# | next:15'1 error: no match found in search range
# | 27: ld.param::func.b64 %rd1, [cp_async_bulk_tensor_s2g_tile_override_addr_param_0];
# | next:15'2 ? possible intended match
# | 28: ld.param::func.b64 %rd2, [cp_async_bulk_tensor_s2g_tile_override_addr_param_1];
# | 29: ld.param::func.b64 %rd3, [cp_async_bulk_tensor_s2g_tile_override_addr_param_2];
# | 30: ld.param::func.b32 %r1, [cp_async_bulk_tensor_s2g_tile_override_addr_param_3];
# | 31: cp.async.bulk.tensor.1d.global.shared::cta.tile.bulk_group.override::global_address [%rd2, %rd3, {%r1}], [%rd1];
# | 32: ld.param::func.b32 %r2, [cp_async_bulk_tensor_s2g_tile_override_addr_param_4];
# | .
# | .
# | .
# | 45: cp.async.bulk.tensor.5d.global.shared::cta.tile.bulk_group.override::global_address.L2::cache_hint [%rd2, %rd3, {%r1, %r2, %r3, %r4, %r5}], [%rd1], %rd4;
# | 46: ret;
# | 47: // -- End function
# | 48: }
# | 49: // .globl cp_async_bulk_tensor_s2g_im2col_override_addr // -- Begin function cp_async_bulk_tensor_s2g_im2col_override_addr
# | 50: .visible .func cp_async_bulk_tensor_s2g_im2col_override_addr(
# | next:15'3 } search range end (exclusive)
# | 51: .param .b64 cp_async_bulk_tensor_s2g_im2col_override_addr_param_0,
# | 52: .param .b64 cp_async_bulk_tensor_s2g_im2col_override_addr_param_1,
# | 53: .param .b64 cp_async_bulk_tensor_s2g_im2col_override_addr_param_2,
# | 54: .param .b32 cp_async_bulk_tensor_s2g_im2col_override_addr_param_3,
# | 55: .param .b32 cp_async_bulk_tensor_s2g_im2col_override_addr_param_4,
# | .
# | .
# | .
# | 60: ) // @cp_async_bulk_tensor_s2g_im2col_override_addr
# | 61: {
# | 62: .reg .b32 %r<6>;
# | 63: .reg .b64 %rd<5>;
# | 64:
# | 65: // %bb.0:
# | next:71'0 { search range start (exclusive)
# | next:71'1 error: no match found in search range
# | 66: ld.param::func.b64 %rd1, [cp_async_bulk_tensor_s2g_im2col_override_addr_param_0];
# | next:71'2 ? possible intended match
# | 67: ld.param::func.b64 %rd2, [cp_async_bulk_tensor_s2g_im2col_override_addr_param_1];
# | 68: ld.param::func.b64 %rd3, [cp_async_bulk_tensor_s2g_im2col_override_addr_param_2];
# | 69: ld.param::func.b32 %r1, [cp_async_bulk_tensor_s2g_im2col_override_addr_param_3];
# | 70: ld.param::func.b32 %r2, [cp_async_bulk_tensor_s2g_im2col_override_addr_param_4];
# | 71: ld.param::func.b32 %r3, [cp_async_bulk_tensor_s2g_im2col_override_addr_param_5];
# | .
# | .
# | .
# | 80: cp.async.bulk.tensor.5d.global.shared::cta.im2col_no_offs.bulk_group.override::global_address.L2::cache_hint [%rd2, %rd3, {%r1, %r2, %r3, %r4, %r5}], [%rd1], %rd4;
# | 81: ret;
# | 82: // -- End function
# | 83: }
# | 84: // .globl cp_async_bulk_tensor_s2g_im2col_w_override_addr // -- Begin function cp_async_bulk_tensor_s2g_im2col_w_override_addr
# | 85: .visible .func cp_async_bulk_tensor_s2g_im2col_w_override_addr(
# | next:71'3 } search range end (exclusive)
# | 86: .param .b64 cp_async_bulk_tensor_s2g_im2col_w_override_addr_param_0,
# | 87: .param .b64 cp_async_bulk_tensor_s2g_im2col_w_override_addr_param_1,
# | 88: .param .b64 cp_async_bulk_tensor_s2g_im2col_w_override_addr_param_2,
# | 89: .param .b32 cp_async_bulk_tensor_s2g_im2col_w_override_addr_param_3,
# | 90: .param .b32 cp_async_bulk_tensor_s2g_im2col_w_override_addr_param_4,
# | .
# | .
# | .
# | 95: ) // @cp_async_bulk_tensor_s2g_im2col_w_override_addr
# | 96: {
# | 97: .reg .b32 %r<6>;
# | 98: .reg .b64 %rd<5>;
# | 99:
# | 100: // %bb.0:
# | next:115'0 { search range start (exclusive)
# | next:115'1 error: no match found in search range
# | 101: ld.param::func.b64 %rd1, [cp_async_bulk_tensor_s2g_im2col_w_override_addr_param_0];
# | next:115'2 ? possible intended match
# | 102: ld.param::func.b64 %rd2, [cp_async_bulk_tensor_s2g_im2col_w_override_addr_param_1];
# | 103: ld.param::func.b64 %rd3, [cp_async_bulk_tensor_s2g_im2col_w_override_addr_param_2];
# | 104: ld.param::func.b32 %r1, [cp_async_bulk_tensor_s2g_im2col_w_override_addr_param_3];
# | 105: ld.param::func.b32 %r2, [cp_async_bulk_tensor_s2g_im2col_w_override_addr_param_4];
# | 106: ld.param::func.b32 %r3, [cp_async_bulk_tensor_s2g_im2col_w_override_addr_param_5];
# | .
# | .
# | .
# | 115: cp.async.bulk.tensor.5d.global.shared::cta.im2col_no_offs::w.bulk_group.override::global_address.L2::cache_hint [%rd2, %rd3, {%r1, %r2, %r3, %r4, %r5}], [%rd1], %rd4;
# | 116: ret;
# | 117: // -- End function
# | 118: }
# | 119: // .globl cp_async_bulk_tensor_s2g_tile_override_addr_dim_stride // -- Begin function cp_async_bulk_tensor_s2g_tile_override_addr_dim_stride
# | 120: .visible .func cp_async_bulk_tensor_s2g_tile_override_addr_dim_stride(
# | next:115'3 } search range end (exclusive)
# | 121: .param .b64 cp_async_bulk_tensor_s2g_tile_override_addr_dim_stride_param_0,
# | 122: .param .b64 cp_async_bulk_tensor_s2g_tile_override_addr_dim_stride_param_1,
# | 123: .param .b64 cp_async_bulk_tensor_s2g_tile_override_addr_dim_stride_param_2,
# | 124: .param .b32 cp_async_bulk_tensor_s2g_tile_override_addr_dim_stride_param_3,
# | 125: .param .b32 cp_async_bulk_tensor_s2g_tile_override_addr_dim_stride_param_4,
# | .
# | .
# | .
# | 141: {
# | 142: .reg .b16 %rs<7>;
# | 143: .reg .b32 %r<10>;
# | 144: .reg .b64 %rd<5>;
# | 145:
# | 146: // %bb.0:
# | next:160'0 { search range start (exclusive)
# | next:160'1 error: no match found in search range
# | 147: ld.param::func.b64 %rd1, [cp_async_bulk_tensor_s2g_tile_override_addr_dim_stride_param_0];
# | next:160'2 ? possible intended match
# | 148: ld.param::func.b64 %rd2, [cp_async_bulk_tensor_s2g_tile_override_addr_dim_stride_param_1];
# | 149: ld.param::func.b64 %rd3, [cp_async_bulk_tensor_s2g_tile_override_addr_dim_stride_param_2];
# | 150: ld.param::func.b16 %rs1, [cp_async_bulk_tensor_s2g_tile_override_addr_dim_stride_param_3];
# | 151: ld.param::func.b32 %r1, [cp_async_bulk_tensor_s2g_tile_override_addr_dim_stride_param_13];
# | 152: cp.async.bulk.tensor.1d.global.shared::cta.tile.bulk_group.override::global_address.override::gl
# | ...
# `---data was truncated (10240/13175) (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/215503
More information about the llvm-commits
mailing list