shylie/llvm-project
1
0
Fork 0
Code Issues Pull Requests Actions 6 Packages Projects Releases Wiki Activity

AMDGPU: support s_monitor_sleep on gfx1250 (#146293)

Browse Source 
Co-Authored-by: Stanislav Mekhanoshin <Stanislav.Mekhanoshin@amd.com>
This commit is contained in:
Changpeng Fang 2025-06-29 21:01:47 -07:00 committed by GitHub
parent a17f63590a
commit bb982e733c
No known key found for this signature in database
GPG Key ID: B5690EEEBB952194
8 changed files with 76 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

1
clang/include/clang/Basic/BuiltinsAMDGPU.def
View File

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

9
clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250.cl
View File

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

4
clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1250-param.cl
View File

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

9
llvm/include/llvm/IR/IntrinsicsAMDGPU.td
View File

@ -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.

12
llvm/lib/Target/AMDGPU/SOPInstructions.td
View File

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

20
llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.monitor.sleep.ll Normal file
View File

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

12
llvm/test/MC/AMDGPU/gfx1250_asm_sopp.s
View File

@ -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

9
llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_sopp.txt
View File

@ -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