[Mlir-commits] [mlir] c9e07d9 - [MLIR][NVVM] Shard generated operation definitions (NFC) (#221812)
llvmlistbot at llvm.org
llvmlistbot at llvm.org
Thu Sep 10 02:51:39 PDT 2026
Author: Mehdi Amini
Date: 2026-09-10T09:51:33Z
New Revision: c9e07d91eb9338f1922adbc120f68cf071bd9c2c
URL: https://github.com/llvm/llvm-project/commit/c9e07d91eb9338f1922adbc120f68cf071bd9c2c
DIFF: https://github.com/llvm/llvm-project/commit/c9e07d91eb9338f1922adbc120f68cf071bd9c2c.diff
LOG: [MLIR][NVVM] Shard generated operation definitions (NFC) (#221812)
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
Added:
mlir/lib/Dialect/LLVMIR/IR/NVVMOps.cpp
mlir/lib/Dialect/LLVMIR/IR/NVVMOps.h
Modified:
mlir/include/mlir/Dialect/LLVMIR/CMakeLists.txt
mlir/include/mlir/Dialect/LLVMIR/NVVMDialect.h
mlir/lib/Dialect/LLVMIR/CMakeLists.txt
mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
Removed:
################################################################################
diff --git a/mlir/include/mlir/Dialect/LLVMIR/CMakeLists.txt b/mlir/include/mlir/Dialect/LLVMIR/CMakeLists.txt
index a81288495f625..1fc410ef075b8 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/CMakeLists.txt
+++ b/mlir/include/mlir/Dialect/LLVMIR/CMakeLists.txt
@@ -67,7 +67,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 6bd582d66ed25..551949b415db0 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 aa65ae820c817..ae78dc9a2810f 100644
--- a/mlir/lib/Dialect/LLVMIR/CMakeLists.txt
+++ b/mlir/lib/Dialect/LLVMIR/CMakeLists.txt
@@ -2,6 +2,7 @@ add_subdirectory(Transforms)
set(LLVM_OPTIONAL_SOURCES
IR/LLVMOps.cpp
+ IR/NVVMOps.cpp
IR/ROCDLOps.cpp
)
@@ -13,7 +14,6 @@ foreach(index RANGE 7)
-op-shard-index=${index})
endforeach()
add_public_tablegen_target(MLIRLLVMOpsShardGen)
-
add_mlir_dialect_library(MLIRLLVMDialect
IR/FunctionCallUtils.cpp
IR/LLVMAttrs.cpp
@@ -60,10 +60,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
@@ -71,6 +81,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 6705f423e3332..688b3b0cd468a 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());
}
@@ -5836,10 +5838,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()});
@@ -5850,9 +5852,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();
@@ -6975,10 +6976,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"
@@ -7196,8 +7194,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 0000000000000..8bfa736721ca0
--- /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 0000000000000..19ba277b8fcc0
--- /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
More information about the Mlir-commits
mailing list