[Mlir-commits] [llvm] [mlir] [mlir][nvvm]: separate NVVM targets from MLIRLLVMIRTransforms (PR #205640)

[Jeremy Kun]

https://github.com/j2kun created https://github.com/llvm/llvm-project/pull/205640

None

>From e37fc73910d55d0f0e506323b9bf5d676ed1b7a9 Mon Sep 17 00:00:00 2001
From: Jeremy Kun <jkun at google.com>
Date: Wed, 24 Jun 2026 11:51:14 -0700
Subject: [PATCH] [mlir][nvvm]: separate NVVM targets from MLIRLLVMIRTransforms

---
 .../include/mlir/Dialect/LLVMIR/NVVMDialect.h |  4 ++++
 .../LLVMIR/Transforms/InlinerInterfaceImpl.h  |  6 ------
 mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp    | 19 +++++++++++++++++++
 .../Dialect/LLVMIR/Transforms/CMakeLists.txt  | 15 +++++++++++++--
 .../Transforms/InlinerInterfaceImpl.cpp       |  7 -------
 .../llvm-project-overlay/mlir/BUILD.bazel     | 19 ++++++++++++++++---
 6 files changed, 52 insertions(+), 18 deletions(-)

diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMDialect.h b/mlir/include/mlir/Dialect/LLVMIR/NVVMDialect.h
index 6bd582d66ed25..fbf1dee00139d 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMDialect.h
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMDialect.h
@@ -60,6 +60,10 @@ std::pair<mlir::Type, unsigned> inferMMAType(mlir::NVVM::MMATypes type,
                                              mlir::NVVM::MMAFrag frag, int nRow,
                                              int nCol,
                                              mlir::MLIRContext *context);
+
+/// Register the `NVVMInlinerInterface` implementation of
+/// `DialectInlinerInterface` with the NVVM dialect.
+void registerInlinerInterface(DialectRegistry &registry);
 } // namespace NVVM
 } // namespace mlir
 
diff --git a/mlir/include/mlir/Dialect/LLVMIR/Transforms/InlinerInterfaceImpl.h b/mlir/include/mlir/Dialect/LLVMIR/Transforms/InlinerInterfaceImpl.h
index 69cc2e32285b6..472a049d567fd 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/Transforms/InlinerInterfaceImpl.h
+++ b/mlir/include/mlir/Dialect/LLVMIR/Transforms/InlinerInterfaceImpl.h
@@ -24,12 +24,6 @@ void registerInlinerInterface(DialectRegistry &registry);
 
 } // namespace LLVM
 
-namespace NVVM {
-/// Register the `NVVMInlinerInterface` implementation of
-/// `DialectInlinerInterface` with the NVVM dialect.
-void registerInlinerInterface(DialectRegistry &registry);
-} // namespace NVVM
-
 } // namespace mlir
 
 #endif // MLIR_DIALECT_LLVMIR_TRANSFORMS_INLINERINTERFACEIMPL_H
diff --git a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
index 13de1e7f2b2ad..379dbb2bb811e 100644
--- a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
+++ b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
@@ -16,6 +16,9 @@
 
 #include "mlir/Dialect/LLVMIR/NVVMDialect.h"
 
+#include "mlir/IR/IRMapping.h"
+#include "mlir/Transforms/InliningUtils.h"
+
 #include "mlir/Conversion/ConvertToLLVM/ToLLVMInterface.h"
 #include "mlir/Dialect/GPU/IR/CompilationInterfaces.h"
 #include "mlir/Dialect/GPU/IR/GPUDialect.h"
