RPCS3/llvm-mirror
1
0
Fork 0
mirror of https://github.com/RPCS3/llvm-mirror.git synced 2025-01-31 20:51:52 +01:00
Code Issues Projects Releases Wiki Activity

AMDGPU: Add new name for barrier intrinsic

Browse Source 
llvm-svn: 258558
This commit is contained in:
Matt Arsenault 2016-01-22 21:30:43 +00:00
parent 7a5e15697d
commit 3913a77bb9
3 changed files with 38 additions and 1 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

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

@ -68,6 +68,9 @@ def int_AMDGPU_ldexp : Intrinsic<
let TargetPrefix = "amdgcn" in {
def int_amdgcn_s_barrier : GCCBuiltin<"__builtin_amdgcn_s_barrier">,
Intrinsic<[], [], [IntrConvergent]>;
def int_amdgcn_div_scale : GCCBuiltin<"__builtin_amdgcn_div_scale">,
// 1st parameter: Numerator
// 2nd parameter: Denominator

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

@ -481,7 +481,7 @@ def S_CBRANCH_EXECNZ : SOPP <
let hasSideEffects = 1 in {
def S_BARRIER : SOPP <0x0000000a, (ins), "s_barrier",
[(int_AMDGPU_barrier_local)]
[(int_amdgcn_s_barrier)]
> {
let SchedRW = [WriteBarrier];
let simm16 = 0;
@ -2167,11 +2167,17 @@ def : Pat <
// SOPP Patterns
//===----------------------------------------------------------------------===//
// FIXME: These should be removed eventually
def : Pat <
(int_AMDGPU_barrier_global),
(S_BARRIER)
>;
def : Pat <
(int_AMDGPU_barrier_local),
(S_BARRIER)
>;
//===----------------------------------------------------------------------===//
// VOP1 Patterns
//===----------------------------------------------------------------------===//

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

@ -0,0 +1,28 @@
; RUN: llc -march=amdgcn -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s
; GCN-LABEL: {{^}}test_barrier:
; GCN: buffer_store_dword
; GCN: s_waitcnt
; GCN: s_barrier
define void @test_barrier(i32 addrspace(1)* %out) #0 {
entry:
%tmp = call i32 @llvm.r600.read.tidig.x()
%tmp1 = getelementptr i32, i32 addrspace(1)* %out, i32 %tmp
store i32 %tmp, i32 addrspace(1)* %tmp1
call void @llvm.amdgcn.s.barrier()
%tmp2 = call i32 @llvm.r600.read.local.size.x()
%tmp3 = sub i32 %tmp2, 1
%tmp4 = sub i32 %tmp3, %tmp
%tmp5 = getelementptr i32, i32 addrspace(1)* %out, i32 %tmp4
%tmp6 = load i32, i32 addrspace(1)* %tmp5
store i32 %tmp6, i32 addrspace(1)* %tmp1
ret void
}
declare void @llvm.amdgcn.s.barrier() #1
declare i32 @llvm.r600.read.tidig.x() #2
declare i32 @llvm.r600.read.local.size.x() #2
attributes #0 = { nounwind }
attributes #1 = { convergent nounwind }
attributes #2 = { nounwind readnone }