This is an archive of the discontinued LLVM Phabricator instance.

Differential D16391

AMDGPU: Add new name for barrier intrinsic
ClosedPublic

Authored by arsenm on Jan 20 2016, 6:25 PM.

Details

Reviewers
• tstellarAMD
nhaehnle

Diff Detail

Event Timeline

arsenm updated this revision to Diff 45480.Jan 20 2016, 6:25 PM
arsenm retitled this revision from to AMDGPU: Add new name for barrier intrinsic.
arsenm updated this object.
arsenm added a reviewer: • tstellarAMD.
arsenm added a subscriber: llvm-commits.
Herald added a subscriber: arsenm. · View Herald TranscriptJan 20 2016, 6:25 PM
nhaehnle accepted this revision.Jan 21 2016, 8:08 AM
nhaehnle added a reviewer: nhaehnle.
nhaehnle added a subscriber: nhaehnle.

LGTM.

This revision is now accepted and ready to land.Jan 21 2016, 8:08 AM
arsenm closed this revision.Jan 22 2016, 1:35 PM

r258558

Revision Contents

PathSize
include/
llvm/
IR/
3 lines
lib/
Target/
AMDGPU/
8 lines
test/
CodeGen/
AMDGPU/
28 lines

Diff 45480

include/llvm/IR/IntrinsicsAMDGPU.td

Show First 20 Lines • Show All 59 Lines • ▼ Show 20 Lines
def int_AMDGPU_rsq_clamped : Intrinsic< def int_AMDGPU_rsq_clamped : Intrinsic<
[llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem] [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem]
>; >;
} }
let TargetPrefix = "amdgcn" in { 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">, def int_amdgcn_div_scale : GCCBuiltin<"__builtin_amdgcn_div_scale">,
// 1st parameter: Numerator // 1st parameter: Numerator
// 2nd parameter: Denominator // 2nd parameter: Denominator
// 3rd parameter: Constant to select select between first and // 3rd parameter: Constant to select select between first and
// second. (0 = first, 1 = second). // second. (0 = first, 1 = second).
Intrinsic<[llvm_anyfloat_ty, llvm_i1_ty], Intrinsic<[llvm_anyfloat_ty, llvm_i1_ty],
[LLVMMatchType<0>, LLVMMatchType<0>, llvm_i1_ty], [LLVMMatchType<0>, LLVMMatchType<0>, llvm_i1_ty],
[IntrNoMem]>; [IntrNoMem]>;
▲ Show 20 Lines • Show All 94 Lines • Show Last 20 Lines

lib/Target/AMDGPU/SIInstructions.td

Show First 20 Lines • Show All 475 Lines • ▼ Show 20 Lines
} // End Uses = [EXEC] } // End Uses = [EXEC]
} // End isBranch = 1 } // End isBranch = 1
} // End isTerminator = 1 } // End isTerminator = 1
let hasSideEffects = 1 in { let hasSideEffects = 1 in {
def S_BARRIER : SOPP <0x0000000a, (ins), "s_barrier", def S_BARRIER : SOPP <0x0000000a, (ins), "s_barrier",
[(int_AMDGPU_barrier_local)] [(int_amdgcn_s_barrier)]
> { > {
let SchedRW = [WriteBarrier]; let SchedRW = [WriteBarrier];
let simm16 = 0; let simm16 = 0;
let mayLoad = 1; let mayLoad = 1;
let mayStore = 1; let mayStore = 1;
let isConvergent = 1; let isConvergent = 1;
} }
▲ Show 20 Lines • Show All 1,669 Lines • ▼ Show 20 Linesdef : Pat <
(i32 (addc i32:$src0, i32:$src1)), (i32 (addc i32:$src0, i32:$src1)),
(S_ADD_U32 $src0, $src1) (S_ADD_U32 $src0, $src1)
>; >;
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
// SOPP Patterns // SOPP Patterns
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
// FIXME: These should be removed eventually
def : Pat < def : Pat <
(int_AMDGPU_barrier_global), (int_AMDGPU_barrier_global),
(S_BARRIER) (S_BARRIER)
>; >;
def : Pat <
(int_AMDGPU_barrier_local),
(S_BARRIER)
>;
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
// VOP1 Patterns // VOP1 Patterns
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
let Predicates = [UnsafeFPMath] in { let Predicates = [UnsafeFPMath] in {
//def : RcpPat<V_RCP_F64_e32, f64>; //def : RcpPat<V_RCP_F64_e32, f64>;
//defm : RsqPat<V_RSQ_F64_e32, f64>; //defm : RsqPat<V_RSQ_F64_e32, f64>;
▲ Show 20 Lines • Show All 1,097 Lines • Show Last 20 Lines

test/CodeGen/AMDGPU/llvm.amdgcn.s.barrier.ll

  • This file was added.
; 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 }