[Mlir-commits] [mlir] [mlir][nvvm] Introduce `setmaxregister.sync.aligned` Op (PR #73780)

[llvmlistbot at llvm.org]

================
@@ -400,6 +400,28 @@ def NVVM_FenceScClusterOp : NVVM_Op<"fence.sc.cluster"> {
   let assemblyFormat = "attr-dict";
 }
 
+def SetMaxRegisterActionIncrease : I32EnumAttrCase<"increase", 0>;
+def SetMaxRegisterActionDecrease   : I32EnumAttrCase<"decrease", 1>;
+def SetMaxRegisterAction : I32EnumAttr<"SetMaxRegisterAction", "NVVM set max register action",
+  [SetMaxRegisterActionDecrease, SetMaxRegisterActionIncrease]> {
+  let genSpecializedAttr = 0;
+  let cppNamespace = "::mlir::NVVM";
+}
+def SetMaxRegisterActionAttr : EnumAttr<NVVM_Dialect, SetMaxRegisterAction, "action">;
+
+def NVVM_SetMaxRegisterOp : NVVM_PTXBuilder_Op<"setmaxregister"> {
+  let arguments = (ins I32Attr:$regCount, SetMaxRegisterActionAttr:$action);
+  let assemblyFormat = "$action $regCount attr-dict";
+  let extraClassDefinition = [{        
+    std::string $cppClass::getPtx() {
----------------
durga4github wrote:

Thanks for the explanation, Guray!
I am with you on a).

On b): Yes, I was looking at the mbarrier.init as an example, and wanted to confirm if both the intrinsic and the inline PTX forms can co-exist. It seems they can, without any issues.
(I was running the mbarrier.init Op example and was always generating the intrinsic version.
Later realized I was using mlir-translate and not "mlir-opt with conversion-to-llvm" to see the inline-PTX impl.)

Thanks for all the clarifications!


https://github.com/llvm/llvm-project/pull/73780