AMDGPU: support s_monitor_sleep on gfx1250 (llvm#146293)

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

Author: changpeng

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")

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)

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}}
+}

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.

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
 //===----------------------------------------------------------------------===//

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
+}

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

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