replica
/
hanchenye-llvm-project
2
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity

AMDGPU: Add kernarg.segment.ptr intrinsic 

llvm-svn: 268105
This commit is contained in:
Matt Arsenault 2016-04-29 21:16:52 +00:00
parent cd62bf5821
commit dc4ebad6d4
3 changed files with 30 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

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

@ -318,6 +318,10 @@ def int_amdgcn_queue_ptr :
GCCBuiltin<"__builtin_amdgcn_queue_ptr">,
Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 2>], [], [IntrNoMem]>;
def int_amdgcn_kernarg_segment_ptr :
GCCBuiltin<"__builtin_amdgcn_kernarg_segment_ptr">,
Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 2>], [], [IntrNoMem]>;
// __builtin_amdgcn_interp_p1 <i>, <attr_chan>, <attr>, <m0>
def int_amdgcn_interp_p1 :
GCCBuiltin<"__builtin_amdgcn_interp_p1">,

5
llvm/lib/Target/AMDGPU/SIISelLowering.cpp
View File

@ -1559,6 +1559,11 @@ SDValue SITargetLowering::LowerINTRINSIC_WO_CHAIN(SDValue Op,
return CreateLiveInRegister(DAG, &AMDGPU::SReg_64RegClass,
TRI->getPreloadedValue(MF, Reg), VT);
}
case Intrinsic::amdgcn_kernarg_segment_ptr: {
unsigned Reg
= TRI->getPreloadedValue(MF, SIRegisterInfo::KERNARG_SEGMENT_PTR);
return CreateLiveInRegister(DAG, &AMDGPU::SReg_64RegClass, Reg, VT);
}
case Intrinsic::amdgcn_rcp:
return DAG.getNode(AMDGPUISD::RCP, DL, VT, Op.getOperand(1));
case Intrinsic::amdgcn_rsq:

21
llvm/test/CodeGen/AMDGPU/llvm.amdgcn.kernarg.segment.ptr.ll Normal file
View File

@ -0,0 +1,21 @@
; RUN: llc -mtriple=amdgcn--amdhsa -mcpu=kaveri -verify-machineinstrs < %s | FileCheck -check-prefix=HSA -check-prefix=ALL %s
; RUN: llc -mtriple=amdgcn-mesa-mesa3d -verify-machineinstrs < %s | FileCheck -check-prefix=MESA -check-prefix=ALL %s
; ALL-LABEL: {{^}}test:
; HSA: enable_sgpr_kernarg_segment_ptr = 1
; HSA: s_load_dword s{{[0-9]+}}, s[4:5], 0xa
; MESA: s_load_dword s{{[0-9]+}}, s[0:1], 0xa
define void @test(i32 addrspace(1)* %out) #1 {
%kernarg.segment.ptr = call noalias i8 addrspace(2)* @llvm.amdgcn.kernarg.segment.ptr()
%header.ptr = bitcast i8 addrspace(2)* %kernarg.segment.ptr to i32 addrspace(2)*
%gep = getelementptr i32, i32 addrspace(2)* %header.ptr, i64 10
%value = load i32, i32 addrspace(2)* %gep
store i32 %value, i32 addrspace(1)* %out
ret void
}
declare i8 addrspace(2)* @llvm.amdgcn.kernarg.segment.ptr() #0
attributes #0 = { nounwind readnone }
attributes #1 = { nounwind }