[llvm] r254116 - AMDGPU: Add llvm.amdgcn.dispatch.ptr intrinsic
Tom Stellard via llvm-commits
llvm-commits at lists.llvm.org
Wed Nov 25 16:43:30 PST 2015
Author: tstellar
Date: Wed Nov 25 18:43:29 2015
New Revision: 254116
URL: http://llvm.org/viewvc/llvm-project?rev=254116&view=rev
Log:
AMDGPU: Add llvm.amdgcn.dispatch.ptr intrinsic
Summary:
This returns a pointer to the dispatch packet, which can be used to load
information about the kernel dispach.
Reviewers: arsenm
Subscribers: arsenm, llvm-commits
Differential Revision: http://reviews.llvm.org/D14898
Added:
llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.dispatch.ptr.ll
Modified:
llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td
llvm/trunk/lib/Target/AMDGPU/AMDGPUAnnotateKernelFeatures.cpp
llvm/trunk/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp
llvm/trunk/lib/Target/AMDGPU/SIISelLowering.cpp
llvm/trunk/lib/Target/AMDGPU/SIRegisterInfo.cpp
llvm/trunk/lib/Target/AMDGPU/SIRegisterInfo.h
Modified: llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td?rev=254116&r1=254115&r2=254116&view=diff
==============================================================================
--- llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td (original)
+++ llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td Wed Nov 25 18:43:29 2015
@@ -127,4 +127,8 @@ def int_amdgcn_s_dcache_wb_vol :
GCCBuiltin<"__builtin_amdgcn_s_dcache_wb_vol">,
Intrinsic<[], [], []>;
+def int_amdgcn_dispatch_ptr :
+ GCCBuiltin<"__builtin_amdgcn_disptch_ptr">,
+ Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 2>], [], [IntrNoMem]>;
+
}
Modified: llvm/trunk/lib/Target/AMDGPU/AMDGPUAnnotateKernelFeatures.cpp
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/AMDGPU/AMDGPUAnnotateKernelFeatures.cpp?rev=254116&r1=254115&r2=254116&view=diff
==============================================================================
--- llvm/trunk/lib/Target/AMDGPU/AMDGPUAnnotateKernelFeatures.cpp (original)
+++ llvm/trunk/lib/Target/AMDGPU/AMDGPUAnnotateKernelFeatures.cpp Wed Nov 25 18:43:29 2015
@@ -105,7 +105,8 @@ bool AMDGPUAnnotateKernelFeatures::runOn
{ "llvm.r600.read.global.size.x", "amdgpu-dispatch-ptr" },
{ "llvm.r600.read.global.size.y", "amdgpu-dispatch-ptr" },
- { "llvm.r600.read.global.size.z", "amdgpu-dispatch-ptr" }
+ { "llvm.r600.read.global.size.z", "amdgpu-dispatch-ptr" },
+ { "llvm.amdgcn.dispatch.ptr", "amdgpu-dispatch-ptr" }
};
// TODO: Intrinsics that require queue ptr.
Modified: llvm/trunk/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp?rev=254116&r1=254115&r2=254116&view=diff
==============================================================================
--- llvm/trunk/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp (original)
+++ llvm/trunk/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp Wed Nov 25 18:43:29 2015
@@ -528,6 +528,9 @@ void AMDGPUAsmPrinter::EmitAmdKernelCode
AMD_CODE_PROPERTY_ENABLE_SGPR_KERNARG_SEGMENT_PTR |
AMD_CODE_PROPERTY_IS_PTR64;
+ if (MFI->hasDispatchPtr())
+ header.code_properties |= AMD_CODE_PROPERTY_ENABLE_SGPR_DISPATCH_PTR;
+
header.kernarg_segment_byte_size = MFI->ABIArgOffset;
header.wavefront_sgpr_count = KernelInfo.NumSGPR;
header.workitem_vgpr_count = KernelInfo.NumVGPR;
Modified: llvm/trunk/lib/Target/AMDGPU/SIISelLowering.cpp
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/AMDGPU/SIISelLowering.cpp?rev=254116&r1=254115&r2=254116&view=diff
==============================================================================
--- llvm/trunk/lib/Target/AMDGPU/SIISelLowering.cpp (original)
+++ llvm/trunk/lib/Target/AMDGPU/SIISelLowering.cpp Wed Nov 25 18:43:29 2015
@@ -646,6 +646,18 @@ SDValue SITargetLowering::LowerFormalArg
CCInfo.AllocateReg(ScratchPtrRegHi);
MF.addLiveIn(InputPtrReg, &AMDGPU::SReg_64RegClass);
MF.addLiveIn(ScratchPtrReg, &AMDGPU::SReg_64RegClass);
+ SIMachineFunctionInfo *MFI = MF.getInfo<SIMachineFunctionInfo>();
+ if (Subtarget->isAmdHsaOS() && MFI->hasDispatchPtr()) {
+ unsigned DispatchPtrReg =
+ TRI->getPreloadedValue(MF, SIRegisterInfo::DISPATCH_PTR);
+ unsigned DispatchPtrRegLo =
+ TRI->getPhysRegSubReg(DispatchPtrReg, &AMDGPU::SReg_32RegClass, 0);
+ unsigned DispatchPtrRegHi =
+ TRI->getPhysRegSubReg(DispatchPtrReg, &AMDGPU::SReg_32RegClass, 1);
+ CCInfo.AllocateReg(DispatchPtrRegLo);
+ CCInfo.AllocateReg(DispatchPtrRegHi);
+ MF.addLiveIn(DispatchPtrReg, &AMDGPU::SReg_64RegClass);
+ }
}
if (Info->getShaderType() == ShaderType::COMPUTE) {
@@ -1053,6 +1065,10 @@ SDValue SITargetLowering::LowerINTRINSIC
// TODO: Should this propagate fast-math-flags?
switch (IntrinsicID) {
+ case Intrinsic::amdgcn_dispatch_ptr:
+ return CreateLiveInRegister(DAG, &AMDGPU::SReg_64RegClass,
+ TRI->getPreloadedValue(MF, SIRegisterInfo::DISPATCH_PTR), VT);
+
case Intrinsic::r600_read_ngroups_x:
return LowerParameter(DAG, VT, VT, DL, DAG.getEntryNode(),
SI::KernelInputOffsets::NGROUPS_X, false);
Modified: llvm/trunk/lib/Target/AMDGPU/SIRegisterInfo.cpp
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/AMDGPU/SIRegisterInfo.cpp?rev=254116&r1=254115&r2=254116&view=diff
==============================================================================
--- llvm/trunk/lib/Target/AMDGPU/SIRegisterInfo.cpp (original)
+++ llvm/trunk/lib/Target/AMDGPU/SIRegisterInfo.cpp Wed Nov 25 18:43:29 2015
@@ -510,6 +510,7 @@ bool SIRegisterInfo::opCanUseInlineConst
unsigned SIRegisterInfo::getPreloadedValue(const MachineFunction &MF,
enum PreloadedValue Value) const {
+ const AMDGPUSubtarget &ST = MF.getSubtarget<AMDGPUSubtarget>();
const SIMachineFunctionInfo *MFI = MF.getInfo<SIMachineFunctionInfo>();
switch (Value) {
case SIRegisterInfo::TGID_X:
@@ -525,6 +526,11 @@ unsigned SIRegisterInfo::getPreloadedVal
case SIRegisterInfo::SCRATCH_PTR:
return AMDGPU::SGPR2_SGPR3;
case SIRegisterInfo::INPUT_PTR:
+ if (ST.isAmdHsaOS())
+ return MFI->hasDispatchPtr() ? AMDGPU::SGPR2_SGPR3 : AMDGPU::SGPR0_SGPR1;
+ return AMDGPU::SGPR0_SGPR1;
+ case SIRegisterInfo::DISPATCH_PTR:
+ assert(MFI->hasDispatchPtr());
return AMDGPU::SGPR0_SGPR1;
case SIRegisterInfo::TIDIG_X:
return AMDGPU::VGPR0;
Modified: llvm/trunk/lib/Target/AMDGPU/SIRegisterInfo.h
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/AMDGPU/SIRegisterInfo.h?rev=254116&r1=254115&r2=254116&view=diff
==============================================================================
--- llvm/trunk/lib/Target/AMDGPU/SIRegisterInfo.h (original)
+++ llvm/trunk/lib/Target/AMDGPU/SIRegisterInfo.h Wed Nov 25 18:43:29 2015
@@ -99,6 +99,7 @@ public:
enum PreloadedValue {
// SGPRS:
SCRATCH_PTR = 0,
+ DISPATCH_PTR = 1,
INPUT_PTR = 3,
TGID_X = 10,
TGID_Y = 11,
Added: llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.dispatch.ptr.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.dispatch.ptr.ll?rev=254116&view=auto
==============================================================================
--- llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.dispatch.ptr.ll (added)
+++ llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.dispatch.ptr.ll Wed Nov 25 18:43:29 2015
@@ -0,0 +1,16 @@
+; RUN: llc -mtriple=amdgcn--amdhsa -mcpu=kaveri -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s
+
+; GCN-LABEL: {{^}}test:
+; GCN: enable_sgpr_dispatch_ptr = 1
+; GCN: s_load_dword s{{[0-9]+}}, s[0:1], 0x0
+define void @test(i32 addrspace(1)* %out) {
+ %dispatch_ptr = call noalias i8 addrspace(2)* @llvm.amdgcn.dispatch.ptr() #0
+ %header_ptr = bitcast i8 addrspace(2)* %dispatch_ptr to i32 addrspace(2)*
+ %value = load i32, i32 addrspace(2)* %header_ptr
+ store i32 %value, i32 addrspace(1)* %out
+ ret void
+}
+
+declare noalias i8 addrspace(2)* @llvm.amdgcn.dispatch.ptr() #0
+
+attributes #0 = { readnone }
More information about the llvm-commits
mailing list