[clang] [llvm] AMDGPU: support s_monitor_sleep on gfx1250 (PR #146293)
Changpeng Fang via llvm-commits
llvm-commits at lists.llvm.org
Sun Jun 29 16:12:39 PDT 2025
https://github.com/changpeng created https://github.com/llvm/llvm-project/pull/146293
None
>From d6f55a31a5b9cc0716149424e219d219c2970d0c Mon Sep 17 00:00:00 2001
From: Changpeng Fang <changpeng.fang at amd.com>
Date: Sun, 29 Jun 2025 16:03:46 -0700
Subject: [PATCH] AMDGPU: support s_monitor_sleep on gfx1250
Co-Authored-by: Stanislav Mekhanoshin <Stanislav.Mekhanoshin at amd.com>
---
clang/include/clang/Basic/BuiltinsAMDGPU.def | 1 +
.../CodeGenOpenCL/builtins-amdgcn-gfx1250.cl | 9 +++++++++
.../builtins-amdgcn-error-gfx1250-param.cl | 4 ++++
llvm/include/llvm/IR/IntrinsicsAMDGPU.td | 9 +++++++++
llvm/lib/Target/AMDGPU/SOPInstructions.td | 12 +++++++++++
.../AMDGPU/llvm.amdgcn.s.monitor.sleep.ll | 20 +++++++++++++++++++
llvm/test/MC/AMDGPU/gfx1250_asm_sopp.s | 12 +++++++++++
.../Disassembler/AMDGPU/gfx1250_dasm_sopp.txt | 9 +++++++++
8 files changed, 76 insertions(+)
create mode 100644 llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.monitor.sleep.ll
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def
index 4e28f3bb7ef81..948da2c99e47c 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.def
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def
@@ -654,6 +654,7 @@ TARGET_BUILTIN(__builtin_amdgcn_ds_load_tr16_b128_v8f16, "V8hV8h*3", "nc", "gfx1
TARGET_BUILTIN(__builtin_amdgcn_ds_load_tr16_b128_v8bf16, "V8yV8y*3", "nc", "gfx1250-insts,wavefrontsize32")
TARGET_BUILTIN(__builtin_amdgcn_s_setprio_inc_wg, "vIs", "n", "setprio-inc-wg-inst")
+TARGET_BUILTIN(__builtin_amdgcn_s_monitor_sleep, "vIs", "n", "gfx1250-insts")
TARGET_BUILTIN(__builtin_amdgcn_cvt_pk_f16_fp8, "V2hs", "nc", "gfx1250-insts")
TARGET_BUILTIN(__builtin_amdgcn_cvt_pk_f16_bf8, "V2hs", "nc", "gfx1250-insts")
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250.cl
index 864e301859682..569df2f1fb4e6 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250.cl
@@ -15,6 +15,15 @@ void test_setprio_inc_wg() {
__builtin_amdgcn_s_setprio_inc_wg(10);
}
+// CHECK-LABEL: @test_s_monitor_sleep(
+// CHECK-NEXT: entry:
+// CHECK-NEXT: call void @llvm.amdgcn.s.monitor.sleep(i16 10)
+// CHECK-NEXT: ret void
+//
+void test_s_monitor_sleep() {
+ __builtin_amdgcn_s_monitor_sleep(10);
+}
+
// CHECK-LABEL: @test_cvt_pk_f16_fp8(
// CHECK-NEXT: entry:
// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5)
diff --git a/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1250-param.cl b/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1250-param.cl
index b69fcb5f445bc..771ae555508c4 100644
--- a/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1250-param.cl
+++ b/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1250-param.cl
@@ -4,3 +4,7 @@
void test_setprio_inc_wg(short a) {
__builtin_amdgcn_s_setprio_inc_wg(a); // expected-error {{'__builtin_amdgcn_s_setprio_inc_wg' must be a constant integer}}
}
+
+void test_s_monitor_sleep(short a) {
+ __builtin_amdgcn_s_monitor_sleep(a); // expected-error {{'__builtin_amdgcn_s_monitor_sleep' must be a constant integer}}
+}
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index ce37702b91486..b3e937a2d3d9f 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -3500,6 +3500,15 @@ def int_amdgcn_ashr_pk_u8_i32 : ClangBuiltin<"__builtin_amdgcn_ashr_pk_u8_i32">,
DefaultAttrsIntrinsic<[llvm_i16_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
[IntrNoMem, IntrSpeculatable]>;
+//===----------------------------------------------------------------------===//
+// gfx1250 intrinsics
+// ===----------------------------------------------------------------------===//
+
+def int_amdgcn_s_monitor_sleep :
+ ClangBuiltin<"__builtin_amdgcn_s_monitor_sleep">,
+ DefaultAttrsIntrinsic<[], [llvm_i16_ty], [ImmArg<ArgIndex<0>>, IntrNoMem,
+ IntrHasSideEffects]>;
+
//===----------------------------------------------------------------------===//
// Special Intrinsics for backend internal use only. No frontend
// should emit calls to these.
diff --git a/llvm/lib/Target/AMDGPU/SOPInstructions.td b/llvm/lib/Target/AMDGPU/SOPInstructions.td
index de217cc602c98..c7c4276e0e252 100644
--- a/llvm/lib/Target/AMDGPU/SOPInstructions.td
+++ b/llvm/lib/Target/AMDGPU/SOPInstructions.td
@@ -1680,6 +1680,12 @@ def S_SET_GPR_IDX_OFF : SOPP_Pseudo<"s_set_gpr_idx_off", (ins) > {
let Uses = [MODE];
}
}
+
+def S_MONITOR_SLEEP : SOPP_Pseudo <"s_monitor_sleep", (ins i16imm:$simm16), "$simm16",
+ [(int_amdgcn_s_monitor_sleep timm:$simm16)]> {
+ let SubtargetPredicate = isGFX1250Plus;
+}
+
} // End hasSideEffects
let SubtargetPredicate = HasVGPRIndexMode in {
@@ -2692,6 +2698,12 @@ defm S_ICACHE_INV : SOPP_Real_32_gfx11_gfx12<0x03c>;
defm S_BARRIER : SOPP_Real_32_gfx11<0x03d>;
+//===----------------------------------------------------------------------===//
+// SOPP - GFX1250.
+//===----------------------------------------------------------------------===//
+
+defm S_MONITOR_SLEEP : SOPP_Real_32_gfx12<0x004>;
+
//===----------------------------------------------------------------------===//
// SOPP - GFX6, GFX7, GFX8, GFX9, GFX10
//===----------------------------------------------------------------------===//
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.monitor.sleep.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.monitor.sleep.ll
new file mode 100644
index 0000000000000..706f470a6285a
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.monitor.sleep.ll
@@ -0,0 +1,20 @@
+; RUN: llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx1250 < %s | FileCheck --check-prefix=GCN %s
+; RUN: llc -global-isel=1 -mtriple=amdgcn -mcpu=gfx1250 < %s | FileCheck --check-prefix=GCN %s
+
+declare void @llvm.amdgcn.s.monitor.sleep(i16)
+
+; GCN-LABEL: {{^}}test_monitor_sleep_1:
+; GCN: s_monitor_sleep 1
+define amdgpu_ps void @test_monitor_sleep_1() {
+ call void @llvm.amdgcn.s.monitor.sleep(i16 1)
+ ret void
+}
+
+; FIXME: 0x8000 would look better
+
+; GCN-LABEL: {{^}}test_monitor_sleep_forever:
+; GCN: s_monitor_sleep 0xffff8000
+define amdgpu_ps void @test_monitor_sleep_forever() {
+ call void @llvm.amdgcn.s.monitor.sleep(i16 32768)
+ ret void
+}
diff --git a/llvm/test/MC/AMDGPU/gfx1250_asm_sopp.s b/llvm/test/MC/AMDGPU/gfx1250_asm_sopp.s
index 48ec44b410c2c..6ebc17468eed6 100644
--- a/llvm/test/MC/AMDGPU/gfx1250_asm_sopp.s
+++ b/llvm/test/MC/AMDGPU/gfx1250_asm_sopp.s
@@ -16,3 +16,15 @@ s_wait_xcnt 0xf
s_setprio_inc_wg 100
// GFX1250: [0x64,0x00,0xbe,0xbf]
// GFX12-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
+
+s_monitor_sleep 1
+// GFX1250: s_monitor_sleep 1 ; encoding: [0x01,0x00,0x84,0xbf]
+// GFX12-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
+
+s_monitor_sleep 32768
+// GFX1250: s_monitor_sleep 0x8000 ; encoding: [0x00,0x80,0x84,0xbf]
+// GFX12-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
+
+s_monitor_sleep 0
+// GFX1250: s_monitor_sleep 0 ; encoding: [0x00,0x00,0x84,0xbf]
+// GFX12-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
diff --git a/llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_sopp.txt b/llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_sopp.txt
index 55f74d3a31bf7..220f9e5084f0e 100644
--- a/llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_sopp.txt
+++ b/llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_sopp.txt
@@ -11,3 +11,12 @@
# GFX1250: s_setprio_inc_wg 0x64 ; encoding: [0x64,0x00,0xbe,0xbf]
0x64,0x00,0xbe,0xbf
+
+# GFX1250: s_monitor_sleep 0 ; encoding: [0x00,0x00,0x84,0xbf]
+0x00,0x00,0x84,0xbf
+
+# GFX1250: s_monitor_sleep 0x8000 ; encoding: [0x00,0x80,0x84,0xbf]
+0x00,0x80,0x84,0xbf
+
+# GFX1250: s_monitor_sleep 1 ; encoding: [0x01,0x00,0x84,0xbf]
+0x01,0x00,0x84,0xbf
More information about the llvm-commits
mailing list