[Mlir-commits] [mlir] [MLIR][NVVM] Add S2G and Reduce override NVVM Dialect ops (PR #216481)
Rajat Bajpai
llvmlistbot at llvm.org
Sat Aug 22 06:46:10 PDT 2026
https://github.com/rajatbajpai updated https://github.com/llvm/llvm-project/pull/216481
>From 5f9c091cf18f5f82065664f099d1c9ddadb28007 Mon Sep 17 00:00:00 2001
From: rbajpai <rbajpai at nvidia.com>
Date: Fri, 14 Aug 2026 13:42:34 +0530
Subject: [PATCH 1/3] [MLIR][NVVM] Add S2G and Reduce override ops
This change adds S2G and Reduction NVVM Dialect operations with tensor map
override capability.
---
mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td | 212 ++++++++++
mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp | 200 ++++++++++
.../tma_reduce_override_invalid.mlir | 34 ++
.../tma_store_override_invalid.mlir | 50 +++
.../LLVMIR/nvvm/tma_store_override.mlir | 154 ++++++++
.../nvvm/tma_store_reduce_override.mlir | 372 ++++++++++++++++++
6 files changed, 1022 insertions(+)
create mode 100644 mlir/test/Target/LLVMIR/nvvm/nvvmir-invalid/tma_reduce_override_invalid.mlir
create mode 100644 mlir/test/Target/LLVMIR/nvvm/nvvmir-invalid/tma_store_override_invalid.mlir
create mode 100644 mlir/test/Target/LLVMIR/nvvm/tma_store_override.mlir
create mode 100644 mlir/test/Target/LLVMIR/nvvm/tma_store_reduce_override.mlir
diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index 6ac489f9ad2bc..0009a7facda6a 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -4518,6 +4518,10 @@ def NVVM_CpAsyncBulkTensorGlobalToSharedClusterOp :
}];
}
+//===----------------------------------------------------------------------===//
+// NVVM S2G Ops
+//===----------------------------------------------------------------------===//
+
def NVVM_CpAsyncBulkTensorSharedCTAToGlobalOp :
NVVM_PTXBuilder_Op<"cp.async.bulk.tensor.global.shared.cta",
[AttrSizedOperandSegments]>,
@@ -4580,6 +4584,107 @@ def NVVM_CpAsyncBulkTensorSharedCTAToGlobalOp :
}];
}
+def NVVM_CpAsyncBulkTensorSharedCTAToGlobalOverrideAddrOp :
+ NVVM_VoidIntrinsicOp<"cp.async.bulk.tensor.global.shared.cta.override",
+ [AttrSizedOperandSegments]> {
+ let arguments = (ins
+ LLVM_AnyPointer:$tmaDesc,
+ LLVM_PointerShared:$srcMem,
+ LLVM_PointerGlobal:$overrideAdrr,
+ Variadic<I32>:$coordinates,
+ Variadic<I16>:$tensorSize,
+ Variadic<I32>:$lowerStride,
+ Optional<I16>:$upperStride,
+ Optional<I64>:$l2CacheHint,
+ DefaultValuedAttr<TMAStoreModeAttr, "TMAStoreMode::TILE">:$mode);
+
+ let summary = "Async bulk tensor copy from shared::cta to global memory with "
+ "tensor-map field overrides";
+ let description = [{
+ Initiates an asynchronous copy of tensor data from shared::cta memory to global
+ memory while overriding specific fields of the opaque tensor-map object with
+ explicit operands. It corresponds to the `cp.async.bulk.tensor.[1-5]d.*` PTX
+ instructions qualified with `.override::*`.
+
+ The `mode` attribute selects the store mode. The override variant is selected by
+ which optional operands are provided:
+
+ - `override.addr` (`.override::global_address`): only `overrideAdrr` is given; the
+ global base address from the tensor-map is replaced. Supported in `TILE` (1D–5D),
+ `IM2COL`/`IM2COL_W` (3D–5D), and `TILE_SCATTER4` (2D, requires 5 coordinates).
+ - `override.addr.dim` (1D, `TILE` only): `overrideAdrr` plus `tensorSize` (one
+ element); also overrides the tensor global dimension.
+ - `override.addr.dim.stride` (2D–5D, `TILE` only): `overrideAdrr` plus `tensorSize`,
+ `lowerStride`, and `upperStride`; also overrides the global dimensions and strides.
+ The effective global stride is
+ `global_stride[i] = ((%stride{i} + (%upper_stride{i} << 32)) << 4)`.
+
+ `overrideAdrr` must be 16B aligned (a runtime error is raised otherwise) and the
+ memory range `[%override_addr, %override_addr + 128 KiB)` must be allocated and
+ accessible during execution. When overriding dimensions/strides, the base-address
+ override is mandatory and the tensor start coordinates must be zero; otherwise the
+ behavior is undefined.
+
+ The optional `l2CacheHint` specifies a cache-eviction policy for the access.
+
+ Examples:
+
+ // override.addr (TILE, 1D)
+ ```mlir
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0] {
+ mode = #nvvm.tma_store_mode<tile>
+ } : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ ```
+
+ // override.addr (IM2COL, 3D) with L2 cache hint
+ ```mlir
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2]
+ l2_cache_hint = %ch {
+ mode = #nvvm.tma_store_mode<im2col>
+ } : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ ```
+
+ // override.addr (TILE_SCATTER4, 2D) — 5 coordinates: x0, y0, y1, y2, y3
+ ```mlir
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%x0, %y0, %y1, %y2, %y3] {
+ mode = #nvvm.tma_store_mode<tile_scatter4>
+ } : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ ```
+
+ // override.addr.dim (TILE, 1D)
+ ```mlir
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr,
+ box[%d0] tensor_size[%ts0] {
+ mode = #nvvm.tma_store_mode<tile>
+ } : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ ```
+
+ // override.addr.dim.stride (TILE, 2D)
+ ```mlir
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr,
+ box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd] {
+ mode = #nvvm.tma_store_mode<tile>
+ } : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ ```
+
+ [For more information, see PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-bulk-tensor)
+ }];
+
+ let assemblyFormat = [{
+ $tmaDesc `,`
+ $srcMem `,`
+ $overrideAdrr `,`
+ `box` `[`$coordinates `]`
+ (`tensor_size` `[`$tensorSize^ `]`)?
+ (`lower_stride` `[`$lowerStride^ `]`)?
+ (`upper_stride` `[`$upperStride^ `]`)?
+ (`l2_cache_hint` `=` $l2CacheHint^)?
+ attr-dict `:` type($tmaDesc) `,` type($srcMem) `,` type($overrideAdrr)
+ }];
+
+ let hasVerifier = 1;
+}
+
//===----------------------------------------------------------------------===//
// NVVM Prefetch Op
//===----------------------------------------------------------------------===//
@@ -4722,6 +4827,10 @@ def NVVM_CpAsyncBulkTensorPrefetchOp :
let hasVerifier = 1;
}
+//===----------------------------------------------------------------------===//
+// NVVM Reduction Ops
+//===----------------------------------------------------------------------===//
+
// List of Reduction Ops supported with TMA Store
def TMAReduxKindAdd : I32EnumAttrCase<"ADD", 0, "add">;
def TMAReduxKindMin : I32EnumAttrCase<"MIN", 1, "min">;
@@ -4792,6 +4901,109 @@ def NVVM_CpAsyncBulkTensorReduceOp :
}];
}
+def NVVM_CpAsyncBulkTensorReduceOverrideAddrOp :
+ NVVM_VoidIntrinsicOp<"cp.async.bulk.tensor.reduce.override",
+ [AttrSizedOperandSegments]> {
+ let arguments = (ins
+ LLVM_AnyPointer:$tmaDesc,
+ LLVM_PointerShared:$srcMem,
+ LLVM_PointerGlobal:$overrideAdrr,
+ Variadic<I32>:$coordinates,
+ Variadic<I16>:$tensorSize,
+ Variadic<I32>:$lowerStride,
+ Optional<I16>:$upperStride,
+ Optional<I64>:$l2CacheHint,
+ TMAReduxKindAttr:$redKind,
+ DefaultValuedAttr<TMAStoreModeAttr, "TMAStoreMode::TILE">:$mode
+ );
+
+ let summary = "Async bulk tensor reduction from shared::cta to global memory with "
+ "tensor-map field overrides";
+ let description = [{
+ Initiates an asynchronous reduction of tensor data in global memory with the tensor
+ data in shared::cta memory, while overriding specific fields of the opaque tensor-map
+ object with explicit operands. It corresponds to the
+ `cp.reduce.async.bulk.tensor.[1-5]d.global.shared::cta.*` PTX instructions qualified
+ with `.override::*`.
+
+ The `mode` attribute selects the store mode and `redKind` selects the reduction
+ operation (`ADD`, `MIN`, `MAX`, `INC`, `DEC`, `AND`, `OR`, `XOR`) that combines the
+ source data in shared memory with the destination data in global memory. The override
+ variant is selected by which optional operands are provided:
+
+ - `override.addr` (`.override::global_address`): only `overrideAdrr` is given; the
+ global base address from the tensor-map is replaced. Supported in `TILE` (1D–5D),
+ `IM2COL` and `IM2COL_W` (3D–5D). `TILE_SCATTER4` is not supported.
+ - `override.addr.dim` (1D, `TILE` only): `overrideAdrr` plus `tensorSize` (one
+ element); also overrides the tensor global dimension.
+ - `override.addr.dim.stride` (2D–5D, `TILE` only): `overrideAdrr` plus `tensorSize`,
+ `lowerStride`, and `upperStride`; also overrides the global dimensions and strides.
+ The effective global stride is
+ `global_stride[i] = ((%stride{i} + (%upper_stride{i} << 32)) << 4)`.
+
+ `overrideAdrr` must be 16B aligned (a runtime error is raised otherwise) and the
+ memory range `[%override_addr, %override_addr + 128 KiB)` must be allocated and
+ accessible during execution. When overriding dimensions/strides, the base-address
+ override is mandatory and the tensor start coordinates must be zero; otherwise the
+ behavior is undefined.
+
+ The optional `l2CacheHint` specifies a cache-eviction policy for the access.
+
+ Examples:
+
+ // override.addr (TILE, 1D) with ADD reduction
+ ```mlir
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] {
+ redKind = #nvvm.tma_redux_kind<add>,
+ mode = #nvvm.tma_store_mode<tile>
+ } : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ ```
+
+ // override.addr (IM2COL, 3D) with L2 cache hint and MIN reduction
+ ```mlir
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2]
+ l2_cache_hint = %ch {
+ redKind = #nvvm.tma_redux_kind<min>,
+ mode = #nvvm.tma_store_mode<im2col>
+ } : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ ```
+
+ // override.addr.dim (TILE, 1D)
+ ```mlir
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr,
+ box[%d0] tensor_size[%ts0] {
+ redKind = #nvvm.tma_redux_kind<add>,
+ mode = #nvvm.tma_store_mode<tile>
+ } : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ ```
+
+ // override.addr.dim.stride (TILE, 2D)
+ ```mlir
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr,
+ box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd] {
+ redKind = #nvvm.tma_redux_kind<and>,
+ mode = #nvvm.tma_store_mode<tile>
+ } : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ ```
+
+ [For more information, see PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-reduce-async-bulk-tensor)
+ }];
+
+ let assemblyFormat = [{
+ $tmaDesc `,`
+ $srcMem `,`
+ $overrideAdrr `,`
+ `box` `[`$coordinates `]`
+ (`tensor_size` `[`$tensorSize^ `]`)?
+ (`lower_stride` `[`$lowerStride^ `]`)?
+ (`upper_stride` `[`$upperStride^ `]`)?
+ (`l2_cache_hint` `=` $l2CacheHint^)?
+ attr-dict `:` type($tmaDesc) `,` type($srcMem) `,` type($overrideAdrr)
+ }];
+
+ let hasVerifier = 1;
+}
+
def NVVM_CpAsyncBulkGlobalToSharedClusterOp :
NVVM_Op<"cp.async.bulk.shared.cluster.global", [AttrSizedOperandSegments]> {
let summary = "Async bulk copy from global to Shared {cta or cluster} memory";
diff --git a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
index 4dddb33fd050e..3da38367aafd4 100644
--- a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
+++ b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
@@ -154,6 +154,32 @@ static LogicalResult cpAsyncBulkTensorCommonVerifier(size_t tensorDims,
return success();
}
+LogicalResult CpAsyncBulkTensorOverrideAddrCommonVerifier(
+ OperandRange coordinates, OperandRange tensorSize, OperandRange lowerStride,
+ Value upperStride, bool isTile, Location loc) {
+ LogicalResult res = success();
+ if (!tensorSize.empty() && coordinates.size() != tensorSize.size())
+ res =
+ emitError(loc, "Expected coordinates size to be equal to tensor size");
+
+ if (!lowerStride.empty() && lowerStride.size() != tensorSize.size() - 1)
+ res = emitError(
+ loc,
+ "Expected lower_stride size to be equal to one less than tensor size");
+
+ if (!lowerStride.empty() != static_cast<bool>(upperStride))
+ res = emitError(loc,
+ "Expected lower_stride and upper_stride to be either both "
+ "present or both absent");
+
+ bool isDimStride = tensorSize.size() > 0;
+ if (!isTile && isDimStride)
+ res = emitError(
+ loc, "Only tile mode supports override address with dim and stride");
+
+ return res;
+}
+
LogicalResult CpAsyncBulkTensorSharedCTAToGlobalOp::verify() {
TMAStoreMode mode = getMode();
// We lower through inline-ptx when getPredicate() is true.
@@ -180,6 +206,24 @@ LogicalResult CpAsyncBulkTensorSharedCTAToGlobalOp::verify() {
return success();
}
+LogicalResult CpAsyncBulkTensorSharedCTAToGlobalOverrideAddrOp::verify() {
+ TMAStoreMode mode = getMode();
+ bool isIm2Col = mode == TMAStoreMode::IM2COL;
+ bool isTile = mode == TMAStoreMode::TILE;
+
+ LogicalResult commonRes = cpAsyncBulkTensorCommonVerifier(
+ getCoordinates().size(), isIm2Col, 0, getLoc());
+
+ LogicalResult overrideAddrRes = CpAsyncBulkTensorOverrideAddrCommonVerifier(
+ getCoordinates(), getTensorSize(), getLowerStride(), getUpperStride(),
+ isTile, getLoc());
+
+ if (mode == TMAStoreMode::TILE_SCATTER4 && getCoordinates().size() != 5)
+ overrideAddrRes = emitError("Mode tile scatter4 expects 5 coordinates");
+
+ return failed(commonRes) || failed(overrideAddrRes) ? failure() : success();
+}
+
LogicalResult CpAsyncOp::verify() {
if (getModifier() != LoadCacheModifierKind::CG &&
getModifier() != LoadCacheModifierKind::CA)
@@ -279,6 +323,20 @@ LogicalResult CpAsyncBulkTensorReduceOp::verify() {
return success();
}
+LogicalResult CpAsyncBulkTensorReduceOverrideAddrOp::verify() {
+ bool isIm2Col = getMode() == TMAStoreMode::IM2COL;
+ bool isTile = getMode() == TMAStoreMode::TILE;
+
+ LogicalResult commonRes = cpAsyncBulkTensorCommonVerifier(
+ getCoordinates().size(), isIm2Col, 0, getLoc());
+
+ LogicalResult overrideAddrRes = CpAsyncBulkTensorOverrideAddrCommonVerifier(
+ getCoordinates(), getTensorSize(), getLowerStride(), getUpperStride(),
+ isTile, getLoc());
+
+ return failed(commonRes) || failed(overrideAddrRes) ? failure() : success();
+}
+
LogicalResult CpAsyncBulkGlobalToSharedClusterOp::verify() {
bool isSharedCTA = isPtrInSharedCTASpace(getDstMem());
if (isSharedCTA && getMulticastMask())
@@ -4737,6 +4795,78 @@ CpAsyncBulkTensorSharedCTAToGlobalOp::getIntrinsicIDAndArgs(
return {id, std::move(args)};
}
+NVVM::IDArgPair
+CpAsyncBulkTensorSharedCTAToGlobalOverrideAddrOp::getIntrinsicIDAndArgs(
+ Operation &op, LLVM::ModuleTranslation &mt, llvm::IRBuilderBase &builder) {
+ auto thisOp =
+ cast<NVVM::CpAsyncBulkTensorSharedCTAToGlobalOverrideAddrOp>(op);
+
+ llvm::SmallVector<llvm::Value *> args;
+ args.push_back(mt.lookupValue(thisOp.getSrcMem()));
+ args.push_back(mt.lookupValue(thisOp.getTmaDesc()));
+ args.push_back(mt.lookupValue(thisOp.getOverrideAdrr()));
+ for (Value v : thisOp.getTensorSize())
+ args.push_back(mt.lookupValue(v));
+ for (Value v : thisOp.getLowerStride())
+ args.push_back(mt.lookupValue(v));
+ if (thisOp.getUpperStride())
+ args.push_back(mt.lookupValue(thisOp.getUpperStride()));
+ for (Value v : thisOp.getCoordinates())
+ args.push_back(mt.lookupValue(v));
+
+ mlir::Value cacheHint = thisOp.getL2CacheHint();
+ const bool hasCacheHint = static_cast<bool>(cacheHint);
+ args.push_back(hasCacheHint ? mt.lookupValue(cacheHint)
+ : builder.getInt64(0));
+ args.push_back(builder.getInt1(hasCacheHint));
+
+ using namespace llvm::Intrinsic;
+ const unsigned NI = not_intrinsic;
+ // clang-format off
+ // override_addr variants, indexed [mode][dim].
+ static constexpr ID IDTable[][6] = {
+ {NI, nvvm_cp_async_bulk_tensor_s2g_tile_override_addr_1d,
+ nvvm_cp_async_bulk_tensor_s2g_tile_override_addr_2d,
+ nvvm_cp_async_bulk_tensor_s2g_tile_override_addr_3d,
+ nvvm_cp_async_bulk_tensor_s2g_tile_override_addr_4d,
+ nvvm_cp_async_bulk_tensor_s2g_tile_override_addr_5d},
+ {NI, NI, NI, nvvm_cp_async_bulk_tensor_s2g_im2col_override_addr_3d,
+ nvvm_cp_async_bulk_tensor_s2g_im2col_override_addr_4d,
+ nvvm_cp_async_bulk_tensor_s2g_im2col_override_addr_5d},
+ {NI, NI, NI, NI, NI,
+ nvvm_cp_async_bulk_tensor_s2g_tile_scatter4_override_addr_2d},
+ {NI, NI, NI, nvvm_cp_async_bulk_tensor_s2g_im2col_w_override_addr_3d,
+ nvvm_cp_async_bulk_tensor_s2g_im2col_w_override_addr_4d,
+ nvvm_cp_async_bulk_tensor_s2g_im2col_w_override_addr_5d}};
+
+ // Tile-only override_addr_dim (1D) / override_addr_dim_stride (2D-5D)
+ // variants, indexed [dim].
+ static constexpr ID dimStrideIDTable[] = {
+ NI, nvvm_cp_async_bulk_tensor_s2g_tile_override_addr_dim_1d,
+ nvvm_cp_async_bulk_tensor_s2g_tile_override_addr_dim_stride_2d,
+ nvvm_cp_async_bulk_tensor_s2g_tile_override_addr_dim_stride_3d,
+ nvvm_cp_async_bulk_tensor_s2g_tile_override_addr_dim_stride_4d,
+ nvvm_cp_async_bulk_tensor_s2g_tile_override_addr_dim_stride_5d};
+ // clang-format on
+
+ size_t mode = static_cast<size_t>(thisOp.getMode());
+ size_t dim = thisOp.getCoordinates().size();
+ bool isDimStride = !thisOp.getTensorSize().empty();
+
+ assert(mode < std::size(IDTable) &&
+ "Invalid mode for CpAsyncBulkTensorSharedCTAToGlobalOverrideAddrOp");
+ assert(dim < std::size(IDTable[mode]) &&
+ "Invalid dim for CpAsyncBulkTensorSharedCTAToGlobalOverrideAddrOp");
+ assert(dim < std::size(dimStrideIDTable) &&
+ "Invalid dim for CpAsyncBulkTensorSharedCTAToGlobalOverrideAddrOp");
+
+ ID intrinsicID = isDimStride ? dimStrideIDTable[dim] : IDTable[mode][dim];
+ assert(
+ intrinsicID != NI &&
+ "Invalid intrinsic for CpAsyncBulkTensorSharedCTAToGlobalOverrideAddrOp");
+ return {intrinsicID, std::move(args)};
+}
+
NVVM::IDArgPair CpAsyncBulkTensorReduceOp::getIntrinsicIDAndArgs(
Operation &op, LLVM::ModuleTranslation &mt, llvm::IRBuilderBase &builder) {
auto thisOp = cast<NVVM::CpAsyncBulkTensorReduceOp>(op);
@@ -4783,6 +4913,76 @@ NVVM::IDArgPair CpAsyncBulkTensorReduceOp::getIntrinsicIDAndArgs(
return {intrinsicID, std::move(args)};
}
+NVVM::IDArgPair CpAsyncBulkTensorReduceOverrideAddrOp::getIntrinsicIDAndArgs(
+ Operation &op, LLVM::ModuleTranslation &mt, llvm::IRBuilderBase &builder) {
+ auto thisOp = cast<NVVM::CpAsyncBulkTensorReduceOverrideAddrOp>(op);
+
+ llvm::SmallVector<llvm::Value *> args;
+ args.push_back(mt.lookupValue(thisOp.getSrcMem()));
+ args.push_back(mt.lookupValue(thisOp.getTmaDesc()));
+ args.push_back(mt.lookupValue(thisOp.getOverrideAdrr()));
+
+ for (Value v : thisOp.getTensorSize())
+ args.push_back(mt.lookupValue(v));
+ for (Value v : thisOp.getLowerStride())
+ args.push_back(mt.lookupValue(v));
+ if (thisOp.getUpperStride())
+ args.push_back(mt.lookupValue(thisOp.getUpperStride()));
+ for (Value v : thisOp.getCoordinates())
+ args.push_back(mt.lookupValue(v));
+
+ mlir::Value cacheHint = thisOp.getL2CacheHint();
+ const bool hasCacheHint = static_cast<bool>(cacheHint);
+ args.push_back(hasCacheHint ? mt.lookupValue(cacheHint)
+ : builder.getInt64(0));
+ args.push_back(builder.getInt32(static_cast<uint32_t>(thisOp.getRedKind())));
+ args.push_back(builder.getInt1(hasCacheHint));
+
+ using namespace llvm::Intrinsic;
+ const unsigned NI = not_intrinsic;
+ // clang-format off
+// override_addr variants, indexed [mode][dim].
+static constexpr ID IDTable[][6] = {
+ {NI, nvvm_cp_async_bulk_tensor_reduce_tile_override_addr_1d,
+ nvvm_cp_async_bulk_tensor_reduce_tile_override_addr_2d,
+ nvvm_cp_async_bulk_tensor_reduce_tile_override_addr_3d,
+ nvvm_cp_async_bulk_tensor_reduce_tile_override_addr_4d,
+ nvvm_cp_async_bulk_tensor_reduce_tile_override_addr_5d},
+ {NI, NI, NI, nvvm_cp_async_bulk_tensor_reduce_im2col_override_addr_3d,
+ nvvm_cp_async_bulk_tensor_reduce_im2col_override_addr_4d,
+ nvvm_cp_async_bulk_tensor_reduce_im2col_override_addr_5d},
+ {NI, NI, NI, NI, NI, NI}, // scatter4 not supported for reduce
+ {NI, NI, NI, nvvm_cp_async_bulk_tensor_reduce_im2col_w_override_addr_3d,
+ nvvm_cp_async_bulk_tensor_reduce_im2col_w_override_addr_4d,
+ nvvm_cp_async_bulk_tensor_reduce_im2col_w_override_addr_5d}};
+
+// Tile-only override_addr_dim (1D) / override_addr_dim_stride (2D-5D)
+// variants, indexed [dim].
+static constexpr ID dimStrideIDTable[] = {
+ NI, nvvm_cp_async_bulk_tensor_reduce_tile_override_addr_dim_1d,
+ nvvm_cp_async_bulk_tensor_reduce_tile_override_addr_dim_stride_2d,
+ nvvm_cp_async_bulk_tensor_reduce_tile_override_addr_dim_stride_3d,
+ nvvm_cp_async_bulk_tensor_reduce_tile_override_addr_dim_stride_4d,
+ nvvm_cp_async_bulk_tensor_reduce_tile_override_addr_dim_stride_5d};
+ // clang-format on
+
+ size_t mode = static_cast<size_t>(thisOp.getMode());
+ size_t dim = thisOp.getCoordinates().size();
+ bool isDimStride = !thisOp.getTensorSize().empty();
+
+ assert(mode < std::size(IDTable) &&
+ "Invalid mode for CpAsyncBulkTensorReduceOverrideAddrOp");
+ assert(dim < std::size(IDTable[mode]) &&
+ "Invalid dim for CpAsyncBulkTensorReduceOverrideAddrOp");
+ assert(dim < std::size(dimStrideIDTable) &&
+ "Invalid dim for CpAsyncBulkTensorReduceOverrideAddrOp");
+
+ ID intrinsicID = isDimStride ? dimStrideIDTable[dim] : IDTable[mode][dim];
+ assert(intrinsicID != NI &&
+ "Invalid intrinsic for CpAsyncBulkTensorReduceOverrideAddrOp");
+ return {intrinsicID, std::move(args)};
+}
+
#define _none
#define CVT_F2TF32_ID_IMPL(rnd, relu, sf) \
diff --git a/mlir/test/Target/LLVMIR/nvvm/nvvmir-invalid/tma_reduce_override_invalid.mlir b/mlir/test/Target/LLVMIR/nvvm/nvvmir-invalid/tma_reduce_override_invalid.mlir
new file mode 100644
index 0000000000000..a723acc1fccc2
--- /dev/null
+++ b/mlir/test/Target/LLVMIR/nvvm/nvvmir-invalid/tma_reduce_override_invalid.mlir
@@ -0,0 +1,34 @@
+// RUN: mlir-translate -verify-diagnostics -split-input-file -mlir-to-llvmir %s
+
+// -----
+
+llvm.func @tma_store_reduce_tile_override_addr_coord_tensor_size_mismatch(%src : !llvm.ptr<3>, %tma_desc : !llvm.ptr, %override_addr : !llvm.ptr<1>, %d0 : i32, %d1 : i32, %d2 : i32, %d3 : i32, %d4 : i32, %ts0 : i16, %ts1 : i16, %ts2 : i16, %ts3 : i16, %ts4 : i16, %lstrd0 : i32, %lstrd1 : i32, %lstrd2 : i32, %lstrd3 : i32, %ustrd : i16, %ch : i64) {
+ // expected-error @below {{Expected coordinates size to be equal to tensor size}}
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0] {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ llvm.return
+}
+
+// -----
+
+llvm.func @tma_store_reduce_tile_override_addr_lower_stride_tensor_size_mismatch(%src : !llvm.ptr<3>, %tma_desc : !llvm.ptr, %override_addr : !llvm.ptr<1>, %d0 : i32, %d1 : i32, %d2 : i32, %d3 : i32, %d4 : i32, %ts0 : i16, %ts1 : i16, %ts2 : i16, %ts3 : i16, %ts4 : i16, %lstrd0 : i32, %lstrd1 : i32, %lstrd2 : i32, %lstrd3 : i32, %ustrd : i16, %ch : i64) {
+ // expected-error @below {{Expected lower_stride size to be equal to one less than tensor size}}
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0, %lstrd1] upper_stride[%ustrd] {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ llvm.return
+}
+
+// -----
+
+llvm.func @tma_store_reduce_tile_override_addr_stride_mismatch(%src : !llvm.ptr<3>, %tma_desc : !llvm.ptr, %override_addr : !llvm.ptr<1>, %d0 : i32, %d1 : i32, %d2 : i32, %d3 : i32, %d4 : i32, %ts0 : i16, %ts1 : i16, %ts2 : i16, %ts3 : i16, %ts4 : i16, %lstrd0 : i32, %lstrd1 : i32, %lstrd2 : i32, %lstrd3 : i32, %ustrd : i16, %ch : i64) {
+ // expected-error @below {{Expected lower_stride and upper_stride to be either both present or both absent}}
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ llvm.return
+}
+
+// -----
+
+llvm.func @tma_store_reduce_tile_override_addr_im2col_dim_stride(%src : !llvm.ptr<3>, %tma_desc : !llvm.ptr, %override_addr : !llvm.ptr<1>, %d0 : i32, %d1 : i32, %d2 : i32, %d3 : i32, %d4 : i32, %ts0 : i16, %ts1 : i16, %ts2 : i16, %ts3 : i16, %ts4 : i16, %lstrd0 : i32, %lstrd1 : i32, %lstrd2 : i32, %lstrd3 : i32, %ustrd : i16, %ch : i64) {
+ // expected-error @below {{Only tile mode supports override address with dim and stride}}
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] tensor_size[%ts0, %ts1, %ts2] {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<im2col>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ llvm.return
+}
+
diff --git a/mlir/test/Target/LLVMIR/nvvm/nvvmir-invalid/tma_store_override_invalid.mlir b/mlir/test/Target/LLVMIR/nvvm/nvvmir-invalid/tma_store_override_invalid.mlir
new file mode 100644
index 0000000000000..4a16af610da82
--- /dev/null
+++ b/mlir/test/Target/LLVMIR/nvvm/nvvmir-invalid/tma_store_override_invalid.mlir
@@ -0,0 +1,50 @@
+// RUN: mlir-translate -verify-diagnostics -split-input-file -mlir-to-llvmir %s
+
+// -----
+
+llvm.func @tma_shared_to_global_tile_override_addr_coord_tensor_size_mismatch(%src : !llvm.ptr<3>, %tma_desc : !llvm.ptr, %override_addr : !llvm.ptr<1>, %d0 : i32, %d1 : i32, %d2 : i32, %d3 : i32, %d4 : i32, %ts0 : i16, %ts1 : i16, %ts2 : i16, %ts3 : i16, %ts4 : i16, %lstrd0 : i32, %lstrd1 : i32, %lstrd2 : i32, %lstrd3 : i32, %ustrd : i16, %ch : i64) {
+ // expected-error @below {{Expected coordinates size to be equal to tensor size}}
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0] {mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ llvm.return
+}
+
+// -----
+
+llvm.func @tma_shared_to_global_tile_override_addr_lower_stride_tensor_size_mismatch(%src : !llvm.ptr<3>, %tma_desc : !llvm.ptr, %override_addr : !llvm.ptr<1>, %d0 : i32, %d1 : i32, %d2 : i32, %d3 : i32, %d4 : i32, %ts0 : i16, %ts1 : i16, %ts2 : i16, %ts3 : i16, %ts4 : i16, %lstrd0 : i32, %lstrd1 : i32, %lstrd2 : i32, %lstrd3 : i32, %ustrd : i16, %ch : i64) {
+ // expected-error @below {{Expected lower_stride size to be equal to one less than tensor size}}
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0, %lstrd1] upper_stride[%ustrd] {mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ llvm.return
+}
+
+// -----
+
+llvm.func @tma_shared_to_global_tile_override_addr_stride_mismatch(%src : !llvm.ptr<3>, %tma_desc : !llvm.ptr, %override_addr : !llvm.ptr<1>, %d0 : i32, %d1 : i32, %d2 : i32, %d3 : i32, %d4 : i32, %ts0 : i16, %ts1 : i16, %ts2 : i16, %ts3 : i16, %ts4 : i16, %lstrd0 : i32, %lstrd1 : i32, %lstrd2 : i32, %lstrd3 : i32, %ustrd : i16, %ch : i64) {
+ // expected-error @below {{Expected lower_stride and upper_stride to be either both present or both absent}}
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] {mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ llvm.return
+}
+
+// -----
+
+llvm.func @tma_shared_to_global_tile_override_addr_im2col_dim_stride(%src : !llvm.ptr<3>, %tma_desc : !llvm.ptr, %override_addr : !llvm.ptr<1>, %d0 : i32, %d1 : i32, %d2 : i32, %d3 : i32, %d4 : i32, %ts0 : i16, %ts1 : i16, %ts2 : i16, %ts3 : i16, %ts4 : i16, %lstrd0 : i32, %lstrd1 : i32, %lstrd2 : i32, %lstrd3 : i32, %ustrd : i16, %ch : i64) {
+ // expected-error @below {{Only tile mode supports override address with dim and stride}}
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] tensor_size[%ts0, %ts1, %ts2] {mode = #nvvm.tma_store_mode<im2col>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ llvm.return
+}
+
+// -----
+
+llvm.func @tma_shared_to_global_tile_override_addr_tile_scatter4_dim_stride(%src : !llvm.ptr<3>, %tma_desc : !llvm.ptr, %override_addr : !llvm.ptr<1>, %d0 : i32, %d1 : i32, %d2 : i32, %d3 : i32, %d4 : i32, %ts0 : i16, %ts1 : i16, %ts2 : i16, %ts3 : i16, %ts4 : i16, %lstrd0 : i32, %lstrd1 : i32, %lstrd2 : i32, %lstrd3 : i32, %ustrd : i16, %ch : i64) {
+ // expected-error @below {{Only tile mode supports override address with dim and stride}}
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3, %d4] tensor_size[%ts0, %ts1, %ts2, %ts3, %ts4] {mode = #nvvm.tma_store_mode<tile_scatter4>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ llvm.return
+}
+
+// -----
+
+llvm.func @tma_shared_to_global_tile_override_addr_tile_scatter4_coord_mismatch(%src : !llvm.ptr<3>, %tma_desc : !llvm.ptr, %override_addr : !llvm.ptr<1>, %d0 : i32, %d1 : i32, %d2 : i32, %d3 : i32, %d4 : i32, %ts0 : i16, %ts1 : i16, %ts2 : i16, %ts3 : i16, %ts4 : i16, %lstrd0 : i32, %lstrd1 : i32, %lstrd2 : i32, %lstrd3 : i32, %ustrd : i16, %ch : i64) {
+ // expected-error @below {{Mode tile scatter4 expects 5 coordinates}}
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3] {mode = #nvvm.tma_store_mode<tile_scatter4>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ llvm.return
+}
+
diff --git a/mlir/test/Target/LLVMIR/nvvm/tma_store_override.mlir b/mlir/test/Target/LLVMIR/nvvm/tma_store_override.mlir
new file mode 100644
index 0000000000000..5b57d375612b3
--- /dev/null
+++ b/mlir/test/Target/LLVMIR/nvvm/tma_store_override.mlir
@@ -0,0 +1,154 @@
+// RUN: mlir-translate -mlir-to-llvmir %s | FileCheck %s
+
+llvm.func @tma_shared_cta_global_tile_override_addr(%src : !llvm.ptr<3>, %tma_desc : !llvm.ptr, %override_addr : !llvm.ptr<1>, %d0 : i32, %d1 : i32, %d2 : i32, %d3 : i32, %d4 : i32, %ts0 : i16, %ts1 : i16, %ts2 : i16, %ts3 : i16, %ts4 : i16, %lstrd0 : i32, %lstrd1 : i32, %lstrd2 : i32, %lstrd3 : i32, %ustrd : i16, %ch : i64) {
+// CHECK-LABEL: define void @tma_shared_cta_global_tile_override_addr(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i32 %6, i32 %7, i16 %8, i16 %9, i16 %10, i16 %11, i16 %12, i32 %13, i32 %14, i32 %15, i32 %16, i16 %17, i64 %18) {
+// CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.s2g.tile.override.addr.1d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i64 0, /* flag_cache_hint= */ i1 false)
+// CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.s2g.tile.override.addr.2d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i64 0, /* flag_cache_hint= */ i1 false)
+// CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.s2g.tile.override.addr.3d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i64 0, /* flag_cache_hint= */ i1 false)
+// CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.s2g.tile.override.addr.4d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i32 %6, i64 0, /* flag_cache_hint= */ i1 false)
+// CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.s2g.tile.override.addr.5d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i32 %6, i32 %7, i64 0, /* flag_cache_hint= */ i1 false)
+// CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.s2g.tile.override.addr.1d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i64 %18, /* flag_cache_hint= */ i1 true)
+// CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.s2g.tile.override.addr.2d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i64 %18, /* flag_cache_hint= */ i1 true)
+// CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.s2g.tile.override.addr.3d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i64 %18, /* flag_cache_hint= */ i1 true)
+// CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.s2g.tile.override.addr.4d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i32 %6, i64 %18, /* flag_cache_hint= */ i1 true)
+// CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.s2g.tile.override.addr.5d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i32 %6, i32 %7, i64 %18, /* flag_cache_hint= */ i1 true)
+// CHECK-NEXT: ret void
+// CHECK-NEXT: }
+
+ // without cache hint
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0] {mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1] {mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] {mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3] {mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3, %d4] {mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ // with cache hint
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0] l2_cache_hint = %ch {mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1] l2_cache_hint = %ch {mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch {mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3] l2_cache_hint = %ch {mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3, %d4] l2_cache_hint = %ch {mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ llvm.return
+}
+
+llvm.func @tma_shared_cta_global_tile_scatter4_override_addr(%src : !llvm.ptr<3>, %tma_desc : !llvm.ptr, %override_addr : !llvm.ptr<1>, %d0 : i32, %d1 : i32, %d2 : i32, %d3 : i32, %d4 : i32, %ts0 : i16, %ts1 : i16, %ts2 : i16, %ts3 : i16, %ts4 : i16, %lstrd0 : i32, %lstrd1 : i32, %lstrd2 : i32, %lstrd3 : i32, %ustrd : i16, %ch : i64) {
+// CHECK-LABEL: define void @tma_shared_cta_global_tile_scatter4_override_addr(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i32 %6, i32 %7, i16 %8, i16 %9, i16 %10, i16 %11, i16 %12, i32 %13, i32 %14, i32 %15, i32 %16, i16 %17, i64 %18) {
+// CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.s2g.tile.scatter4.override.addr.2d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i32 %6, i32 %7, i64 0, /* flag_cache_hint= */ i1 false)
+// CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.s2g.tile.scatter4.override.addr.2d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i32 %6, i32 %7, i64 %18, /* flag_cache_hint= */ i1 true)
+// CHECK-NEXT: ret void
+// CHECK-NEXT: }
+
+ // without cache hint
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3, %d4] {mode = #nvvm.tma_store_mode<tile_scatter4>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ // with cache hint
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3, %d4] l2_cache_hint = %ch {mode = #nvvm.tma_store_mode<tile_scatter4>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ llvm.return
+}
+
+llvm.func @tma_shared_cta_global_im2col(%src : !llvm.ptr<3>, %tma_desc : !llvm.ptr, %override_addr : !llvm.ptr<1>, %d0 : i32, %d1 : i32, %d2 : i32, %d3 : i32, %d4 : i32, %ch : i64) {
+// CHECK-LABEL: define void @tma_shared_cta_global_im2col(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i32 %6, i32 %7, i64 %8) {
+// CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.s2g.im2col.override.addr.3d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i64 0, /* flag_cache_hint= */ i1 false)
+// CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.s2g.im2col.override.addr.4d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i32 %6, i64 0, /* flag_cache_hint= */ i1 false)
+// CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.s2g.im2col.override.addr.5d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i32 %6, i32 %7, i64 0, /* flag_cache_hint= */ i1 false)
+// CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.s2g.im2col.override.addr.3d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i64 %8, /* flag_cache_hint= */ i1 true)
+// CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.s2g.im2col.override.addr.4d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i32 %6, i64 %8, /* flag_cache_hint= */ i1 true)
+// CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.s2g.im2col.override.addr.5d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i32 %6, i32 %7, i64 %8, /* flag_cache_hint= */ i1 true)
+// CHECK-NEXT: ret void
+// CHECK-NEXT: }
+
+ // without cache hint
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] {mode = #nvvm.tma_store_mode<im2col>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3] {mode = #nvvm.tma_store_mode<im2col>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3, %d4] {mode = #nvvm.tma_store_mode<im2col>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ // with cache hint
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch {mode = #nvvm.tma_store_mode<im2col>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3] l2_cache_hint = %ch {mode = #nvvm.tma_store_mode<im2col>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3, %d4] l2_cache_hint = %ch {mode = #nvvm.tma_store_mode<im2col>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ llvm.return
+}
+
+llvm.func @tma_shared_cta_global_im2col_w(%src : !llvm.ptr<3>, %tma_desc : !llvm.ptr, %override_addr : !llvm.ptr<1>, %d0 : i32, %d1 : i32, %d2 : i32, %d3 : i32, %d4 : i32, %ch : i64) {
+// CHECK-LABEL: define void @tma_shared_cta_global_im2col_w(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i32 %6, i32 %7, i64 %8) {
+// CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.s2g.im2col.w.override.addr.3d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i64 0, /* flag_cache_hint= */ i1 false)
+// CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.s2g.im2col.w.override.addr.4d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i32 %6, i64 0, /* flag_cache_hint= */ i1 false)
+// CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.s2g.im2col.w.override.addr.5d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i32 %6, i32 %7, i64 0, /* flag_cache_hint= */ i1 false)
+// CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.s2g.im2col.w.override.addr.3d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i64 %8, /* flag_cache_hint= */ i1 true)
+// CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.s2g.im2col.w.override.addr.4d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i32 %6, i64 %8, /* flag_cache_hint= */ i1 true)
+// CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.s2g.im2col.w.override.addr.5d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i32 %6, i32 %7, i64 %8, /* flag_cache_hint= */ i1 true)
+// CHECK-NEXT: ret void
+// CHECK-NEXT: }
+
+ // without cache hint
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] {mode = #nvvm.tma_store_mode<im2col_w>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3] {mode = #nvvm.tma_store_mode<im2col_w>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3, %d4] {mode = #nvvm.tma_store_mode<im2col_w>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ // with cache hint
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch {mode = #nvvm.tma_store_mode<im2col_w>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3] l2_cache_hint = %ch {mode = #nvvm.tma_store_mode<im2col_w>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3, %d4] l2_cache_hint = %ch {mode = #nvvm.tma_store_mode<im2col_w>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ llvm.return
+}
+
+llvm.func @tma_shared_cta_global_tile_override_addr_dim_stride(%src : !llvm.ptr<3>, %tma_desc : !llvm.ptr, %override_addr : !llvm.ptr<1>, %d0 : i32, %d1 : i32, %d2 : i32, %d3 : i32, %d4 : i32, %ts0 : i16, %ts1 : i16, %ts2 : i16, %ts3 : i16, %ts4 : i16, %lstrd0 : i32, %lstrd1 : i32, %lstrd2 : i32, %lstrd3 : i32, %ustrd : i16, %ch : i64) {
+// CHECK-LABEL: define void @tma_shared_cta_global_tile_override_addr_dim_stride(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i32 %6, i32 %7, i16 %8, i16 %9, i16 %10, i16 %11, i16 %12, i32 %13, i32 %14, i32 %15, i32 %16, i16 %17, i64 %18) {
+// CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.s2g.tile.override.addr.dim.1d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i16 %8, i32 %3, i64 0, /* flag_cache_hint= */ i1 false)
+// CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.s2g.tile.override.addr.dim.stride.2d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i16 %8, i16 %9, i32 %13, i16 %17, i32 %3, i32 %4, i64 0, /* flag_cache_hint= */ i1 false)
+// CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.s2g.tile.override.addr.dim.stride.3d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i16 %8, i16 %9, i16 %10, i32 %13, i32 %14, i16 %17, i32 %3, i32 %4, i32 %5, i64 0, /* flag_cache_hint= */ i1 false)
+// CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.s2g.tile.override.addr.dim.stride.4d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i16 %8, i16 %9, i16 %10, i16 %11, i32 %13, i32 %14, i32 %15, i16 %17, i32 %3, i32 %4, i32 %5, i32 %6, i64 0, /* flag_cache_hint= */ i1 false)
+// CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.s2g.tile.override.addr.dim.stride.5d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i16 %8, i16 %9, i16 %10, i16 %11, i16 %12, i32 %13, i32 %14, i32 %15, i32 %16, i16 %17, i32 %3, i32 %4, i32 %5, i32 %6, i32 %7, i64 0, /* flag_cache_hint= */ i1 false)
+// CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.s2g.tile.override.addr.dim.1d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i16 %8, i32 %3, i64 %18, /* flag_cache_hint= */ i1 true)
+// CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.s2g.tile.override.addr.dim.stride.2d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i16 %8, i16 %9, i32 %13, i16 %17, i32 %3, i32 %4, i64 %18, /* flag_cache_hint= */ i1 true)
+// CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.s2g.tile.override.addr.dim.stride.3d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i16 %8, i16 %9, i16 %10, i32 %13, i32 %14, i16 %17, i32 %3, i32 %4, i32 %5, i64 %18, /* flag_cache_hint= */ i1 true)
+// CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.s2g.tile.override.addr.dim.stride.4d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i16 %8, i16 %9, i16 %10, i16 %11, i32 %13, i32 %14, i32 %15, i16 %17, i32 %3, i32 %4, i32 %5, i32 %6, i64 %18, /* flag_cache_hint= */ i1 true)
+// CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.s2g.tile.override.addr.dim.stride.5d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i16 %8, i16 %9, i16 %10, i16 %11, i16 %12, i32 %13, i32 %14, i32 %15, i32 %16, i16 %17, i32 %3, i32 %4, i32 %5, i32 %6, i32 %7, i64 %18, /* flag_cache_hint= */ i1 true)
+// CHECK-NEXT: ret void
+// CHECK-NEXT: }
+
+ // without cache hint
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0] tensor_size[%ts0] {mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd] {mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] tensor_size[%ts0, %ts1, %ts2] lower_stride[%lstrd0, %lstrd1] upper_stride[%ustrd] {mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3] tensor_size[%ts0, %ts1, %ts2, %ts3] lower_stride[%lstrd0, %lstrd1, %lstrd2] upper_stride[%ustrd] {mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3, %d4] tensor_size[%ts0, %ts1, %ts2, %ts3, %ts4] lower_stride[%lstrd0, %lstrd1, %lstrd2, %lstrd3] upper_stride[%ustrd] {mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ // with cache hint
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0] tensor_size[%ts0] l2_cache_hint = %ch {mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd] l2_cache_hint = %ch {mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] tensor_size[%ts0, %ts1, %ts2] lower_stride[%lstrd0, %lstrd1] upper_stride[%ustrd] l2_cache_hint = %ch {mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3] tensor_size[%ts0, %ts1, %ts2, %ts3] lower_stride[%lstrd0, %lstrd1, %lstrd2] upper_stride[%ustrd] l2_cache_hint = %ch {mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3, %d4] tensor_size[%ts0, %ts1, %ts2, %ts3, %ts4] lower_stride[%lstrd0, %lstrd1, %lstrd2, %lstrd3] upper_stride[%ustrd] l2_cache_hint = %ch {mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ llvm.return
+}
+
diff --git a/mlir/test/Target/LLVMIR/nvvm/tma_store_reduce_override.mlir b/mlir/test/Target/LLVMIR/nvvm/tma_store_reduce_override.mlir
new file mode 100644
index 0000000000000..1236155f9df29
--- /dev/null
+++ b/mlir/test/Target/LLVMIR/nvvm/tma_store_reduce_override.mlir
@@ -0,0 +1,372 @@
+// RUN: mlir-translate -mlir-to-llvmir %s | FileCheck %s
+
+llvm.func @tma_store_reduce_tile_override_addr(%src : !llvm.ptr<3>, %tma_desc : !llvm.ptr, %override_addr : !llvm.ptr<1>, %d0 : i32, %d1 : i32, %d2 : i32, %d3 : i32, %d4 : i32, %ts0 : i16, %ts1 : i16, %ts2 : i16, %ts3 : i16, %ts4 : i16, %lstrd0 : i32, %lstrd1 : i32, %lstrd2 : i32, %lstrd3 : i32, %ustrd : i16, %ch : i64) {
+ // CHECK-LABEL: define void @tma_store_reduce_tile_override_addr(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i32 %6, i32 %7, i16 %8, i16 %9, i16 %10, i16 %11, i16 %12, i32 %13, i32 %14, i32 %15, i32 %16, i16 %17, i64 %18) {
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.1d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i64 0, /* red_op=add */ i32 0, /* flag_cache_hint= */ i1 false)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.2d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i64 0, /* red_op=add */ i32 0, /* flag_cache_hint= */ i1 false)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.3d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i64 0, /* red_op=add */ i32 0, /* flag_cache_hint= */ i1 false)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.4d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i32 %6, i64 0, /* red_op=add */ i32 0, /* flag_cache_hint= */ i1 false)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.5d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i32 %6, i32 %7, i64 0, /* red_op=add */ i32 0, /* flag_cache_hint= */ i1 false)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.1d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i64 %18, /* red_op=add */ i32 0, /* flag_cache_hint= */ i1 true)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.2d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i64 %18, /* red_op=add */ i32 0, /* flag_cache_hint= */ i1 true)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.3d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i64 %18, /* red_op=add */ i32 0, /* flag_cache_hint= */ i1 true)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.4d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i32 %6, i64 %18, /* red_op=add */ i32 0, /* flag_cache_hint= */ i1 true)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.5d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i32 %6, i32 %7, i64 %18, /* red_op=add */ i32 0, /* flag_cache_hint= */ i1 true)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.1d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i64 0, /* red_op=min */ i32 1, /* flag_cache_hint= */ i1 false)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.1d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i64 %18, /* red_op=min */ i32 1, /* flag_cache_hint= */ i1 true)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.1d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i64 0, /* red_op=max */ i32 2, /* flag_cache_hint= */ i1 false)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.1d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i64 %18, /* red_op=max */ i32 2, /* flag_cache_hint= */ i1 true)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.1d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i64 0, /* red_op=inc */ i32 3, /* flag_cache_hint= */ i1 false)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.1d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i64 %18, /* red_op=inc */ i32 3, /* flag_cache_hint= */ i1 true)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.1d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i64 0, /* red_op=dec */ i32 4, /* flag_cache_hint= */ i1 false)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.1d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i64 %18, /* red_op=dec */ i32 4, /* flag_cache_hint= */ i1 true)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.1d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i64 0, /* red_op=and */ i32 5, /* flag_cache_hint= */ i1 false)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.1d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i64 %18, /* red_op=and */ i32 5, /* flag_cache_hint= */ i1 true)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.1d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i64 0, /* red_op=or */ i32 6, /* flag_cache_hint= */ i1 false)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.1d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i64 %18, /* red_op=or */ i32 6, /* flag_cache_hint= */ i1 true)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.1d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i64 0, /* red_op=xor */ i32 7, /* flag_cache_hint= */ i1 false)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.1d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i64 %18, /* red_op=xor */ i32 7, /* flag_cache_hint= */ i1 true)
+ // CHECK-NEXT: ret void
+ // CHECK-NEXT: }
+ // without cache hint
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3] {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3, %d4] {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ // with cache hint
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3, %d4] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ // Test min reduction
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] {redKind = #nvvm.tma_redux_kind<min>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<min>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ // Test max reduction
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] {redKind = #nvvm.tma_redux_kind<max>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<max>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ // Test inc reduction
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] {redKind = #nvvm.tma_redux_kind<inc>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<inc>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ // Test dec reduction
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] {redKind = #nvvm.tma_redux_kind<dec>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<dec>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ // Test and reduction
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] {redKind = #nvvm.tma_redux_kind<and>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<and>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ // Test or reduction
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] {redKind = #nvvm.tma_redux_kind<or>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<or>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ // Test xor reduction
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] {redKind = #nvvm.tma_redux_kind<xor>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<xor>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ llvm.return
+}
+
+llvm.func @tma_store_reduce_im2col_override_addr(%src : !llvm.ptr<3>, %tma_desc : !llvm.ptr, %override_addr : !llvm.ptr<1>, %d0 : i32, %d1 : i32, %d2 : i32, %d3 : i32, %d4 : i32, %ts0 : i16, %ts1 : i16, %ts2 : i16, %ts3 : i16, %ts4 : i16, %lstrd0 : i32, %lstrd1 : i32, %lstrd2 : i32, %lstrd3 : i32, %ustrd : i16, %ch : i64) {
+ // CHECK-LABEL: define void @tma_store_reduce_im2col_override_addr(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i32 %6, i32 %7, i16 %8, i16 %9, i16 %10, i16 %11, i16 %12, i32 %13, i32 %14, i32 %15, i32 %16, i16 %17, i64 %18) {
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.override.addr.3d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i64 0, /* red_op=add */ i32 0, /* flag_cache_hint= */ i1 false)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.override.addr.4d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i32 %6, i64 0, /* red_op=add */ i32 0, /* flag_cache_hint= */ i1 false)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.override.addr.5d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i32 %6, i32 %7, i64 0, /* red_op=add */ i32 0, /* flag_cache_hint= */ i1 false)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.override.addr.3d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i64 %18, /* red_op=add */ i32 0, /* flag_cache_hint= */ i1 true)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.override.addr.4d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i32 %6, i64 %18, /* red_op=add */ i32 0, /* flag_cache_hint= */ i1 true)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.override.addr.5d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i32 %6, i32 %7, i64 %18, /* red_op=add */ i32 0, /* flag_cache_hint= */ i1 true)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.override.addr.3d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i64 0, /* red_op=min */ i32 1, /* flag_cache_hint= */ i1 false)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.override.addr.3d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i64 %18, /* red_op=min */ i32 1, /* flag_cache_hint= */ i1 true)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.override.addr.3d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i64 0, /* red_op=max */ i32 2, /* flag_cache_hint= */ i1 false)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.override.addr.3d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i64 %18, /* red_op=max */ i32 2, /* flag_cache_hint= */ i1 true)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.override.addr.3d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i64 0, /* red_op=inc */ i32 3, /* flag_cache_hint= */ i1 false)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.override.addr.3d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i64 %18, /* red_op=inc */ i32 3, /* flag_cache_hint= */ i1 true)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.override.addr.3d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i64 0, /* red_op=dec */ i32 4, /* flag_cache_hint= */ i1 false)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.override.addr.3d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i64 %18, /* red_op=dec */ i32 4, /* flag_cache_hint= */ i1 true)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.override.addr.3d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i64 0, /* red_op=and */ i32 5, /* flag_cache_hint= */ i1 false)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.override.addr.3d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i64 %18, /* red_op=and */ i32 5, /* flag_cache_hint= */ i1 true)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.override.addr.3d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i64 0, /* red_op=or */ i32 6, /* flag_cache_hint= */ i1 false)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.override.addr.3d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i64 %18, /* red_op=or */ i32 6, /* flag_cache_hint= */ i1 true)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.override.addr.3d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i64 0, /* red_op=xor */ i32 7, /* flag_cache_hint= */ i1 false)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.override.addr.3d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i64 %18, /* red_op=xor */ i32 7, /* flag_cache_hint= */ i1 true)
+ // CHECK-NEXT: ret void
+ // CHECK-NEXT: }
+ // without cache hint
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<im2col>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3] {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<im2col>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3, %d4] {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<im2col>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ // with cache hint
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<im2col>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<im2col>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3, %d4] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<im2col>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ // Test min reduction
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] {redKind = #nvvm.tma_redux_kind<min>, mode = #nvvm.tma_store_mode<im2col>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<min>, mode = #nvvm.tma_store_mode<im2col>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ // Test max reduction
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] {redKind = #nvvm.tma_redux_kind<max>, mode = #nvvm.tma_store_mode<im2col>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<max>, mode = #nvvm.tma_store_mode<im2col>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ // Test inc reduction
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] {redKind = #nvvm.tma_redux_kind<inc>, mode = #nvvm.tma_store_mode<im2col>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<inc>, mode = #nvvm.tma_store_mode<im2col>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ // Test dec reduction
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] {redKind = #nvvm.tma_redux_kind<dec>, mode = #nvvm.tma_store_mode<im2col>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<dec>, mode = #nvvm.tma_store_mode<im2col>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ // Test and reduction
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] {redKind = #nvvm.tma_redux_kind<and>, mode = #nvvm.tma_store_mode<im2col>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<and>, mode = #nvvm.tma_store_mode<im2col>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ // Test or reduction
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] {redKind = #nvvm.tma_redux_kind<or>, mode = #nvvm.tma_store_mode<im2col>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<or>, mode = #nvvm.tma_store_mode<im2col>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ // Test xor reduction
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] {redKind = #nvvm.tma_redux_kind<xor>, mode = #nvvm.tma_store_mode<im2col>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<xor>, mode = #nvvm.tma_store_mode<im2col>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ llvm.return
+}
+
+llvm.func @tma_store_reduce_im2col_w_override_addr(%src : !llvm.ptr<3>, %tma_desc : !llvm.ptr, %override_addr : !llvm.ptr<1>, %d0 : i32, %d1 : i32, %d2 : i32, %d3 : i32, %d4 : i32, %ts0 : i16, %ts1 : i16, %ts2 : i16, %ts3 : i16, %ts4 : i16, %lstrd0 : i32, %lstrd1 : i32, %lstrd2 : i32, %lstrd3 : i32, %ustrd : i16, %ch : i64) {
+ // CHECK-LABEL: define void @tma_store_reduce_im2col_w_override_addr(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i32 %6, i32 %7, i16 %8, i16 %9, i16 %10, i16 %11, i16 %12, i32 %13, i32 %14, i32 %15, i32 %16, i16 %17, i64 %18) {
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.w.override.addr.3d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i64 0, /* red_op=add */ i32 0, /* flag_cache_hint= */ i1 false)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.w.override.addr.4d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i32 %6, i64 0, /* red_op=add */ i32 0, /* flag_cache_hint= */ i1 false)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.w.override.addr.5d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i32 %6, i32 %7, i64 0, /* red_op=add */ i32 0, /* flag_cache_hint= */ i1 false)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.w.override.addr.3d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i64 %18, /* red_op=add */ i32 0, /* flag_cache_hint= */ i1 true)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.w.override.addr.4d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i32 %6, i64 %18, /* red_op=add */ i32 0, /* flag_cache_hint= */ i1 true)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.w.override.addr.5d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i32 %6, i32 %7, i64 %18, /* red_op=add */ i32 0, /* flag_cache_hint= */ i1 true)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.w.override.addr.3d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i64 0, /* red_op=min */ i32 1, /* flag_cache_hint= */ i1 false)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.w.override.addr.3d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i64 %18, /* red_op=min */ i32 1, /* flag_cache_hint= */ i1 true)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.w.override.addr.3d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i64 0, /* red_op=max */ i32 2, /* flag_cache_hint= */ i1 false)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.w.override.addr.3d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i64 %18, /* red_op=max */ i32 2, /* flag_cache_hint= */ i1 true)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.w.override.addr.3d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i64 0, /* red_op=inc */ i32 3, /* flag_cache_hint= */ i1 false)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.w.override.addr.3d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i64 %18, /* red_op=inc */ i32 3, /* flag_cache_hint= */ i1 true)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.w.override.addr.3d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i64 0, /* red_op=dec */ i32 4, /* flag_cache_hint= */ i1 false)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.w.override.addr.3d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i64 %18, /* red_op=dec */ i32 4, /* flag_cache_hint= */ i1 true)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.w.override.addr.3d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i64 0, /* red_op=and */ i32 5, /* flag_cache_hint= */ i1 false)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.w.override.addr.3d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i64 %18, /* red_op=and */ i32 5, /* flag_cache_hint= */ i1 true)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.w.override.addr.3d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i64 0, /* red_op=or */ i32 6, /* flag_cache_hint= */ i1 false)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.w.override.addr.3d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i64 %18, /* red_op=or */ i32 6, /* flag_cache_hint= */ i1 true)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.w.override.addr.3d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i64 0, /* red_op=xor */ i32 7, /* flag_cache_hint= */ i1 false)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.w.override.addr.3d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i64 %18, /* red_op=xor */ i32 7, /* flag_cache_hint= */ i1 true)
+ // CHECK-NEXT: ret void
+ // CHECK-NEXT: }
+ // without cache hint
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<im2col_w>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3] {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<im2col_w>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3, %d4] {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<im2col_w>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ // with cache hint
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<im2col_w>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<im2col_w>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3, %d4] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<im2col_w>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ // Test min reduction
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] {redKind = #nvvm.tma_redux_kind<min>, mode = #nvvm.tma_store_mode<im2col_w>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<min>, mode = #nvvm.tma_store_mode<im2col_w>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ // Test max reduction
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] {redKind = #nvvm.tma_redux_kind<max>, mode = #nvvm.tma_store_mode<im2col_w>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<max>, mode = #nvvm.tma_store_mode<im2col_w>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ // Test inc reduction
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] {redKind = #nvvm.tma_redux_kind<inc>, mode = #nvvm.tma_store_mode<im2col_w>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<inc>, mode = #nvvm.tma_store_mode<im2col_w>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ // Test dec reduction
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] {redKind = #nvvm.tma_redux_kind<dec>, mode = #nvvm.tma_store_mode<im2col_w>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<dec>, mode = #nvvm.tma_store_mode<im2col_w>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ // Test and reduction
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] {redKind = #nvvm.tma_redux_kind<and>, mode = #nvvm.tma_store_mode<im2col_w>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<and>, mode = #nvvm.tma_store_mode<im2col_w>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ // Test or reduction
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] {redKind = #nvvm.tma_redux_kind<or>, mode = #nvvm.tma_store_mode<im2col_w>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<or>, mode = #nvvm.tma_store_mode<im2col_w>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ // Test xor reduction
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] {redKind = #nvvm.tma_redux_kind<xor>, mode = #nvvm.tma_store_mode<im2col_w>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<xor>, mode = #nvvm.tma_store_mode<im2col_w>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ llvm.return
+}
+
+llvm.func @tma_store_reduce_tile_override_addr_dim_stride(%src : !llvm.ptr<3>, %tma_desc : !llvm.ptr, %override_addr : !llvm.ptr<1>, %d0 : i32, %d1 : i32, %d2 : i32, %d3 : i32, %d4 : i32, %ts0 : i16, %ts1 : i16, %ts2 : i16, %ts3 : i16, %ts4 : i16, %lstrd0 : i32, %lstrd1 : i32, %lstrd2 : i32, %lstrd3 : i32, %ustrd : i16, %ch : i64) {
+ // CHECK-LABEL: define void @tma_store_reduce_tile_override_addr_dim_stride(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i32 %6, i32 %7, i16 %8, i16 %9, i16 %10, i16 %11, i16 %12, i32 %13, i32 %14, i32 %15, i32 %16, i16 %17, i64 %18) {
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.dim.1d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i16 %8, i32 %3, i64 0, /* red_op=add */ i32 0, /* flag_cache_hint= */ i1 false)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.dim.stride.2d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i16 %8, i16 %9, i32 %13, i16 %17, i32 %3, i32 %4, i64 0, /* red_op=add */ i32 0, /* flag_cache_hint= */ i1 false)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.dim.stride.3d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i16 %8, i16 %9, i16 %10, i32 %13, i32 %14, i16 %17, i32 %3, i32 %4, i32 %5, i64 0, /* red_op=add */ i32 0, /* flag_cache_hint= */ i1 false)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.dim.stride.4d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i16 %8, i16 %9, i16 %10, i16 %11, i32 %13, i32 %14, i32 %15, i16 %17, i32 %3, i32 %4, i32 %5, i32 %6, i64 0, /* red_op=add */ i32 0, /* flag_cache_hint= */ i1 false)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.dim.stride.5d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i16 %8, i16 %9, i16 %10, i16 %11, i16 %12, i32 %13, i32 %14, i32 %15, i32 %16, i16 %17, i32 %3, i32 %4, i32 %5, i32 %6, i32 %7, i64 0, /* red_op=add */ i32 0, /* flag_cache_hint= */ i1 false)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.dim.1d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i16 %8, i32 %3, i64 %18, /* red_op=add */ i32 0, /* flag_cache_hint= */ i1 true)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.dim.stride.2d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i16 %8, i16 %9, i32 %13, i16 %17, i32 %3, i32 %4, i64 %18, /* red_op=add */ i32 0, /* flag_cache_hint= */ i1 true)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.dim.stride.3d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i16 %8, i16 %9, i16 %10, i32 %13, i32 %14, i16 %17, i32 %3, i32 %4, i32 %5, i64 %18, /* red_op=add */ i32 0, /* flag_cache_hint= */ i1 true)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.dim.stride.4d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i16 %8, i16 %9, i16 %10, i16 %11, i32 %13, i32 %14, i32 %15, i16 %17, i32 %3, i32 %4, i32 %5, i32 %6, i64 %18, /* red_op=add */ i32 0, /* flag_cache_hint= */ i1 true)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.dim.stride.5d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i16 %8, i16 %9, i16 %10, i16 %11, i16 %12, i32 %13, i32 %14, i32 %15, i32 %16, i16 %17, i32 %3, i32 %4, i32 %5, i32 %6, i32 %7, i64 %18, /* red_op=add */ i32 0, /* flag_cache_hint= */ i1 true)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.dim.1d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i16 %8, i32 %3, i64 0, /* red_op=min */ i32 1, /* flag_cache_hint= */ i1 false)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.dim.1d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i16 %8, i32 %3, i64 %18, /* red_op=min */ i32 1, /* flag_cache_hint= */ i1 true)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.dim.stride.2d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i16 %8, i16 %9, i32 %13, i16 %17, i32 %3, i32 %4, i64 0, /* red_op=min */ i32 1, /* flag_cache_hint= */ i1 false)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.dim.stride.2d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i16 %8, i16 %9, i32 %13, i16 %17, i32 %3, i32 %4, i64 %18, /* red_op=min */ i32 1, /* flag_cache_hint= */ i1 true)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.dim.1d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i16 %8, i32 %3, i64 0, /* red_op=max */ i32 2, /* flag_cache_hint= */ i1 false)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.dim.1d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i16 %8, i32 %3, i64 %18, /* red_op=max */ i32 2, /* flag_cache_hint= */ i1 true)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.dim.stride.2d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i16 %8, i16 %9, i32 %13, i16 %17, i32 %3, i32 %4, i64 0, /* red_op=max */ i32 2, /* flag_cache_hint= */ i1 false)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.dim.stride.2d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i16 %8, i16 %9, i32 %13, i16 %17, i32 %3, i32 %4, i64 %18, /* red_op=max */ i32 2, /* flag_cache_hint= */ i1 true)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.dim.1d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i16 %8, i32 %3, i64 0, /* red_op=inc */ i32 3, /* flag_cache_hint= */ i1 false)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.dim.1d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i16 %8, i32 %3, i64 %18, /* red_op=inc */ i32 3, /* flag_cache_hint= */ i1 true)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.dim.stride.2d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i16 %8, i16 %9, i32 %13, i16 %17, i32 %3, i32 %4, i64 0, /* red_op=inc */ i32 3, /* flag_cache_hint= */ i1 false)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.dim.stride.2d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i16 %8, i16 %9, i32 %13, i16 %17, i32 %3, i32 %4, i64 %18, /* red_op=inc */ i32 3, /* flag_cache_hint= */ i1 true)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.dim.1d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i16 %8, i32 %3, i64 0, /* red_op=dec */ i32 4, /* flag_cache_hint= */ i1 false)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.dim.1d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i16 %8, i32 %3, i64 %18, /* red_op=dec */ i32 4, /* flag_cache_hint= */ i1 true)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.dim.stride.2d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i16 %8, i16 %9, i32 %13, i16 %17, i32 %3, i32 %4, i64 0, /* red_op=dec */ i32 4, /* flag_cache_hint= */ i1 false)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.dim.stride.2d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i16 %8, i16 %9, i32 %13, i16 %17, i32 %3, i32 %4, i64 %18, /* red_op=dec */ i32 4, /* flag_cache_hint= */ i1 true)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.dim.1d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i16 %8, i32 %3, i64 0, /* red_op=and */ i32 5, /* flag_cache_hint= */ i1 false)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.dim.1d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i16 %8, i32 %3, i64 %18, /* red_op=and */ i32 5, /* flag_cache_hint= */ i1 true)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.dim.stride.2d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i16 %8, i16 %9, i32 %13, i16 %17, i32 %3, i32 %4, i64 0, /* red_op=and */ i32 5, /* flag_cache_hint= */ i1 false)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.dim.stride.2d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i16 %8, i16 %9, i32 %13, i16 %17, i32 %3, i32 %4, i64 %18, /* red_op=and */ i32 5, /* flag_cache_hint= */ i1 true)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.dim.1d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i16 %8, i32 %3, i64 0, /* red_op=or */ i32 6, /* flag_cache_hint= */ i1 false)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.dim.1d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i16 %8, i32 %3, i64 %18, /* red_op=or */ i32 6, /* flag_cache_hint= */ i1 true)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.dim.stride.2d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i16 %8, i16 %9, i32 %13, i16 %17, i32 %3, i32 %4, i64 0, /* red_op=or */ i32 6, /* flag_cache_hint= */ i1 false)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.dim.stride.2d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i16 %8, i16 %9, i32 %13, i16 %17, i32 %3, i32 %4, i64 %18, /* red_op=or */ i32 6, /* flag_cache_hint= */ i1 true)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.dim.1d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i16 %8, i32 %3, i64 0, /* red_op=xor */ i32 7, /* flag_cache_hint= */ i1 false)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.dim.1d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i16 %8, i32 %3, i64 %18, /* red_op=xor */ i32 7, /* flag_cache_hint= */ i1 true)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.dim.stride.2d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i16 %8, i16 %9, i32 %13, i16 %17, i32 %3, i32 %4, i64 0, /* red_op=xor */ i32 7, /* flag_cache_hint= */ i1 false)
+ // CHECK-NEXT: call void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.override.addr.dim.stride.2d(ptr addrspace(3) %0, ptr %1, ptr addrspace(1) %2, i16 %8, i16 %9, i32 %13, i16 %17, i32 %3, i32 %4, i64 %18, /* red_op=xor */ i32 7, /* flag_cache_hint= */ i1 true)
+ // CHECK-NEXT: ret void
+ // CHECK-NEXT: }
+ // without cache hint
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] tensor_size[%ts0] {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd] {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] tensor_size[%ts0, %ts1, %ts2] lower_stride[%lstrd0, %lstrd1] upper_stride[%ustrd] {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3] tensor_size[%ts0, %ts1, %ts2, %ts3] lower_stride[%lstrd0, %lstrd1, %lstrd2] upper_stride[%ustrd] {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3, %d4] tensor_size[%ts0, %ts1, %ts2, %ts3, %ts4] lower_stride[%lstrd0, %lstrd1, %lstrd2, %lstrd3] upper_stride[%ustrd] {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ // with cache hint
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] tensor_size[%ts0] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] tensor_size[%ts0, %ts1, %ts2] lower_stride[%lstrd0, %lstrd1] upper_stride[%ustrd] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3] tensor_size[%ts0, %ts1, %ts2, %ts3] lower_stride[%lstrd0, %lstrd1, %lstrd2] upper_stride[%ustrd] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3, %d4] tensor_size[%ts0, %ts1, %ts2, %ts3, %ts4] lower_stride[%lstrd0, %lstrd1, %lstrd2, %lstrd3] upper_stride[%ustrd] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ // Test min reduction
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] tensor_size[%ts0] {redKind = #nvvm.tma_redux_kind<min>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] tensor_size[%ts0] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<min>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd] {redKind = #nvvm.tma_redux_kind<min>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<min>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ // Test max reduction
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] tensor_size[%ts0] {redKind = #nvvm.tma_redux_kind<max>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] tensor_size[%ts0] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<max>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd] {redKind = #nvvm.tma_redux_kind<max>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<max>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ // Test inc reduction
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] tensor_size[%ts0] {redKind = #nvvm.tma_redux_kind<inc>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] tensor_size[%ts0] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<inc>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd] {redKind = #nvvm.tma_redux_kind<inc>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<inc>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ // Test dec reduction
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] tensor_size[%ts0] {redKind = #nvvm.tma_redux_kind<dec>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] tensor_size[%ts0] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<dec>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd] {redKind = #nvvm.tma_redux_kind<dec>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<dec>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ // Test and reduction
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] tensor_size[%ts0] {redKind = #nvvm.tma_redux_kind<and>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] tensor_size[%ts0] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<and>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd] {redKind = #nvvm.tma_redux_kind<and>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<and>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ // Test or reduction
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] tensor_size[%ts0] {redKind = #nvvm.tma_redux_kind<or>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] tensor_size[%ts0] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<or>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd] {redKind = #nvvm.tma_redux_kind<or>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<or>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ // Test xor reduction
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] tensor_size[%ts0] {redKind = #nvvm.tma_redux_kind<xor>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] tensor_size[%ts0] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<xor>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd] {redKind = #nvvm.tma_redux_kind<xor>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<xor>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+
+ llvm.return
+}
+
>From 21f64e2bc80df3033fc8609f0ca0754b0709cee6 Mon Sep 17 00:00:00 2001
From: Rajat Bajpai <rbajpai at nvidia.com>
Date: Sat, 22 Aug 2026 12:32:17 +0000
Subject: [PATCH 2/3] Addressed review comments
---
mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td | 48 +++++++--------------
mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp | 28 ++++++------
2 files changed, 30 insertions(+), 46 deletions(-)
diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index 0009a7facda6a..51156e91d1328 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -4519,7 +4519,7 @@ def NVVM_CpAsyncBulkTensorGlobalToSharedClusterOp :
}
//===----------------------------------------------------------------------===//
-// NVVM S2G Ops
+// NVVM TMA Shared to Global Ops
//===----------------------------------------------------------------------===//
def NVVM_CpAsyncBulkTensorSharedCTAToGlobalOp :
@@ -4588,9 +4588,9 @@ def NVVM_CpAsyncBulkTensorSharedCTAToGlobalOverrideAddrOp :
NVVM_VoidIntrinsicOp<"cp.async.bulk.tensor.global.shared.cta.override",
[AttrSizedOperandSegments]> {
let arguments = (ins
- LLVM_AnyPointer:$tmaDesc,
+ LLVM_PointerGeneric:$tmaDescriptor,
LLVM_PointerShared:$srcMem,
- LLVM_PointerGlobal:$overrideAdrr,
+ LLVM_PointerGlobal:$overrideAddr,
Variadic<I32>:$coordinates,
Variadic<I16>:$tensorSize,
Variadic<I32>:$lowerStride,
@@ -4609,17 +4609,17 @@ def NVVM_CpAsyncBulkTensorSharedCTAToGlobalOverrideAddrOp :
The `mode` attribute selects the store mode. The override variant is selected by
which optional operands are provided:
- - `override.addr` (`.override::global_address`): only `overrideAdrr` is given; the
+ - `override.addr` (`.override::global_address`): only `overrideAddr` is given; the
global base address from the tensor-map is replaced. Supported in `TILE` (1D–5D),
`IM2COL`/`IM2COL_W` (3D–5D), and `TILE_SCATTER4` (2D, requires 5 coordinates).
- - `override.addr.dim` (1D, `TILE` only): `overrideAdrr` plus `tensorSize` (one
+ - `override.addr.dim` (1D, `TILE` only): `overrideAddr` plus `tensorSize` (one
element); also overrides the tensor global dimension.
- - `override.addr.dim.stride` (2D–5D, `TILE` only): `overrideAdrr` plus `tensorSize`,
+ - `override.addr.dim.stride` (2D–5D, `TILE` only): `overrideAddr` plus `tensorSize`,
`lowerStride`, and `upperStride`; also overrides the global dimensions and strides.
The effective global stride is
`global_stride[i] = ((%stride{i} + (%upper_stride{i} << 32)) << 4)`.
- `overrideAdrr` must be 16B aligned (a runtime error is raised otherwise) and the
+ `overrideAddr` must be 16B aligned (a runtime error is raised otherwise) and the
memory range `[%override_addr, %override_addr + 128 KiB)` must be allocated and
accessible during execution. When overriding dimensions/strides, the base-address
override is mandatory and the tensor start coordinates must be zero; otherwise the
@@ -4671,15 +4671,15 @@ def NVVM_CpAsyncBulkTensorSharedCTAToGlobalOverrideAddrOp :
}];
let assemblyFormat = [{
- $tmaDesc `,`
+ $tmaDescriptor `,`
$srcMem `,`
- $overrideAdrr `,`
+ $overrideAddr `,`
`box` `[`$coordinates `]`
(`tensor_size` `[`$tensorSize^ `]`)?
(`lower_stride` `[`$lowerStride^ `]`)?
(`upper_stride` `[`$upperStride^ `]`)?
(`l2_cache_hint` `=` $l2CacheHint^)?
- attr-dict `:` type($tmaDesc) `,` type($srcMem) `,` type($overrideAdrr)
+ attr-dict `:` type($tmaDescriptor) `,` type($srcMem) `,` type($overrideAddr)
}];
let hasVerifier = 1;
@@ -4905,9 +4905,9 @@ def NVVM_CpAsyncBulkTensorReduceOverrideAddrOp :
NVVM_VoidIntrinsicOp<"cp.async.bulk.tensor.reduce.override",
[AttrSizedOperandSegments]> {
let arguments = (ins
- LLVM_AnyPointer:$tmaDesc,
+ LLVM_PointerGeneric:$tmaDescriptor,
LLVM_PointerShared:$srcMem,
- LLVM_PointerGlobal:$overrideAdrr,
+ LLVM_PointerGlobal:$overrideAddr,
Variadic<I32>:$coordinates,
Variadic<I16>:$tensorSize,
Variadic<I32>:$lowerStride,
@@ -4929,23 +4929,7 @@ def NVVM_CpAsyncBulkTensorReduceOverrideAddrOp :
The `mode` attribute selects the store mode and `redKind` selects the reduction
operation (`ADD`, `MIN`, `MAX`, `INC`, `DEC`, `AND`, `OR`, `XOR`) that combines the
source data in shared memory with the destination data in global memory. The override
- variant is selected by which optional operands are provided:
-
- - `override.addr` (`.override::global_address`): only `overrideAdrr` is given; the
- global base address from the tensor-map is replaced. Supported in `TILE` (1D–5D),
- `IM2COL` and `IM2COL_W` (3D–5D). `TILE_SCATTER4` is not supported.
- - `override.addr.dim` (1D, `TILE` only): `overrideAdrr` plus `tensorSize` (one
- element); also overrides the tensor global dimension.
- - `override.addr.dim.stride` (2D–5D, `TILE` only): `overrideAdrr` plus `tensorSize`,
- `lowerStride`, and `upperStride`; also overrides the global dimensions and strides.
- The effective global stride is
- `global_stride[i] = ((%stride{i} + (%upper_stride{i} << 32)) << 4)`.
-
- `overrideAdrr` must be 16B aligned (a runtime error is raised otherwise) and the
- memory range `[%override_addr, %override_addr + 128 KiB)` must be allocated and
- accessible during execution. When overriding dimensions/strides, the base-address
- override is mandatory and the tensor start coordinates must be zero; otherwise the
- behavior is undefined.
+ variant is similar to the `cp.async.bulk.tensor.global.shared.cta.override` op described above.
The optional `l2CacheHint` specifies a cache-eviction policy for the access.
@@ -4990,15 +4974,15 @@ def NVVM_CpAsyncBulkTensorReduceOverrideAddrOp :
}];
let assemblyFormat = [{
- $tmaDesc `,`
+ $tmaDescriptor `,`
$srcMem `,`
- $overrideAdrr `,`
+ $overrideAddr `,`
`box` `[`$coordinates `]`
(`tensor_size` `[`$tensorSize^ `]`)?
(`lower_stride` `[`$lowerStride^ `]`)?
(`upper_stride` `[`$upperStride^ `]`)?
(`l2_cache_hint` `=` $l2CacheHint^)?
- attr-dict `:` type($tmaDesc) `,` type($srcMem) `,` type($overrideAdrr)
+ attr-dict `:` type($tmaDescriptor) `,` type($srcMem) `,` type($overrideAddr)
}];
let hasVerifier = 1;
diff --git a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
index 3da38367aafd4..ec619a23475aa 100644
--- a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
+++ b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
@@ -158,24 +158,28 @@ LogicalResult CpAsyncBulkTensorOverrideAddrCommonVerifier(
OperandRange coordinates, OperandRange tensorSize, OperandRange lowerStride,
Value upperStride, bool isTile, Location loc) {
LogicalResult res = success();
- if (!tensorSize.empty() && coordinates.size() != tensorSize.size())
+ if (!tensorSize.empty() && coordinates.size() != tensorSize.size()) {
res =
emitError(loc, "Expected coordinates size to be equal to tensor size");
+ }
- if (!lowerStride.empty() && lowerStride.size() != tensorSize.size() - 1)
+ if (!lowerStride.empty() && lowerStride.size() != tensorSize.size() - 1) {
res = emitError(
loc,
"Expected lower_stride size to be equal to one less than tensor size");
+ }
- if (!lowerStride.empty() != static_cast<bool>(upperStride))
+ if (!lowerStride.empty() != static_cast<bool>(upperStride)) {
res = emitError(loc,
"Expected lower_stride and upper_stride to be either both "
"present or both absent");
+ }
bool isDimStride = tensorSize.size() > 0;
- if (!isTile && isDimStride)
+ if (!isTile && isDimStride) {
res = emitError(
loc, "Only tile mode supports override address with dim and stride");
+ }
return res;
}
@@ -4803,8 +4807,8 @@ CpAsyncBulkTensorSharedCTAToGlobalOverrideAddrOp::getIntrinsicIDAndArgs(
llvm::SmallVector<llvm::Value *> args;
args.push_back(mt.lookupValue(thisOp.getSrcMem()));
- args.push_back(mt.lookupValue(thisOp.getTmaDesc()));
- args.push_back(mt.lookupValue(thisOp.getOverrideAdrr()));
+ args.push_back(mt.lookupValue(thisOp.getTmaDescriptor()));
+ args.push_back(mt.lookupValue(thisOp.getOverrideAddr()));
for (Value v : thisOp.getTensorSize())
args.push_back(mt.lookupValue(v));
for (Value v : thisOp.getLowerStride())
@@ -4855,9 +4859,7 @@ CpAsyncBulkTensorSharedCTAToGlobalOverrideAddrOp::getIntrinsicIDAndArgs(
assert(mode < std::size(IDTable) &&
"Invalid mode for CpAsyncBulkTensorSharedCTAToGlobalOverrideAddrOp");
- assert(dim < std::size(IDTable[mode]) &&
- "Invalid dim for CpAsyncBulkTensorSharedCTAToGlobalOverrideAddrOp");
- assert(dim < std::size(dimStrideIDTable) &&
+ assert(dim < std::size(IDTable[mode]) && dim < std::size(dimStrideIDTable) &&
"Invalid dim for CpAsyncBulkTensorSharedCTAToGlobalOverrideAddrOp");
ID intrinsicID = isDimStride ? dimStrideIDTable[dim] : IDTable[mode][dim];
@@ -4919,8 +4921,8 @@ NVVM::IDArgPair CpAsyncBulkTensorReduceOverrideAddrOp::getIntrinsicIDAndArgs(
llvm::SmallVector<llvm::Value *> args;
args.push_back(mt.lookupValue(thisOp.getSrcMem()));
- args.push_back(mt.lookupValue(thisOp.getTmaDesc()));
- args.push_back(mt.lookupValue(thisOp.getOverrideAdrr()));
+ args.push_back(mt.lookupValue(thisOp.getTmaDescriptor()));
+ args.push_back(mt.lookupValue(thisOp.getOverrideAddr()));
for (Value v : thisOp.getTensorSize())
args.push_back(mt.lookupValue(v));
@@ -4972,9 +4974,7 @@ static constexpr ID dimStrideIDTable[] = {
assert(mode < std::size(IDTable) &&
"Invalid mode for CpAsyncBulkTensorReduceOverrideAddrOp");
- assert(dim < std::size(IDTable[mode]) &&
- "Invalid dim for CpAsyncBulkTensorReduceOverrideAddrOp");
- assert(dim < std::size(dimStrideIDTable) &&
+ assert(dim < std::size(IDTable[mode]) && dim < std::size(dimStrideIDTable) &&
"Invalid dim for CpAsyncBulkTensorReduceOverrideAddrOp");
ID intrinsicID = isDimStride ? dimStrideIDTable[dim] : IDTable[mode][dim];
>From aadbdfb64f233584aef26f8f92922178b5b5185b Mon Sep 17 00:00:00 2001
From: Rajat Bajpai <rbajpai at nvidia.com>
Date: Sat, 22 Aug 2026 13:42:45 +0000
Subject: [PATCH 3/3] Incorporate assembly format changes
---
mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td | 55 ++---
.../tma_reduce_override_invalid.mlir | 9 +-
.../tma_store_override_invalid.mlir | 13 +-
.../LLVMIR/nvvm/tma_store_override.mlir | 69 +++---
.../nvvm/tma_store_reduce_override.mlir | 205 +++++++++---------
5 files changed, 165 insertions(+), 186 deletions(-)
diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index 51156e91d1328..3226782937d0c 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -4631,40 +4631,31 @@ def NVVM_CpAsyncBulkTensorSharedCTAToGlobalOverrideAddrOp :
// override.addr (TILE, 1D)
```mlir
- nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0] {
- mode = #nvvm.tma_store_mode<tile>
- } : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0] : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
```
// override.addr (IM2COL, 3D) with L2 cache hint
```mlir
nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2]
- l2_cache_hint = %ch {
- mode = #nvvm.tma_store_mode<im2col>
- } : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ l2_cache_hint = %ch mode = im2col : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
```
// override.addr (TILE_SCATTER4, 2D) — 5 coordinates: x0, y0, y1, y2, y3
```mlir
- nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%x0, %y0, %y1, %y2, %y3] {
- mode = #nvvm.tma_store_mode<tile_scatter4>
- } : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%x0, %y0, %y1, %y2, %y3]
+ mode = tile_scatter4 : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
```
// override.addr.dim (TILE, 1D)
```mlir
nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr,
- box[%d0] tensor_size[%ts0] {
- mode = #nvvm.tma_store_mode<tile>
- } : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ box[%d0] tensor_size[%ts0] : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
```
// override.addr.dim.stride (TILE, 2D)
```mlir
nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr,
- box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd] {
- mode = #nvvm.tma_store_mode<tile>
- } : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd] : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
```
[For more information, see PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-bulk-tensor)
@@ -4675,11 +4666,14 @@ def NVVM_CpAsyncBulkTensorSharedCTAToGlobalOverrideAddrOp :
$srcMem `,`
$overrideAddr `,`
`box` `[`$coordinates `]`
- (`tensor_size` `[`$tensorSize^ `]`)?
- (`lower_stride` `[`$lowerStride^ `]`)?
- (`upper_stride` `[`$upperStride^ `]`)?
- (`l2_cache_hint` `=` $l2CacheHint^)?
- attr-dict `:` type($tmaDescriptor) `,` type($srcMem) `,` type($overrideAddr)
+ oilist(
+ `tensor_size` `[` $tensorSize `]`
+ | `lower_stride` `[` $lowerStride `]`
+ | `upper_stride` `[` $upperStride `]`
+ | `l2_cache_hint` `=` $l2CacheHint
+ | `mode` `=` $mode
+ )
+ attr-dict `:` type($tmaDescriptor) `,` type($srcMem) `,` type($overrideAddr)
}];
let hasVerifier = 1;
@@ -4937,37 +4931,25 @@ def NVVM_CpAsyncBulkTensorReduceOverrideAddrOp :
// override.addr (TILE, 1D) with ADD reduction
```mlir
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] {
- redKind = #nvvm.tma_redux_kind<add>,
- mode = #nvvm.tma_store_mode<tile>
- } : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0], reduction = add : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
```
// override.addr (IM2COL, 3D) with L2 cache hint and MIN reduction
```mlir
nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2]
- l2_cache_hint = %ch {
- redKind = #nvvm.tma_redux_kind<min>,
- mode = #nvvm.tma_store_mode<im2col>
- } : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ l2_cache_hint = %ch, reduction = min mode = im2col : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
```
// override.addr.dim (TILE, 1D)
```mlir
nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr,
- box[%d0] tensor_size[%ts0] {
- redKind = #nvvm.tma_redux_kind<add>,
- mode = #nvvm.tma_store_mode<tile>
- } : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ box[%d0] tensor_size[%ts0], reduction = add : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
```
// override.addr.dim.stride (TILE, 2D)
```mlir
nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr,
- box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd] {
- redKind = #nvvm.tma_redux_kind<and>,
- mode = #nvvm.tma_store_mode<tile>
- } : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd], reduction = and : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
```
[For more information, see PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-reduce-async-bulk-tensor)
@@ -4982,6 +4964,7 @@ def NVVM_CpAsyncBulkTensorReduceOverrideAddrOp :
(`lower_stride` `[`$lowerStride^ `]`)?
(`upper_stride` `[`$upperStride^ `]`)?
(`l2_cache_hint` `=` $l2CacheHint^)?
+ `,` `reduction` `=` $redKind oilist(`mode` `=` $mode)
attr-dict `:` type($tmaDescriptor) `,` type($srcMem) `,` type($overrideAddr)
}];
diff --git a/mlir/test/Target/LLVMIR/nvvm/nvvmir-invalid/tma_reduce_override_invalid.mlir b/mlir/test/Target/LLVMIR/nvvm/nvvmir-invalid/tma_reduce_override_invalid.mlir
index a723acc1fccc2..a714e7be0f575 100644
--- a/mlir/test/Target/LLVMIR/nvvm/nvvmir-invalid/tma_reduce_override_invalid.mlir
+++ b/mlir/test/Target/LLVMIR/nvvm/nvvmir-invalid/tma_reduce_override_invalid.mlir
@@ -4,7 +4,7 @@
llvm.func @tma_store_reduce_tile_override_addr_coord_tensor_size_mismatch(%src : !llvm.ptr<3>, %tma_desc : !llvm.ptr, %override_addr : !llvm.ptr<1>, %d0 : i32, %d1 : i32, %d2 : i32, %d3 : i32, %d4 : i32, %ts0 : i16, %ts1 : i16, %ts2 : i16, %ts3 : i16, %ts4 : i16, %lstrd0 : i32, %lstrd1 : i32, %lstrd2 : i32, %lstrd3 : i32, %ustrd : i16, %ch : i64) {
// expected-error @below {{Expected coordinates size to be equal to tensor size}}
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0] {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0], reduction = add : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
llvm.return
}
@@ -12,7 +12,7 @@ llvm.func @tma_store_reduce_tile_override_addr_coord_tensor_size_mismatch(%src :
llvm.func @tma_store_reduce_tile_override_addr_lower_stride_tensor_size_mismatch(%src : !llvm.ptr<3>, %tma_desc : !llvm.ptr, %override_addr : !llvm.ptr<1>, %d0 : i32, %d1 : i32, %d2 : i32, %d3 : i32, %d4 : i32, %ts0 : i16, %ts1 : i16, %ts2 : i16, %ts3 : i16, %ts4 : i16, %lstrd0 : i32, %lstrd1 : i32, %lstrd2 : i32, %lstrd3 : i32, %ustrd : i16, %ch : i64) {
// expected-error @below {{Expected lower_stride size to be equal to one less than tensor size}}
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0, %lstrd1] upper_stride[%ustrd] {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0, %lstrd1] upper_stride[%ustrd], reduction = add : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
llvm.return
}
@@ -20,7 +20,7 @@ llvm.func @tma_store_reduce_tile_override_addr_lower_stride_tensor_size_mismatch
llvm.func @tma_store_reduce_tile_override_addr_stride_mismatch(%src : !llvm.ptr<3>, %tma_desc : !llvm.ptr, %override_addr : !llvm.ptr<1>, %d0 : i32, %d1 : i32, %d2 : i32, %d3 : i32, %d4 : i32, %ts0 : i16, %ts1 : i16, %ts2 : i16, %ts3 : i16, %ts4 : i16, %lstrd0 : i32, %lstrd1 : i32, %lstrd2 : i32, %lstrd3 : i32, %ustrd : i16, %ch : i64) {
// expected-error @below {{Expected lower_stride and upper_stride to be either both present or both absent}}
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0], reduction = add : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
llvm.return
}
@@ -28,7 +28,6 @@ llvm.func @tma_store_reduce_tile_override_addr_stride_mismatch(%src : !llvm.ptr<
llvm.func @tma_store_reduce_tile_override_addr_im2col_dim_stride(%src : !llvm.ptr<3>, %tma_desc : !llvm.ptr, %override_addr : !llvm.ptr<1>, %d0 : i32, %d1 : i32, %d2 : i32, %d3 : i32, %d4 : i32, %ts0 : i16, %ts1 : i16, %ts2 : i16, %ts3 : i16, %ts4 : i16, %lstrd0 : i32, %lstrd1 : i32, %lstrd2 : i32, %lstrd3 : i32, %ustrd : i16, %ch : i64) {
// expected-error @below {{Only tile mode supports override address with dim and stride}}
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] tensor_size[%ts0, %ts1, %ts2] {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<im2col>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] tensor_size[%ts0, %ts1, %ts2], reduction = add mode = im2col : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
llvm.return
}
-
diff --git a/mlir/test/Target/LLVMIR/nvvm/nvvmir-invalid/tma_store_override_invalid.mlir b/mlir/test/Target/LLVMIR/nvvm/nvvmir-invalid/tma_store_override_invalid.mlir
index 4a16af610da82..7d4c2a0a50cb4 100644
--- a/mlir/test/Target/LLVMIR/nvvm/nvvmir-invalid/tma_store_override_invalid.mlir
+++ b/mlir/test/Target/LLVMIR/nvvm/nvvmir-invalid/tma_store_override_invalid.mlir
@@ -4,7 +4,7 @@
llvm.func @tma_shared_to_global_tile_override_addr_coord_tensor_size_mismatch(%src : !llvm.ptr<3>, %tma_desc : !llvm.ptr, %override_addr : !llvm.ptr<1>, %d0 : i32, %d1 : i32, %d2 : i32, %d3 : i32, %d4 : i32, %ts0 : i16, %ts1 : i16, %ts2 : i16, %ts3 : i16, %ts4 : i16, %lstrd0 : i32, %lstrd1 : i32, %lstrd2 : i32, %lstrd3 : i32, %ustrd : i16, %ch : i64) {
// expected-error @below {{Expected coordinates size to be equal to tensor size}}
- nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0] {mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0] : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
llvm.return
}
@@ -12,7 +12,7 @@ llvm.func @tma_shared_to_global_tile_override_addr_coord_tensor_size_mismatch(%s
llvm.func @tma_shared_to_global_tile_override_addr_lower_stride_tensor_size_mismatch(%src : !llvm.ptr<3>, %tma_desc : !llvm.ptr, %override_addr : !llvm.ptr<1>, %d0 : i32, %d1 : i32, %d2 : i32, %d3 : i32, %d4 : i32, %ts0 : i16, %ts1 : i16, %ts2 : i16, %ts3 : i16, %ts4 : i16, %lstrd0 : i32, %lstrd1 : i32, %lstrd2 : i32, %lstrd3 : i32, %ustrd : i16, %ch : i64) {
// expected-error @below {{Expected lower_stride size to be equal to one less than tensor size}}
- nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0, %lstrd1] upper_stride[%ustrd] {mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0, %lstrd1] upper_stride[%ustrd] : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
llvm.return
}
@@ -20,7 +20,7 @@ llvm.func @tma_shared_to_global_tile_override_addr_lower_stride_tensor_size_mism
llvm.func @tma_shared_to_global_tile_override_addr_stride_mismatch(%src : !llvm.ptr<3>, %tma_desc : !llvm.ptr, %override_addr : !llvm.ptr<1>, %d0 : i32, %d1 : i32, %d2 : i32, %d3 : i32, %d4 : i32, %ts0 : i16, %ts1 : i16, %ts2 : i16, %ts3 : i16, %ts4 : i16, %lstrd0 : i32, %lstrd1 : i32, %lstrd2 : i32, %lstrd3 : i32, %ustrd : i16, %ch : i64) {
// expected-error @below {{Expected lower_stride and upper_stride to be either both present or both absent}}
- nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] {mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
llvm.return
}
@@ -28,7 +28,7 @@ llvm.func @tma_shared_to_global_tile_override_addr_stride_mismatch(%src : !llvm.
llvm.func @tma_shared_to_global_tile_override_addr_im2col_dim_stride(%src : !llvm.ptr<3>, %tma_desc : !llvm.ptr, %override_addr : !llvm.ptr<1>, %d0 : i32, %d1 : i32, %d2 : i32, %d3 : i32, %d4 : i32, %ts0 : i16, %ts1 : i16, %ts2 : i16, %ts3 : i16, %ts4 : i16, %lstrd0 : i32, %lstrd1 : i32, %lstrd2 : i32, %lstrd3 : i32, %ustrd : i16, %ch : i64) {
// expected-error @below {{Only tile mode supports override address with dim and stride}}
- nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] tensor_size[%ts0, %ts1, %ts2] {mode = #nvvm.tma_store_mode<im2col>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] tensor_size[%ts0, %ts1, %ts2] mode = im2col : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
llvm.return
}
@@ -36,7 +36,7 @@ llvm.func @tma_shared_to_global_tile_override_addr_im2col_dim_stride(%src : !llv
llvm.func @tma_shared_to_global_tile_override_addr_tile_scatter4_dim_stride(%src : !llvm.ptr<3>, %tma_desc : !llvm.ptr, %override_addr : !llvm.ptr<1>, %d0 : i32, %d1 : i32, %d2 : i32, %d3 : i32, %d4 : i32, %ts0 : i16, %ts1 : i16, %ts2 : i16, %ts3 : i16, %ts4 : i16, %lstrd0 : i32, %lstrd1 : i32, %lstrd2 : i32, %lstrd3 : i32, %ustrd : i16, %ch : i64) {
// expected-error @below {{Only tile mode supports override address with dim and stride}}
- nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3, %d4] tensor_size[%ts0, %ts1, %ts2, %ts3, %ts4] {mode = #nvvm.tma_store_mode<tile_scatter4>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3, %d4] tensor_size[%ts0, %ts1, %ts2, %ts3, %ts4] mode = tile_scatter4 : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
llvm.return
}
@@ -44,7 +44,6 @@ llvm.func @tma_shared_to_global_tile_override_addr_tile_scatter4_dim_stride(%src
llvm.func @tma_shared_to_global_tile_override_addr_tile_scatter4_coord_mismatch(%src : !llvm.ptr<3>, %tma_desc : !llvm.ptr, %override_addr : !llvm.ptr<1>, %d0 : i32, %d1 : i32, %d2 : i32, %d3 : i32, %d4 : i32, %ts0 : i16, %ts1 : i16, %ts2 : i16, %ts3 : i16, %ts4 : i16, %lstrd0 : i32, %lstrd1 : i32, %lstrd2 : i32, %lstrd3 : i32, %ustrd : i16, %ch : i64) {
// expected-error @below {{Mode tile scatter4 expects 5 coordinates}}
- nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3] {mode = #nvvm.tma_store_mode<tile_scatter4>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3] mode = tile_scatter4 : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
llvm.return
}
-
diff --git a/mlir/test/Target/LLVMIR/nvvm/tma_store_override.mlir b/mlir/test/Target/LLVMIR/nvvm/tma_store_override.mlir
index 5b57d375612b3..1378524343a45 100644
--- a/mlir/test/Target/LLVMIR/nvvm/tma_store_override.mlir
+++ b/mlir/test/Target/LLVMIR/nvvm/tma_store_override.mlir
@@ -16,26 +16,26 @@ llvm.func @tma_shared_cta_global_tile_override_addr(%src : !llvm.ptr<3>, %tma_de
// CHECK-NEXT: }
// without cache hint
- nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0] {mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0] : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1] {mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1] : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] {mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3] {mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3] : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3, %d4] {mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3, %d4] : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
// with cache hint
- nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0] l2_cache_hint = %ch {mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0] l2_cache_hint = %ch : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1] l2_cache_hint = %ch {mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1] l2_cache_hint = %ch : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch {mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3] l2_cache_hint = %ch {mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3] l2_cache_hint = %ch : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3, %d4] l2_cache_hint = %ch {mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3, %d4] l2_cache_hint = %ch : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
llvm.return
}
@@ -48,10 +48,10 @@ llvm.func @tma_shared_cta_global_tile_scatter4_override_addr(%src : !llvm.ptr<3>
// CHECK-NEXT: }
// without cache hint
- nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3, %d4] {mode = #nvvm.tma_store_mode<tile_scatter4>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3, %d4] mode = tile_scatter4 : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
// with cache hint
- nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3, %d4] l2_cache_hint = %ch {mode = #nvvm.tma_store_mode<tile_scatter4>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3, %d4] l2_cache_hint = %ch mode = tile_scatter4 : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
llvm.return
}
@@ -68,18 +68,18 @@ llvm.func @tma_shared_cta_global_im2col(%src : !llvm.ptr<3>, %tma_desc : !llvm.p
// CHECK-NEXT: }
// without cache hint
- nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] {mode = #nvvm.tma_store_mode<im2col>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] mode = im2col : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3] {mode = #nvvm.tma_store_mode<im2col>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3] mode = im2col : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3, %d4] {mode = #nvvm.tma_store_mode<im2col>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3, %d4] mode = im2col : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
// with cache hint
- nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch {mode = #nvvm.tma_store_mode<im2col>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch mode = im2col : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3] l2_cache_hint = %ch {mode = #nvvm.tma_store_mode<im2col>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3] l2_cache_hint = %ch mode = im2col : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3, %d4] l2_cache_hint = %ch {mode = #nvvm.tma_store_mode<im2col>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3, %d4] l2_cache_hint = %ch mode = im2col : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
llvm.return
}
@@ -96,18 +96,18 @@ llvm.func @tma_shared_cta_global_im2col_w(%src : !llvm.ptr<3>, %tma_desc : !llvm
// CHECK-NEXT: }
// without cache hint
- nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] {mode = #nvvm.tma_store_mode<im2col_w>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] mode = im2col_w : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3] {mode = #nvvm.tma_store_mode<im2col_w>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3] mode = im2col_w : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3, %d4] {mode = #nvvm.tma_store_mode<im2col_w>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3, %d4] mode = im2col_w : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
// with cache hint
- nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch {mode = #nvvm.tma_store_mode<im2col_w>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch mode = im2col_w : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3] l2_cache_hint = %ch {mode = #nvvm.tma_store_mode<im2col_w>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3] l2_cache_hint = %ch mode = im2col_w : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3, %d4] l2_cache_hint = %ch {mode = #nvvm.tma_store_mode<im2col_w>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3, %d4] l2_cache_hint = %ch mode = im2col_w : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
llvm.return
}
@@ -128,27 +128,26 @@ llvm.func @tma_shared_cta_global_tile_override_addr_dim_stride(%src : !llvm.ptr<
// CHECK-NEXT: }
// without cache hint
- nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0] tensor_size[%ts0] {mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0] tensor_size[%ts0] : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd] {mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd] : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] tensor_size[%ts0, %ts1, %ts2] lower_stride[%lstrd0, %lstrd1] upper_stride[%ustrd] {mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] tensor_size[%ts0, %ts1, %ts2] lower_stride[%lstrd0, %lstrd1] upper_stride[%ustrd] : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3] tensor_size[%ts0, %ts1, %ts2, %ts3] lower_stride[%lstrd0, %lstrd1, %lstrd2] upper_stride[%ustrd] {mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3] tensor_size[%ts0, %ts1, %ts2, %ts3] lower_stride[%lstrd0, %lstrd1, %lstrd2] upper_stride[%ustrd] : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3, %d4] tensor_size[%ts0, %ts1, %ts2, %ts3, %ts4] lower_stride[%lstrd0, %lstrd1, %lstrd2, %lstrd3] upper_stride[%ustrd] {mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3, %d4] tensor_size[%ts0, %ts1, %ts2, %ts3, %ts4] lower_stride[%lstrd0, %lstrd1, %lstrd2, %lstrd3] upper_stride[%ustrd] : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
// with cache hint
- nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0] tensor_size[%ts0] l2_cache_hint = %ch {mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0] tensor_size[%ts0] l2_cache_hint = %ch : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd] l2_cache_hint = %ch {mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd] l2_cache_hint = %ch : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] tensor_size[%ts0, %ts1, %ts2] lower_stride[%lstrd0, %lstrd1] upper_stride[%ustrd] l2_cache_hint = %ch {mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] tensor_size[%ts0, %ts1, %ts2] lower_stride[%lstrd0, %lstrd1] upper_stride[%ustrd] l2_cache_hint = %ch : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3] tensor_size[%ts0, %ts1, %ts2, %ts3] lower_stride[%lstrd0, %lstrd1, %lstrd2] upper_stride[%ustrd] l2_cache_hint = %ch {mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3] tensor_size[%ts0, %ts1, %ts2, %ts3] lower_stride[%lstrd0, %lstrd1, %lstrd2] upper_stride[%ustrd] l2_cache_hint = %ch : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3, %d4] tensor_size[%ts0, %ts1, %ts2, %ts3, %ts4] lower_stride[%lstrd0, %lstrd1, %lstrd2, %lstrd3] upper_stride[%ustrd] l2_cache_hint = %ch {mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.global.shared.cta.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3, %d4] tensor_size[%ts0, %ts1, %ts2, %ts3, %ts4] lower_stride[%lstrd0, %lstrd1, %lstrd2, %lstrd3] upper_stride[%ustrd] l2_cache_hint = %ch : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
llvm.return
}
-
diff --git a/mlir/test/Target/LLVMIR/nvvm/tma_store_reduce_override.mlir b/mlir/test/Target/LLVMIR/nvvm/tma_store_reduce_override.mlir
index 1236155f9df29..d5a18bf588ccc 100644
--- a/mlir/test/Target/LLVMIR/nvvm/tma_store_reduce_override.mlir
+++ b/mlir/test/Target/LLVMIR/nvvm/tma_store_reduce_override.mlir
@@ -29,61 +29,61 @@ llvm.func @tma_store_reduce_tile_override_addr(%src : !llvm.ptr<3>, %tma_desc :
// CHECK-NEXT: ret void
// CHECK-NEXT: }
// without cache hint
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0], reduction = add : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1], reduction = add : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2], reduction = add : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3] {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3], reduction = add : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3, %d4] {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3, %d4], reduction = add : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
// with cache hint
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] l2_cache_hint = %ch, reduction = add : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] l2_cache_hint = %ch, reduction = add : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch, reduction = add : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3] l2_cache_hint = %ch, reduction = add : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3, %d4] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3, %d4] l2_cache_hint = %ch, reduction = add : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
// Test min reduction
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] {redKind = #nvvm.tma_redux_kind<min>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0], reduction = min : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<min>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] l2_cache_hint = %ch, reduction = min : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
// Test max reduction
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] {redKind = #nvvm.tma_redux_kind<max>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0], reduction = max : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<max>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] l2_cache_hint = %ch, reduction = max : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
// Test inc reduction
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] {redKind = #nvvm.tma_redux_kind<inc>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0], reduction = inc : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<inc>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] l2_cache_hint = %ch, reduction = inc : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
// Test dec reduction
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] {redKind = #nvvm.tma_redux_kind<dec>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0], reduction = dec : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<dec>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] l2_cache_hint = %ch, reduction = dec : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
// Test and reduction
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] {redKind = #nvvm.tma_redux_kind<and>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0], reduction = and : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<and>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] l2_cache_hint = %ch, reduction = and : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
// Test or reduction
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] {redKind = #nvvm.tma_redux_kind<or>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0], reduction = or : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<or>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] l2_cache_hint = %ch, reduction = or : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
// Test xor reduction
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] {redKind = #nvvm.tma_redux_kind<xor>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0], reduction = xor : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<xor>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] l2_cache_hint = %ch, reduction = xor : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
llvm.return
}
@@ -113,53 +113,53 @@ llvm.func @tma_store_reduce_im2col_override_addr(%src : !llvm.ptr<3>, %tma_desc
// CHECK-NEXT: ret void
// CHECK-NEXT: }
// without cache hint
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<im2col>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2], reduction = add mode = im2col : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3] {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<im2col>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3], reduction = add mode = im2col : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3, %d4] {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<im2col>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3, %d4], reduction = add mode = im2col : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
// with cache hint
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<im2col>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch, reduction = add mode = im2col : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<im2col>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3] l2_cache_hint = %ch, reduction = add mode = im2col : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3, %d4] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<im2col>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3, %d4] l2_cache_hint = %ch, reduction = add mode = im2col : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
// Test min reduction
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] {redKind = #nvvm.tma_redux_kind<min>, mode = #nvvm.tma_store_mode<im2col>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2], reduction = min mode = im2col : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<min>, mode = #nvvm.tma_store_mode<im2col>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch, reduction = min mode = im2col : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
// Test max reduction
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] {redKind = #nvvm.tma_redux_kind<max>, mode = #nvvm.tma_store_mode<im2col>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2], reduction = max mode = im2col : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<max>, mode = #nvvm.tma_store_mode<im2col>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch, reduction = max mode = im2col : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
// Test inc reduction
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] {redKind = #nvvm.tma_redux_kind<inc>, mode = #nvvm.tma_store_mode<im2col>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2], reduction = inc mode = im2col : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<inc>, mode = #nvvm.tma_store_mode<im2col>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch, reduction = inc mode = im2col : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
// Test dec reduction
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] {redKind = #nvvm.tma_redux_kind<dec>, mode = #nvvm.tma_store_mode<im2col>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2], reduction = dec mode = im2col : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<dec>, mode = #nvvm.tma_store_mode<im2col>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch, reduction = dec mode = im2col : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
// Test and reduction
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] {redKind = #nvvm.tma_redux_kind<and>, mode = #nvvm.tma_store_mode<im2col>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2], reduction = and mode = im2col : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<and>, mode = #nvvm.tma_store_mode<im2col>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch, reduction = and mode = im2col : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
// Test or reduction
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] {redKind = #nvvm.tma_redux_kind<or>, mode = #nvvm.tma_store_mode<im2col>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2], reduction = or mode = im2col : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<or>, mode = #nvvm.tma_store_mode<im2col>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch, reduction = or mode = im2col : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
// Test xor reduction
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] {redKind = #nvvm.tma_redux_kind<xor>, mode = #nvvm.tma_store_mode<im2col>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2], reduction = xor mode = im2col : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<xor>, mode = #nvvm.tma_store_mode<im2col>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch, reduction = xor mode = im2col : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
llvm.return
}
@@ -189,53 +189,53 @@ llvm.func @tma_store_reduce_im2col_w_override_addr(%src : !llvm.ptr<3>, %tma_des
// CHECK-NEXT: ret void
// CHECK-NEXT: }
// without cache hint
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<im2col_w>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2], reduction = add mode = im2col_w : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3] {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<im2col_w>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3], reduction = add mode = im2col_w : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3, %d4] {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<im2col_w>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3, %d4], reduction = add mode = im2col_w : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
// with cache hint
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<im2col_w>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch, reduction = add mode = im2col_w : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<im2col_w>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3] l2_cache_hint = %ch, reduction = add mode = im2col_w : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3, %d4] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<im2col_w>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3, %d4] l2_cache_hint = %ch, reduction = add mode = im2col_w : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
// Test min reduction
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] {redKind = #nvvm.tma_redux_kind<min>, mode = #nvvm.tma_store_mode<im2col_w>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2], reduction = min mode = im2col_w : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<min>, mode = #nvvm.tma_store_mode<im2col_w>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch, reduction = min mode = im2col_w : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
// Test max reduction
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] {redKind = #nvvm.tma_redux_kind<max>, mode = #nvvm.tma_store_mode<im2col_w>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2], reduction = max mode = im2col_w : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<max>, mode = #nvvm.tma_store_mode<im2col_w>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch, reduction = max mode = im2col_w : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
// Test inc reduction
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] {redKind = #nvvm.tma_redux_kind<inc>, mode = #nvvm.tma_store_mode<im2col_w>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2], reduction = inc mode = im2col_w : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<inc>, mode = #nvvm.tma_store_mode<im2col_w>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch, reduction = inc mode = im2col_w : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
// Test dec reduction
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] {redKind = #nvvm.tma_redux_kind<dec>, mode = #nvvm.tma_store_mode<im2col_w>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2], reduction = dec mode = im2col_w : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<dec>, mode = #nvvm.tma_store_mode<im2col_w>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch, reduction = dec mode = im2col_w : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
// Test and reduction
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] {redKind = #nvvm.tma_redux_kind<and>, mode = #nvvm.tma_store_mode<im2col_w>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2], reduction = and mode = im2col_w : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<and>, mode = #nvvm.tma_store_mode<im2col_w>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch, reduction = and mode = im2col_w : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
// Test or reduction
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] {redKind = #nvvm.tma_redux_kind<or>, mode = #nvvm.tma_store_mode<im2col_w>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2], reduction = or mode = im2col_w : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<or>, mode = #nvvm.tma_store_mode<im2col_w>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch, reduction = or mode = im2col_w : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
// Test xor reduction
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] {redKind = #nvvm.tma_redux_kind<xor>, mode = #nvvm.tma_store_mode<im2col_w>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2], reduction = xor mode = im2col_w : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<xor>, mode = #nvvm.tma_store_mode<im2col_w>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] l2_cache_hint = %ch, reduction = xor mode = im2col_w : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
llvm.return
}
@@ -283,90 +283,89 @@ llvm.func @tma_store_reduce_tile_override_addr_dim_stride(%src : !llvm.ptr<3>, %
// CHECK-NEXT: ret void
// CHECK-NEXT: }
// without cache hint
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] tensor_size[%ts0] {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] tensor_size[%ts0], reduction = add : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd] {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd], reduction = add : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] tensor_size[%ts0, %ts1, %ts2] lower_stride[%lstrd0, %lstrd1] upper_stride[%ustrd] {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] tensor_size[%ts0, %ts1, %ts2] lower_stride[%lstrd0, %lstrd1] upper_stride[%ustrd], reduction = add : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3] tensor_size[%ts0, %ts1, %ts2, %ts3] lower_stride[%lstrd0, %lstrd1, %lstrd2] upper_stride[%ustrd] {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3] tensor_size[%ts0, %ts1, %ts2, %ts3] lower_stride[%lstrd0, %lstrd1, %lstrd2] upper_stride[%ustrd], reduction = add : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3, %d4] tensor_size[%ts0, %ts1, %ts2, %ts3, %ts4] lower_stride[%lstrd0, %lstrd1, %lstrd2, %lstrd3] upper_stride[%ustrd] {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3, %d4] tensor_size[%ts0, %ts1, %ts2, %ts3, %ts4] lower_stride[%lstrd0, %lstrd1, %lstrd2, %lstrd3] upper_stride[%ustrd], reduction = add : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
// with cache hint
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] tensor_size[%ts0] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] tensor_size[%ts0] l2_cache_hint = %ch, reduction = add : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd] l2_cache_hint = %ch, reduction = add : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] tensor_size[%ts0, %ts1, %ts2] lower_stride[%lstrd0, %lstrd1] upper_stride[%ustrd] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2] tensor_size[%ts0, %ts1, %ts2] lower_stride[%lstrd0, %lstrd1] upper_stride[%ustrd] l2_cache_hint = %ch, reduction = add : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3] tensor_size[%ts0, %ts1, %ts2, %ts3] lower_stride[%lstrd0, %lstrd1, %lstrd2] upper_stride[%ustrd] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3] tensor_size[%ts0, %ts1, %ts2, %ts3] lower_stride[%lstrd0, %lstrd1, %lstrd2] upper_stride[%ustrd] l2_cache_hint = %ch, reduction = add : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3, %d4] tensor_size[%ts0, %ts1, %ts2, %ts3, %ts4] lower_stride[%lstrd0, %lstrd1, %lstrd2, %lstrd3] upper_stride[%ustrd] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1, %d2, %d3, %d4] tensor_size[%ts0, %ts1, %ts2, %ts3, %ts4] lower_stride[%lstrd0, %lstrd1, %lstrd2, %lstrd3] upper_stride[%ustrd] l2_cache_hint = %ch, reduction = add : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
// Test min reduction
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] tensor_size[%ts0] {redKind = #nvvm.tma_redux_kind<min>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] tensor_size[%ts0], reduction = min : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] tensor_size[%ts0] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<min>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] tensor_size[%ts0] l2_cache_hint = %ch, reduction = min : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd] {redKind = #nvvm.tma_redux_kind<min>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd], reduction = min : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<min>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd] l2_cache_hint = %ch, reduction = min : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
// Test max reduction
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] tensor_size[%ts0] {redKind = #nvvm.tma_redux_kind<max>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] tensor_size[%ts0], reduction = max : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] tensor_size[%ts0] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<max>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] tensor_size[%ts0] l2_cache_hint = %ch, reduction = max : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd] {redKind = #nvvm.tma_redux_kind<max>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd], reduction = max : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<max>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd] l2_cache_hint = %ch, reduction = max : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
// Test inc reduction
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] tensor_size[%ts0] {redKind = #nvvm.tma_redux_kind<inc>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] tensor_size[%ts0], reduction = inc : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] tensor_size[%ts0] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<inc>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] tensor_size[%ts0] l2_cache_hint = %ch, reduction = inc : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd] {redKind = #nvvm.tma_redux_kind<inc>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd], reduction = inc : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<inc>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd] l2_cache_hint = %ch, reduction = inc : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
// Test dec reduction
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] tensor_size[%ts0] {redKind = #nvvm.tma_redux_kind<dec>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] tensor_size[%ts0], reduction = dec : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] tensor_size[%ts0] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<dec>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] tensor_size[%ts0] l2_cache_hint = %ch, reduction = dec : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd] {redKind = #nvvm.tma_redux_kind<dec>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd], reduction = dec : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<dec>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd] l2_cache_hint = %ch, reduction = dec : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
// Test and reduction
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] tensor_size[%ts0] {redKind = #nvvm.tma_redux_kind<and>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] tensor_size[%ts0], reduction = and : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] tensor_size[%ts0] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<and>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] tensor_size[%ts0] l2_cache_hint = %ch, reduction = and : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd] {redKind = #nvvm.tma_redux_kind<and>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd], reduction = and : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<and>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd] l2_cache_hint = %ch, reduction = and : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
// Test or reduction
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] tensor_size[%ts0] {redKind = #nvvm.tma_redux_kind<or>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] tensor_size[%ts0], reduction = or : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] tensor_size[%ts0] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<or>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] tensor_size[%ts0] l2_cache_hint = %ch, reduction = or : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd] {redKind = #nvvm.tma_redux_kind<or>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd], reduction = or : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<or>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd] l2_cache_hint = %ch, reduction = or : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
// Test xor reduction
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] tensor_size[%ts0] {redKind = #nvvm.tma_redux_kind<xor>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] tensor_size[%ts0], reduction = xor : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] tensor_size[%ts0] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<xor>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0] tensor_size[%ts0] l2_cache_hint = %ch, reduction = xor : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd] {redKind = #nvvm.tma_redux_kind<xor>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd], reduction = xor : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
- nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd] l2_cache_hint = %ch {redKind = #nvvm.tma_redux_kind<xor>, mode = #nvvm.tma_store_mode<tile>} : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
+ nvvm.cp.async.bulk.tensor.reduce.override %tma_desc, %src, %override_addr, box[%d0, %d1] tensor_size[%ts0, %ts1] lower_stride[%lstrd0] upper_stride[%ustrd] l2_cache_hint = %ch, reduction = xor : !llvm.ptr, !llvm.ptr<3>, !llvm.ptr<1>
llvm.return
}
-
More information about the Mlir-commits
mailing list