[AMDGPU] Add agpr_count to metadata and AsmParser

[lamb-j]

Summary

gfx90a allows the number of ACC registers (AGPRs) to be set
independently to the VGPR registers. For both HSA and PAL metadata, we
now include an "agpr_count" key to report the number of AGPRs set for
supported devices (gfx90a, gfx908, as determined by hasMAIInsts()).
This is collected from SIProgramInfo.NumAccVGPR for both HSA and PAL.
The AsmParser also now recognizes ".kernel.agpr_count" for supported
devices.

[kzhuravl]

Add @tpr to look at pal metadata changes

[arsenm]

llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp

1158

Formatting and weird consta placement

1162

Ditto

1179

Fallthrough here for other targets

[lamb-j]

Removing const from MCSymbol variable and creating fall through case for getTotalNumVgprs()

[lamb-j]

llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp

1158

The const does seem to be unnecessary. Should I remove from lines 1134 and 1144 as well, or leave that for another patch?

[arsenm]

llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp

1168

Really should avoid repeating this then. Usually things that need to be shared between codegen and MC go in AMDGPUBaseInfo

[lamb-j]

llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp

1168

Ok perfect, I was looking for a good way to do this. So would you recommend lifting the implementation of getTotalNumVGPRs from AMDGPUResourceUsageAnalysis.cpp into Utils/AMDGPUBaseInfo.cpp, possibly as a separate patch?

Moreover, there’s actually already a function defined as “getTotalNumVGPRs” in Utils/AMDGPUBaseInfo.cpp, but its implementation is different than what we’re looking for in the AsmParser. Here are the two overlapping functions.

AMDGPUResourceUsageAnalysis.cpp
Returns a variable value that represents the number of VGPRs and AGPRs employed by a code region.

int32_t AMDGPUResourceUsageAnalysis::SIFunctionResourceInfo::getTotalNumVGPRs(
const GCNSubtarget &ST, int32_t ArgNumAGPR, int32_t ArgNumVGPR) const {
  if (ST.hasGFX90AInsts() && ArgNumAGPR)
    return alignTo(ArgNumVGPR, 4) + ArgNumAGPR;
  return std::max(ArgNumVGPR, ArgNumAGPR);
}

Utils/AMDGPUBaseInfo.cpp
Returns a constant value indicating the total number of physical VGPRs accessible on a device.

unsigned getTotalNumVGPRs(const MCSubtargetInfo *STI) {
  if (STI->getFeatureBits().test(FeatureGFX90AInsts))
    return 512;
  if (!isGFX10Plus(*STI))
    return 256;
  return STI->getFeatureBits().test(FeatureWavefrontSize32) ? 1024 : 512;
}

Are the overlapping names here intentional? Would it make sense to rename one of these in a separate patch, for example the AMDGPUResourceUsageAnalysis one to “getTotalNumUsedVGPRs” or something similar?

[lamb-j]

Fixing MCSymbol const declaration formatting

[lamb-j]

Switching to use new getTotalNumVGPRs from AMDGPUBaseInfo

[lamb-j]

Updating commit message

llvm/docs/AMDGPUUsage.rst

".vgpr_count"                       integer        Required  Number of vector
                                                                   registers required by
                                                                   each work-item for
                                                                   GFX6-GFX9. A register
                                                                   is required if it is
                                                                   used explicitly, or
                                                                   if a higher numbered
                                                                   register is used
                                                                   explicitly.
+     ".agpr_count"                       integer        Required  Number of accumulator
+                                                                  registers required by
+                                                                  each work-item for
+                                                                  GFX90A, GFX908.
      ".max_flat_workgroup_size"          integer        Required  Maximum flat
                                                                   work-group size
                                                                   supported by the
                                                                   kernel in work-items.
                                                                   Must be >=1 and
                                                                   consistent with
                                                                   ReqdWorkGroupSize if
                                                                   not 0, 0, 0.
.. table:: AMDPAL Code Object Hardware Stage Metadata Map
      ========================== ============== ========= ===============================================================
      String Key                 Value Type     Required? Description
      ========================== ============== ========= ===============================================================
      ".entry_point"             string                   The ELF symbol pointing to this pipeline's stage entry point.
      ".scratch_memory_size"     integer                  Scratch memory size in bytes.
      ".lds_size"                integer                  Local Data Share size in bytes.
      ".perf_data_buffer_size"   integer                  Performance data buffer size in bytes.
      ".vgpr_count"              integer                  Number of VGPRs used.
