[llvm] [AMDGPU][WIP] Added mask intrinsics (PR #193538)

[Ravil Dorozhinskii via llvm-commits]

https://github.com/ravil-mobile created https://github.com/llvm/llvm-project/pull/193538

None

>From 8c9c1cf54ccca07b24966c30a6681f9ecc58521a Mon Sep 17 00:00:00 2001
From: ravil-mobile <ravil.aviva.com at gmail.com>
Date: Wed, 22 Apr 2026 16:25:30 +0000
Subject: [PATCH] [AMDGPU][WIP] Added mask intrinsics

---
 llvm/include/llvm/IR/IntrinsicsAMDGPU.td  | 10 +++++++
 llvm/lib/Target/AMDGPU/SIISelLowering.cpp | 34 +++++++++++++++++++++++
 2 files changed, 44 insertions(+)

diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index 9528fb2b446bc..a9f77d9c4ad07 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -3844,6 +3844,16 @@ def int_amdgcn_cooperative_atomic_store_16x8B : AMDGPUCooperativeAtomicStore<llv
 def int_amdgcn_cooperative_atomic_load_8x16B : AMDGPUCooperativeAtomicLoad<llvm_v4i32_ty>;
 def int_amdgcn_cooperative_atomic_store_8x16B : AMDGPUCooperativeAtomicStore<llvm_v4i32_ty>;
 
+def int_amdgcn_mask : ClangBuiltin<"__builtin_amdgcn_mask">,
+  Intrinsic<[llvm_anyint_ty],
+  [llvm_i1_ty], [IntrWillReturn, IntrNoCallback, IntrNoFree]
+>;
+
+def int_amdgcn_reset_mask : ClangBuiltin<"__builtin_amdgcn_reset_mask">,
+  Intrinsic<[],
+  [llvm_anyint_ty], [IntrWillReturn, IntrNoCallback, IntrNoFree]
+>;
+
 //===----------------------------------------------------------------------===//
 // Special Intrinsics for backend internal use only. No frontend
 // should emit calls to these.
diff --git a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
index 8bb072a9626a9..1efb0a0078d60 100644
--- a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
+++ b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
@@ -11742,6 +11742,30 @@ SDValue SITargetLowering::LowerINTRINSIC_W_CHAIN(SDValue Op,
                                    Op->getVTList(), {Chain, Ptr},
                                    MII->getMemoryVT(), MII->getMemOperand());
   }
+  case Intrinsic::amdgcn_mask: {
+    SDNode *N = Op.getNode();
+    SDLoc SL(N);
+    SDValue Pred = N->getOperand(2);
+
+    // generate a new EXEC mask based on the predicate value
+    SmallVector<SDValue, 2> Operands;
+    Operands.push_back(
+    DAG.getTargetConstant(Intrinsic::amdgcn_ballot, SL, MVT::i32));
+    Operands.push_back(Pred);
+    SDValue NewMask = DAG.getNode(ISD::INTRINSIC_WO_CHAIN, SL, MVT::i32, Operands);
+
+    // save the original EXEC mask
+    EVT ExecVT = Subtarget->isWave32() ? MVT::i32 : MVT::i64;
+    Register ExecReg = Subtarget->isWave32() ? AMDGPU::EXEC_LO : AMDGPU::EXEC;
+    SDValue OrigMask = DAG.getCopyFromReg(DAG.getEntryNode(), SL, ExecReg, ExecVT);
+    SDValue Chain = OrigMask.getValue(1);
+
+    // change the EXEC mask
+    SDValue Res = DAG.getCopyToReg(Chain, SL, ExecReg, NewMask, SDValue());
+
+    return DAG.getMergeValues({OrigMask.getValue(0), Res.getValue(0)}, SL);
+  }
+
   default:
 
     if (const AMDGPU::ImageDimIntrinsicInfo *ImageDimIntr =
@@ -12425,6 +12449,16 @@ SDValue SITargetLowering::LowerINTRINSIC_VOID(SDValue Op,
     return DAG.getAtomic(ISD::ATOMIC_STORE, DL, MII->getMemoryVT(), Chain, Val,
                          Ptr, MII->getMemOperand());
   }
+  case Intrinsic::amdgcn_reset_mask: {
+    SDNode *N = Op.getNode();
+    SDLoc SL(N);
+
+    SDValue Chain = N->getOperand(0);
+    SDValue Mask = N->getOperand(2);
+    const GCNSubtarget *ST = this->getSubtarget();
+    Register ExecReg = ST->isWave32() ? AMDGPU::EXEC_LO : AMDGPU::EXEC;
+    return DAG.getCopyToReg(Chain, SL, ExecReg, Mask.getValue(0), SDValue());
+  }
   default: {
     if (const AMDGPU::ImageDimIntrinsicInfo *ImageDimIntr =
             AMDGPU::getImageDimIntrinsicInfo(IntrinsicID))