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

[Guray Ozen]

================
@@ -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() {
----------------
grypp wrote:

FWIW, I think we should add LLVM intrinsic (this is "the right thing"). 

a) We can (should) keep the Op's interface same. So `nvgpu->nvvm` lowering remains unchanged (and other lowerings we have internally). We can use LLVM intrinsic when we have it. 

Let's say we have LLVM intrinsic for this Op, what we need to is to add `llvmBuilder` part and remove `getPtx`. 

```
def NVVM_SetMaxRegisterOp : NVVM_Op<"setmaxregister"> {
  let arguments = (ins I32Attr:$regCount, SetMaxRegisterActionAttr:$action);
  let assemblyFormat = "$action $regCount attr-dict";
  string llvmBuilder = [{
    if(getAction() == NVVM::SetMaxRegisterAction::increase)
      createIntrinsicCall(builder, llvm::Intrinsic::nvvm_set_max_register_increase);
    else createIntrinsicCall(builder,   llvm::Intrinsic::nvvm_set_max_register_descreate);
  }];
  let hasVerifier = 1;
}
```


b) I've added a `hasIntrinsic()` in `BasicPtxBuilderOpInterface`  interface to use inline PTX and LLVM intrinsic together. If LLVM contains all necessary intrinsics for the Op, we don't need `BasicPtxBuilderOpInterface` and `hasIntrinsic()`.

For example you can take a look at [`mbarrier.init` Op](https://github.com/llvm/llvm-project/blob/main/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td#L213-L225C2)


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