[clang] [llvm] [NVPTX] Add pm_event intrinsics (PR #141278)

[Durgadoss R via cfe-commits]

https://github.com/durga4github updated https://github.com/llvm/llvm-project/pull/141278

>From 0fc21a165a6f9202b441d1d8c4afa1252f9d6cc6 Mon Sep 17 00:00:00 2001
From: Durgadoss R <durgadossr at nvidia.com>
Date: Fri, 23 May 2025 20:43:18 +0530
Subject: [PATCH] [NVPTX] Add pm_event intrinsics

This patch adds pm_event.mask intrinsic and its
clang-builtin.

Signed-off-by: Durgadoss R <durgadossr at nvidia.com>
---
 clang/include/clang/Basic/BuiltinsNVPTX.td |  1 +
 clang/test/CodeGen/builtins-nvptx.c        |  7 +++++++
 llvm/docs/NVPTXUsage.rst                   | 23 ++++++++++++++++++++++
 llvm/include/llvm/IR/IntrinsicsNVVM.td     |  5 +++++
 llvm/lib/Target/NVPTX/NVPTXIntrinsics.td   | 10 ++++++++++
 llvm/test/CodeGen/NVPTX/pm-event.ll        | 15 ++++++++++++++
 6 files changed, 61 insertions(+)
 create mode 100644 llvm/test/CodeGen/NVPTX/pm-event.ll

diff --git a/clang/include/clang/Basic/BuiltinsNVPTX.td b/clang/include/clang/Basic/BuiltinsNVPTX.td
index 2cea44e224674..3e479a3d62dd8 100644
--- a/clang/include/clang/Basic/BuiltinsNVPTX.td
+++ b/clang/include/clang/Basic/BuiltinsNVPTX.td
@@ -177,6 +177,7 @@ let Attributes = [NoReturn] in {
 }
 let Attributes = [NoThrow] in {
   def __nvvm_nanosleep : NVPTXBuiltinSMAndPTX<"void(unsigned int)", SM_70, PTX63>;
+  def __nvvm_pm_event_mask : NVPTXBuiltin<"void(unsigned short)">;
 }
 
 // Min Max
diff --git a/clang/test/CodeGen/builtins-nvptx.c b/clang/test/CodeGen/builtins-nvptx.c
index cef529163bb39..f994adb14e457 100644
--- a/clang/test/CodeGen/builtins-nvptx.c
+++ b/clang/test/CodeGen/builtins-nvptx.c
@@ -883,6 +883,13 @@ __device__ void nvvm_vote(int pred) {
   // CHECK: ret void
 }
 
+// CHECK-LABEL: nvvm_pm_event_mask
+__device__ void nvvm_pm_event_mask() {
+  // CHECK: call void @llvm.nvvm.pm.event.mask(i16 255)
+  __nvvm_pm_event_mask(255);
+  // CHECK: ret void
+}
+
 // CHECK-LABEL: nvvm_nanosleep
 __device__ void nvvm_nanosleep(int d) {
 #if __CUDA_ARCH__ >= 700
diff --git a/llvm/docs/NVPTXUsage.rst b/llvm/docs/NVPTXUsage.rst
index 8bb0f2ed17c32..d51686c0b830c 100644
--- a/llvm/docs/NVPTXUsage.rst
+++ b/llvm/docs/NVPTXUsage.rst
@@ -1868,6 +1868,29 @@ If the request failed, the behavior of these intrinsics is undefined.
 
 For more information, refer `PTX ISA <https://docs.nvidia.com/cuda/parallel-thread-execution/?a#parallel-synchronization-and-communication-instructions-clusterlaunchcontrol-query-cancel>`__.
 
+Perf Monitor Event Intrinsics
+-----------------------------
+
+'``llvm.nvvm.pm.event.mask``' Intrinsic
+^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
+Syntax:
+"""""""
+
+.. code-block:: llvm
+
+    declare void @llvm.nvvm.pm.event.mask(i16 immarg %mask_val)
+
+Overview:
+"""""""""
+
+The '``llvm.nvvm.pm.event.mask``' intrinsic triggers one or more
+performance monitor events. Each bit in the 16-bit immediate operand
+``%mask_val`` controls an event.
+
+For more information on the pmevent instructions, refer to the PTX ISA
+`<https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#miscellaneous-instructions-pmevent>`_.
+
 Other Intrinsics
 ----------------
 
diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index 91e7d188c8533..8c8e778b57061 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -768,6 +768,11 @@ let TargetPrefix = "nvvm" in {
       DefaultAttrsIntrinsic<[], [llvm_i32_ty],
                             [IntrConvergent, IntrNoMem, IntrHasSideEffects]>;
 
+  // Performance Monitor Events (pm events) intrinsics
+  def int_nvvm_pm_event_mask : NVVMBuiltin,
+      DefaultAttrsIntrinsic<[], [llvm_i16_ty],
+                [IntrConvergent, IntrNoMem, IntrHasSideEffects,
+                 ImmArg<ArgIndex<0>>]>;
 //
 // Min Max
 //
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index 8fb5884fa2a20..71da857841c95 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -1052,6 +1052,16 @@ def INT_NVVM_NANOSLEEP_I : NVPTXInst<(outs), (ins i32imm:$i), "nanosleep.u32 \t$
 def INT_NVVM_NANOSLEEP_R : NVPTXInst<(outs), (ins Int32Regs:$i), "nanosleep.u32 \t$i;",
                              [(int_nvvm_nanosleep i32:$i)]>,
         Requires<[hasPTX<63>, hasSM<70>]>;
+
+let hasSideEffects = 1 in {
+// Performance Monitor events
+def INT_PM_EVENT_MASK : BasicNVPTXInst<(outs),
+                        (ins i16imm:$mask),
+                        "pmevent.mask",
+                        [(int_nvvm_pm_event_mask timm:$mask)]>,
+                        Requires<[hasSM<20>, hasPTX<30>]>;
+} // hasSideEffects
+
 //
 // Min Max
 //
diff --git a/llvm/test/CodeGen/NVPTX/pm-event.ll b/llvm/test/CodeGen/NVPTX/pm-event.ll
new file mode 100644
index 0000000000000..871da6d414978
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/pm-event.ll
@@ -0,0 +1,15 @@
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 | FileCheck %s
+; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
+
+declare void @llvm.nvvm.pm.event.mask(i16 %mask)
+
+; CHECK-LABEL: test_pm_event
+define void @test_pm_event() {
+  ; CHECK: pmevent.mask 255;
+  call void @llvm.nvvm.pm.event.mask(i16 u0xff)
+
+  ; CHECK: pmevent.mask 4096;
+  call void @llvm.nvvm.pm.event.mask(i16 u0x1000)
+
+  ret void
+}