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

[llvmlistbot at llvm.org]

Author: Rajat Bajpai
Date: 2026-02-11T16:23:22+05:30
New Revision: 435a162b01532080dc6cfc2579e8fdfe91ecaa3a

URL: https://github.com/llvm/llvm-project/commit/435a162b01532080dc6cfc2579e8fdfe91ecaa3a
DIFF: https://github.com/llvm/llvm-project/commit/435a162b01532080dc6cfc2579e8fdfe91ecaa3a.diff

LOG: [MLIR][NVVM][NFC] Fix PTX builder class api (#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.

Added: 
    

Modified: 
    mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td

Removed: 
    


################################################################################
diff  --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index a6ffc81c68688..9725cdb6a4c4d 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,