RPCS3/llvm-mirror
1
0
Fork 0
mirror of https://github.com/RPCS3/llvm-mirror.git synced 2024-11-26 04:32:44 +01:00
Code Issues Projects Releases Wiki Activity

[AMDGPU] New intrinsic void llvm.amdgcn.s.sethalt(i32)

Browse Source 
The expected use case is for frontends to insert this into
shaders that are to be run under a debugger. The shader can
then be resumed or single stepped from the point of the call
under debugger control.

Differential Revision: https://reviews.llvm.org/D97670
This commit is contained in:
Jay Foad 2021-03-01 09:45:55 +00:00
parent 82b9132694
commit d90cb44457
3 changed files with 34 additions and 1 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

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

@ -1283,6 +1283,10 @@ def int_amdgcn_s_decperflevel :
IntrHasSideEffects, IntrWillReturn]> {
}
def int_amdgcn_s_sethalt :
Intrinsic<[], [llvm_i32_ty], [ImmArg<ArgIndex<0>>, IntrNoMem,
IntrHasSideEffects, IntrWillReturn]>;
def int_amdgcn_s_getreg :
GCCBuiltin<"__builtin_amdgcn_s_getreg">,
Intrinsic<[llvm_i32_ty], [llvm_i32_ty],

3
lib/Target/AMDGPU/SOPInstructions.td
View File

@ -1228,7 +1228,8 @@ def S_WAKEUP : SOPP_Pseudo <"s_wakeup", (ins) > {
let mayLoad = 0, mayStore = 0, hasSideEffects = 1 in
def S_WAITCNT : SOPP_Pseudo <"s_waitcnt" , (ins WAIT_FLAG:$simm16), "$simm16",
[(int_amdgcn_s_waitcnt timm:$simm16)]>;
def S_SETHALT : SOPP_Pseudo <"s_sethalt" , (ins i16imm:$simm16), "$simm16">;
def S_SETHALT : SOPP_Pseudo <"s_sethalt" , (ins i32imm:$simm16), "$simm16",
[(int_amdgcn_s_sethalt timm:$simm16)]>;
def S_SETKILL : SOPP_Pseudo <"s_setkill" , (ins i16imm:$simm16), "$simm16">;
// On SI the documentation says sleep for approximately 64 * low 2

28
test/CodeGen/AMDGPU/llvm.amdgcn.s.sethalt.ll Normal file
View File

@ -0,0 +1,28 @@
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py
; RUN: llc -march=amdgcn -mcpu=gfx900 -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s
; RUN: llc -global-isel -march=amdgcn -mcpu=gfx900 -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s
define amdgpu_kernel void @test_s_sethalt() {
; GCN-LABEL: test_s_sethalt:
; GCN: ; %bb.0:
; GCN-NEXT: s_sethalt 0
; GCN-NEXT: s_sethalt 1
; GCN-NEXT: s_sethalt 2
; GCN-NEXT: s_sethalt 3
; GCN-NEXT: s_sethalt 4
; GCN-NEXT: s_sethalt 5
; GCN-NEXT: s_sethalt 6
; GCN-NEXT: s_sethalt 7
; GCN-NEXT: s_endpgm
call void @llvm.amdgcn.s.sethalt(i32 0)
call void @llvm.amdgcn.s.sethalt(i32 1)
call void @llvm.amdgcn.s.sethalt(i32 2)
call void @llvm.amdgcn.s.sethalt(i32 3)
call void @llvm.amdgcn.s.sethalt(i32 4)
call void @llvm.amdgcn.s.sethalt(i32 5)
call void @llvm.amdgcn.s.sethalt(i32 6)
call void @llvm.amdgcn.s.sethalt(i32 7)
ret void
}
declare void @llvm.amdgcn.s.sethalt(i32)