@@ -6552,3 +6555,19 @@ LogicalResult NVVMTargetAttr::verifyTarget(Operation *gpuModule) {
 
 #define GET_ATTRDEF_CLASSES
 #include "mlir/Dialect/LLVMIR/NVVMOpsAttributes.cpp.inc"
+
+namespace {
+struct NVVMInlinerInterface : public DialectInlinerInterface {
+  using DialectInlinerInterface::DialectInlinerInterface;
+
+  bool isLegalToInline(Operation *, Region *, bool, IRMapping &) const final {
+    return true;
+  }
+};
+} // namespace
+
+void mlir::NVVM::registerInlinerInterface(DialectRegistry &registry) {
+  registry.addExtension(+[](MLIRContext *ctx, NVVM::NVVMDialect *dialect) {
+    dialect->addInterfaces<NVVMInlinerInterface>();
+  });
+}
diff --git a/mlir/lib/Dialect/LLVMIR/Transforms/CMakeLists.txt b/mlir/lib/Dialect/LLVMIR/Transforms/CMakeLists.txt
index 6ffa2d9a8a9e8..bc3dfdbc9887c 100644
--- a/mlir/lib/Dialect/LLVMIR/Transforms/CMakeLists.txt
+++ b/mlir/lib/Dialect/LLVMIR/Transforms/CMakeLists.txt
@@ -5,7 +5,6 @@ add_mlir_dialect_library(MLIRLLVMIRTransforms
   DIScopeForLLVMFuncOp.cpp
   InlinerInterfaceImpl.cpp
   LegalizeForExport.cpp
-  OptimizeForNVVM.cpp
   RequestCWrappers.cpp
   UseDefaultVisibilityPass.cpp
 
@@ -18,6 +17,18 @@ add_mlir_dialect_library(MLIRLLVMIRTransforms
   MLIRLLVMDialect
   MLIRPass
   MLIRTransforms
-  MLIRNVVMDialect
   MLIROpenMPDialect
   )
+
+add_mlir_dialect_library(MLIRLLVMIRNVVMTransforms
+  OptimizeForNVVM.cpp
+
+  DEPENDS
+  MLIRLLVMPassIncGen
+
+  LINK_LIBS PUBLIC
+  MLIRLLVMIRTransforms
+  MLIRNVVMDialect
+  MLIRPass
+  MLIRTransforms
+  )
diff --git a/mlir/lib/Dialect/LLVMIR/Transforms/InlinerInterfaceImpl.cpp b/mlir/lib/Dialect/LLVMIR/Transforms/InlinerInterfaceImpl.cpp
index 680fafcc099f2..657835ae54149 100644
--- a/mlir/lib/Dialect/LLVMIR/Transforms/InlinerInterfaceImpl.cpp
+++ b/mlir/lib/Dialect/LLVMIR/Transforms/InlinerInterfaceImpl.cpp
@@ -14,7 +14,6 @@
 #include "mlir/Dialect/LLVMIR/Transforms/InlinerInterfaceImpl.h"
 #include "mlir/Analysis/SliceWalk.h"
 #include "mlir/Dialect/LLVMIR/LLVMDialect.h"
-#include "mlir/Dialect/LLVMIR/NVVMDialect.h"
 #include "mlir/IR/Matchers.h"
 #include "mlir/Interfaces/DataLayoutInterfaces.h"
 #include "mlir/Interfaces/ViewLikeInterface.h"
@@ -861,9 +860,3 @@ void mlir::LLVM::registerInlinerInterface(DialectRegistry &registry) {
     dialect->addInterfaces<LLVMInlinerInterface>();
   });
 }
-
-void mlir::NVVM::registerInlinerInterface(DialectRegistry &registry) {
-  registry.addExtension(+[](MLIRContext *ctx, NVVM::NVVMDialect *dialect) {
-    dialect->addInterfaces<LLVMInlinerInterface>();
-  });
-}
diff --git a/utils/bazel/llvm-project-overlay/mlir/BUILD.bazel b/utils/bazel/llvm-project-overlay/mlir/BUILD.bazel
index fd0bf0bcce72c..f7d6359d5a135 100644
--- a/utils/bazel/llvm-project-overlay/mlir/BUILD.bazel
+++ b/utils/bazel/llvm-project-overlay/mlir/BUILD.bazel
@@ -5543,6 +5543,7 @@ cc_library(
             "lib/Dialect/LLVMIR/Transforms/DIExpressionLegalization.cpp",
             "lib/Dialect/LLVMIR/Transforms/DIExpressionRewriter.cpp",
             "lib/Dialect/LLVMIR/Transforms/LegalizeForExport.cpp",
+            "lib/Dialect/LLVMIR/Transforms/OptimizeForNVVM.cpp",
         ],
     ),
     hdrs = glob(
@@ -5563,12 +5564,24 @@ cc_library(
         ":LLVMDialect",
         ":LLVMIRTransformsLegalizeForExport",
         ":LLVMPassIncGen",
-        ":NVVMDialect",
+        "//llvm:BinaryFormat",
+        "//llvm:Support",
         ":Pass",
         ":TransformUtils",
         ":ViewLikeInterface",
-        "//llvm:BinaryFormat",
-        "//llvm:Support",
+    ],
+)
+
+cc_library(
+    name = "LLVMIRNVVMTransforms",
+    srcs = ["lib/Dialect/LLVMIR/Transforms/OptimizeForNVVM.cpp"],
+    hdrs = ["include/mlir/Dialect/LLVMIR/Transforms/OptimizeForNVVM.h"],
+    includes = ["include"],
+    deps = [
+        ":LLVMIRTransforms",
+        ":NVVMDialect",
+        ":Pass",
+        ":Transforms",
     ],
 )