+     ".agpr_count"              integer                  Number of AGPRs used.
      ".sgpr_count"              integer                  Number of SGPRs used.
      ".vgpr_limit"              integer                  If non-zero, indicates the shader was compiled with a
                                                          directive to instruct the compiler to limit the VGPR usage to
                                                          be less than or equal to the specified value (only set if
                                                          different from HW default).
      ".sgpr_limit"              integer                  SGPR count upper limit (only set if different from HW
                                                          default).
      ".threadgroup_dimensions"  sequence of              Thread-group X/Y/Z dimensions (Compute only).

llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp

 void AMDGPUAsmPrinter::EmitPALMetadata(const MachineFunction &MF,
        const SIProgramInfo &CurrentProgramInfo) {
   const SIMachineFunctionInfo *MFI = MF.getInfo<SIMachineFunctionInfo>();
   auto CC = MF.getFunction().getCallingConv();
   auto MD = getTargetStreamer()->getPALMetadata();
 
   MD->setEntryPoint(CC, MF.getFunction().getName());
   MD->setNumUsedVgprs(CC, CurrentProgramInfo.NumVGPRsForWavesPerEU);
+
+  // Only set AGPRs for supported devices
+  const GCNSubtarget &STM = MF.getSubtarget<GCNSubtarget>();
+  if (STM.hasMAIInsts()) {
+    MD->setNumUsedAgprs(CC, CurrentProgramInfo.NumAccVGPR);
+  }
+
   MD->setNumUsedSgprs(CC, CurrentProgramInfo.NumSGPRsForWavesPerEU);
   MD->setRsrc1(CC, CurrentProgramInfo.getPGMRSrc1(CC));
   if (AMDGPU::isCompute(CC)) {
     MD->setRsrc2(CC, CurrentProgramInfo.ComputePGMRSrc2);
   } else {
     if (CurrentProgramInfo.ScratchBlocks > 0)
       MD->setRsrc2(CC, S_00B84C_SCRATCH_EN(1));
   }
   // ScratchSize is in bytes, 16 aligned.
   MD->setScratchSize(CC, alignTo(CurrentProgramInfo.ScratchSize, 16));
   if (MF.getFunction().getCallingConv() == CallingConv::AMDGPU_PS) {
     MD->setRsrc2(CC, S_00B02C_EXTRA_LDS_SIZE(CurrentProgramInfo.LDSBlocks));
     MD->setSpiPsInputEna(MFI->getPSInputEnable());
     MD->setSpiPsInputAddr(MFI->getPSInputAddr());
   }
 
-  const GCNSubtarget &STM = MF.getSubtarget<GCNSubtarget>();
   if (STM.isWave32())
     MD->setWave32(MF.getFunction().getCallingConv());
 }
 
 void AMDGPUAsmPrinter::emitPALFunctionMetadata(const MachineFunction &MF) {
   auto *MD = getTargetStreamer()->getPALMetadata();
   const MachineFrameInfo &MFI = MF.getFrameInfo();
   MD->setFunctionScratchSize(MF, MFI.getStackSize());

llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp

MetadataStreamerV3::getHSAKernelProps(const MachineFunction &MF,
 
   // 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);
+
+  // Only add AGPR count to metadata for supported devices
+  if (STM.hasMAIInsts()) {
+    Kern[".agpr_count"] = Kern.getDocument()->getNode(ProgramInfo.NumAccVGPR);
+  }
+
   Kern[".max_flat_workgroup_size"] =
       Kern.getDocument()->getNode(MFI.getMaxFlatWorkGroupSize());
   Kern[".sgpr_spill_count"] =
       Kern.getDocument()->getNode(MFI.getNumSpilledSGPRs());
   Kern[".vgpr_spill_count"] =
       Kern.getDocument()->getNode(MFI.getNumSpilledVGPRs());
 
   return Kern;

llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp

 //===----------------------------------------------------------------------===//
 
 // Holds info related to the current kernel, e.g. count of SGPRs used.
 // Kernel scope begins at .amdgpu_hsa_kernel directive, ends at next
 // .amdgpu_hsa_kernel or at EOF.
 class KernelScopeInfo {
   int SgprIndexUnusedMin = -1;
   int VgprIndexUnusedMin = -1;
+  int AgprIndexUnusedMin = -1;
   MCContext *Ctx = nullptr;
+  MCSubtargetInfo const *MSTI = nullptr;
 
   void usesSgprAt(int i) {
     if (i >= SgprIndexUnusedMin) {
       SgprIndexUnusedMin = ++i;
       if (Ctx) {
         MCSymbol * const Sym = Ctx->getOrCreateSymbol(Twine(".kernel.sgpr_count"));
         Sym->setVariableValue(MCConstantExpr::create(SgprIndexUnusedMin, *Ctx));
       }
     }
   }
 
   void usesVgprAt(int i) {
     if (i >= VgprIndexUnusedMin) {
       VgprIndexUnusedMin = ++i;
       if (Ctx) {
         MCSymbol * const Sym = Ctx->getOrCreateSymbol(Twine(".kernel.vgpr_count"));
-        Sym->setVariableValue(MCConstantExpr::create(VgprIndexUnusedMin, *Ctx));
+        Sym->setVariableValue(MCConstantExpr::create(getTotalNumVgprs(), *Ctx));
       }
     }
   }
 
+  void usesAgprAt(int i) {
+    // Instruction will be flagged/error in AMDGPUAsmParser::MatchAndEmitInstruction
+    if (!hasMAIInsts(*MSTI))
+      return;
+
+    if (i >= AgprIndexUnusedMin) {
+      AgprIndexUnusedMin = ++i;
+      if (Ctx) {
+        MCSymbol * const Sym = Ctx->getOrCreateSymbol(Twine(".kernel.agpr_count"));
+        Sym->setVariableValue(MCConstantExpr::create(AgprIndexUnusedMin, *Ctx));
+
+        // Also update vgpr_count, which is dependent on agpr_count (on gfx908, gfx90a)
+        MCSymbol * const vSym = Ctx->getOrCreateSymbol(Twine(".kernel.vgpr_count"));
+        vSym->setVariableValue(MCConstantExpr::create(getTotalNumVgprs(), *Ctx));
+      }
+    }
+  }
+
+  // Sourced from AMDGPUResourceUsageAnalysis::SIFunctionResourceInfo::getTotalNumVGPRs()
+  int getTotalNumVgprs() {
+    // Default
+    if (!hasMAIInsts(*MSTI))
+      return VgprIndexUnusedMin;
+    // 908, vgpr_count == max(used_vgprs, used_agprs)
+    if (!isGFX90A(*MSTI))
+      return std::max(VgprIndexUnusedMin, AgprIndexUnusedMin);
+    // 90a, vgpr_count == sum(used_vgprs, used_agprs)
+    if (isGFX90A(*MSTI))
+      return alignTo(VgprIndexUnusedMin, 4) + AgprIndexUnusedMin;
+  }
+
 public:
   KernelScopeInfo() = default;
 
   void initialize(MCContext &Context) {
     Ctx = &Context;
+    MSTI = Ctx->getSubtargetInfo();
+
     usesSgprAt(SgprIndexUnusedMin = -1);
     usesVgprAt(VgprIndexUnusedMin = -1);
+    if (hasMAIInsts(*MSTI)) {
+      usesAgprAt(AgprIndexUnusedMin = -1);
+    }
   }
 
   void usesRegister(RegisterKind RegKind, unsigned DwordRegIndex, unsigned RegWidth) {
     switch (RegKind) {
       case IS_SGPR: usesSgprAt(DwordRegIndex + RegWidth - 1); break;
-      case IS_AGPR: // fall through
+      case IS_AGPR: usesAgprAt(DwordRegIndex + RegWidth - 1); break;
       case IS_VGPR: usesVgprAt(DwordRegIndex + RegWidth - 1); break;
       default: break;
     }
   }
 };
 
 class AMDGPUAsmParser : public MCTargetAsmParser {
   MCAsmParser &Parser;

llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h

 bool isGFX10(const MCSubtargetInfo &STI);
 bool isGFX10Plus(const MCSubtargetInfo &STI);
 bool isGCN3Encoding(const MCSubtargetInfo &STI);
 bool isGFX10_AEncoding(const MCSubtargetInfo &STI);
 bool isGFX10_BEncoding(const MCSubtargetInfo &STI);
 bool hasGFX10_3Insts(const MCSubtargetInfo &STI);
 bool isGFX90A(const MCSubtargetInfo &STI);
 bool hasArchitectedFlatScratch(const MCSubtargetInfo &STI);
+bool hasMAIInsts(const MCSubtargetInfo &STI);
 
 /// Is Reg - scalar register
 bool isSGPR(unsigned Reg, const MCRegisterInfo* TRI);
 
 /// Is there any intersection between registers
 bool isRegIntersect(unsigned Reg0, unsigned Reg1, const MCRegisterInfo* TRI);
 
 /// If \p Reg is a pseudo reg, return the correct hardware register given

llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp

 bool isGFX90A(const MCSubtargetInfo &STI) {
   return STI.getFeatureBits()[AMDGPU::FeatureGFX90AInsts];
 }
 
 bool hasArchitectedFlatScratch(const MCSubtargetInfo &STI) {
   return STI.getFeatureBits()[AMDGPU::FeatureArchitectedFlatScratch];
 }
 
+bool hasMAIInsts(const MCSubtargetInfo &STI) {
+  return STI.getFeatureBits()[AMDGPU::FeatureMAIInsts];
+}
+
 bool isSGPR(unsigned Reg, const MCRegisterInfo* TRI) {
   const MCRegisterClass SGPRClass = TRI->getRegClass(AMDGPU::SReg_32RegClassID);
   const unsigned FirstSubReg = TRI->getSubReg(Reg, AMDGPU::sub0);
   return SGPRClass.contains(FirstSubReg != 0 ? FirstSubReg : Reg) ||
     Reg == AMDGPU::SCC;
 }
 
 bool isRegIntersect(unsigned Reg0, unsigned Reg1, const MCRegisterInfo* TRI) {

llvm/lib/Target/AMDGPU/Utils/AMDGPUPALMetadata.h

public:
   // Set the entry point name for one shader.
   void setEntryPoint(unsigned CC, StringRef Name);
 
   // Set the number of used vgprs in the metadata. This is an optional advisory
   // record for logging etc; wave dispatch actually uses the rsrc1 register for
   // the shader stage to determine the number of vgprs to allocate.
   void setNumUsedVgprs(unsigned CC, unsigned Val);
 
+  // Set the number of used agprs in the metadata. This is an optional advisory
+  // record for logging etc;
+  void setNumUsedAgprs(unsigned CC, unsigned Val);
+
   // Set the number of used sgprs in the metadata. This is an optional advisory
   // record for logging etc; wave dispatch actually uses the rsrc1 register for
   // the shader stage to determine the number of sgprs to allocate.
   void setNumUsedSgprs(unsigned CC, unsigned Val);
 
   // Set the scratch size in the metadata.
   void setScratchSize(unsigned CC, unsigned Val);
 

llvm/lib/Target/AMDGPU/Utils/AMDGPUPALMetadata.cpp

unsigned NumUsedVgprsKey = getScratchSizeKey(CC) +
                                PALMD::Key::VS_SCRATCH_SIZE;
     setRegister(NumUsedVgprsKey, Val);
     return;
   }
   // Msgpack format.
   getHwStage(CC)[".vgpr_count"] = MsgPackDoc.getNode(Val);
 }
 
+// Set the number of used agprs in the metadata.
+void AMDGPUPALMetadata::setNumUsedAgprs(CallingConv::ID CC, unsigned Val) {
+  getHwStage(CC)[".agpr_count"] = Val;
+}
+
 // Set the number of used sgprs in the metadata. This is an optional advisory
 // record for logging etc; wave dispatch actually uses the rsrc1 register for
 // the shader stage to determine the number of sgprs to allocate.
 void AMDGPUPALMetadata::setNumUsedSgprs(CallingConv::ID CC, unsigned Val) {
   if (isLegacy()) {
     // Old non-msgpack format.
     unsigned NumUsedSgprsKey = getScratchSizeKey(CC) +
                                PALMD::Key::VS_NUM_USED_SGPRS -

llvm/test/CodeGen/AMDGPU/amdpal-metadata-agpr-register-count.ll

+; RUN: llc -mtriple=amdgcn--amdpal -mcpu=gfx90a < %s | FileCheck -check-prefixes=CHECK,GFX90A %s
+; RUN: llc -mtriple=amdgcn--amdpal -mcpu=gfx908 < %s | FileCheck -check-prefixes=CHECK,GFX908 %s
+
+; COM: Adapted from agpr-register-count.ll
+; COM: GFX900 and below should not have .agpr_count present in the metadata
+
+
+; CHECK:      .type          kernel_32_agprs
+; CHECK:      NumAgprs:       32
+define amdgpu_kernel void @kernel_32_agprs() #0 {
+bb:
+  call void asm sideeffect "", "~{v8}" ()
+  call void asm sideeffect "", "~{a31}" ()
+  ret void
+}
+
+; CHECK:      .type          kernel_0_agprs
+; CHECK:      NumAgprs:       0
+define amdgpu_kernel void @kernel_0_agprs() #0 {
+bb:
+  call void asm sideeffect "", "~{v0}" ()
+  ret void
+}
+
+; CHECK:      .type           kernel_40_vgprs
+; CHECK:      NumAgprs:       16
+define amdgpu_kernel void @kernel_40_vgprs() #0 {
+bb:
+  call void asm sideeffect "", "~{v39}" ()
+  call void asm sideeffect "", "~{a15}" ()
+  ret void
+}
+
+; CHECK:      .type          kernel_max_gprs
+; CHECK:      NumAgprs:       256
+define amdgpu_kernel void @kernel_max_gprs() #0 {
+bb:
+  call void asm sideeffect "", "~{v255}" ()
+  call void asm sideeffect "", "~{a255}" ()
+  ret void
+}
+
+; CHECK:      .type          func_32_agprs
+; CHECK:      NumAgprs:       32
+define void @func_32_agprs() #0 {
+bb:
+  call void asm sideeffect "", "~{v8}" ()
+  call void asm sideeffect "", "~{a31}" ()
+  ret void
+}
+
+; CHECK:      .type          kernel_call_func_32_agprs
+; CHECK:      NumAgprs:       32
+define amdgpu_kernel void @kernel_call_func_32_agprs() #0 {
+bb:
+  call void @func_32_agprs() #0
+  ret void
+}
+
+declare void @undef_func()
+
+; CHECK:      .type          kernel_call_undef_func
+; CHECK:      NumAgprs:       32
+define amdgpu_kernel void @kernel_call_undef_func() #0 {
+bb:
+  call void @undef_func()
+  ret void
+}
+
+; CHECK: ---
+; CHECK:  amdpal.pipelines:
+; GFX90A: agpr_count:  0x20
+; GFX90A: vgpr_count:  0x40
+
+; GFX908: agpr_count:  0x20
+; GFX908: vgpr_count:  0x20
+
+attributes #0 = { nounwind noinline "amdgpu-flat-work-group-size"="1,512" }

