AMDGPU: support s_monitor_sleep on gfx1250 #146293

changpeng · 2025-06-29T23:12:32Z

Co-Authored-by: Stanislav Mekhanoshin [email protected]

Co-Authored-by: Stanislav Mekhanoshin <[email protected]>

llvmbot · 2025-06-29T23:13:00Z

@llvm/pr-subscribers-llvm-ir

@llvm/pr-subscribers-mc

Author: Changpeng Fang (changpeng)

Changes

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

8 Files Affected:

(modified) clang/include/clang/Basic/BuiltinsAMDGPU.def (+1)
(modified) clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250.cl (+9)
(modified) clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1250-param.cl (+4)
(modified) llvm/include/llvm/IR/IntrinsicsAMDGPU.td (+9)
(modified) llvm/lib/Target/AMDGPU/SOPInstructions.td (+12)
(added) llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.monitor.sleep.ll (+20)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_sopp.s (+12)
(modified) llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_sopp.txt (+9)

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

llvmbot · 2025-06-29T23:13:01Z

@llvm/pr-subscribers-backend-amdgpu

Author: Changpeng Fang (changpeng)

Changes

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

8 Files Affected:

(modified) clang/include/clang/Basic/BuiltinsAMDGPU.def (+1)
(modified) clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250.cl (+9)
(modified) clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1250-param.cl (+4)
(modified) llvm/include/llvm/IR/IntrinsicsAMDGPU.td (+9)
(modified) llvm/lib/Target/AMDGPU/SOPInstructions.td (+12)
(added) llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.monitor.sleep.ll (+20)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_sopp.s (+12)
(modified) llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_sopp.txt (+9)

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

llvmbot · 2025-06-29T23:13:01Z

@llvm/pr-subscribers-clang

Author: Changpeng Fang (changpeng)

Changes

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

8 Files Affected:

(modified) clang/include/clang/Basic/BuiltinsAMDGPU.def (+1)
(modified) clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250.cl (+9)
(modified) clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1250-param.cl (+4)
(modified) llvm/include/llvm/IR/IntrinsicsAMDGPU.td (+9)
(modified) llvm/lib/Target/AMDGPU/SOPInstructions.td (+12)
(added) llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.monitor.sleep.ll (+20)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_sopp.s (+12)
(modified) llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_sopp.txt (+9)

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

Co-Authored-by: Stanislav Mekhanoshin <[email protected]>

AMDGPU: support s_monitor_sleep on gfx1250

d6f55a3

Co-Authored-by: Stanislav Mekhanoshin <[email protected]>

changpeng requested review from arsenm, shiltian and rampitec June 29, 2025 23:12

llvmbot added clang Clang issues not falling into any other category backend:AMDGPU clang:frontend Language frontend issues, e.g. anything involving "Sema" mc Machine (object) code llvm:ir labels Jun 29, 2025

arsenm approved these changes Jun 30, 2025

View reviewed changes

changpeng merged commit bb982e7 into llvm:main Jun 30, 2025
13 checks passed

changpeng deleted the sleep branch June 30, 2025 04:02

rlavaee pushed a commit to rlavaee/llvm-project that referenced this pull request Jul 1, 2025

AMDGPU: support s_monitor_sleep on gfx1250 (llvm#146293)

af3c470

Co-Authored-by: Stanislav Mekhanoshin <[email protected]>

rlavaee pushed a commit to rlavaee/llvm-project that referenced this pull request Jul 1, 2025

AMDGPU: support s_monitor_sleep on gfx1250 (llvm#146293)

913b4cd

Co-Authored-by: Stanislav Mekhanoshin <[email protected]>

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Uh oh!

AMDGPU: support s_monitor_sleep on gfx1250 #146293

AMDGPU: support s_monitor_sleep on gfx1250 #146293

Uh oh!

changpeng commented Jun 29, 2025 •

edited

Loading

Uh oh!

llvmbot commented Jun 29, 2025 •

edited

Loading

Uh oh!

llvmbot commented Jun 29, 2025

Uh oh!

llvmbot commented Jun 29, 2025

Uh oh!

Uh oh!

Uh oh!

AMDGPU: support s_monitor_sleep on gfx1250 #146293

AMDGPU: support s_monitor_sleep on gfx1250 #146293

Uh oh!

Conversation

changpeng commented Jun 29, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

llvmbot commented Jun 29, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

llvmbot commented Jun 29, 2025

Uh oh!

llvmbot commented Jun 29, 2025

Uh oh!

Uh oh!

Uh oh!

changpeng commented Jun 29, 2025 •

edited

Loading

llvmbot commented Jun 29, 2025 •

edited

Loading