RPCSX/llvm
1
0
Fork 0
mirror of https://github.com/RPCSX/llvm.git synced 2024-11-30 23:20:54 +00:00
Code Issues Actions Packages Projects Releases Wiki Activity

AMDGPU: Add s_sleep intrinsic

Browse Source 
git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@262120 91177308-0d34-0410-b5e6-96231b3b80d8
This commit is contained in:
Matt Arsenault 2016-02-27 08:53:52 +00:00
parent a164276e20
commit 788be52946
4 changed files with 67 additions and 1 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

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

@ -192,6 +192,11 @@ def int_amdgcn_s_memtime :
GCCBuiltin<"__builtin_amdgcn_s_memtime">,
Intrinsic<[llvm_i64_ty], [], []>;
def int_amdgcn_s_sleep :
GCCBuiltin<"__builtin_amdgcn_s_sleep">,
Intrinsic<[], [llvm_i32_ty], []> {
}
def int_amdgcn_dispatch_ptr :
GCCBuiltin<"__builtin_amdgcn_dispatch_ptr">,
Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 2>], [], [IntrNoMem]>;

4
lib/Target/AMDGPU/SIInstrInfo.td
View File

@ -371,6 +371,10 @@ def IMM16bit : PatLeaf <(imm),
[{return isUInt<16>(N->getZExtValue());}]
>;
def SIMM16bit : PatLeaf <(imm),
[{return isInt<16>(N->getSExtValue());}]
>;
def IMM20bit : PatLeaf <(imm),
[{return isUInt<20>(N->getZExtValue());}]
>;

14
lib/Target/AMDGPU/SIInstructions.td
View File

@ -501,10 +501,22 @@ def S_BARRIER : SOPP <0x0000000a, (ins), "s_barrier",
def S_WAITCNT : SOPP <0x0000000c, (ins WAIT_FLAG:$simm16), "s_waitcnt $simm16">;
def S_SETHALT : SOPP <0x0000000d, (ins i16imm:$simm16), "s_sethalt $simm16">;
def S_SLEEP : SOPP <0x0000000e, (ins i16imm:$simm16), "s_sleep $simm16">;
// On SI the documentation says sleep for approximately 64 * low 2
// bits, consistent with the reported maximum of 448. On VI the
// maximum reported is 960 cycles, so 960 / 64 = 15 max, so is the
// maximum really 15 on VI?
def S_SLEEP : SOPP <0x0000000e, (ins i32imm:$simm16),
"s_sleep $simm16", [(int_amdgcn_s_sleep SIMM16bit:$simm16)]> {
let hasSideEffects = 1;
let mayLoad = 1;
let mayStore = 1;
}
def S_SETPRIO : SOPP <0x0000000f, (ins i16imm:$sim16), "s_setprio $sim16">;
let Uses = [EXEC, M0] in {
// FIXME: Should this be mayLoad+mayStore?
def S_SENDMSG : SOPP <0x00000010, (ins SendMsgImm:$simm16), "s_sendmsg $simm16",
[(AMDGPUsendmsg (i32 imm:$simm16))]
>;

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

@ -0,0 +1,45 @@
; RUN: llc -march=amdgcn -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s
; RUN: llc -march=amdgcn -mcpu=tonga -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s
declare void @llvm.amdgcn.s.sleep(i32) #0
; GCN-LABEL: {{^}}test_s_sleep:
; GCN: s_sleep 0{{$}}
; GCN: s_sleep 1{{$}}
; GCN: s_sleep 2{{$}}
; GCN: s_sleep 3{{$}}
; GCN: s_sleep 4{{$}}
; GCN: s_sleep 5{{$}}
; GCN: s_sleep 6{{$}}
; GCN: s_sleep 7{{$}}
; GCN: s_sleep 8{{$}}
; GCN: s_sleep 9{{$}}
; GCN: s_sleep 10{{$}}
; GCN: s_sleep 11{{$}}
; GCN: s_sleep 12{{$}}
; GCN: s_sleep 13{{$}}
; GCN: s_sleep 14{{$}}
; GCN: s_sleep 15{{$}}
define void @test_s_sleep(i32 %x) #0 {
call void @llvm.amdgcn.s.sleep(i32 0)
call void @llvm.amdgcn.s.sleep(i32 1)
call void @llvm.amdgcn.s.sleep(i32 2)
call void @llvm.amdgcn.s.sleep(i32 3)
call void @llvm.amdgcn.s.sleep(i32 4)
call void @llvm.amdgcn.s.sleep(i32 5)
call void @llvm.amdgcn.s.sleep(i32 6)
call void @llvm.amdgcn.s.sleep(i32 7)
; Values that might only work on VI
call void @llvm.amdgcn.s.sleep(i32 8)
call void @llvm.amdgcn.s.sleep(i32 9)
call void @llvm.amdgcn.s.sleep(i32 10)
call void @llvm.amdgcn.s.sleep(i32 11)
call void @llvm.amdgcn.s.sleep(i32 12)
call void @llvm.amdgcn.s.sleep(i32 13)
call void @llvm.amdgcn.s.sleep(i32 14)
call void @llvm.amdgcn.s.sleep(i32 15)
ret void
}
attributes #0 = { nounwind }