llvm/test/CodeGen/AMDGPU/hsa-metadata-agpr-register-count.ll

+; RUN: llc -mtriple=amdgcn--amdhsa -mcpu=gfx90a -verify-machineinstrs < %s | FileCheck -check-prefixes=CHECK,GFX90A %s
+; RUN: llc -mtriple=amdgcn--amdhsa -mcpu=gfx908 -verify-machineinstrs < %s | FileCheck -check-prefixes=CHECK,GFX908 %s
+; RUN: llc -mtriple=amdgcn--amdhsa -mcpu=gfx801 -verify-machineinstrs < %s | FileCheck -check-prefixes=CHECK,GFX801 %s
+
+; COM: Adapted from agpr-register-count.ll
+; COM: GFX900 and below should not have .agpr_count present in the metadata
+
+; CHECK: ---
+; CHECK:  amdhsa.kernels:
+
+; GFX90A:    - .agpr_count:    32
+; GFX908:    - .agpr_count:    32
+; GFX801-NOT:    - .agpr_count:
+; CHECK:      .name:          kernel_32_agprs
+; GFX90A:      .vgpr_count:    44
+; GFX908:      .vgpr_count:    32
+; GFX801:      .vgpr_count:    9
+define amdgpu_kernel void @kernel_32_agprs() #0 {
+bb:
+  call void asm sideeffect "", "~{v8}" ()
+  call void asm sideeffect "", "~{a31}" ()
+  ret void
+}
+
+; GFX90A:    - .agpr_count:    0
+; GFX908:    - .agpr_count:    0
+; GFX801-NOT:    - .agpr_count:
+; CHECK:      .name:          kernel_0_agprs
+; GFX90A:      .vgpr_count:    1
+; GFX908:      .vgpr_count:    1
+; GFX801:      .vgpr_count:    1
+define amdgpu_kernel void @kernel_0_agprs() #0 {
+bb:
+  call void asm sideeffect "", "~{v0}" ()
+  ret void
+}
+
+; GFX90A:    - .agpr_count:    16
+; GFX908:    - .agpr_count:    16
+; GFX801-NOT:    - .agpr_count:
+; CHECK:      .name:          kernel_40_vgprs
+; GFX90A:      .vgpr_count:    56
+; GFX908:      .vgpr_count:    40
+; GFX801:      .vgpr_count:    40
+define amdgpu_kernel void @kernel_40_vgprs() #0 {
+bb:
+  call void asm sideeffect "", "~{v39}" ()
+  call void asm sideeffect "", "~{a15}" ()
+  ret void
+}
+
+; GFX90A:    - .agpr_count:    256
+; GFX908:    - .agpr_count:    256
+; GFX801-NOT:    - .agpr_count:
+; CHECK:      .name:          kernel_max_gprs
+; GFX90A:      .vgpr_count:    512
+; GFX908:      .vgpr_count:    256
+; GFX801:      .vgpr_count:    256
+define amdgpu_kernel void @kernel_max_gprs() #0 {
+bb:
+  call void asm sideeffect "", "~{v255}" ()
+  call void asm sideeffect "", "~{a255}" ()
+  ret void
+}
+
+define void @func_32_agprs() #0 {
+bb:
+  call void asm sideeffect "", "~{v8}" ()
+  call void asm sideeffect "", "~{a31}" ()
+  ret void
+}
+
+; GFX90A:    - .agpr_count:    32
+; GFX908:    - .agpr_count:    32
+; GFX801-NOT:    - .agpr_count:
+; CHECK:      .name:          kernel_call_func_32_agprs
+; GFX90A:      .vgpr_count:    44
+; GFX908:      .vgpr_count:    32
+; GFX801:      .vgpr_count:    9
+define amdgpu_kernel void @kernel_call_func_32_agprs() #0 {
+bb:
+  call void @func_32_agprs() #0
+  ret void
+}
+
+declare void @undef_func()
+
+; GFX90A:    - .agpr_count:    32
+; GFX908:    - .agpr_count:    32
+; GFX801-NOT:    - .agpr_count:
+; CHECK:      .name:          kernel_call_undef_func
+; GFX90A:      .vgpr_count:    64
+; GFX908:      .vgpr_count:    32
+; GFX801:      .vgpr_count:    32
+define amdgpu_kernel void @kernel_call_undef_func() #0 {
+bb:
+  call void @undef_func()
+  ret void
+}
+
+attributes #0 = { nounwind noinline "amdgpu-flat-work-group-size"="1,512" }

