[llvm] [NVVM][NVPTX] Add im2col_w support for S2G and reduction intrinsics (PR #214436)
via llvm-commits
llvm-commits at lists.llvm.org
Thu Aug 6 02:16:06 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-llvm-ir
Author: Rajat Bajpai (rajatbajpai)
<details>
<summary>Changes</summary>
PTX ISA 9.4 adds the im2col_no_offs::w mode to shared-to-global tensor copy and reduction instructions for Rubin family targets.
This change adds the corresponding NVVM intrinsics and NVPTX lowering.
---
Patch is 56.62 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/214436.diff
6 Files Affected:
- (modified) llvm/docs/NVPTXUsage.md (+50-2)
- (modified) llvm/include/llvm/IR/IntrinsicsNVVM.td (+1-1)
- (modified) llvm/lib/Target/NVPTX/NVPTXIntrinsics.td (+19-6)
- (modified) llvm/lib/Target/NVPTX/NVPTXSubtarget.h (+4)
- (added) llvm/test/CodeGen/NVPTX/cp-async-bulk-tensor-reduce-im2colw.ll (+303)
- (added) llvm/test/CodeGen/NVPTX/cp-async-bulk-tensor-s2g-im2colw.ll (+122)
``````````diff
diff --git a/llvm/docs/NVPTXUsage.md b/llvm/docs/NVPTXUsage.md
index f0a1439b30334..c8de49780a581 100644
--- a/llvm/docs/NVPTXUsage.md
+++ b/llvm/docs/NVPTXUsage.md
@@ -2051,8 +2051,8 @@ declare void @llvm.nvvm.cp.async.bulk.tensor.s2g.im2col.5d(..., i32 %d0, i32 %d1
##### Overview:
-The '`@llvm.nvvm.cp.async.bulk.tensor.s2g.im2col.[1-5]d`' intrinsics
-correspond to the `cp.async.bulk.tensor.[1-5]d.*` set of PTX instructions.
+The '`@llvm.nvvm.cp.async.bulk.tensor.s2g.im2col.[3-5]d`' intrinsics
+correspond to the `cp.async.bulk.tensor.[3-5]d.*` set of PTX instructions.
These instructions initiate an asynchronous copy of tensor data from
shared::cta to global memory (indicated by the `s2g` prefix) in `im2col`
mode. In this mode, the tensor has to be at least three-dimensional. Unlike the
@@ -2062,6 +2062,30 @@ described in the `s2g.tile` mode intrinsics above.
For more information, refer [PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-bulk-tensor).
+
+#### '`llvm.nvvm.cp.async.bulk.tensor.s2g.im2col.w.[3-5]d`'
+
+##### Syntax:
+
+```llvm
+declare void @llvm.nvvm.cp.async.bulk.tensor.s2g.im2col.w.3d(ptr addrspace(3) %src, ptr %tensor_map, i32 %d0, i32 %d1, i32 %d2, i64 %ch, i1 %flag_ch)
+declare void @llvm.nvvm.cp.async.bulk.tensor.s2g.im2col.w.4d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, ...)
+declare void @llvm.nvvm.cp.async.bulk.tensor.s2g.im2col.w.5d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, i32 %d4, ...)
+```
+
+##### Overview:
+
+The '`@llvm.nvvm.cp.async.bulk.tensor.s2g.im2col.w.[3-5]d`' intrinsics
+correspond to the `cp.async.bulk.tensor.[3-5]d.*` set of PTX instructions.
+These instructions initiate an asynchronous copy of tensor data from
+shared::cta to global memory (indicated by the `s2g` prefix) in `im2col_w`
+mode. In this mode, the tensor has to be at least three-dimensional. Unlike the
+`g2s` variants, there are no im2col_offsets for these intrinsics. The last
+argument to these intrinsics is a boolean flag, with the same functionality as
+described in the `s2g.tile` mode intrinsics above.
+
+For more information, refer [PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-bulk-tensor).
+
#### '`llvm.nvvm.cp.async.bulk.tensor.reduce.tile.[1-5]d`'
##### Syntax:
@@ -2131,6 +2155,30 @@ intrinsics above.
For more information, refer [PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-reduce-async-bulk-tensor).
+#### '`llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.w.[3-5]d`'
+
+##### Syntax:
+
+```llvm
+declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.w.3d(ptr addrspace(3) %src, ptr %tensor_map, i32 %d0, i32 %d1, i32 %d2, i64 %ch, i32 %red_op, i1 %flag_ch)
+declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.w.4d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, ...)
+declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.w.5d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, i32 %d4, ...)
+```
+
+##### Overview:
+
+The '`@llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.w.[3-5]d`'
+intrinsics correspond to the `cp.reduce.async.bulk.tensor.[3-5]d.*` set of PTX
+instructions. These instructions initiate an asynchronous reduction operation of
+tensor data in global memory with the tensor data in shared\{::cta} memory, using
+`im2col_w` mode. In this mode, the tensor has to be at least three-dimensional.
+The supported reduction operations are the same as the ones
+in the `tile` mode. The `i32 %red_op` argument and the last boolean flag
+argument have the same functionality as described in the `tile` mode
+intrinsics above.
+
+For more information, refer [PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-reduce-async-bulk-tensor).
+
#### '`llvm.nvvm.cp.async.bulk.tensor.prefetch.tile.[1-5]d`'
##### Syntax:
diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index fc04a5dc64e78..da9740a57634b 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -2823,7 +2823,7 @@ class DefaultAttrsIntrinsicFlags<list<LLVMType> ret_types,
// TMA Tensor Copy Intrinsics: S2G -> From Shared to Global memory variants
foreach dim = 1...5 in {
defvar tensor_dim_args = !listsplat(llvm_i32_ty, dim);
- foreach mode = !if(!ge(dim, 3), ["tile", "im2col"], ["tile"]) in {
+ foreach mode = !if(!ge(dim, 3), ["tile", "im2col", "im2col_w"], ["tile"]) in {
def int_nvvm_cp_async_bulk_tensor_s2g_ # mode # _ # dim # d :
DefaultAttrsIntrinsicFlags<[],
!listconcat([llvm_shared_ptr_ty, // src_smem_ptr
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index ea47ddcc02ae7..ca2742f41ee00 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -824,6 +824,7 @@ multiclass TMA_TENSOR_S2G_INTR<int dim, string mode,
// Fix-up the asm_str when it is im2col/scatter4.
defvar mode_asm_str = !cond(
!eq(mode, "im2col") : "im2col_no_offs",
+ !eq(mode, "im2col_w") : "im2col_no_offs::w",
!eq(mode, "tile_scatter4") : "tile::scatter4",
true : mode);
defvar prefix = "cp.async.bulk.tensor"
@@ -849,6 +850,11 @@ foreach dim = 1...5 in {
defm TMA_TENSOR_S2G_ # suffix : TMA_TENSOR_S2G_INTR<dim, mode>;
}
}
+foreach dim = 3...5 in {
+ defm TMA_TENSOR_S2G_IM2COL_W_ # dim # D
+ : TMA_TENSOR_S2G_INTR<dim, "im2col_w",
+ [callSubtarget<"hasRubinFamilySupport">]>;
+}
defm TMA_S2G_TILE_SCATTER4_2D : TMA_TENSOR_S2G_INTR<5, "tile_scatter4",
[callSubtarget<"hasTMABlackwellSupport">]>;
@@ -860,14 +866,15 @@ def tma_tensor_reduction_imm :
TImmLeaf<i32, [{ return Imm >= 0 && Imm < 8; }]>;
// TMA Copy from Shared to Global memory with Reduction
-multiclass CP_ASYNC_BULK_TENSOR_REDUCE_INTR<int dim, string mode> {
+multiclass CP_ASYNC_BULK_TENSOR_REDUCE_INTR<int dim, string mode,
+ list<Predicate> pred = [hasPTX<80>, hasSM<90>]> {
defvar dims_dag = TMA_DIMS_UTIL<dim>.ins_dag;
defvar dims_str = TMA_DIMS_UTIL<dim>.base_str;
defvar asm_str = " [$tmap, {{" # dims_str # "}}], [$src]";
- // For im2col mode, the actual asm_str is "im2col_no_offs"
- defvar mode_asm_str = !if(!eq(mode, "im2col"),
- "im2col_no_offs", mode);
+ defvar mode_asm_str = !cond(!eq(mode, "im2col") : "im2col_no_offs",
+ !eq(mode, "im2col_w") : "im2col_no_offs::w",
+ true : mode);
defvar prefix = "cp.reduce.async.bulk.tensor"
# "." # dim # "d"
# ".global.shared::cta";
@@ -894,14 +901,14 @@ multiclass CP_ASYNC_BULK_TENSOR_REDUCE_INTR<int dim, string mode> {
(ins TMAReductionFlags:$red_op)),
!strconcat(prefix, "${red_op}", suffix, asm_str, ";"),
[intr_dag]>,
- Requires<[hasPTX<80>, hasSM<90>]>;
+ Requires<pred>;
def _CH : NVPTXInst<(outs),
!con((ins ADDR:$src, B64:$tmap), dims_dag,
(ins B64:$ch, TMAReductionFlags:$red_op)),
!strconcat(prefix, "${red_op}", suffix,
".L2::cache_hint", asm_str, ", $ch;"),
[intr_dag_with_ch]>,
- Requires<[hasPTX<80>, hasSM<90>]>;
+ Requires<pred>;
}
foreach dim = 1...5 in {
@@ -912,6 +919,12 @@ foreach dim = 1...5 in {
}
}
+foreach dim = 3...5 in {
+ defm CP_ASYNC_BULK_TENSOR_RED_ # dim # D_IM2COL_W
+ : CP_ASYNC_BULK_TENSOR_REDUCE_INTR<dim, "im2col_w",
+ [callSubtarget<"hasRubinFamilySupport">]>;
+}
+
// TMA Prefetch from Global memory to L2 cache
multiclass TMA_TENSOR_PREFETCH_INTR<int dim, string mode,
list<Predicate> pred = [hasPTX<80>, hasSM<90>]> {
diff --git a/llvm/lib/Target/NVPTX/NVPTXSubtarget.h b/llvm/lib/Target/NVPTX/NVPTXSubtarget.h
index 22affe4f40759..f6faa022be8cc 100644
--- a/llvm/lib/Target/NVPTX/NVPTXSubtarget.h
+++ b/llvm/lib/Target/NVPTX/NVPTXSubtarget.h
@@ -140,6 +140,10 @@ class NVPTXSubtarget : public NVPTXGenSubtargetInfo {
hasPTXWithAccelSMs(86, {100, 101});
}
+ // Checks Rubin family extensions support.
+ // - TMA S2G im2col_w mode support
+ bool hasRubinFamilySupport() const { return hasPTXWithFamilySMs(94, {107}); }
+
// Checks tcgen05.shift instruction support.
bool hasTcgen05ShiftSupport() const {
// sm_101 renamed to sm_110 in PTX 9.0
diff --git a/llvm/test/CodeGen/NVPTX/cp-async-bulk-tensor-reduce-im2colw.ll b/llvm/test/CodeGen/NVPTX/cp-async-bulk-tensor-reduce-im2colw.ll
new file mode 100644
index 0000000000000..dbc240415be15
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/cp-async-bulk-tensor-reduce-im2colw.ll
@@ -0,0 +1,303 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
+
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_107f | FileCheck --check-prefixes=CHECK-PTX64 %s
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_107f --nvptx-short-ptr| FileCheck --check-prefixes=CHECK-PTX-SHARED32 %s
+; RUN: llvm-as < %s | llvm-dis | FileCheck --check-prefixes=CHECK-FORMAT %s
+; RUN: %if ptxas-sm_107f && ptxas-isa-9.4 %{ llc < %s -march=nvptx64 -mcpu=sm_107f -mattr=+ptx94 | %ptxas-verify -arch=sm_107f %}
+
+define void @cp_async_bulk_tensor_reduce_im2colw_3d(ptr addrspace(3) %src, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i64 %ch) {
+; CHECK-PTX64-LABEL: cp_async_bulk_tensor_reduce_im2colw_3d(
+; CHECK-PTX64: {
+; CHECK-PTX64-NEXT: .reg .b32 %r<4>;
+; CHECK-PTX64-NEXT: .reg .b64 %rd<4>;
+; CHECK-PTX64-EMPTY:
+; CHECK-PTX64-NEXT: // %bb.0:
+; CHECK-PTX64-NEXT: ld.param.b64 %rd1, [cp_async_bulk_tensor_reduce_im2colw_3d_param_0];
+; CHECK-PTX64-NEXT: ld.param.b64 %rd2, [cp_async_bulk_tensor_reduce_im2colw_3d_param_1];
+; CHECK-PTX64-NEXT: ld.param.b32 %r1, [cp_async_bulk_tensor_reduce_im2colw_3d_param_2];
+; CHECK-PTX64-NEXT: ld.param.b32 %r2, [cp_async_bulk_tensor_reduce_im2colw_3d_param_3];
+; CHECK-PTX64-NEXT: ld.param.b32 %r3, [cp_async_bulk_tensor_reduce_im2colw_3d_param_4];
+; CHECK-PTX64-NEXT: ld.param.b64 %rd3, [cp_async_bulk_tensor_reduce_im2colw_3d_param_5];
+; CHECK-PTX64-NEXT: cp.reduce.async.bulk.tensor.3d.global.shared::cta.add.im2col_no_offs::w.bulk_group.L2::cache_hint [%rd2, {%r1, %r2, %r3}], [%rd1], %rd3;
+; CHECK-PTX64-NEXT: cp.reduce.async.bulk.tensor.3d.global.shared::cta.min.im2col_no_offs::w.bulk_group.L2::cache_hint [%rd2, {%r1, %r2, %r3}], [%rd1], %rd3;
+; CHECK-PTX64-NEXT: cp.reduce.async.bulk.tensor.3d.global.shared::cta.max.im2col_no_offs::w.bulk_group.L2::cache_hint [%rd2, {%r1, %r2, %r3}], [%rd1], %rd3;
+; CHECK-PTX64-NEXT: cp.reduce.async.bulk.tensor.3d.global.shared::cta.inc.im2col_no_offs::w.bulk_group.L2::cache_hint [%rd2, {%r1, %r2, %r3}], [%rd1], %rd3;
+; CHECK-PTX64-NEXT: cp.reduce.async.bulk.tensor.3d.global.shared::cta.dec.im2col_no_offs::w.bulk_group.L2::cache_hint [%rd2, {%r1, %r2, %r3}], [%rd1], %rd3;
+; CHECK-PTX64-NEXT: cp.reduce.async.bulk.tensor.3d.global.shared::cta.and.im2col_no_offs::w.bulk_group.L2::cache_hint [%rd2, {%r1, %r2, %r3}], [%rd1], %rd3;
+; CHECK-PTX64-NEXT: cp.reduce.async.bulk.tensor.3d.global.shared::cta.or.im2col_no_offs::w.bulk_group.L2::cache_hint [%rd2, {%r1, %r2, %r3}], [%rd1], %rd3;
+; CHECK-PTX64-NEXT: cp.reduce.async.bulk.tensor.3d.global.shared::cta.xor.im2col_no_offs::w.bulk_group.L2::cache_hint [%rd2, {%r1, %r2, %r3}], [%rd1], %rd3;
+; CHECK-PTX64-NEXT: cp.reduce.async.bulk.tensor.3d.global.shared::cta.add.im2col_no_offs::w.bulk_group [%rd2, {%r1, %r2, %r3}], [%rd1];
+; CHECK-PTX64-NEXT: cp.reduce.async.bulk.tensor.3d.global.shared::cta.min.im2col_no_offs::w.bulk_group [%rd2, {%r1, %r2, %r3}], [%rd1];
+; CHECK-PTX64-NEXT: cp.reduce.async.bulk.tensor.3d.global.shared::cta.max.im2col_no_offs::w.bulk_group [%rd2, {%r1, %r2, %r3}], [%rd1];
+; CHECK-PTX64-NEXT: cp.reduce.async.bulk.tensor.3d.global.shared::cta.inc.im2col_no_offs::w.bulk_group [%rd2, {%r1, %r2, %r3}], [%rd1];
+; CHECK-PTX64-NEXT: cp.reduce.async.bulk.tensor.3d.global.shared::cta.dec.im2col_no_offs::w.bulk_group [%rd2, {%r1, %r2, %r3}], [%rd1];
+; CHECK-PTX64-NEXT: cp.reduce.async.bulk.tensor.3d.global.shared::cta.and.im2col_no_offs::w.bulk_group [%rd2, {%r1, %r2, %r3}], [%rd1];
+; CHECK-PTX64-NEXT: cp.reduce.async.bulk.tensor.3d.global.shared::cta.or.im2col_no_offs::w.bulk_group [%rd2, {%r1, %r2, %r3}], [%rd1];
+; CHECK-PTX64-NEXT: cp.reduce.async.bulk.tensor.3d.global.shared::cta.xor.im2col_no_offs::w.bulk_group [%rd2, {%r1, %r2, %r3}], [%rd1];
+; CHECK-PTX64-NEXT: ret;
+;
+; CHECK-PTX-SHARED32-LABEL: cp_async_bulk_tensor_reduce_im2colw_3d(
+; CHECK-PTX-SHARED32: {
+; CHECK-PTX-SHARED32-NEXT: .reg .b32 %r<5>;
+; CHECK-PTX-SHARED32-NEXT: .reg .b64 %rd<3>;
+; CHECK-PTX-SHARED32-EMPTY:
+; CHECK-PTX-SHARED32-NEXT: // %bb.0:
+; CHECK-PTX-SHARED32-NEXT: ld.param.b32 %r1, [cp_async_bulk_tensor_reduce_im2colw_3d_param_0];
+; CHECK-PTX-SHARED32-NEXT: ld.param.b64 %rd1, [cp_async_bulk_tensor_reduce_im2colw_3d_param_1];
+; CHECK-PTX-SHARED32-NEXT: ld.param.b32 %r2, [cp_async_bulk_tensor_reduce_im2colw_3d_param_2];
+; CHECK-PTX-SHARED32-NEXT: ld.param.b32 %r3, [cp_async_bulk_tensor_reduce_im2colw_3d_param_3];
+; CHECK-PTX-SHARED32-NEXT: ld.param.b32 %r4, [cp_async_bulk_tensor_reduce_im2colw_3d_param_4];
+; CHECK-PTX-SHARED32-NEXT: ld.param.b64 %rd2, [cp_async_bulk_tensor_reduce_im2colw_3d_param_5];
+; CHECK-PTX-SHARED32-NEXT: cp.reduce.async.bulk.tensor.3d.global.shared::cta.add.im2col_no_offs::w.bulk_group.L2::cache_hint [%rd1, {%r2, %r3, %r4}], [%r1], %rd2;
+; CHECK-PTX-SHARED32-NEXT: cp.reduce.async.bulk.tensor.3d.global.shared::cta.min.im2col_no_offs::w.bulk_group.L2::cache_hint [%rd1, {%r2, %r3, %r4}], [%r1], %rd2;
+; CHECK-PTX-SHARED32-NEXT: cp.reduce.async.bulk.tensor.3d.global.shared::cta.max.im2col_no_offs::w.bulk_group.L2::cache_hint [%rd1, {%r2, %r3, %r4}], [%r1], %rd2;
+; CHECK-PTX-SHARED32-NEXT: cp.reduce.async.bulk.tensor.3d.global.shared::cta.inc.im2col_no_offs::w.bulk_group.L2::cache_hint [%rd1, {%r2, %r3, %r4}], [%r1], %rd2;
+; CHECK-PTX-SHARED32-NEXT: cp.reduce.async.bulk.tensor.3d.global.shared::cta.dec.im2col_no_offs::w.bulk_group.L2::cache_hint [%rd1, {%r2, %r3, %r4}], [%r1], %rd2;
+; CHECK-PTX-SHARED32-NEXT: cp.reduce.async.bulk.tensor.3d.global.shared::cta.and.im2col_no_offs::w.bulk_group.L2::cache_hint [%rd1, {%r2, %r3, %r4}], [%r1], %rd2;
+; CHECK-PTX-SHARED32-NEXT: cp.reduce.async.bulk.tensor.3d.global.shared::cta.or.im2col_no_offs::w.bulk_group.L2::cache_hint [%rd1, {%r2, %r3, %r4}], [%r1], %rd2;
+; CHECK-PTX-SHARED32-NEXT: cp.reduce.async.bulk.tensor.3d.global.shared::cta.xor.im2col_no_offs::w.bulk_group.L2::cache_hint [%rd1, {%r2, %r3, %r4}], [%r1], %rd2;
+; CHECK-PTX-SHARED32-NEXT: cp.reduce.async.bulk.tensor.3d.global.shared::cta.add.im2col_no_offs::w.bulk_group [%rd1, {%r2, %r3, %r4}], [%r1];
+; CHECK-PTX-SHARED32-NEXT: cp.reduce.async.bulk.tensor.3d.global.shared::cta.min.im2col_no_offs::w.bulk_group [%rd1, {%r2, %r3, %r4}], [%r1];
+; CHECK-PTX-SHARED32-NEXT: cp.reduce.async.bulk.tensor.3d.global.shared::cta.max.im2col_no_offs::w.bulk_group [%rd1, {%r2, %r3, %r4}], [%r1];
+; CHECK-PTX-SHARED32-NEXT: cp.reduce.async.bulk.tensor.3d.global.shared::cta.inc.im2col_no_offs::w.bulk_group [%rd1, {%r2, %r3, %r4}], [%r1];
+; CHECK-PTX-SHARED32-NEXT: cp.reduce.async.bulk.tensor.3d.global.shared::cta.dec.im2col_no_offs::w.bulk_group [%rd1, {%r2, %r3, %r4}], [%r1];
+; CHECK-PTX-SHARED32-NEXT: cp.reduce.async.bulk.tensor.3d.global.shared::cta.and.im2col_no_offs::w.bulk_group [%rd1, {%r2, %r3, %r4}], [%r1];
+; CHECK-PTX-SHARED32-NEXT: cp.reduce.async.bulk.tensor.3d.global.shared::cta.or.im2col_no_offs::w.bulk_group [%rd1, {%r2, %r3, %r4}], [%r1];
+; CHECK-PTX-SHARED32-NEXT: cp.reduce.async.bulk.tensor.3d.global.shared::cta.xor.im2col_no_offs::w.bulk_group [%rd1, {%r2, %r3, %r4}], [%r1];
+; CHECK-PTX-SHARED32-NEXT: ret;
+; CHECK-FORMAT-LABEL: define void @cp_async_bulk_tensor_reduce_im2colw_3d(ptr addrspace(3) %src, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i64 %ch) {
+; CHECK-FORMAT: tail call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.w.3d(ptr addrspace(3) %src, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i64 %ch, /* red_op=add */ i32 0, i1 true)
+; CHECK-FORMAT: tail call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.w.3d(ptr addrspace(3) %src, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i64 %ch, /* red_op=min */ i32 1, i1 true)
+; CHECK-FORMAT: tail call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.w.3d(ptr addrspace(3) %src, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i64 %ch, /* red_op=max */ i32 2, i1 true)
+; CHECK-FORMAT: tail call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.w.3d(ptr addrspace(3) %src, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i64 %ch, /* red_op=inc */ i32 3, i1 true)
+; CHECK-FORMAT: tail call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.w.3d(ptr addrspace(3) %src, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i64 %ch, /* red_op=dec */ i32 4, i1 true)
+; CHECK-FORMAT: tail call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.w.3d(ptr addrspace(3) %src, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i64 %ch, /* red_op=and */ i32 5, i1 true)
+; CHECK-FORMAT: tail call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.w.3d(ptr addrspace(3) %src, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i64 %ch, /* red_op=or */ i32 6, i1 true)
+; CHECK-FORMAT: tail call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.w.3d(ptr addrspace(3) %src, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i64 %ch, /* red_op=xor */ i32 7, i1 true)
+; CHECK-FORMAT: tail call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.w.3d(ptr addrspace(3) %src, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i64 %ch, /* red_op=add */ i32 0, i1 false)
+; CHECK-FORMAT: tail call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.w.3d(ptr addrspace(3) %src, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i64 %ch, /* red_op=min */ i32 1, i1 false)
+; CHECK-FORMAT: tail call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.w.3d(ptr addrspace(3) %src, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i64 %ch, /* red_op=max */ i32 2, i1 false)
+; CHECK-FORMAT: tail call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.w.3d(ptr addrspace(3) %src, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i64 %ch, /* red_op=inc */ i32 3, i1 false)
+; CHECK-FORMAT: tail call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.w.3d(ptr addrspace(3) %src, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i64 %ch, /* red_op=dec */ i32 4, i1 false)
+; CHECK-FORMAT: tail call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.w.3d(ptr addrspace(3) %src, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i64 %ch, /* red_op=and */ i32 5, i1 false)
+; CHECK-FORMAT: tail call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.w.3d(ptr addrspace(3) %src, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i64 %ch, /* red_op=or */ i32 6, i1 false)
+; CHECK-FORMAT: tail call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.w.3d(ptr addrspace(3) %src, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i64 %ch, /* red_op=xor */ i32 7, i1 false)
+ tail call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.w.3d(ptr addrspace(3) %src, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i64 %ch, i32 0, i1 1)
+ tail call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.w.3d(ptr addrspace(3) %src, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i64 %ch, i32 1, i1 1)
+ tail call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.w.3d(ptr addrspace(3) %src, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i64 %ch, i32 2, i1 1)
+ tail call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.w.3d(ptr addrspace(3) %src, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i64 %ch, i32 3, i1 1)
+ tail call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.w.3d(ptr addrspace(3) %src, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i64 %ch, i32 4, i1 1)
+ tail call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.w.3d(ptr addrspace(3) %src, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i64 %ch, i32 5, i1 1)
+ tail call void @llvm.nvvm.cp.async.bulk.tensor.redu...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/214436
More information about the llvm-commits
mailing list