[Mlir-commits] [mlir] [MLIR][NVVM] Update PMEvent lowering to intrinsics (PR #171649)
Durgadoss R
llvmlistbot at llvm.org
Wed Dec 10 08:55:37 PST 2025
https://github.com/durga4github created https://github.com/llvm/llvm-project/pull/171649
The patch updates the lowering of `id` based pmevent
also to intrinsics. The mask is simply (1 << event-id).
>From cd18519cc8407dc555db688ab729374969d7423a Mon Sep 17 00:00:00 2001
From: Durgadoss R <durgadossr at nvidia.com>
Date: Wed, 10 Dec 2025 19:11:06 +0530
Subject: [PATCH] [MLIR][NVVM] Update PMEvent lowering to intrinsics
The patch updates the lowering of `id` based pmevent
also to intrinsics. The mask is simply (1 << event-id).
Signed-off-by: Durgadoss R <durgadossr at nvidia.com>
---
mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td | 20 ++++++++--------
mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp | 19 +++++++++++++++
.../Conversion/NVVMToLLVM/nvvm-to-llvm.mlir | 15 ------------
mlir/test/Target/LLVMIR/nvvm/pm_event.mlir | 23 +++++++++++++++++++
.../Target/LLVMIR/nvvm/pm_event_invalid.mlir | 21 +++++++++++++++++
mlir/test/Target/LLVMIR/nvvmir-invalid.mlir | 21 -----------------
mlir/test/Target/LLVMIR/nvvmir.mlir | 11 ---------
7 files changed, 73 insertions(+), 57 deletions(-)
create mode 100644 mlir/test/Target/LLVMIR/nvvm/pm_event.mlir
create mode 100644 mlir/test/Target/LLVMIR/nvvm/pm_event_invalid.mlir
diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index a0a00513d7da5..51d310970fda9 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -543,7 +543,7 @@ def NVVM_NanosleepOp : NVVM_Op<"nanosleep">,
// NVVM Performance Monitor events
//===----------------------------------------------------------------------===//
-def NVVM_PMEventOp : NVVM_PTXBuilder_Op<"pmevent">,
+def NVVM_PMEventOp : NVVM_Op<"pmevent">,
Arguments<(ins OptionalAttr<I16Attr>:$maskedEventId,
OptionalAttr<I32Attr>:$eventId)> {
let summary = "Trigger one or more Performance Monitor events.";
@@ -561,20 +561,20 @@ def NVVM_PMEventOp : NVVM_PTXBuilder_Op<"pmevent">,
[For more information, see PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#miscellaneous-instructions-pmevent)
}];
- string llvmBuilder = [{
- llvm::Value *mId = builder.getInt16(* $maskedEventId);
- createIntrinsicCall(builder, llvm::Intrinsic::nvvm_pm_event_mask, {mId});
- }];
-
let assemblyFormat = "attr-dict (`id` `=` $eventId^)? (`mask` `=` $maskedEventId^)?";
+ let hasVerifier = 1;
let extraClassDeclaration = [{
- bool hasIntrinsic() { return !getEventId(); }
+ static mlir::NVVM::IDArgPair
+ getIntrinsicIDAndArgs(Operation &op, LLVM::ModuleTranslation &mt,
+ llvm::IRBuilderBase& builder);
}];
- let extraClassDefinition = [{
- std::string $cppClass::getPtx() { return std::string("pmevent %0;"); }
+
+ string llvmBuilder = [{
+ auto [id, args] = NVVM::PMEventOp::getIntrinsicIDAndArgs(
+ *op, moduleTranslation, builder);
+ createIntrinsicCall(builder, id, args);
}];
- let hasVerifier = 1;
}
//===----------------------------------------------------------------------===//
diff --git a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
index 5ce56e6399e31..d9c895e6827f9 100644
--- a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
+++ b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
@@ -2476,6 +2476,25 @@ mlir::NVVM::IDArgPair NVVM::BarrierOp::getIntrinsicIDAndArgs(
return {id, std::move(args)};
}
+mlir::NVVM::IDArgPair
+PMEventOp::getIntrinsicIDAndArgs(Operation &op, LLVM::ModuleTranslation &mt,
+ llvm::IRBuilderBase &builder) {
+ auto thisOp = cast<NVVM::PMEventOp>(op);
+ llvm::Type *i16Ty = llvm::Type::getInt16Ty(mt.getLLVMContext());
+
+ // With event-id, mask is generated as (1 << event-id)
+ llvm::Value *maskVal;
+ if (auto eventAttr = thisOp.getEventIdAttr()) {
+ uint16_t mask = static_cast<uint16_t>(1u << eventAttr.getInt());
+ maskVal = llvm::ConstantInt::get(i16Ty, mask);
+ } else {
+ maskVal =
+ llvm::ConstantInt::get(i16Ty, thisOp.getMaskedEventIdAttr().getValue());
+ }
+
+ return {llvm::Intrinsic::nvvm_pm_event_mask, {maskVal}};
+}
+
mlir::NVVM::IDArgPair MBarrierInitOp::getIntrinsicIDAndArgs(
Operation &op, LLVM::ModuleTranslation &mt, llvm::IRBuilderBase &builder) {
auto thisOp = cast<NVVM::MBarrierInitOp>(op);
diff --git a/mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir b/mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir
index 8fb36ace2c463..c4b8e93b6a9f9 100644
--- a/mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir
+++ b/mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir
@@ -678,21 +678,6 @@ llvm.func @inline_ptx_multi_rw_r(%a : i32, %b : i32, %rw_c : f32, %rw_d : f32)
llvm.return %r5 : f32
}
-
-// -----
-
-// CHECK-LABEL: @nvvm_pmevent
-llvm.func @nvvm_pmevent() {
- // CHECK: %[[S0:.+]] = llvm.mlir.constant(10 : i32) : i32
- // CHECK: llvm.inline_asm has_side_effects asm_dialect = att "pmevent $0;", "n" %[[S0]] : (i32) -> ()
-
- nvvm.pmevent id = 10
- // CHECK: %[[S1:.+]] = llvm.mlir.constant(4 : i32) : i32
- // CHECK: llvm.inline_asm has_side_effects asm_dialect = att "pmevent $0;", "n" %[[S1]] : (i32) -> ()
- nvvm.pmevent id = 4
- llvm.return
-}
-
// -----
llvm.func @inline_ptx_pack_4i8(%src : vector<4xi8>, %mask : i32, %zero: i32) {
diff --git a/mlir/test/Target/LLVMIR/nvvm/pm_event.mlir b/mlir/test/Target/LLVMIR/nvvm/pm_event.mlir
new file mode 100644
index 0000000000000..0092d32319a83
--- /dev/null
+++ b/mlir/test/Target/LLVMIR/nvvm/pm_event.mlir
@@ -0,0 +1,23 @@
+// RUN: mlir-translate -mlir-to-llvmir %s | FileCheck %s
+
+llvm.func @nvvm_pmevent_mask() {
+ // CHECK-LABEL: define void @nvvm_pmevent_mask() {
+ // CHECK-NEXT: call void @llvm.nvvm.pm.event.mask(i16 15000)
+ // CHECK-NEXT: call void @llvm.nvvm.pm.event.mask(i16 4)
+ // CHECK-NEXT: ret void
+ // CHECK-NEXT: }
+ nvvm.pmevent mask = 15000
+ nvvm.pmevent mask = 4
+ llvm.return
+}
+
+llvm.func @nvvm_pmevent_id() {
+ // CHECK-LABEL: define void @nvvm_pmevent_id() {
+ // CHECK-NEXT: call void @llvm.nvvm.pm.event.mask(i16 1024)
+ // CHECK-NEXT: call void @llvm.nvvm.pm.event.mask(i16 16)
+ // CHECK-NEXT: ret void
+ // CHECK-NEXT: }
+ nvvm.pmevent id = 10
+ nvvm.pmevent id = 4
+ llvm.return
+}
diff --git a/mlir/test/Target/LLVMIR/nvvm/pm_event_invalid.mlir b/mlir/test/Target/LLVMIR/nvvm/pm_event_invalid.mlir
new file mode 100644
index 0000000000000..783988fb36368
--- /dev/null
+++ b/mlir/test/Target/LLVMIR/nvvm/pm_event_invalid.mlir
@@ -0,0 +1,21 @@
+// RUN: mlir-translate -verify-diagnostics -split-input-file -mlir-to-llvmir %s
+
+llvm.func @pmevent_no_id() {
+ // expected-error @below {{either `id` or `mask` must be set}}
+ nvvm.pmevent
+}
+
+// -----
+
+llvm.func @pmevent_bigger15() {
+ // expected-error @below {{`id` must be between 0 and 15}}
+ nvvm.pmevent id = 16
+}
+
+// -----
+
+llvm.func @pmevent_many_ids() {
+ // expected-error @below {{`id` and `mask` cannot be set at the same time}}
+ nvvm.pmevent id = 1 mask = 1
+}
+
diff --git a/mlir/test/Target/LLVMIR/nvvmir-invalid.mlir b/mlir/test/Target/LLVMIR/nvvmir-invalid.mlir
index d5868ee73cc50..c0fe0fa11f497 100644
--- a/mlir/test/Target/LLVMIR/nvvmir-invalid.mlir
+++ b/mlir/test/Target/LLVMIR/nvvmir-invalid.mlir
@@ -1,26 +1,5 @@
// RUN: mlir-translate -verify-diagnostics -split-input-file -mlir-to-llvmir %s
-llvm.func @pmevent_no_id() {
- // expected-error @below {{either `id` or `mask` must be set}}
- nvvm.pmevent
-}
-
-// -----
-
-llvm.func @pmevent_bigger15() {
- // expected-error @below {{`id` must be between 0 and 15}}
- nvvm.pmevent id = 141
-}
-
-// -----
-
-llvm.func @pmevent_many_ids() {
- // expected-error @below {{`id` and `mask` cannot be set at the same time}}
- nvvm.pmevent id = 1 mask = 1
-}
-
-// -----
-
llvm.func @kernel_func(%numberOfThreads : i32) {
// expected-error @below {{'nvvm.barrier' op barrier id is missing, it should be set between 0 to 15}}
nvvm.barrier number_of_threads = %numberOfThreads
diff --git a/mlir/test/Target/LLVMIR/nvvmir.mlir b/mlir/test/Target/LLVMIR/nvvmir.mlir
index c4a69097692cb..9e4aadac69896 100644
--- a/mlir/test/Target/LLVMIR/nvvmir.mlir
+++ b/mlir/test/Target/LLVMIR/nvvmir.mlir
@@ -903,17 +903,6 @@ llvm.func @nvvm_dot_accumulate_2way(%a: vector<2xi16>, %b: vector<4xi8>, %c: i32
// -----
-// CHECK-LABEL: @nvvm_pmevent
-llvm.func @nvvm_pmevent() {
- // CHECK: call void @llvm.nvvm.pm.event.mask(i16 15000)
- nvvm.pmevent mask = 15000
- // CHECK: call void @llvm.nvvm.pm.event.mask(i16 4)
- nvvm.pmevent mask = 4
- llvm.return
-}
-
-// -----
-
// CHECK-LABEL: @nanosleep
llvm.func @nanosleep(%duration: i32) {
// CHECK: call void @llvm.nvvm.nanosleep(i32 %{{.*}})
More information about the Mlir-commits
mailing list