[Mlir-commits] [mlir] [MLIR][NVVM] Add trap (PR #120247)

[llvmlistbot at llvm.org]

llvmbot wrote:


<!--LLVM PR SUMMARY COMMENT-->

@llvm/pr-subscribers-mlir

Author: Guray Ozen (grypp)

<details>
<summary>Changes</summary>

PR adds `trap` instruction to nvvm dialect.

---
Full diff: https://github.com/llvm/llvm-project/pull/120247.diff


2 Files Affected:

- (modified) mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td (+36-22) 
- (modified) mlir/test/Target/LLVMIR/nvvmir.mlir (+10) 


``````````diff
diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index 5d8772d9d5c5f5..a594d564f5a71c 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -685,27 +685,6 @@ def NVVM_FenceProxyReleaseOp : NVVM_Op<"fence.proxy.release">,
   let hasVerifier = 1;
 }
 
-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_Op<"setmaxregister"> {
-  let arguments = (ins I32Attr:$regCount, SetMaxRegisterActionAttr:$action);
-  let assemblyFormat = "$action $regCount attr-dict";
-  let hasVerifier = 1;
-  string llvmBuilder = [{
-    auto intId = (op.getAction() == NVVM::SetMaxRegisterAction::increase) ?
-      llvm::Intrinsic::nvvm_setmaxnreg_inc_sync_aligned_u32 :
-      llvm::Intrinsic::nvvm_setmaxnreg_dec_sync_aligned_u32;
-
-    createIntrinsicCall(builder, intId, builder.getInt32($regCount));
-  }];
-}
 
 def NVVM_FenceMbarrierInitOp : NVVM_PTXBuilder_Op<"fence.mbarrier.init"> {
   let arguments = (ins );
@@ -2317,7 +2296,7 @@ def NVVM_WgmmaMmaAsyncOp : NVVM_Op<"wgmma.mma_async",
 }
 
 //===----------------------------------------------------------------------===//
-// NVVM breakpoint Op
+// NVVM Miscellaneous instructions
 //===----------------------------------------------------------------------===//
 
 def NVVM_Breakpoint : NVVM_Op<"breakpoint"> {
@@ -2333,6 +2312,41 @@ def NVVM_Breakpoint : NVVM_Op<"breakpoint"> {
   let assemblyFormat = "attr-dict";
 }
 
+def NVVM_Trap : NVVM_Op<"trap"> {
+  let summary = "Trap Op";
+  let description = [{
+    Trap aborts execution and generate an interrupt to the host CPU.
+    [For more information, see PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#miscellaneous-instructions-trap)
+  }];
+  string llvmBuilder = [{
+    createIntrinsicCall(builder, llvm::Intrinsic::trap);
+  }];
+
+  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_Op<"setmaxregister"> {
+  let arguments = (ins I32Attr:$regCount, SetMaxRegisterActionAttr:$action);
+  let assemblyFormat = "$action $regCount attr-dict";
+  let hasVerifier = 1;
+  string llvmBuilder = [{
+    auto intId = (op.getAction() == NVVM::SetMaxRegisterAction::increase) ?
+      llvm::Intrinsic::nvvm_setmaxnreg_inc_sync_aligned_u32 :
+      llvm::Intrinsic::nvvm_setmaxnreg_dec_sync_aligned_u32;
+
+    createIntrinsicCall(builder, intId, builder.getInt32($regCount));
+  }];
+}
+
 //===----------------------------------------------------------------------===//
 // NVVM target attribute.
 //===----------------------------------------------------------------------===//
diff --git a/mlir/test/Target/LLVMIR/nvvmir.mlir b/mlir/test/Target/LLVMIR/nvvmir.mlir
index 2749e42c40bc34..a2c6f31c0c9be5 100644
--- a/mlir/test/Target/LLVMIR/nvvmir.mlir
+++ b/mlir/test/Target/LLVMIR/nvvmir.mlir
@@ -698,9 +698,19 @@ llvm.func @nvvm_fence_proxy_tensormap_generic_acquire(%addr : !llvm.ptr) {
 }
 
 // -----
+
 // CHECK-LABEL: @nvvm_breakpoint
 llvm.func @nvvm_breakpoint() {
   // CHECK: call void @llvm.debugtrap()
   nvvm.breakpoint
   llvm.return
 }
+
+// -----
+
+// CHECK-LABEL: @nvvm_trap
+llvm.func @nvvm_trap() {
+  // CHECK: call void @llvm.trap()
+  nvvm.trap
+  llvm.return
+}

``````````

</details>


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