llvm/test/CodeGen/AMDGPU/hsa-metadata-agpr-small.ll

+; RUN: llc -mtriple=amdgcn--amdhsa -mcpu=gfx908 < %s | FileCheck -check-prefixes=CHECK,GFX908 %s
+; RUN: llc -mtriple=amdgcn--amdhsa -mcpu=gfx90a < %s | FileCheck -check-prefixes=CHECK,GFX90A %s
+; RUN: llc -mtriple=amdgcn--amdhsa -mcpu=gfx801 < %s | FileCheck -check-prefixes=CHECK,GFX801 %s
+
+; COM: Comments for each kernel
+; CHECK: kernel_32_agprs
+; GFX908:   ; NumVgprs: 9
+; GFX908    ; NumAgprs: 32
+; GFX908    ; TotalNumVgprs: 32
+
+; GFX90A:   ; NumVgprs: 9
+; GFX90A    ; NumAgprs: 32
+; GFX90A    ; TotalNumVgprs: 44
+
+; GFX801:   ; NumVgprs: 9
+
+; CHECK: kernel_40_vgprs
+; GFX908:   ; NumVgprs: 40
+; GFX908    ; NumAgprs: 16
+; GFX908    ; TotalNumVgprs: 40
+
+; GFX90A:   ; NumVgprs: 40
+; GFX90A    ; NumAgprs: 16
+; GFX90A    ; TotalNumVgprs: 56
+
+; GFX801:   ; NumVgprs: 40
+
+; COM: Metadata
+; GFX908:    - .agpr_count:    32
+; GFX908:      .vgpr_count:    32
+
+; GFX90A:    - .agpr_count:    32
+; GFX90A:      .vgpr_count:    44
+
+; GFX801:      .vgpr_count:    9
+define amdgpu_kernel void @kernel_32_agprs() #0 {
+bb:
+  call void asm sideeffect "", "~{v8}" ()
+  call void asm sideeffect "", "~{a31}" ()
+  ret void
+}
+
+; GFX908:    - .agpr_count:    16
+; GFX908:      .vgpr_count:    40
+
+; GFX90A:    - .agpr_count:    16
+; GFX90A:      .vgpr_count:    56
+
+; GFX801:      .vgpr_count:    40
+define amdgpu_kernel void @kernel_40_vgprs() #0 {
+bb:
+  call void asm sideeffect "", "~{v39}" ()
+  call void asm sideeffect "", "~{a15}" ()
+  ret void
+}
+
+attributes #0 = { nounwind noinline "amdgpu-flat-work-group-size"="1,512" }

