AMDGPU: Add llvm.amdgcn.dispatch.ptr intrinsic

[• tstellarAMD]

Summary

This returns a pointer to the dispatch packet, which can be used to load
information about the kernel dispach.

[arsenm]

include/llvm/IR/IntrinsicsAMDGPU.td

131 ↗

(On Diff #40855)

Typo: disaptch

132 ↗

(On Diff #40855)

Is there a way to mark noalias on the return value?

lib/Target/AMDGPU/AMDGPUAnnotateKernelFeatures.cpp

106–109 ↗

(On Diff #40855)

Later this should probably be cleaned up to only check the dispatch.ptr one for HSA, and the read.global.size.x should not set this

[• tstellarAMD]

Added test case, fixed typo.

[• tstellarAMD]

include/llvm/IR/IntrinsicsAMDGPU.td

132 ↗

(On Diff #40855)

I added noalias to the declaration and call in the test case and it still selects correctly. I'll look into this more when adding the clang patch for the builtin.

[arsenm]

LGTM. The intrinsic should probably have a comment that it is only for when targeting HSA

llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td

def int_amdgcn_s_dcache_wb :
   GCCBuiltin<"__builtin_amdgcn_s_dcache_wb">,
   Intrinsic<[], [], []>;
 
 // VI
 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]>;
+
 }

llvm/trunk/lib/Target/AMDGPU/AMDGPUAnnotateKernelFeatures.cpp

bool AMDGPUAnnotateKernelFeatures::runOnModule(Module &M) {
 
   static const StringRef HSAIntrinsicToAttr[][2] = {
     { "llvm.r600.read.local.size.x", "amdgpu-dispatch-ptr" },
     { "llvm.r600.read.local.size.y", "amdgpu-dispatch-ptr" },
     { "llvm.r600.read.local.size.z", "amdgpu-dispatch-ptr" },
 
     { "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.
 
   // We do not need to note the x workitem or workgroup id because they are
   // always initialized.
 
   bool Changed = addAttrsForIntrinsics(M, IntrinsicToAttr);

llvm/trunk/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp

void AMDGPUAsmPrinter::EmitAmdKernelCodeT(const MachineFunction &MF,
 
   header.compute_pgm_resource_registers =
       KernelInfo.ComputePGMRSrc1 |
       (KernelInfo.ComputePGMRSrc2 << 32);
   header.code_properties =
       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;
 
   AMDGPUTargetStreamer *TS =
       static_cast<AMDGPUTargetStreamer *>(OutStreamer->getTargetStreamer());
   TS->EmitAMDKernelCodeT(header);
 }

llvm/trunk/lib/Target/AMDGPU/SIISelLowering.cpp

unsigned ScratchPtrRegHi =
         TRI->getPhysRegSubReg(ScratchPtrReg, &AMDGPU::SReg_32RegClass, 1);
 
     CCInfo.AllocateReg(InputPtrRegLo);
     CCInfo.AllocateReg(InputPtrRegHi);
     CCInfo.AllocateReg(ScratchPtrRegLo);
     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) {
     getOriginalFunctionArgs(DAG, DAG.getMachineFunction().getFunction(), Ins,
                             Splits);
   }
 
   AnalyzeFormalArguments(CCInfo, Splits);
SDValue SITargetLowering::LowerINTRINSIC_WO_CHAIN(SDValue Op,
 
   EVT VT = Op.getValueType();
   SDLoc DL(Op);
   unsigned IntrinsicID = cast<ConstantSDNode>(Op.getOperand(0))->getZExtValue();
 
   // 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);
   case Intrinsic::r600_read_ngroups_y:
     return LowerParameter(DAG, VT, VT, DL, DAG.getEntryNode(),
                           SI::KernelInputOffsets::NGROUPS_Y, false);
   case Intrinsic::r600_read_ngroups_z:
     return LowerParameter(DAG, VT, VT, DL, DAG.getEntryNode(),

llvm/trunk/lib/Target/AMDGPU/SIRegisterInfo.h

public:
   /// \returns True if operands defined with this operand type can accept
   /// an inline constant. i.e. An integer value in the range (-16, 64) or
   /// -4.0f, -2.0f, -1.0f, -0.5f, 0.0f, 0.5f, 1.0f, 2.0f, 4.0f. 
   bool opCanUseInlineConstant(unsigned OpType) const;
 
   enum PreloadedValue {
     // SGPRS:
     SCRATCH_PTR         =  0,
+    DISPATCH_PTR        =  1,
     INPUT_PTR           =  3,
     TGID_X              = 10,
     TGID_Y              = 11,
     TGID_Z              = 12,
     SCRATCH_WAVE_OFFSET = 14,
     // VGPRS:
     FIRST_VGPR_VALUE    = 15,
     TIDIG_X             = FIRST_VGPR_VALUE,

llvm/trunk/lib/Target/AMDGPU/SIRegisterInfo.cpp

if (opCanUseLiteralConstant(OpType))
     return true;
 
   return OpType == AMDGPU::OPERAND_REG_INLINE_C;
 }
 
 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:
     return AMDGPU::SReg_32RegClass.getRegister(MFI->NumUserSGPRs + 0);
   case SIRegisterInfo::TGID_Y:
     return AMDGPU::SReg_32RegClass.getRegister(MFI->NumUserSGPRs + 1);
   case SIRegisterInfo::TGID_Z:
     return AMDGPU::SReg_32RegClass.getRegister(MFI->NumUserSGPRs + 2);
   case SIRegisterInfo::SCRATCH_WAVE_OFFSET:
     if (MFI->getShaderType() != ShaderType::COMPUTE)
       return MFI->ScratchOffsetReg;
     return AMDGPU::SReg_32RegClass.getRegister(MFI->NumUserSGPRs + 4);
   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;
   case SIRegisterInfo::TIDIG_Y:
     return AMDGPU::VGPR1;
   case SIRegisterInfo::TIDIG_Z:
     return AMDGPU::VGPR2;
   }

llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.dispatch.ptr.ll

+; 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 }