AMDGPU: Account for implicit argument alignment for kernarg segment

[arsenm]

Summary

If a kernel had no formal arguments but did have the implicit
arguments, we were reporting a required kernarg alignment of 4. For
some reason we require an 8-byte alignment for this, even though
there's no real advantage and I don't see where this is documented in
the ABI.

The code object header code also claims the minimum alignment is 16,
which is what I thought you always got at runtime anyway so I don't
know why this matters.

[arsenm]

ping

[arsenm]

ping

[kzhuravl]

LGTM

The metadata treats the minimum as 16?

I found some internal docs that say the minimum must be 16, but those docs describe the time when kernarg segment alignment was a part of kernel descriptor. So I believe metadata just inherited that. I will find out if it should stay the same in metadata and update the amdgpu usage accordingly. thanks.

[arsenm]

90ff14871904881fb156a1d4d5fb083ca75998ab

llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp

void AMDGPUAsmPrinter::getAmdKernelCode(amd_kernel_code_t &Out,
   Out.kernarg_segment_byte_size = STM.getKernArgSegmentSize(F, MaxKernArgAlign);
   Out.wavefront_sgpr_count = CurrentProgramInfo.NumSGPR;
   Out.workitem_vgpr_count = CurrentProgramInfo.NumVGPR;
   Out.workitem_private_segment_byte_size = CurrentProgramInfo.ScratchSize;
   Out.workgroup_group_segment_byte_size = CurrentProgramInfo.LDSSize;
 
   // kernarg_segment_alignment is specified as log of the alignment.
   // The minimum alignment is 16.
+  // FIXME: The metadata treats the minimum as 4?
   Out.kernarg_segment_alignment = Log2(std::max(Align(16), MaxKernArgAlign));
 }
 
 bool AMDGPUAsmPrinter::PrintAsmOperand(const MachineInstr *MI, unsigned OpNo,
                                        const char *ExtraCode, raw_ostream &O) {
   // First try the generic code, which knows about modifiers like 'c' and 'n'.
   if (!AsmPrinter::PrintAsmOperand(MI, OpNo, ExtraCode, O))
     return false;

llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp

MetadataStreamerV2::getHSACodeProps(const MachineFunction &MF,
   const Function &F = MF.getFunction();
 
   assert(F.getCallingConv() == CallingConv::AMDGPU_KERNEL ||
          F.getCallingConv() == CallingConv::SPIR_KERNEL);
 
   Align MaxKernArgAlign;
   HSACodeProps.mKernargSegmentSize = STM.getKernArgSegmentSize(F,
                                                                MaxKernArgAlign);
-  HSACodeProps.mGroupSegmentFixedSize = ProgramInfo.LDSSize;
-  HSACodeProps.mPrivateSegmentFixedSize = ProgramInfo.ScratchSize;
   HSACodeProps.mKernargSegmentAlign =
-      std::max(MaxKernArgAlign, Align(4)).value();
+
+  HSACodeProps.mGroupSegmentFixedSize = ProgramInfo.LDSSize;
+  HSACodeProps.mPrivateSegmentFixedSize = ProgramInfo.ScratchSize;
   HSACodeProps.mWavefrontSize = STM.getWavefrontSize();
   HSACodeProps.mNumSGPRs = ProgramInfo.NumSGPR;
   HSACodeProps.mNumVGPRs = ProgramInfo.NumVGPR;
   HSACodeProps.mMaxFlatWorkGroupSize = MFI.getMaxFlatWorkGroupSize();
   HSACodeProps.mIsDynamicCallStack = ProgramInfo.DynamicCallStack;
   HSACodeProps.mIsXNACKEnabled = STM.isXNACKEnabled();
   HSACodeProps.mNumSpilledSGPRs = MFI.getNumSpilledSGPRs();
   HSACodeProps.mNumSpilledVGPRs = MFI.getNumSpilledVGPRs();
MetadataStreamerV3::getHSAKernelProps(const MachineFunction &MF,
 
   Align MaxKernArgAlign;
   Kern[".kernarg_segment_size"] = Kern.getDocument()->getNode(
       STM.getKernArgSegmentSize(F, MaxKernArgAlign));
   Kern[".group_segment_fixed_size"] =
       Kern.getDocument()->getNode(ProgramInfo.LDSSize);
   Kern[".private_segment_fixed_size"] =
       Kern.getDocument()->getNode(ProgramInfo.ScratchSize);
+
+  // FIXME: The metadata treats the minimum as 16?
   Kern[".kernarg_segment_align"] =
       Kern.getDocument()->getNode(std::max(Align(4), MaxKernArgAlign).value());
   Kern[".wavefront_size"] =
       Kern.getDocument()->getNode(STM.getWavefrontSize());
   Kern[".sgpr_count"] = Kern.getDocument()->getNode(ProgramInfo.NumSGPR);
   Kern[".vgpr_count"] = Kern.getDocument()->getNode(ProgramInfo.NumVGPR);
   Kern[".max_flat_workgroup_size"] =
       Kern.getDocument()->getNode(MFI.getMaxFlatWorkGroupSize());

llvm/lib/Target/AMDGPU/AMDGPUSubtarget.cpp

unsigned AMDGPUSubtarget::getKernArgSegmentSize(const Function &F,
 
   unsigned ExplicitOffset = getExplicitKernelArgOffset(F);
 
   uint64_t TotalSize = ExplicitOffset + ExplicitArgBytes;
   unsigned ImplicitBytes = getImplicitArgNumBytes(F);
   if (ImplicitBytes != 0) {
     const Align Alignment = getAlignmentForImplicitArgPtr();
     TotalSize = alignTo(ExplicitArgBytes, Alignment) + ImplicitBytes;
+    MaxAlign = std::max(MaxAlign, Alignment);
   }
 
   // Being able to dereference past the end is useful for emitting scalar loads.
   return alignTo(TotalSize, 4);
 }
 
 AMDGPUDwarfFlavour AMDGPUSubtarget::getAMDGPUDwarfFlavour() const {
   return getWavefrontSize() == 32 ? AMDGPUDwarfFlavour::Wave32

llvm/test/CodeGen/AMDGPU/llvm.amdgcn.implicitarg.ptr.ll

define amdgpu_kernel void @kernel_implicitarg_no_struct_align_padding(<16 x i32>, i32) #1 {
   ret void
 }
 
 
 ; HSA-LABEL: Kernels:
 ; HSA-LABEL: - Name:            kernel_implicitarg_ptr_empty
 ; HSA: CodeProps:
 ; HSA: KernargSegmentSize: 56
-; HSA: KernargSegmentAlign: 4
+; HSA: KernargSegmentAlign: 8
 
 ; HSA-LABEL: - Name:            kernel_implicitarg_ptr_empty_0implicit
 ; HSA: KernargSegmentSize: 0
 ; HSA: KernargSegmentAlign: 4
 
 ; HSA-LABEL: - Name:            opencl_kernel_implicitarg_ptr_empty
 ; HSA: KernargSegmentSize: 48
-; HSA: KernargSegmentAlign: 4
+; HSA: KernargSegmentAlign: 8
 
 ; HSA-LABEL: - Name:            kernel_implicitarg_ptr
 ; HSA: KernargSegmentSize: 168
-; HSA: KernargSegmentAlign: 4
+; HSA: KernargSegmentAlign: 8
 
 ; HSA-LABEL: - Name:            opencl_kernel_implicitarg_ptr
 ; HSA: KernargSegmentSize: 160
-; HSA: KernargSegmentAlign: 4
+; HSA: KernargSegmentAlign: 8
 
 ; HSA-LABEL: - Name:            kernel_call_implicitarg_ptr_func_empty
 ; HSA: KernargSegmentSize: 56
-; HSA: KernargSegmentAlign: 4
+; HSA: KernargSegmentAlign: 8
 
 ; HSA-LABEL: - Name:            kernel_call_implicitarg_ptr_func_empty_implicit0
 ; HSA: KernargSegmentSize: 0
 ; HSA: KernargSegmentAlign: 4
 
 ; HSA-LABEL:  - Name:            opencl_kernel_call_implicitarg_ptr_func_empty
 ; HSA: KernargSegmentSize: 48
-; HSA: KernargSegmentAlign: 4
+; HSA: KernargSegmentAlign: 8
 
 ; HSA-LABEL:  - Name:            kernel_call_implicitarg_ptr_func
 ; HSA: KernargSegmentSize: 168
-; HSA: KernargSegmentAlign: 4
+; HSA: KernargSegmentAlign: 8
 
 ; HSA-LABEL:  - Name:            opencl_kernel_call_implicitarg_ptr_func
 ; HSA: KernargSegmentSize: 160
-; HSA: KernargSegmentAlign: 4
+; HSA: KernargSegmentAlign: 8
 
 ; HSA-LABEL: - Name:            kernel_call_kernarg_implicitarg_ptr_func
 ; HSA: KernargSegmentSize: 168
-; HSA: KernargSegmentAlign: 4
+; HSA: KernargSegmentAlign: 8
 
 ; HSA-LABEL: - Name:            kernel_implicitarg_no_struct_align_padding
 ; HSA: KernargSegmentSize: 120
 ; HSA: KernargSegmentAlign: 64
 
 declare i8 addrspace(4)* @llvm.amdgcn.implicitarg.ptr() #2
 declare i8 addrspace(4)* @llvm.amdgcn.kernarg.segment.ptr() #2
 
 attributes #0 = { nounwind noinline }
 attributes #1 = { nounwind noinline "amdgpu-implicitarg-num-bytes"="48" }
 attributes #2 = { nounwind readnone speculatable }
 attributes #3 = { nounwind noinline "amdgpu-implicitarg-num-bytes"="0" }