llvm/test/MC/AMDGPU/sym_kernel_scope_agpr.s

+// RUN: llvm-mc -arch=amdgcn -mcpu=gfx90a %s 2>&1 | FileCheck -check-prefixes=GFX90A %s
+// RUN: llvm-mc -arch=amdgcn -mcpu=gfx908 %s 2>&1 | FileCheck -check-prefixes=GFX908 %s
+// Based on sym_kernel_scope.s
+
+.byte .kernel.agpr_count
+// CHECK: .byte 0
+.byte .kernel.vgpr_count
+// CHECK: .byte 0
+
+    v_accvgpr_write_b32 a0, v6
+    v_accvgpr_read_b32 v3, a3
+    s_endpgm
+.byte .kernel.agpr_count
+// GFX90A: .byte 4
+// GFX908: .byte 4
+.byte .kernel.vgpr_count
+// GFX90A: .byte 12
+// GFX908: .byte 7
+
+.amdgpu_hsa_kernel K1
+K1:
+.byte .kernel.agpr_count
+// CHECK: .byte 0
+.byte .kernel.vgpr_count
+// CHECK: .byte 0
+    v_accvgpr_write_b32 a44, v6
+    s_endpgm
+.byte .kernel.agpr_count
+// GFX90A: .byte 45
+// GFX908: .byte 45
+.byte .kernel.vgpr_count
+// GFX90A: .byte 53
+// GFX908: .byte 45
+
+.amdgpu_hsa_kernel K2
+.byte .kernel.agpr_count
+// CHECK: .byte 0
+.byte .kernel.vgpr_count
+// CHECK: .byte 0
+K2:
+    v_mfma_f32_4x4x1f32 a[0:3], v1, v0, a[0:3] cbsz:1 abid:2 blgp:3
+    s_endpgm
+.byte .kernel.agpr_count
+// GFX90A: .byte 4
+// GFX908: .byte 4
+.byte .kernel.vgpr_count
+// GFX90A: .byte 8
+// GFX908: .byte 4
+
+.text
+.amdgpu_hsa_kernel K3
+K3:
+    v_accvgpr_read_b32 v[0], a0
+    v_mfma_f32_16x16x1f32 a[0:15], v1, v0, a[0:15] cbsz:1 abid:2 blgp:3
+    s_endpgm
+
+.byte .kernel.agpr_count
+// GFX90A: .byte 16
+// GFX908: .byte 16
+.byte .kernel.vgpr_count
+// GFX90A: .byte 20
+// GFX908: .byte 16