[Mlir-commits] [mlir] [MLIR][NVVM] Shard generated operation definitions (NFC) (PR #221812)
llvmlistbot at llvm.org
llvmlistbot at llvm.org
Mon Sep 7 12:49:14 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-mlir
Author: Mehdi Amini (joker-eph)
<details>
<summary>Changes</summary>
Generate the NVVM operation definitions in eight shards, use the generated registration hook, and expose only the four parser and range helpers referenced by generated code.
Assisted-by: Codex
---
Full diff: https://github.com/llvm/llvm-project/pull/221812.diff
6 Files Affected:
- (modified) mlir/include/mlir/Dialect/LLVMIR/CMakeLists.txt (+8-1)
- (modified) mlir/include/mlir/Dialect/LLVMIR/NVVMDialect.h (+2-2)
- (modified) mlir/lib/Dialect/LLVMIR/CMakeLists.txt (+15)
- (modified) mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp (+13-18)
- (added) mlir/lib/Dialect/LLVMIR/IR/NVVMOps.cpp (+16)
- (added) mlir/lib/Dialect/LLVMIR/IR/NVVMOps.h (+33)
``````````diff
diff --git a/mlir/include/mlir/Dialect/LLVMIR/CMakeLists.txt b/mlir/include/mlir/Dialect/LLVMIR/CMakeLists.txt
index de9a3e1c932823..3a39ebdf6137be 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/CMakeLists.txt
+++ b/mlir/include/mlir/Dialect/LLVMIR/CMakeLists.txt
@@ -62,7 +62,14 @@ mlir_tablegen(NVVMRequiresSMTraits.h.inc -gen-op-interface-decls)
mlir_tablegen(NVVMRequiresSMTraits.cpp.inc -gen-op-interface-defs)
add_mlir_dialect_tablegen_target(MLIRNVVMRequiresSMTraitsIncGen)
-add_mlir_dialect(NVVMOps nvvm)
+set(LLVM_TARGET_DEFINITIONS NVVMOps.td)
+mlir_tablegen(NVVMOps.h.inc -gen-op-decls -op-shard-count=8)
+mlir_tablegen(NVVMOps.cpp.inc -gen-op-defs -op-shard-count=8)
+mlir_tablegen(NVVMOpsTypes.h.inc -gen-typedef-decls -typedefs-dialect=nvvm)
+mlir_tablegen(NVVMOpsTypes.cpp.inc -gen-typedef-defs -typedefs-dialect=nvvm)
+mlir_tablegen(NVVMOpsDialect.h.inc -gen-dialect-decls -dialect=nvvm)
+mlir_tablegen(NVVMOpsDialect.cpp.inc -gen-dialect-defs -dialect=nvvm)
+add_mlir_dialect_tablegen_target(MLIRNVVMOpsIncGen)
add_mlir_doc(NVVMOps NVVMOps Dialects/ -gen-op-doc -dialect=nvvm)
set(LLVM_TARGET_DEFINITIONS NVVMOps.td)
mlir_tablegen(NVVMConversions.inc -gen-llvmir-conversions)
diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMDialect.h b/mlir/include/mlir/Dialect/LLVMIR/NVVMDialect.h
index 6bd582d66ed252..551949b415db05 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMDialect.h
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMDialect.h
@@ -67,9 +67,9 @@ std::pair<mlir::Type, unsigned> inferMMAType(mlir::NVVM::MMATypes type,
#define GET_ATTRDEF_CLASSES
#include "mlir/Dialect/LLVMIR/NVVMOpsAttributes.h.inc"
+#include "mlir/Dialect/LLVMIR/NVVMOpsDialect.h.inc"
+
#define GET_OP_CLASSES
#include "mlir/Dialect/LLVMIR/NVVMOps.h.inc"
-#include "mlir/Dialect/LLVMIR/NVVMOpsDialect.h.inc"
-
#endif /* MLIR_DIALECT_LLVMIR_NVVMDIALECT_H_ */
diff --git a/mlir/lib/Dialect/LLVMIR/CMakeLists.txt b/mlir/lib/Dialect/LLVMIR/CMakeLists.txt
index 494a1e4c3ee975..a93dc0635a9e5d 100644
--- a/mlir/lib/Dialect/LLVMIR/CMakeLists.txt
+++ b/mlir/lib/Dialect/LLVMIR/CMakeLists.txt
@@ -1,5 +1,9 @@
add_subdirectory(Transforms)
+set(LLVM_OPTIONAL_SOURCES
+ IR/NVVMOps.cpp
+ )
+
add_mlir_dialect_library(MLIRLLVMDialect
IR/FunctionCallUtils.cpp
IR/LLVMAttrs.cpp
@@ -41,10 +45,20 @@ add_mlir_dialect_library(MLIRLLVMDialect
MLIRSupport
)
+set(LLVM_TARGET_DEFINITIONS IR/NVVMOps.cpp)
+foreach(index RANGE 7)
+ set(NVVM_SHARDED_SRC NVVMOps.${index}.cpp)
+ list(APPEND NVVM_SHARDED_SRCS ${NVVM_SHARDED_SRC})
+ tablegen(MLIR_SRC_SHARDER ${NVVM_SHARDED_SRC}
+ -op-shard-index=${index})
+endforeach()
+add_public_tablegen_target(MLIRNVVMOpsShardGen)
+
add_mlir_dialect_library(MLIRNVVMDialect
IR/NVVMDialect.cpp
IR/BasicPtxBuilderInterface.cpp
IR/NVVMRequiresSMTraits.cpp
+ ${NVVM_SHARDED_SRCS}
ADDITIONAL_HEADER_DIRS
${MLIR_MAIN_INCLUDE_DIR}/mlir/Dialect/LLVMIR
@@ -52,6 +66,7 @@ add_mlir_dialect_library(MLIRNVVMDialect
DEPENDS
MLIRGPUCompilationAttrInterfacesIncGen
MLIRNVVMOpsIncGen
+ MLIRNVVMOpsShardGen
MLIRNVVMConversionsIncGen
MLIRBasicPtxBuilderInterfaceIncGen
MLIRNVVMRequiresSMTraitsIncGen
diff --git a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
index c6ed0a840a1b05..b201c09712810b 100644
--- a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
+++ b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
@@ -16,6 +16,8 @@
#include "mlir/Dialect/LLVMIR/NVVMDialect.h"
+#include "IR/NVVMOps.h"
+
#include "mlir/Conversion/ConvertToLLVM/ToLLVMInterface.h"
#include "mlir/Dialect/GPU/IR/CompilationInterfaces.h"
#include "mlir/Dialect/GPU/IR/GPUDialect.h"
@@ -93,8 +95,8 @@ getNVVMCtaGroupKind(NVVM::CTAGroupKind ctaGroup) {
llvm_unreachable("unsupported cta_group value");
}
-static ParseResult parseCTAGroup(OpAsmParser &parser,
- NVVM::CTAGroupKindAttr &groupAttr) {
+ParseResult mlir::NVVM::parseCTAGroup(OpAsmParser &parser,
+ NVVM::CTAGroupKindAttr &groupAttr) {
StringRef keyword;
if (parser.parseKeyword(&keyword))
return failure();
@@ -106,8 +108,8 @@ static ParseResult parseCTAGroup(OpAsmParser &parser,
return success();
}
-static void printCTAGroup(OpAsmPrinter &printer, Operation *,
- NVVM::CTAGroupKindAttr groupAttr) {
+void mlir::NVVM::printCTAGroup(OpAsmPrinter &printer, Operation *,
+ NVVM::CTAGroupKindAttr groupAttr) {
printer << NVVM::stringifyCTAGroupKind(groupAttr.getValue());
}
@@ -5818,10 +5820,10 @@ LogicalResult Tcgen05StOp::verify() {
/// Infer the result ranges for the NVVM SpecialRangeableRegisterOp that might
/// have ConstantRangeAttr.
-static void nvvmInferResultRanges(std::optional<LLVM::ConstantRangeAttr> range,
- Value result,
- ArrayRef<::mlir::ConstantIntRanges> argRanges,
- SetIntRangeFn setResultRanges) {
+void mlir::NVVM::nvvmInferResultRanges(
+ std::optional<LLVM::ConstantRangeAttr> range, Value result,
+ ArrayRef<::mlir::ConstantIntRanges> argRanges,
+ SetIntRangeFn setResultRanges) {
if (range) {
setResultRanges(result, {range->getLower(), range->getUpper(),
range->getLower(), range->getUpper()});
@@ -5832,9 +5834,8 @@ static void nvvmInferResultRanges(std::optional<LLVM::ConstantRangeAttr> range,
/// Verify the range attribute satisfies LLVM ConstantRange constructor
/// requirements for NVVM SpecialRangeableRegisterOp.
-static LogicalResult
-verifyConstantRangeAttr(Operation *op,
- std::optional<LLVM::ConstantRangeAttr> rangeAttr) {
+LogicalResult mlir::NVVM::verifyConstantRangeAttr(
+ Operation *op, std::optional<LLVM::ConstantRangeAttr> rangeAttr) {
if (!rangeAttr)
return success();
@@ -6957,10 +6958,7 @@ struct NVVMInlinerInterface final : DialectInlinerInterface {
// TODO: This should be the llvm.nvvm dialect once this is supported.
void NVVMDialect::initialize() {
- addOperations<
-#define GET_OP_LIST
-#include "mlir/Dialect/LLVMIR/NVVMOps.cpp.inc"
- >();
+ registerNVVMDialectOperations(this);
addAttributes<
#define GET_ATTRDEF_LIST
#include "mlir/Dialect/LLVMIR/NVVMOpsAttributes.cpp.inc"
@@ -7173,8 +7171,5 @@ LogicalResult NVVMTargetAttr::verifyTarget(Operation *gpuModule) {
return success();
}
-#define GET_OP_CLASSES
-#include "mlir/Dialect/LLVMIR/NVVMOps.cpp.inc"
-
#define GET_ATTRDEF_CLASSES
#include "mlir/Dialect/LLVMIR/NVVMOpsAttributes.cpp.inc"
diff --git a/mlir/lib/Dialect/LLVMIR/IR/NVVMOps.cpp b/mlir/lib/Dialect/LLVMIR/IR/NVVMOps.cpp
new file mode 100644
index 00000000000000..8bfa736721ca07
--- /dev/null
+++ b/mlir/lib/Dialect/LLVMIR/IR/NVVMOps.cpp
@@ -0,0 +1,16 @@
+//===- NVVMOps.cpp - NVVM operation definitions --------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include "mlir/Dialect/LLVMIR/NVVMDialect.h"
+
+#include "IR/NVVMOps.h"
+
+using namespace mlir;
+using namespace NVVM;
+
+#include "mlir/Dialect/LLVMIR/NVVMOps.cpp.inc"
diff --git a/mlir/lib/Dialect/LLVMIR/IR/NVVMOps.h b/mlir/lib/Dialect/LLVMIR/IR/NVVMOps.h
new file mode 100644
index 00000000000000..19ba277b8fcc0d
--- /dev/null
+++ b/mlir/lib/Dialect/LLVMIR/IR/NVVMOps.h
@@ -0,0 +1,33 @@
+//===- NVVMOps.h - NVVM operation implementation helpers -------*- C++ -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef MLIR_LIB_DIALECT_LLVMIR_IR_NVVMOPS_H
+#define MLIR_LIB_DIALECT_LLVMIR_IR_NVVMOPS_H
+
+#include "mlir/Dialect/LLVMIR/NVVMDialect.h"
+
+#include <optional>
+
+namespace mlir::NVVM {
+
+mlir::ParseResult parseCTAGroup(mlir::OpAsmParser &parser,
+ mlir::NVVM::CTAGroupKindAttr &groupAttr);
+void printCTAGroup(mlir::OpAsmPrinter &printer, mlir::Operation *,
+ mlir::NVVM::CTAGroupKindAttr groupAttr);
+
+void nvvmInferResultRanges(std::optional<mlir::LLVM::ConstantRangeAttr> range,
+ mlir::Value result,
+ mlir::ArrayRef<mlir::ConstantIntRanges> argRanges,
+ mlir::SetIntRangeFn setResultRanges);
+mlir::LogicalResult
+verifyConstantRangeAttr(mlir::Operation *op,
+ std::optional<mlir::LLVM::ConstantRangeAttr> rangeAttr);
+
+} // namespace mlir::NVVM
+
+#endif // MLIR_LIB_DIALECT_LLVMIR_IR_NVVMOPS_H
``````````
</details>
https://github.com/llvm/llvm-project/pull/221812
More information about the Mlir-commits
mailing list