[Mlir-commits] [llvm] [mlir] [NVVM][NVPTX] Change S2G reduction ops to use flag for reduction ops (PR #213638)
llvmlistbot at llvm.org
llvmlistbot at llvm.org
Mon Aug 3 03:04:19 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-backend-nvptx
Author: Rajat Bajpai (rajatbajpai)
<details>
<summary>Changes</summary>
Currently, TMA S2G reduction intrinsics use reduction operation in the name. This leads to large number of intrinsics without any difference between them besides the operation. This change move reduction operation from name into an immediate flag. This reduces number of intrinsics from 64 to 8.
---
Patch is 224.02 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/213638.diff
12 Files Affected:
- (modified) llvm/docs/NVPTXUsage.md (+69-61)
- (modified) llvm/include/llvm/IR/IntrinsicsNVVM.td (+16-10)
- (modified) llvm/include/llvm/IR/NVVMIntrinsicUtils.h (+25)
- (modified) llvm/lib/IR/AutoUpgrade.cpp (+63)
- (modified) llvm/lib/IR/NVVMIntrinsicUtils.cpp (+9)
- (modified) llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.cpp (+3-30)
- (modified) llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp (-185)
- (modified) llvm/lib/Target/NVPTX/NVPTXIntrinsics.td (+40-19)
- (modified) llvm/test/Assembler/auto_upgrade_nvvm_intrinsics.ll (+56)
- (modified) llvm/test/CodeGen/NVPTX/cp-async-bulk-tensor-reduce.ll (+517-372)
- (modified) mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp (+20-122)
- (modified) mlir/test/Target/LLVMIR/nvvm/tma_store_reduce.mlir (+128-128)
``````````diff
diff --git a/llvm/docs/NVPTXUsage.md b/llvm/docs/NVPTXUsage.md
index 810938303c591..58753e76a40c3 100644
--- a/llvm/docs/NVPTXUsage.md
+++ b/llvm/docs/NVPTXUsage.md
@@ -2062,6 +2062,75 @@ described in the `s2g.tile` mode intrinsics above.
For more information, refer [PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-bulk-tensor).
+#### '`llvm.nvvm.cp.async.bulk.tensor.reduce.tile.[1-5]d`'
+
+##### Syntax:
+
+```llvm
+declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.1d(ptr addrspace(3) %src, ptr %tensor_map, i32 %d0, i64 %ch, i32 %red_op, i1 %flag_ch)
+declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.2d(..., i32 %d0, i32 %d1, ...)
+declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.3d(..., i32 %d0, i32 %d1, i32 %d2, ...)
+declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.4d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, ...)
+declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.tile.5d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, i32 %d4, ...)
+```
+
+##### Overview:
+
+The '`@llvm.nvvm.cp.async.bulk.tensor.reduce.tile.[1-5]d`' intrinsics
+correspond to the `cp.reduce.async.bulk.tensor.[1-5]d.global.shared::cta.*`
+set of PTX instructions. These instructions initiate an asynchronous reduction
+operation of tensor data in global memory with the tensor data in shared::cta
+memory, using `tile` mode. The dimension of the tensor data ranges from 1d to
+5d with the coordinates specified by the `i32 %d0 ... i32 %d4` arguments. The
+`i32 %red_op` argument selects the reduction operation to perform. It must be
+a compile-time constant in the half-open range `[0, 8)`, with the following
+encoding:
+
+| `red_op` | Reduction |
+|---------:|-----------|
+| 0 | add |
+| 1 | min |
+| 2 | max |
+| 3 | inc |
+| 4 | dec |
+| 5 | and |
+| 6 | or |
+| 7 | xor |
+
+The symbolic LLVM IR annotation for `red_op` and the PTX reduction suffix use
+the same canonical operator spelling.
+
+- The last argument to these intrinsics is a boolean flag indicating support
+ for cache_hint. This flag argument must be a compile-time constant. When
+ set, it indicates a valid cache_hint (`i64 %ch`) and generates the
+ `.L2::cache_hint` variant of the PTX instruction.
+
+For more information, refer [PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-reduce-async-bulk-tensor).
+
+#### '`llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.[3-5]d`'
+
+##### Syntax:
+
+```llvm
+declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.3d(ptr addrspace(3) %src, ptr %tensor_map, i32 %d0, i32 %d1, i32 %d2, i64 %ch, i32 %red_op, i1 %flag_ch)
+declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.4d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, ...)
+declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.5d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, i32 %d4, ...)
+```
+
+##### Overview:
+
+The '`@llvm.nvvm.cp.async.bulk.tensor.reduce.im2col.[3-5]d`' intrinsics
+correspond to the `cp.reduce.async.bulk.tensor.[3-5]d.global.shared::cta.*`
+set of PTX instructions. These instructions initiate an asynchronous reduction
+operation of tensor data in global memory with the tensor data in shared::cta
+memory, using `im2col` mode. In this mode, the tensor has to be at least
+three-dimensional. The supported reduction operations are the same as the ones
+in the `tile` mode. The `i32 %red_op` argument and the last boolean flag
+argument have the same functionality as described in the `tile` mode
+intrinsics above.
+
+For more information, refer [PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-reduce-async-bulk-tensor).
+
#### '`llvm.nvvm.cp.async.bulk.tensor.prefetch.tile.[1-5]d`'
##### Syntax:
@@ -2137,67 +2206,6 @@ functionality as described in the `tile` mode intrinsics above.
For more information, refer [PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-cp-async-bulk-prefetch-tensor).
-#### '`llvm.nvvm.cp.async.bulk.tensor.reduce.[red_op].tile.[1-5]d`'
-
-##### Syntax:
-
-```llvm
-declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.add.tile.1d(ptr addrspace(3) %src, ptr %tensor_map, i32 %d0, i64 %ch, i1 %flag_ch)
-declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.min.tile.1d(ptr addrspace(3) %src, ptr %tensor_map, i32 %d0, i64 %ch, i1 %flag_ch)
-declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.max.tile.1d(ptr addrspace(3) %src, ptr %tensor_map, i32 %d0, i64 %ch, i1 %flag_ch)
-declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.inc.tile.1d(ptr addrspace(3) %src, ptr %tensor_map, i32 %d0, i64 %ch, i1 %flag_ch)
-declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.dec.tile.1d(ptr addrspace(3) %src, ptr %tensor_map, i32 %d0, i64 %ch, i1 %flag_ch)
-declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.and.tile.1d(ptr addrspace(3) %src, ptr %tensor_map, i32 %d0, i64 %ch, i1 %flag_ch)
-declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.or.tile.1d(ptr addrspace(3) %src, ptr %tensor_map, i32 %d0, i64 %ch, i1 %flag_ch)
-declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.xor.tile.1d(ptr addrspace(3) %src, ptr %tensor_map, i32 %d0, i64 %ch, i1 %flag_ch)
-
-declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.<red_op>.tile.2d(..., i32 %d0, i32 %d1, ...)
-declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.<red_op>.tile.3d(..., i32 %d0, i32 %d1, i32 %d2, ...)
-declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.<red_op>.tile.4d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, ...)
-declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.<red_op>.tile.5d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, i32 %d4, ...)
-```
-
-##### Overview:
-
-The '`@llvm.nvvm.cp.async.bulk.tensor.reduce.<red_op>.tile.[1-5]d`' intrinsics
-correspond to the `cp.reduce.async.bulk.tensor.[1-5]d.*` set of PTX
-instructions. These instructions initiate an asynchronous reduction operation of
-tensor data in global memory with the tensor data in shared\{::cta} memory, using
-`tile` mode. The dimension of the tensor data ranges from 1d to 5d with the
-coordinates specified by the `i32 %d0 ... i32 %d4` arguments. The supported
-reduction operations are {add, min, max, inc, dec, and, or, xor} as described in
-the `tile.1d` intrinsics.
-
-- The last argument to these intrinsics is a boolean flag indicating support for
- cache_hint. This flag argument must be a compile-time constant. When set, it
- indicates a valid cache_hint (`i64 %ch`) and generates the
- `.L2::cache_hint` variant of the PTX instruction.
-
-For more information, refer [PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-reduce-async-bulk-tensor).
-
-#### '`llvm.nvvm.cp.async.bulk.tensor.reduce.[red_op].im2col.[3-5]d`'
-
-##### Syntax:
-
-```llvm
-declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.<red_op>.im2col.3d(ptr addrspace(3) %src, ptr %tensor_map, i32 %d0, i32 %d1, i32 %d2, i64 %ch, i1 %flag_ch)
-declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.<red_op>.im2col.4d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, ...)
-declare void @llvm.nvvm.cp.async.bulk.tensor.reduce.<red_op>.im2col.5d(..., i32 %d0, i32 %d1, i32 %d2, i32 %d3, i32 %d4, ...)
-```
-
-##### Overview:
-
-The '`@llvm.nvvm.cp.async.bulk.tensor.reduce.<red_op>.im2col.[3-5]d`'
-intrinsics correspond to the `cp.reduce.async.bulk.tensor.[3-5]d.*` set of PTX
-instructions. These instructions initiate an asynchronous reduction operation of
-tensor data in global memory with the tensor data in shared\{::cta} memory, using
-`im2col` mode. In this mode, the tensor has to be at least three-dimensional.
-The supported reduction operations supported are the same as the ones in the
-tile mode. The last argument to these intrinsics is a boolean flag, with the
-same functionality as described in the `tile` mode intrinsics above.
-
-For more information, refer [PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-reduce-async-bulk-tensor).
-
### Warp Group Intrinsics
#### '`llvm.nvvm.wgmma.fence.sync.aligned`'
diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index 40f3c40a76bc6..6af8e6563ac09 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -2833,16 +2833,22 @@ foreach dim = 1...5 in {
[llvm_i1_ty], // Flag for cache_hint
[IntrConvergent, ReadOnly<ArgIndex<0>>, ReadOnly<ArgIndex<1>>]>;
- // Intrinsics for TMA Copy with reduction
- foreach red_op = ["add", "min", "max", "inc", "dec", "and", "or", "xor"] in
- def int_nvvm_cp_async_bulk_tensor_reduce_ # red_op # _ # mode # _ # dim # d :
- DefaultAttrsIntrinsicFlags<[],
- !listconcat([llvm_shared_ptr_ty, // src_smem_ptr
- llvm_ptr_ty], // tensormap_ptr
- tensor_dim_args, // actual tensor dims
- [llvm_i64_ty]), // cache_hint
- [llvm_i1_ty], // Flag for cache_hint
- [IntrConvergent, ReadOnly<ArgIndex<0>>, ReadOnly<ArgIndex<1>>]>;
+ defvar reduce_params =
+ !listconcat([llvm_shared_ptr_ty, // src_smem_ptr
+ llvm_ptr_ty], // tensormap_ptr
+ tensor_dim_args, // actual tensor dims
+ [llvm_i64_ty]); // cache_hint
+ defvar red_op_idx = !size(reduce_params);
+ def int_nvvm_cp_async_bulk_tensor_reduce_ # mode # _ # dim # d :
+ DefaultAttrsIntrinsicFlags<[],
+ reduce_params,
+ [llvm_i32_ty, // reduction operation
+ llvm_i1_ty], // Flag for cache_hint
+ [IntrConvergent, ReadOnly<ArgIndex<0>>, ReadOnly<ArgIndex<1>>,
+ // Allowed values for red_op are {0..7} i.e. [0, 8).
+ Range<ArgIndex<red_op_idx>, 0, 8>,
+ ArgInfo<ArgIndex<red_op_idx>, [ArgName<"red_op">,
+ ImmArgPrinter<"printTMAReductionOp">]>]>;
}
}
diff --git a/llvm/include/llvm/IR/NVVMIntrinsicUtils.h b/llvm/include/llvm/IR/NVVMIntrinsicUtils.h
index 083be6d4247c1..a19a93aae08a6 100644
--- a/llvm/include/llvm/IR/NVVMIntrinsicUtils.h
+++ b/llvm/include/llvm/IR/NVVMIntrinsicUtils.h
@@ -19,6 +19,7 @@
#include "llvm/ADT/APFloat.h"
#include "llvm/ADT/APInt.h"
+#include "llvm/ADT/StringRef.h"
#include "llvm/IR/Constants.h"
#include "llvm/IR/Intrinsics.h"
#include "llvm/IR/IntrinsicsNVPTX.h"
@@ -41,6 +42,28 @@ enum class TMAReductionOp : uint8_t {
XOR = 7,
};
+inline StringRef getTMAReductionOpName(TMAReductionOp Op) {
+ switch (Op) {
+ case TMAReductionOp::ADD:
+ return "add";
+ case TMAReductionOp::MIN:
+ return "min";
+ case TMAReductionOp::MAX:
+ return "max";
+ case TMAReductionOp::INC:
+ return "inc";
+ case TMAReductionOp::DEC:
+ return "dec";
+ case TMAReductionOp::AND:
+ return "and";
+ case TMAReductionOp::OR:
+ return "or";
+ case TMAReductionOp::XOR:
+ return "xor";
+ }
+ llvm_unreachable("invalid TMA reduction operation");
+}
+
// Enum to represent the cta_group::1 and
// cta_group::2 variants in TMA/TCGEN05 family of
// PTX instructions.
@@ -106,6 +129,8 @@ enum class TensormapFillMode : uint8_t {
LLVM_ABI void printTcgen05MMAKind(raw_ostream &OS, const Constant *ImmArgVal);
+LLVM_ABI void printTMAReductionOp(raw_ostream &OS, const Constant *ImmArgVal);
+
LLVM_ABI void printTcgen05CollectorUsageOp(raw_ostream &OS,
const Constant *ImmArgVal);
diff --git a/llvm/lib/IR/AutoUpgrade.cpp b/llvm/lib/IR/AutoUpgrade.cpp
index f2a2ddbfc2e31..d2216e4072337 100644
--- a/llvm/lib/IR/AutoUpgrade.cpp
+++ b/llvm/lib/IR/AutoUpgrade.cpp
@@ -43,6 +43,7 @@
#include "llvm/IR/MDBuilder.h"
#include "llvm/IR/Metadata.h"
#include "llvm/IR/Module.h"
+#include "llvm/IR/NVVMIntrinsicUtils.h"
#include "llvm/IR/Value.h"
#include "llvm/IR/Verifier.h"
#include "llvm/Support/AMDGPUAddrSpace.h"
@@ -1192,6 +1193,42 @@ static Intrinsic::ID shouldUpgradeNVPTXTMAG2SIntrinsics(Function *F,
return Intrinsic::not_intrinsic;
}
+// The legacy TMA reduction intrinsics encode the reduction operator in their
+// name, while the current ones take it as an immediate argument. Map the
+// operator part of a legacy name to the corresponding immediate value.
+static std::optional<unsigned> getNVPTXTMAReductionOp(StringRef Name) {
+ return StringSwitch<std::optional<unsigned>>(Name)
+ .Case("add", static_cast<unsigned>(nvvm::TMAReductionOp::ADD))
+ .Case("min", static_cast<unsigned>(nvvm::TMAReductionOp::MIN))
+ .Case("max", static_cast<unsigned>(nvvm::TMAReductionOp::MAX))
+ .Case("inc", static_cast<unsigned>(nvvm::TMAReductionOp::INC))
+ .Case("dec", static_cast<unsigned>(nvvm::TMAReductionOp::DEC))
+ .Case("and", static_cast<unsigned>(nvvm::TMAReductionOp::AND))
+ .Case("or", static_cast<unsigned>(nvvm::TMAReductionOp::OR))
+ .Case("xor", static_cast<unsigned>(nvvm::TMAReductionOp::XOR))
+ .Default(std::nullopt);
+}
+
+static Intrinsic::ID shouldUpgradeNVPTXTMAReductionIntrinsics(StringRef Name) {
+ if (!Name.consume_front("cp.async.bulk.tensor.reduce."))
+ return Intrinsic::not_intrinsic;
+
+ auto [RedOpName, ShapeName] = Name.split('.');
+ if (!getNVPTXTMAReductionOp(RedOpName))
+ return Intrinsic::not_intrinsic;
+
+ return StringSwitch<Intrinsic::ID>(ShapeName)
+ .Case("tile.1d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_1d)
+ .Case("tile.2d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_2d)
+ .Case("tile.3d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_3d)
+ .Case("tile.4d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_4d)
+ .Case("tile.5d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_5d)
+ .Case("im2col.3d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_3d)
+ .Case("im2col.4d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_4d)
+ .Case("im2col.5d", Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_5d)
+ .Default(Intrinsic::not_intrinsic);
+}
+
static Intrinsic::ID shouldUpgradeNVPTXSharedClusterIntrinsic(Function *F,
StringRef Name) {
if (Name.consume_front("mapa.shared.cluster"))
@@ -1718,6 +1755,15 @@ static bool upgradeIntrinsicFunction1(Function *F, Function *&NewFn,
return true;
}
+ // Upgrade TMA reduction intrinsics
+ // llvm.nvvm.cp.async.bulk.tensor.reduce.<red_op>* =>
+ // llvm.nvvm.cp.async.bulk.tensor.reduce.<shape>*
+ IID = shouldUpgradeNVPTXTMAReductionIntrinsics(Name);
+ if (IID != Intrinsic::not_intrinsic) {
+ NewFn = Intrinsic::getOrInsertDeclaration(F->getParent(), IID);
+ return true;
+ }
+
// Upgrade TMA copy G2S Intrinsics
IID = shouldUpgradeNVPTXTMAG2SIntrinsics(F, Name);
if (IID != Intrinsic::not_intrinsic) {
@@ -5617,6 +5663,23 @@ void llvm::UpgradeIntrinsicCall(CallBase *CI, Function *NewFn) {
CI->eraseFromParent();
return;
}
+ case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_1d:
+ case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_2d:
+ case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_3d:
+ case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_4d:
+ case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_tile_5d:
+ case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_3d:
+ case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_4d:
+ case Intrinsic::nvvm_cp_async_bulk_tensor_reduce_im2col_5d: {
+ StringRef Name = F->getName();
+ Name.consume_front("llvm.nvvm.cp.async.bulk.tensor.reduce.");
+ auto RedOp = getNVPTXTMAReductionOp(Name.split('.').first);
+
+ SmallVector<Value *, 16> Args(CI->args());
+ Args.insert(Args.end() - 1, Builder.getInt32(*RedOp));
+ NewCall = Builder.CreateCall(NewFn, Args);
+ break;
+ }
case Intrinsic::riscv_sha256sig0:
case Intrinsic::riscv_sha256sig1:
case Intrinsic::riscv_sha256sum0:
diff --git a/llvm/lib/IR/NVVMIntrinsicUtils.cpp b/llvm/lib/IR/NVVMIntrinsicUtils.cpp
index d745a1cfb72cd..04df25c79dcf5 100644
--- a/llvm/lib/IR/NVVMIntrinsicUtils.cpp
+++ b/llvm/lib/IR/NVVMIntrinsicUtils.cpp
@@ -16,6 +16,15 @@
using namespace llvm;
using namespace nvvm;
+void nvvm::printTMAReductionOp(raw_ostream &OS, const Constant *ImmArgVal) {
+ const auto *CI = dyn_cast<ConstantInt>(ImmArgVal);
+ if (!CI || CI->getZExtValue() > static_cast<uint64_t>(TMAReductionOp::XOR))
+ llvm_unreachable(
+ "printTMAReductionOp called with invalid value for immediate argument");
+
+ OS << getTMAReductionOpName(static_cast<TMAReductionOp>(CI->getZExtValue()));
+}
+
void nvvm::printTcgen05MMAKind(raw_ostream &OS, const Constant *ImmArgVal) {
if (const auto *CI = dyn_cast<ConstantInt>(ImmArgVal)) {
uint64_t Val = CI->getZExtValue();
diff --git a/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.cpp b/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.cpp
index 9db838523fdcd..0b0b000107332 100644
--- a/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.cpp
+++ b/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.cpp
@@ -468,36 +468,9 @@ void NVPTXInstPrinter::printTmaReductionMode(const MCInst *MI, int OpNum,
const MCSubtargetInfo &,
raw_ostream &O) {
const MCOperand &MO = MI->getOperand(OpNum);
- using RedTy = nvvm::TMAReductionOp;
-
- switch (static_cast<RedTy>(MO.getImm())) {
- case RedTy::ADD:
- O << ".add";
- return;
- case RedTy::MIN:
- O << ".min";
- return;
- case RedTy::MAX:
- O << ".max";
- return;
- case RedTy::INC:
- O << ".inc";
- return;
- case RedTy::DEC:
- O << ".dec";
- return;
- case RedTy::AND:
- O << ".and";
- return;
- case RedTy::OR:
- O << ".or";
- return;
- case RedTy::XOR:
- O << ".xor";
- return;
- }
- llvm_unreachable(
- "Invalid Reduction Op in printCpAsyncBulkTensorReductionMode");
+ O << '.'
+ << nvvm::getTMAReductionOpName(
+ static_cast<nvvm::TMAReductionOp>(MO.getImm()));
}
void NVPTXInstPrinter::printCTAGroup(const MCInst *MI, int OpNum,
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
index a0e3b1cc8e47f..2ca1db5d77300 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
@@ -28,8 +28,6 @@
#include "llvm/IR/Instructions.h"
#include "llvm/IR/Intrinsics.h"
#include "llvm/IR/IntrinsicsNVPTX.h"
-#include "llvm/IR/LLVMContext.h"
-#include "llvm/IR/NVVMIntrinsicUtils.h"
#include "llvm/Support/AtomicOrdering.h"
#include "llvm/Support/CommandLine.h"
#include "llvm/Support/ErrorHandling.h"
@@ -1909,82 +1907,6 @@ NVPTX::Scope NVPTXScopes::operator[](SyncScope::ID ID) const {
bool NVPTXScopes::empty() const { return Scopes.size() == 0; }
-#define CP_ASYNC_BULK_TENSOR_OPCODE(dir, dim, mode, is_s32, suffix) \
- (is_s32 \
- ? NVPTX::CP_ASYNC_BULK_TENSOR_##dir##_##dim##_SHARED32_##mode##suffix \
- : NVPTX::CP_ASYNC_BULK_TENSOR_##dir##_##dim##_##mode##suffix)
-
-#define GET_CP_ASYNC_BULK_TENSOR_OPCODE_S2G_RED(dim, mode, is_ch, is_s32) \
- (is_ch ? (CP_ASYNC_BULK_TENSOR_OPCODE(RED, dim, mode, is_s32, _CH)) \
- : (CP_ASYNC_BULK_TENSOR_OPCODE(RED, dim, mode, is_s32, )))
-
-static unsigned GetCpAsyncBulkTensorS2GReductionOpcode(size_t Dim,
- bool IsShared32,
- bool IsCacheHint,
- bool IsIm2Col) {
- if (IsIm2Col) {
- switch (Dim) {
- case 3:
- return GET_CP_ASYNC_BULK_TENSOR_OPCODE_S2G_RED(3D, IM2COL, IsCacheHint,
- IsShared32);
- case 4:
- return GET_CP_ASYNC_BULK_TENSOR_OPCODE_S2G_RED(4D, IM2COL, IsCacheHint,
- IsShared32);
- case 5:
- return GET_CP_ASYNC_BULK_TENSOR_OPCODE_S2G_RED(5D, IM2COL, IsCacheHint,
- IsShared32);
- default:
- llvm_unreachable("Invalid Dimension in...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/213638
More information about the Mlir-commits
mailing list