[Mlir-commits] [mlir] [NFC][MLIR][NVVM] Fix PTX builder class api (PR #180787)

[Rajat Bajpai]

https://github.com/rajatbajpai created https://github.com/llvm/llvm-project/pull/180787

Previously, `NVVM_PTXBuilder_Op` included `BasicPtxBuilderOpInterface` as part of the default value of the `traits` parameter. This meant any subclass that provided an explicit traits list would silently replace the default and lose the interface, defeating the purpose of the base class. Callers had to redundantly re-specify the interface.

>From 0263a7ca27bebd778107903b52ee3b81b1efa5e3 Mon Sep 17 00:00:00 2001
From: rbajpai <rbajpai at nvidia.com>
Date: Tue, 10 Feb 2026 17:21:23 +0530
Subject: [PATCH] [NFC][MLIR][NVVM] Fix PTX builder class api
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

Previously, `NVVM_PTXBuilder_Op` included `BasicPtxBuilderOpInterface` as part
of the default value of the `traits` parameter. This meant any subclass that
provided an explicit traits list would silently replace the default and lose the
interface — defeating the purpose of the base class. Callers had to redundantly
re-specify the interface.
---
 mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td | 10 +++++-----
 1 file changed, 5 insertions(+), 5 deletions(-)

diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index 5cd4a1af1f881..e6b870fe1461b 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -117,9 +117,10 @@ class NVVM_Op<string mnemonic, list<Trait> traits = []> :
 }
 
 /// Base class that defines BasicPtxBuilderOpInterface. 
-class NVVM_PTXBuilder_Op<string mnemonic, 
-  list<Trait> traits = [DeclareOpInterfaceMethods<BasicPtxBuilderOpInterface>]> :
-  LLVM_OpBase<NVVM_Dialect, mnemonic, traits> {
+class NVVM_PTXBuilder_Op<string mnemonic, list<Trait> extraTraits = []> :
+  LLVM_OpBase<NVVM_Dialect, mnemonic,
+              !listconcat([DeclareOpInterfaceMethods<BasicPtxBuilderOpInterface>],
+                          extraTraits)> {
 }
 
 //===----------------------------------------------------------------------===//
@@ -4049,8 +4050,7 @@ def NVVM_CpAsyncBulkTensorGlobalToSharedClusterOp :
 
 def NVVM_CpAsyncBulkTensorSharedCTAToGlobalOp : 
   NVVM_PTXBuilder_Op<"cp.async.bulk.tensor.global.shared.cta",
-  [DeclareOpInterfaceMethods<BasicPtxBuilderOpInterface>,
-  AttrSizedOperandSegments]>,
+  [AttrSizedOperandSegments]>,
   Arguments<(ins LLVM_PointerGeneric:$tmaDescriptor,
                  LLVM_PointerShared:$srcMem,
                  Variadic<I32>:$coordinates,