[llvm] [AMDGPU] MCExpr-ify AMDGPU HSAMetadata (PR #94788)
Janek van Oirschot via llvm-commits
llvm-commits at lists.llvm.org
Fri Jun 7 11:49:50 PDT 2024
https://github.com/JanekvO created https://github.com/llvm/llvm-project/pull/94788
Do note that this PR should be considered a stacked PR and depends on #93236
The only commit to review is https://github.com/llvm/llvm-project/commit/73814b6b2d7c552f7b00a7dbcd5c7b1fd72f4bd5
Enables MCExpr for HSAMetadata, particularly msgpack format.
>From f3656539152691294233ad20ca27d97b49d571cb Mon Sep 17 00:00:00 2001
From: Janek van Oirschot <janek.vanoirschot at amd.com>
Date: Thu, 23 May 2024 12:34:37 -0700
Subject: [PATCH 1/6] MCExpr-ify AMDGPU PALMetadata
---
llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp | 71 +++--
.../AMDGPU/Utils/AMDGPUDelayedMCExpr.cpp | 64 +++++
.../Target/AMDGPU/Utils/AMDGPUDelayedMCExpr.h | 39 +++
.../Target/AMDGPU/Utils/AMDGPUPALMetadata.cpp | 113 ++++++++
.../Target/AMDGPU/Utils/AMDGPUPALMetadata.h | 24 ++
llvm/lib/Target/AMDGPU/Utils/CMakeLists.txt | 1 +
llvm/test/CodeGen/AMDGPU/amdpal-es.ll | 1 +
llvm/test/CodeGen/AMDGPU/amdpal-gs.ll | 1 +
llvm/test/CodeGen/AMDGPU/amdpal-hs.ll | 1 +
llvm/test/CodeGen/AMDGPU/amdpal-ls.ll | 1 +
llvm/test/CodeGen/AMDGPU/amdpal-vs.ll | 1 +
llvm/unittests/Target/AMDGPU/CMakeLists.txt | 1 +
llvm/unittests/Target/AMDGPU/PALMetadata.cpp | 245 ++++++++++++++++++
13 files changed, 543 insertions(+), 20 deletions(-)
create mode 100644 llvm/lib/Target/AMDGPU/Utils/AMDGPUDelayedMCExpr.cpp
create mode 100644 llvm/lib/Target/AMDGPU/Utils/AMDGPUDelayedMCExpr.h
create mode 100644 llvm/unittests/Target/AMDGPU/PALMetadata.cpp
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp b/llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp
index cad4a3430327b..f4028adc84828 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp
@@ -1194,6 +1194,30 @@ static void EmitPALMetadataCommon(AMDGPUPALMetadata *MD,
getLdsDwGranularity(ST) * sizeof(uint32_t)));
}
+static constexpr std::pair<unsigned, unsigned> getShiftMask(unsigned Value) {
+ unsigned Shift = 0;
+ unsigned Mask = 0;
+
+ Mask = ~Value;
+ for (; !(Mask & 1); Shift++, Mask >>= 1) {
+ }
+
+ return std::make_pair(Shift, Mask);
+}
+
+static const MCExpr *MaskShiftSet(const MCExpr *Val, uint32_t Mask,
+ uint32_t Shift, MCContext &Ctx) {
+ if (Mask) {
+ const MCExpr *MaskExpr = MCConstantExpr::create(Mask, Ctx);
+ Val = MCBinaryExpr::createAnd(Val, MaskExpr, Ctx);
+ }
+ if (Shift) {
+ const MCExpr *ShiftExpr = MCConstantExpr::create(Shift, Ctx);
+ Val = MCBinaryExpr::createShl(Val, ShiftExpr, Ctx);
+ }
+ return Val;
+}
+
// This is the equivalent of EmitProgramInfoSI above, but for when the OS type
// is AMDPAL. It stores each compute/SPI register setting and other PAL
// metadata items into the PALMD::Metadata, combining with any provided by the
@@ -1207,41 +1231,49 @@ void AMDGPUAsmPrinter::EmitPALMetadata(const MachineFunction &MF,
auto &Ctx = MF.getContext();
MD->setEntryPoint(CC, MF.getFunction().getName());
- MD->setNumUsedVgprs(
- CC, getMCExprValue(CurrentProgramInfo.NumVGPRsForWavesPerEU, Ctx));
+ MD->setNumUsedVgprs(CC, CurrentProgramInfo.NumVGPRsForWavesPerEU, Ctx);
// Only set AGPRs for supported devices
const GCNSubtarget &STM = MF.getSubtarget<GCNSubtarget>();
if (STM.hasMAIInsts()) {
- MD->setNumUsedAgprs(CC, getMCExprValue(CurrentProgramInfo.NumAccVGPR, Ctx));
+ MD->setNumUsedAgprs(CC, CurrentProgramInfo.NumAccVGPR);
}
- MD->setNumUsedSgprs(
- CC, getMCExprValue(CurrentProgramInfo.NumSGPRsForWavesPerEU, Ctx));
+ MD->setNumUsedSgprs(CC, CurrentProgramInfo.NumSGPRsForWavesPerEU, Ctx);
if (MD->getPALMajorVersion() < 3) {
- MD->setRsrc1(CC, CurrentProgramInfo.getPGMRSrc1(CC, STM));
+ MD->setRsrc1(CC, CurrentProgramInfo.getPGMRSrc1(CC, STM, Ctx), Ctx);
if (AMDGPU::isCompute(CC)) {
- MD->setRsrc2(CC, CurrentProgramInfo.getComputePGMRSrc2());
+ MD->setRsrc2(CC, CurrentProgramInfo.getComputePGMRSrc2(Ctx), Ctx);
} else {
- if (getMCExprValue(CurrentProgramInfo.ScratchBlocks, Ctx) > 0)
- MD->setRsrc2(CC, S_00B84C_SCRATCH_EN(1));
+ const MCExpr *HasScratchBlocks =
+ MCBinaryExpr::createGT(CurrentProgramInfo.ScratchBlocks,
+ MCConstantExpr::create(0, Ctx), Ctx);
+ auto [Shift, Mask] = getShiftMask(C_00B84C_SCRATCH_EN);
+ MD->setRsrc2(CC, MaskShiftSet(HasScratchBlocks, Mask, Shift, Ctx), Ctx);
}
} else {
MD->setHwStage(CC, ".debug_mode", (bool)CurrentProgramInfo.DebugMode);
- MD->setHwStage(CC, ".scratch_en",
- (bool)getMCExprValue(CurrentProgramInfo.ScratchEnable, Ctx));
+ MD->setHwStage(CC, ".scratch_en", msgpack::Type::Boolean,
+ CurrentProgramInfo.ScratchEnable);
EmitPALMetadataCommon(MD, CurrentProgramInfo, CC, STM);
}
// ScratchSize is in bytes, 16 aligned.
MD->setScratchSize(
- CC, alignTo(getMCExprValue(CurrentProgramInfo.ScratchSize, Ctx), 16));
+ CC,
+ AMDGPUVariadicMCExpr::createAlignTo(CurrentProgramInfo.ScratchSize,
+ MCConstantExpr::create(16, Ctx), Ctx),
+ Ctx);
+
if (MF.getFunction().getCallingConv() == CallingConv::AMDGPU_PS) {
unsigned ExtraLDSSize = STM.getGeneration() >= AMDGPUSubtarget::GFX11
? divideCeil(CurrentProgramInfo.LDSBlocks, 2)
: CurrentProgramInfo.LDSBlocks;
if (MD->getPALMajorVersion() < 3) {
- MD->setRsrc2(CC, S_00B02C_EXTRA_LDS_SIZE(ExtraLDSSize));
+ MD->setRsrc2(
+ CC,
+ MCConstantExpr::create(S_00B02C_EXTRA_LDS_SIZE(ExtraLDSSize), Ctx),
+ Ctx);
MD->setSpiPsInputEna(MFI->getPSInputEnable());
MD->setSpiPsInputAddr(MFI->getPSInputAddr());
} else {
@@ -1288,20 +1320,19 @@ void AMDGPUAsmPrinter::emitPALFunctionMetadata(const MachineFunction &MF) {
if (MD->getPALMajorVersion() < 3) {
// Set compute registers
- MD->setRsrc1(CallingConv::AMDGPU_CS,
- CurrentProgramInfo.getPGMRSrc1(CallingConv::AMDGPU_CS, ST));
+ MD->setRsrc1(
+ CallingConv::AMDGPU_CS,
+ CurrentProgramInfo.getPGMRSrc1(CallingConv::AMDGPU_CS, ST, Ctx), Ctx);
MD->setRsrc2(CallingConv::AMDGPU_CS,
- CurrentProgramInfo.getComputePGMRSrc2());
+ CurrentProgramInfo.getComputePGMRSrc2(Ctx), Ctx);
} else {
EmitPALMetadataCommon(MD, CurrentProgramInfo, CallingConv::AMDGPU_CS, ST);
}
// Set optional info
MD->setFunctionLdsSize(FnName, CurrentProgramInfo.LDSSize);
- MD->setFunctionNumUsedVgprs(
- FnName, getMCExprValue(CurrentProgramInfo.NumVGPRsForWavesPerEU, Ctx));
- MD->setFunctionNumUsedSgprs(
- FnName, getMCExprValue(CurrentProgramInfo.NumSGPRsForWavesPerEU, Ctx));
+ MD->setFunctionNumUsedVgprs(FnName, CurrentProgramInfo.NumVGPRsForWavesPerEU);
+ MD->setFunctionNumUsedSgprs(FnName, CurrentProgramInfo.NumSGPRsForWavesPerEU);
}
// This is supposed to be log2(Size)
diff --git a/llvm/lib/Target/AMDGPU/Utils/AMDGPUDelayedMCExpr.cpp b/llvm/lib/Target/AMDGPU/Utils/AMDGPUDelayedMCExpr.cpp
new file mode 100644
index 0000000000000..3955f557b9a25
--- /dev/null
+++ b/llvm/lib/Target/AMDGPU/Utils/AMDGPUDelayedMCExpr.cpp
@@ -0,0 +1,64 @@
+//===- AMDGPUDelayedMCExpr.cpp - Delayed MCExpr resolve ---------*- C++ -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include "AMDGPUDelayedMCExpr.h"
+#include "llvm/MC/MCExpr.h"
+#include "llvm/MC/MCValue.h"
+
+using namespace llvm;
+
+static msgpack::DocNode getNode(msgpack::DocNode DN, msgpack::Type Type,
+ MCValue Val) {
+ msgpack::Document *Doc = DN.getDocument();
+ switch (Type) {
+ default:
+ return Doc->getEmptyNode();
+ case msgpack::Type::Int:
+ return Doc->getNode(static_cast<int64_t>(Val.getConstant()));
+ case msgpack::Type::UInt:
+ return Doc->getNode(static_cast<uint64_t>(Val.getConstant()));
+ case msgpack::Type::Boolean:
+ return Doc->getNode(static_cast<bool>(Val.getConstant()));
+ }
+}
+
+void DelayedMCExpr::AssignDocNode(msgpack::DocNode &DN, msgpack::Type Type,
+ const MCExpr *Expr) {
+ MCValue Res;
+ if (Expr->evaluateAsRelocatable(Res, nullptr, nullptr)) {
+ if (Res.isAbsolute()) {
+ DN = getNode(DN, Type, Res);
+ return;
+ }
+ }
+
+ DelayedExprs.push_back(DelayedExpr{DN, Type, Expr});
+}
+
+bool DelayedMCExpr::ResolveDelayedExpressions() {
+ bool Success;
+
+ while (!DelayedExprs.empty()) {
+ DelayedExpr DE = DelayedExprs.front();
+ MCValue Res;
+
+ Success = DE.Expr->evaluateAsRelocatable(Res, nullptr, nullptr);
+ Success &= Res.isAbsolute();
+ if (!Success)
+ return false;
+
+ DelayedExprs.pop_front();
+ DE.DN = getNode(DE.DN, DE.Type, Res);
+ }
+
+ return true;
+}
+
+void DelayedMCExpr::clear() { DelayedExprs.clear(); }
+
+bool DelayedMCExpr::empty() { return DelayedExprs.empty(); }
diff --git a/llvm/lib/Target/AMDGPU/Utils/AMDGPUDelayedMCExpr.h b/llvm/lib/Target/AMDGPU/Utils/AMDGPUDelayedMCExpr.h
new file mode 100644
index 0000000000000..c546660a0d996
--- /dev/null
+++ b/llvm/lib/Target/AMDGPU/Utils/AMDGPUDelayedMCExpr.h
@@ -0,0 +1,39 @@
+//===- AMDGPUDelayedMCExpr.h - Delayed MCExpr resolve -----------*- C++ -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_LIB_TARGET_AMDGPU_UTILS_AMDGPUDELAYEDMCEXPR_H
+#define LLVM_LIB_TARGET_AMDGPU_UTILS_AMDGPUDELAYEDMCEXPR_H
+
+#include "llvm/BinaryFormat/MsgPackDocument.h"
+#include <deque>
+
+namespace llvm {
+class MCExpr;
+
+class DelayedMCExpr {
+ struct DelayedExpr {
+ msgpack::DocNode &DN;
+ msgpack::Type Type;
+ const MCExpr *Expr;
+ DelayedExpr(msgpack::DocNode &DN, msgpack::Type Type, const MCExpr *Expr)
+ : DN(DN), Type(Type), Expr(Expr) {}
+ };
+
+ std::deque<DelayedExpr> DelayedExprs;
+
+public:
+ bool ResolveDelayedExpressions();
+ void AssignDocNode(msgpack::DocNode &DN, msgpack::Type Type,
+ const MCExpr *Expr);
+ void clear();
+ bool empty();
+};
+
+} // end namespace llvm
+
+#endif // LLVM_LIB_TARGET_AMDGPU_UTILS_AMDGPUDELAYEDMCEXPR_H
diff --git a/llvm/lib/Target/AMDGPU/Utils/AMDGPUPALMetadata.cpp b/llvm/lib/Target/AMDGPU/Utils/AMDGPUPALMetadata.cpp
index 0fa67c559cb29..4597dab142470 100644
--- a/llvm/lib/Target/AMDGPU/Utils/AMDGPUPALMetadata.cpp
+++ b/llvm/lib/Target/AMDGPU/Utils/AMDGPUPALMetadata.cpp
@@ -20,6 +20,7 @@
#include "llvm/BinaryFormat/ELF.h"
#include "llvm/IR/Constants.h"
#include "llvm/IR/Module.h"
+#include "llvm/MC/MCExpr.h"
#include "llvm/Support/AMDGPUMetadata.h"
#include "llvm/Support/EndianStream.h"
@@ -137,12 +138,22 @@ void AMDGPUPALMetadata::setRsrc1(CallingConv::ID CC, unsigned Val) {
setRegister(getRsrc1Reg(CC), Val);
}
+void AMDGPUPALMetadata::setRsrc1(CallingConv::ID CC, const MCExpr *Val,
+ MCContext &Ctx) {
+ setRegister(getRsrc1Reg(CC), Val, Ctx);
+}
+
// Set the rsrc2 register in the metadata for a particular shader stage.
// In fact this ORs the value into any previous setting of the register.
void AMDGPUPALMetadata::setRsrc2(CallingConv::ID CC, unsigned Val) {
setRegister(getRsrc1Reg(CC) + 1, Val);
}
+void AMDGPUPALMetadata::setRsrc2(CallingConv::ID CC, const MCExpr *Val,
+ MCContext &Ctx) {
+ setRegister(getRsrc1Reg(CC) + 1, Val, Ctx);
+}
+
// Set the SPI_PS_INPUT_ENA register in the metadata.
// In fact this ORs the value into any previous setting of the register.
void AMDGPUPALMetadata::setSpiPsInputEna(unsigned Val) {
@@ -182,6 +193,40 @@ void AMDGPUPALMetadata::setRegister(unsigned Reg, unsigned Val) {
N = N.getDocument()->getNode(Val);
}
+// Set a register in the metadata.
+// In fact this ORs the value into any previous setting of the register.
+void AMDGPUPALMetadata::setRegister(unsigned Reg, const MCExpr *Val,
+ MCContext &Ctx) {
+ if (!isLegacy()) {
+ // In the new MsgPack format, ignore register numbered >= 0x10000000. It
+ // is a PAL ABI pseudo-register in the old non-MsgPack format.
+ if (Reg >= 0x10000000)
+ return;
+ }
+ auto &N = getRegisters()[MsgPackDoc.getNode(Reg)];
+ bool RegSeenInREM = REM.find(Reg) != REM.end();
+
+ if (RegSeenInREM) {
+ Val = MCBinaryExpr::createOr(Val, REM[Reg], Ctx);
+ // This conditional may be redundant most of the time, but
+ // setRegister(unsigned, unsigned) could've been called while RegSeenInREM
+ // is true.
+ if (N.getKind() == msgpack::Type::UInt) {
+ const MCExpr *NExpr = MCConstantExpr::create(N.getUInt(), Ctx);
+ Val = MCBinaryExpr::createOr(Val, NExpr, Ctx);
+ }
+ REM[Reg] = Val;
+ } else if (N.getKind() == msgpack::Type::UInt) {
+ const MCExpr *NExpr = MCConstantExpr::create(N.getUInt(), Ctx);
+ Val = MCBinaryExpr::createOr(Val, NExpr, Ctx);
+ int64_t Unused;
+ if (!Val->evaluateAsAbsolute(Unused))
+ REM[Reg] = Val;
+ (void)Unused;
+ }
+ DelayedExprs.AssignDocNode(N, msgpack::Type::UInt, Val);
+}
+
// Set the entry point name for one shader.
void AMDGPUPALMetadata::setEntryPoint(unsigned CC, StringRef Name) {
if (isLegacy())
@@ -207,11 +252,29 @@ void AMDGPUPALMetadata::setNumUsedVgprs(CallingConv::ID CC, unsigned Val) {
getHwStage(CC)[".vgpr_count"] = MsgPackDoc.getNode(Val);
}
+void AMDGPUPALMetadata::setNumUsedVgprs(CallingConv::ID CC, const MCExpr *Val,
+ MCContext &Ctx) {
+ if (isLegacy()) {
+ // Old non-msgpack format.
+ unsigned NumUsedVgprsKey = getScratchSizeKey(CC) +
+ PALMD::Key::VS_NUM_USED_VGPRS -
+ PALMD::Key::VS_SCRATCH_SIZE;
+ setRegister(NumUsedVgprsKey, Val, Ctx);
+ return;
+ }
+ // Msgpack format.
+ setHwStage(CC, ".vgpr_count", msgpack::Type::UInt, Val);
+}
+
// Set the number of used agprs in the metadata.
void AMDGPUPALMetadata::setNumUsedAgprs(CallingConv::ID CC, unsigned Val) {
getHwStage(CC)[".agpr_count"] = Val;
}
+void AMDGPUPALMetadata::setNumUsedAgprs(unsigned CC, const MCExpr *Val) {
+ setHwStage(CC, ".agpr_count", msgpack::Type::UInt, 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.
@@ -228,6 +291,20 @@ void AMDGPUPALMetadata::setNumUsedSgprs(CallingConv::ID CC, unsigned Val) {
getHwStage(CC)[".sgpr_count"] = MsgPackDoc.getNode(Val);
}
+void AMDGPUPALMetadata::setNumUsedSgprs(unsigned CC, const MCExpr *Val,
+ MCContext &Ctx) {
+ if (isLegacy()) {
+ // Old non-msgpack format.
+ unsigned NumUsedSgprsKey = getScratchSizeKey(CC) +
+ PALMD::Key::VS_NUM_USED_SGPRS -
+ PALMD::Key::VS_SCRATCH_SIZE;
+ setRegister(NumUsedSgprsKey, Val, Ctx);
+ return;
+ }
+ // Msgpack format.
+ setHwStage(CC, ".sgpr_count", msgpack::Type::UInt, Val);
+}
+
// Set the scratch size in the metadata.
void AMDGPUPALMetadata::setScratchSize(CallingConv::ID CC, unsigned Val) {
if (isLegacy()) {
@@ -239,6 +316,17 @@ void AMDGPUPALMetadata::setScratchSize(CallingConv::ID CC, unsigned Val) {
getHwStage(CC)[".scratch_memory_size"] = MsgPackDoc.getNode(Val);
}
+void AMDGPUPALMetadata::setScratchSize(unsigned CC, const MCExpr *Val,
+ MCContext &Ctx) {
+ if (isLegacy()) {
+ // Old non-msgpack format.
+ setRegister(getScratchSizeKey(CC), Val, Ctx);
+ return;
+ }
+ // Msgpack format.
+ setHwStage(CC, ".scratch_memory_size", msgpack::Type::UInt, Val);
+}
+
// Set the stack frame size of a function in the metadata.
void AMDGPUPALMetadata::setFunctionScratchSize(StringRef FnName, unsigned Val) {
auto Node = getShaderFunction(FnName);
@@ -259,6 +347,12 @@ void AMDGPUPALMetadata::setFunctionNumUsedVgprs(StringRef FnName,
Node[".vgpr_count"] = MsgPackDoc.getNode(Val);
}
+void AMDGPUPALMetadata::setFunctionNumUsedVgprs(StringRef FnName,
+ const MCExpr *Val) {
+ auto Node = getShaderFunction(FnName);
+ DelayedExprs.AssignDocNode(Node[".vgpr_count"], msgpack::Type::UInt, Val);
+}
+
// Set the number of used vgprs in the metadata.
void AMDGPUPALMetadata::setFunctionNumUsedSgprs(StringRef FnName,
unsigned Val) {
@@ -266,6 +360,12 @@ void AMDGPUPALMetadata::setFunctionNumUsedSgprs(StringRef FnName,
Node[".sgpr_count"] = MsgPackDoc.getNode(Val);
}
+void AMDGPUPALMetadata::setFunctionNumUsedSgprs(StringRef FnName,
+ const MCExpr *Val) {
+ auto Node = getShaderFunction(FnName);
+ DelayedExprs.AssignDocNode(Node[".sgpr_count"], msgpack::Type::UInt, Val);
+}
+
// Set the hardware register bit in PAL metadata to enable wave32 on the
// shader of the given calling convention.
void AMDGPUPALMetadata::setWave32(unsigned CC) {
@@ -662,6 +762,7 @@ void AMDGPUPALMetadata::toString(std::string &String) {
String.clear();
if (!BlobType)
return;
+ ResolvedAll = DelayedExprs.ResolveDelayedExpressions();
raw_string_ostream Stream(String);
if (isLegacy()) {
if (MsgPackDoc.getRoot().getKind() == msgpack::Type::Nil)
@@ -711,6 +812,7 @@ void AMDGPUPALMetadata::toString(std::string &String) {
// a .note record of the specified AMD type. Returns an empty blob if
// there is no PAL metadata,
void AMDGPUPALMetadata::toBlob(unsigned Type, std::string &Blob) {
+ ResolvedAll = DelayedExprs.ResolveDelayedExpressions();
if (Type == ELF::NT_AMD_PAL_METADATA)
toLegacyBlob(Blob);
else if (Type)
@@ -906,11 +1008,17 @@ void AMDGPUPALMetadata::setLegacy() {
// Erase all PAL metadata.
void AMDGPUPALMetadata::reset() {
MsgPackDoc.clear();
+ REM.clear();
+ DelayedExprs.clear();
Registers = MsgPackDoc.getEmptyNode();
HwStages = MsgPackDoc.getEmptyNode();
ShaderFunctions = MsgPackDoc.getEmptyNode();
}
+bool AMDGPUPALMetadata::resolvedAllMCExpr() {
+ return ResolvedAll && DelayedExprs.empty();
+}
+
unsigned AMDGPUPALMetadata::getPALVersion(unsigned idx) {
assert(idx < 2 &&
"illegal index to PAL version - should be 0 (major) or 1 (minor)");
@@ -942,6 +1050,11 @@ void AMDGPUPALMetadata::setHwStage(unsigned CC, StringRef field, bool Val) {
getHwStage(CC)[field] = Val;
}
+void AMDGPUPALMetadata::setHwStage(unsigned CC, StringRef field,
+ msgpack::Type Type, const MCExpr *Val) {
+ DelayedExprs.AssignDocNode(getHwStage(CC)[field], Type, Val);
+}
+
void AMDGPUPALMetadata::setComputeRegisters(StringRef field, unsigned Val) {
getComputeRegisters()[field] = Val;
}
diff --git a/llvm/lib/Target/AMDGPU/Utils/AMDGPUPALMetadata.h b/llvm/lib/Target/AMDGPU/Utils/AMDGPUPALMetadata.h
index 158f766d04854..1dcdd4b985142 100644
--- a/llvm/lib/Target/AMDGPU/Utils/AMDGPUPALMetadata.h
+++ b/llvm/lib/Target/AMDGPU/Utils/AMDGPUPALMetadata.h
@@ -13,7 +13,10 @@
#ifndef LLVM_LIB_TARGET_AMDGPU_AMDGPUPALMETADATA_H
#define LLVM_LIB_TARGET_AMDGPU_AMDGPUPALMETADATA_H
+#include "AMDGPUDelayedMCExpr.h"
+#include "llvm/ADT/DenseMap.h"
#include "llvm/BinaryFormat/MsgPackDocument.h"
+#include "llvm/MC/MCContext.h"
namespace llvm {
@@ -21,6 +24,10 @@ class Module;
class StringRef;
class AMDGPUPALMetadata {
+public:
+ using RegisterExprMap = DenseMap<unsigned, const MCExpr *>;
+
+private:
unsigned BlobType = 0;
msgpack::Document MsgPackDoc;
msgpack::DocNode Registers;
@@ -32,6 +39,10 @@ class AMDGPUPALMetadata {
msgpack::DocNode ComputeRegisters;
msgpack::DocNode GraphicsRegisters;
+ DelayedMCExpr DelayedExprs;
+ RegisterExprMap REM;
+ bool ResolvedAll = true;
+
public:
// Read the amdgpu.pal.metadata supplied by the frontend, ready for
// per-function modification.
@@ -45,10 +56,12 @@ class AMDGPUPALMetadata {
// Set the rsrc1 register in the metadata for a particular shader stage.
// In fact this ORs the value into any previous setting of the register.
void setRsrc1(unsigned CC, unsigned Val);
+ void setRsrc1(unsigned CC, const MCExpr *Val, MCContext &Ctx);
// Set the rsrc2 register in the metadata for a particular shader stage.
// In fact this ORs the value into any previous setting of the register.
void setRsrc2(unsigned CC, unsigned Val);
+ void setRsrc2(unsigned CC, const MCExpr *Val, MCContext &Ctx);
// Set the SPI_PS_INPUT_ENA register in the metadata.
// In fact this ORs the value into any previous setting of the register.
@@ -64,6 +77,7 @@ class AMDGPUPALMetadata {
// Set a register in the metadata.
// In fact this ORs the value into any previous setting of the register.
void setRegister(unsigned Reg, unsigned Val);
+ void setRegister(unsigned Reg, const MCExpr *Val, MCContext &Ctx);
// Set the entry point name for one shader.
void setEntryPoint(unsigned CC, StringRef Name);
@@ -72,18 +86,22 @@ class AMDGPUPALMetadata {
// 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);
+ void setNumUsedVgprs(unsigned CC, const MCExpr *Val, MCContext &Ctx);
// Set the number of used agprs in the metadata. This is an optional advisory
// record for logging etc;
void setNumUsedAgprs(unsigned CC, unsigned Val);
+ void setNumUsedAgprs(unsigned CC, const MCExpr *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);
+ void setNumUsedSgprs(unsigned CC, const MCExpr *Val, MCContext &Ctx);
// Set the scratch size in the metadata.
void setScratchSize(unsigned CC, unsigned Val);
+ void setScratchSize(unsigned CC, const MCExpr *Val, MCContext &Ctx);
// Set the stack frame size of a function in the metadata.
void setFunctionScratchSize(StringRef FnName, unsigned Val);
@@ -97,11 +115,13 @@ class AMDGPUPALMetadata {
// record for logging etc; wave dispatch actually uses the rsrc1 register for
// the shader stage to determine the number of vgprs to allocate.
void setFunctionNumUsedVgprs(StringRef FnName, unsigned Val);
+ void setFunctionNumUsedVgprs(StringRef FnName, const MCExpr *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 setFunctionNumUsedSgprs(StringRef FnName, unsigned Val);
+ void setFunctionNumUsedSgprs(StringRef FnName, const MCExpr *Val);
// Set the hardware register bit in PAL metadata to enable wave32 on the
// shader of the given calling convention.
@@ -138,6 +158,8 @@ class AMDGPUPALMetadata {
void setHwStage(unsigned CC, StringRef field, unsigned Val);
void setHwStage(unsigned CC, StringRef field, bool Val);
+ void setHwStage(unsigned CC, StringRef field, msgpack::Type Type,
+ const MCExpr *Val);
void setComputeRegisters(StringRef field, unsigned Val);
void setComputeRegisters(StringRef field, bool Val);
@@ -156,6 +178,8 @@ class AMDGPUPALMetadata {
// Erase all PAL metadata.
void reset();
+ bool resolvedAllMCExpr();
+
private:
// Return whether the blob type is legacy PAL metadata.
bool isLegacy() const;
diff --git a/llvm/lib/Target/AMDGPU/Utils/CMakeLists.txt b/llvm/lib/Target/AMDGPU/Utils/CMakeLists.txt
index 2f4ce8eaf1d60..09b8da9f5dd48 100644
--- a/llvm/lib/Target/AMDGPU/Utils/CMakeLists.txt
+++ b/llvm/lib/Target/AMDGPU/Utils/CMakeLists.txt
@@ -1,6 +1,7 @@
add_llvm_component_library(LLVMAMDGPUUtils
AMDGPUAsmUtils.cpp
AMDGPUBaseInfo.cpp
+ AMDGPUDelayedMCExpr.cpp
AMDGPUMemoryUtils.cpp
AMDGPUPALMetadata.cpp
AMDKernelCodeTUtils.cpp
diff --git a/llvm/test/CodeGen/AMDGPU/amdpal-es.ll b/llvm/test/CodeGen/AMDGPU/amdpal-es.ll
index 679e0858819eb..657fe80be04da 100644
--- a/llvm/test/CodeGen/AMDGPU/amdpal-es.ll
+++ b/llvm/test/CodeGen/AMDGPU/amdpal-es.ll
@@ -11,6 +11,7 @@
; GCN-NEXT: .scratch_memory_size: 0
; GCN: .registers:
; GCN-NEXT: '0x2cca (SPI_SHADER_PGM_RSRC1_ES)': 0
+; GCN-NEXT: '0x2ccb (SPI_SHADER_PGM_RSRC2_ES)': 0
; GCN-NEXT: ...
; GCN-NEXT: .end_amdgpu_pal_metadata
define amdgpu_es half @es_amdpal(half %arg0) {
diff --git a/llvm/test/CodeGen/AMDGPU/amdpal-gs.ll b/llvm/test/CodeGen/AMDGPU/amdpal-gs.ll
index 75f7a1dc266d3..9f5eb3927c489 100644
--- a/llvm/test/CodeGen/AMDGPU/amdpal-gs.ll
+++ b/llvm/test/CodeGen/AMDGPU/amdpal-gs.ll
@@ -12,6 +12,7 @@
; GCN-NEXT: .scratch_memory_size: 0
; GCN: .registers:
; GCN-NEXT: '0x2c8a (SPI_SHADER_PGM_RSRC1_GS)': 0
+; GCN-NEXT: '0x2c8b (SPI_SHADER_PGM_RSRC2_GS)': 0
; GCN-NEXT: ...
; GCN-NEXT: .end_amdgpu_pal_metadata
define amdgpu_gs half @gs_amdpal(half %arg0) {
diff --git a/llvm/test/CodeGen/AMDGPU/amdpal-hs.ll b/llvm/test/CodeGen/AMDGPU/amdpal-hs.ll
index c61578a967b62..7eacedf44d09d 100644
--- a/llvm/test/CodeGen/AMDGPU/amdpal-hs.ll
+++ b/llvm/test/CodeGen/AMDGPU/amdpal-hs.ll
@@ -12,6 +12,7 @@
; GCN-NEXT: .scratch_memory_size: 0
; GCN: .registers:
; GCN-NEXT: '0x2d0a (SPI_SHADER_PGM_RSRC1_HS)': 0
+; GCN-NEXT: '0x2d0b (SPI_SHADER_PGM_RSRC2_HS)': 0
; GCN-NEXT: ...
; GCN-NEXT: .end_amdgpu_pal_metadata
define amdgpu_hs half @hs_amdpal(half %arg0) {
diff --git a/llvm/test/CodeGen/AMDGPU/amdpal-ls.ll b/llvm/test/CodeGen/AMDGPU/amdpal-ls.ll
index 8162c824dc2ce..973eb561a9a3d 100644
--- a/llvm/test/CodeGen/AMDGPU/amdpal-ls.ll
+++ b/llvm/test/CodeGen/AMDGPU/amdpal-ls.ll
@@ -11,6 +11,7 @@
; GCN-NEXT: .scratch_memory_size: 0
; GCN: .registers:
; GCN-NEXT: '0x2d4a (SPI_SHADER_PGM_RSRC1_LS)': 0
+; GCN-NEXT: '0x2d4b (SPI_SHADER_PGM_RSRC2_LS)': 0
; GCN-NEXT: ...
; GCN-NEXT: .end_amdgpu_pal_metadata
define amdgpu_ls half @ls_amdpal(half %arg0) {
diff --git a/llvm/test/CodeGen/AMDGPU/amdpal-vs.ll b/llvm/test/CodeGen/AMDGPU/amdpal-vs.ll
index c300ba187740c..e554bb8980cec 100644
--- a/llvm/test/CodeGen/AMDGPU/amdpal-vs.ll
+++ b/llvm/test/CodeGen/AMDGPU/amdpal-vs.ll
@@ -12,6 +12,7 @@
; GCN-NEXT: .scratch_memory_size: 0
; GCN: .registers:
; GCN-NEXT: '0x2c4a (SPI_SHADER_PGM_RSRC1_VS)': 0
+; GCN-NEXT: '0x2c4b (SPI_SHADER_PGM_RSRC2_VS)': 0
; GCN-NEXT: ...
; GCN-NEXT: .end_amdgpu_pal_metadata
define amdgpu_vs half @vs_amdpal(half %arg0) {
diff --git a/llvm/unittests/Target/AMDGPU/CMakeLists.txt b/llvm/unittests/Target/AMDGPU/CMakeLists.txt
index 2d7a47943df69..502aaaa90c07b 100644
--- a/llvm/unittests/Target/AMDGPU/CMakeLists.txt
+++ b/llvm/unittests/Target/AMDGPU/CMakeLists.txt
@@ -20,6 +20,7 @@ add_llvm_target_unittest(AMDGPUTests
AMDGPUUnitTests.cpp
DwarfRegMappings.cpp
ExecMayBeModifiedBeforeAnyUse.cpp
+ PALMetadata.cpp
)
set_property(TARGET AMDGPUTests PROPERTY FOLDER "Tests/UnitTests/TargetTests")
diff --git a/llvm/unittests/Target/AMDGPU/PALMetadata.cpp b/llvm/unittests/Target/AMDGPU/PALMetadata.cpp
new file mode 100644
index 0000000000000..f58a91c5df892
--- /dev/null
+++ b/llvm/unittests/Target/AMDGPU/PALMetadata.cpp
@@ -0,0 +1,245 @@
+//===- llvm/unittests/MC/AMDGPU/PALMetadata.cpp ---------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include "AMDGPUTargetMachine.h"
+#include "GCNSubtarget.h"
+#include "SIProgramInfo.h"
+#include "Utils/AMDGPUPALMetadata.h"
+#include "llvm/CodeGen/MachineModuleInfo.h"
+#include "llvm/MC/MCContext.h"
+#include "llvm/MC/MCExpr.h"
+#include "llvm/MC/MCStreamer.h"
+#include "llvm/MC/MCSymbol.h"
+#include "llvm/MC/MCTargetOptions.h"
+#include "llvm/MC/TargetRegistry.h"
+#include "llvm/Support/TargetSelect.h"
+#include "llvm/Target/TargetMachine.h"
+#include "gtest/gtest.h"
+
+using namespace llvm;
+
+class PALMetadata : public testing::Test {
+protected:
+ std::unique_ptr<GCNTargetMachine> TM;
+ std::unique_ptr<LLVMContext> Ctx;
+ std::unique_ptr<GCNSubtarget> ST;
+ std::unique_ptr<MachineModuleInfo> MMI;
+ std::unique_ptr<MachineFunction> MF;
+ std::unique_ptr<Module> M;
+ AMDGPUPALMetadata MD;
+
+ static void SetUpTestSuite() {
+ LLVMInitializeAMDGPUTargetInfo();
+ LLVMInitializeAMDGPUTarget();
+ LLVMInitializeAMDGPUTargetMC();
+ }
+
+ PALMetadata() {
+ std::string Triple = "amdgcn--amdpal";
+ std::string CPU = "gfx1010";
+ std::string FS = "";
+
+ std::string Error;
+ const Target *TheTarget = TargetRegistry::lookupTarget(Triple, Error);
+ TargetOptions Options;
+
+ TM.reset(static_cast<GCNTargetMachine *>(TheTarget->createTargetMachine(
+ Triple, CPU, FS, Options, std::nullopt, std::nullopt)));
+
+ Ctx = std::make_unique<LLVMContext>();
+ M = std::make_unique<Module>("Module", *Ctx);
+ M->setDataLayout(TM->createDataLayout());
+ auto *FType = FunctionType::get(Type::getVoidTy(*Ctx), false);
+ auto *F = Function::Create(FType, GlobalValue::ExternalLinkage, "Test", *M);
+ MMI = std::make_unique<MachineModuleInfo>(TM.get());
+
+ ST = std::make_unique<GCNSubtarget>(TM->getTargetTriple(),
+ TM->getTargetCPU(),
+ TM->getTargetFeatureString(), *TM);
+
+ MF = std::make_unique<MachineFunction>(*F, *TM, *ST, 1, *MMI);
+ }
+};
+
+TEST_F(PALMetadata, ResourceRegisterSetORsResolvableUnknown) {
+ std::string yaml = "---\n"
+ "amdpal.pipelines:\n"
+ " - .hardware_stages:\n"
+ " .es:\n"
+ " .entry_point: Test\n"
+ " .scratch_memory_size: 0\n"
+ " .sgpr_count: 0x1\n"
+ " .vgpr_count: 0x1\n"
+ " .registers:\n"
+ " \'0x2c4a (SPI_SHADER_PGM_RSRC1_VS)\': 0x2f0000\n"
+ " \'0x2c4b (SPI_SHADER_PGM_RSRC2_VS)\': 0\n"
+ "...\n";
+
+ MCContext &MCCtx = MF->getContext();
+ auto CC = CallingConv::AMDGPU_VS;
+ MD.setFromString(yaml);
+ MD.setRsrc2(CC, MCConstantExpr::create(42, MCCtx), MCCtx);
+ MCSymbol *Sym = MCCtx.getOrCreateSymbol("Unknown");
+ MD.setRsrc2(CC, MCSymbolRefExpr::create(Sym, MCCtx), MCCtx);
+ EXPECT_FALSE(MD.resolvedAllMCExpr());
+
+ MD.setRsrc2(CC, MCConstantExpr::create(0xff00, MCCtx), MCCtx);
+ Sym->setVariableValue(MCConstantExpr::create(0xffff0000, MCCtx));
+ std::string Output;
+ MD.toString(Output);
+
+ EXPECT_TRUE(MD.resolvedAllMCExpr());
+
+ auto n = Output.find("\'0x2c4b (SPI_SHADER_PGM_RSRC2_VS)\': 0xffffff2a");
+ EXPECT_TRUE(n != std::string::npos);
+}
+
+TEST_F(PALMetadata, ResourceRegisterSetORsResolvableUnknowns) {
+ std::string yaml = "---\n"
+ "amdpal.pipelines:\n"
+ " - .hardware_stages:\n"
+ " .es:\n"
+ " .entry_point: Test\n"
+ " .scratch_memory_size: 0\n"
+ " .sgpr_count: 0x1\n"
+ " .vgpr_count: 0x1\n"
+ " .registers:\n"
+ " \'0x2c4a (SPI_SHADER_PGM_RSRC1_VS)\': 0x2f0000\n"
+ " \'0x2c4b (SPI_SHADER_PGM_RSRC2_VS)\': 0\n"
+ "...\n";
+
+ MCContext &MCCtx = MF->getContext();
+ auto CC = CallingConv::AMDGPU_VS;
+ MD.setFromString(yaml);
+ MCSymbol *SymOne = MCCtx.getOrCreateSymbol("UnknownOne");
+ MD.setRsrc2(CC, MCSymbolRefExpr::create(SymOne, MCCtx), MCCtx);
+
+ MD.setRsrc2(CC, MCConstantExpr::create(42, MCCtx), MCCtx);
+
+ MCSymbol *SymTwo = MCCtx.getOrCreateSymbol("UnknownTwo");
+ MD.setRsrc2(CC, MCSymbolRefExpr::create(SymTwo, MCCtx), MCCtx);
+ EXPECT_FALSE(MD.resolvedAllMCExpr());
+
+ SymOne->setVariableValue(MCConstantExpr::create(0xffff0000, MCCtx));
+ SymTwo->setVariableValue(MCConstantExpr::create(0x0000ff00, MCCtx));
+
+ std::string Output;
+ MD.toString(Output);
+
+ EXPECT_TRUE(MD.resolvedAllMCExpr());
+
+ auto n = Output.find("\'0x2c4b (SPI_SHADER_PGM_RSRC2_VS)\': 0xffffff2a");
+ EXPECT_TRUE(n != std::string::npos);
+}
+
+TEST_F(PALMetadata, ResourceRegisterSetORsPreset) {
+ std::string yaml = "---\n"
+ "amdpal.pipelines:\n"
+ " - .hardware_stages:\n"
+ " .es:\n"
+ " .entry_point: Test\n"
+ " .scratch_memory_size: 0\n"
+ " .sgpr_count: 0x1\n"
+ " .vgpr_count: 0x1\n"
+ " .registers:\n"
+ " \'0x2c4a (SPI_SHADER_PGM_RSRC1_VS)\': 0x2f0000\n"
+ " \'0x2c4b (SPI_SHADER_PGM_RSRC2_VS)\': 0x2a\n"
+ "...\n";
+
+ MCContext &MCCtx = MF->getContext();
+ auto CC = CallingConv::AMDGPU_VS;
+ MD.setFromString(yaml);
+ MCSymbol *Sym = MCCtx.getOrCreateSymbol("Unknown");
+ MD.setRsrc2(CC, MCSymbolRefExpr::create(Sym, MCCtx), MCCtx);
+ MD.setRsrc2(CC, MCConstantExpr::create(0xff00, MCCtx), MCCtx);
+ Sym->setVariableValue(MCConstantExpr::create(0xffff0000, MCCtx));
+ std::string Output;
+ MD.toString(Output);
+
+ auto n = Output.find("\'0x2c4b (SPI_SHADER_PGM_RSRC2_VS)\': 0xffffff2a");
+ EXPECT_TRUE(n != std::string::npos);
+}
+
+TEST_F(PALMetadata, ResourceRegisterSetORs) {
+ std::string yaml = "---\n"
+ "amdpal.pipelines:\n"
+ " - .hardware_stages:\n"
+ " .es:\n"
+ " .entry_point: Test\n"
+ " .scratch_memory_size: 0\n"
+ " .sgpr_count: 0x1\n"
+ " .vgpr_count: 0x1\n"
+ " .registers:\n"
+ " \'0x2c4a (SPI_SHADER_PGM_RSRC1_VS)\': 0x2f0000\n"
+ " \'0x2c4b (SPI_SHADER_PGM_RSRC2_VS)\': 0\n"
+ "...\n";
+
+ MCContext &MCCtx = MF->getContext();
+ auto CC = CallingConv::AMDGPU_VS;
+ MD.setFromString(yaml);
+ MCSymbol *Sym = MCCtx.getOrCreateSymbol("Unknown");
+ MD.setRsrc2(CC, MCSymbolRefExpr::create(Sym, MCCtx), MCCtx);
+ MD.setRsrc2(CC, 42);
+ MD.setRsrc2(CC, MCConstantExpr::create(0xff00, MCCtx), MCCtx);
+ Sym->setVariableValue(MCConstantExpr::create(0xffff0000, MCCtx));
+ std::string Output;
+ MD.toString(Output);
+
+ auto n = Output.find("\'0x2c4b (SPI_SHADER_PGM_RSRC2_VS)\': 0xffffff2a");
+ EXPECT_TRUE(n != std::string::npos);
+}
+
+TEST_F(PALMetadata, ResourceRegisterSetUnresolvedSym) {
+ std::string yaml = "---\n"
+ "amdpal.pipelines:\n"
+ " - .hardware_stages:\n"
+ " .es:\n"
+ " .entry_point: Test\n"
+ " .scratch_memory_size: 0\n"
+ " .sgpr_count: 0x1\n"
+ " .vgpr_count: 0x1\n"
+ " .registers:\n"
+ " \'0x2c4a (SPI_SHADER_PGM_RSRC1_VS)\': 0x2f0000\n"
+ " \'0x2c4b (SPI_SHADER_PGM_RSRC2_VS)\': 0\n"
+ "...\n";
+
+ MCContext &MCCtx = MF->getContext();
+ auto CC = CallingConv::AMDGPU_VS;
+ MD.setFromString(yaml);
+ MCSymbol *Sym = MCCtx.getOrCreateSymbol("Unknown");
+ MD.setRsrc2(CC, MCSymbolRefExpr::create(Sym, MCCtx), MCCtx);
+ MD.setRsrc2(CC, MCConstantExpr::create(0xff00, MCCtx), MCCtx);
+ std::string Output;
+
+ MD.toString(Output);
+ EXPECT_FALSE(MD.resolvedAllMCExpr());
+}
+
+TEST_F(PALMetadata, ResourceRegisterSetNoEmitUnresolved) {
+ std::string yaml = "---\n"
+ "amdpal.pipelines:\n"
+ " - .hardware_stages:\n"
+ " .es:\n"
+ " .entry_point: Test\n"
+ " .scratch_memory_size: 0\n"
+ " .sgpr_count: 0x1\n"
+ " .vgpr_count: 0x1\n"
+ " .registers:\n"
+ " \'0x2c4a (SPI_SHADER_PGM_RSRC1_VS)\': 0x2f0000\n"
+ " \'0x2c4b (SPI_SHADER_PGM_RSRC2_VS)\': 0\n"
+ "...\n";
+
+ MCContext &MCCtx = MF->getContext();
+ auto CC = CallingConv::AMDGPU_VS;
+ MD.setFromString(yaml);
+ MCSymbol *Sym = MCCtx.getOrCreateSymbol("Unknown");
+ MD.setRsrc2(CC, MCSymbolRefExpr::create(Sym, MCCtx), MCCtx);
+ MD.setRsrc2(CC, MCConstantExpr::create(0xff00, MCCtx), MCCtx);
+
+ EXPECT_FALSE(MD.resolvedAllMCExpr());
+}
>From b8234b2890631bec3348df120c9670e5fc5f89d4 Mon Sep 17 00:00:00 2001
From: Janek van Oirschot <janek.vanoirschot at amd.com>
Date: Fri, 31 May 2024 14:21:10 +0100
Subject: [PATCH 2/6] Move the SIDefines helper functions for MCExpr cases to
its own util header, add comments with explanations
---
llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp | 27 +------
.../AMDGPU/Utils/AMDKernelCodeTUtils.cpp | 46 ++---------
llvm/lib/Target/AMDGPU/Utils/SIDefinesUtils.h | 79 +++++++++++++++++++
3 files changed, 86 insertions(+), 66 deletions(-)
create mode 100644 llvm/lib/Target/AMDGPU/Utils/SIDefinesUtils.h
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp b/llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp
index f4028adc84828..f0c7a5efb53a8 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp
@@ -29,6 +29,7 @@
#include "TargetInfo/AMDGPUTargetInfo.h"
#include "Utils/AMDGPUBaseInfo.h"
#include "Utils/AMDKernelCodeTUtils.h"
+#include "Utils/SIDefinesUtils.h"
#include "llvm/Analysis/OptimizationRemarkEmitter.h"
#include "llvm/BinaryFormat/ELF.h"
#include "llvm/CodeGen/MachineFrameInfo.h"
@@ -1194,30 +1195,6 @@ static void EmitPALMetadataCommon(AMDGPUPALMetadata *MD,
getLdsDwGranularity(ST) * sizeof(uint32_t)));
}
-static constexpr std::pair<unsigned, unsigned> getShiftMask(unsigned Value) {
- unsigned Shift = 0;
- unsigned Mask = 0;
-
- Mask = ~Value;
- for (; !(Mask & 1); Shift++, Mask >>= 1) {
- }
-
- return std::make_pair(Shift, Mask);
-}
-
-static const MCExpr *MaskShiftSet(const MCExpr *Val, uint32_t Mask,
- uint32_t Shift, MCContext &Ctx) {
- if (Mask) {
- const MCExpr *MaskExpr = MCConstantExpr::create(Mask, Ctx);
- Val = MCBinaryExpr::createAnd(Val, MaskExpr, Ctx);
- }
- if (Shift) {
- const MCExpr *ShiftExpr = MCConstantExpr::create(Shift, Ctx);
- Val = MCBinaryExpr::createShl(Val, ShiftExpr, Ctx);
- }
- return Val;
-}
-
// This is the equivalent of EmitProgramInfoSI above, but for when the OS type
// is AMDPAL. It stores each compute/SPI register setting and other PAL
// metadata items into the PALMD::Metadata, combining with any provided by the
@@ -1249,7 +1226,7 @@ void AMDGPUAsmPrinter::EmitPALMetadata(const MachineFunction &MF,
MCBinaryExpr::createGT(CurrentProgramInfo.ScratchBlocks,
MCConstantExpr::create(0, Ctx), Ctx);
auto [Shift, Mask] = getShiftMask(C_00B84C_SCRATCH_EN);
- MD->setRsrc2(CC, MaskShiftSet(HasScratchBlocks, Mask, Shift, Ctx), Ctx);
+ MD->setRsrc2(CC, maskShiftSet(HasScratchBlocks, Mask, Shift, Ctx), Ctx);
}
} else {
MD->setHwStage(CC, ".debug_mode", (bool)CurrentProgramInfo.DebugMode);
diff --git a/llvm/lib/Target/AMDGPU/Utils/AMDKernelCodeTUtils.cpp b/llvm/lib/Target/AMDGPU/Utils/AMDKernelCodeTUtils.cpp
index eaee1a2a97399..720d5a1853dbb 100644
--- a/llvm/lib/Target/AMDGPU/Utils/AMDKernelCodeTUtils.cpp
+++ b/llvm/lib/Target/AMDGPU/Utils/AMDKernelCodeTUtils.cpp
@@ -14,6 +14,7 @@
#include "AMDKernelCodeT.h"
#include "SIDefines.h"
#include "Utils/AMDGPUBaseInfo.h"
+#include "Utils/SIDefinesUtils.h"
#include "llvm/ADT/IndexedMap.h"
#include "llvm/ADT/StringRef.h"
#include "llvm/MC/MCContext.h"
@@ -220,43 +221,6 @@ static int get_amd_kernel_code_t_FieldIndex(StringRef name) {
return map.lookup(name) - 1; // returns -1 if not found
}
-static constexpr std::pair<unsigned, unsigned> getShiftMask(unsigned Value) {
- unsigned Shift = 0;
- unsigned Mask = 0;
-
- Mask = ~Value;
- for (; !(Mask & 1); Shift++, Mask >>= 1) {
- }
-
- return std::make_pair(Shift, Mask);
-}
-
-static const MCExpr *MaskShiftSet(const MCExpr *Val, uint32_t Mask,
- uint32_t Shift, MCContext &Ctx) {
- if (Mask) {
- const MCExpr *MaskExpr = MCConstantExpr::create(Mask, Ctx);
- Val = MCBinaryExpr::createAnd(Val, MaskExpr, Ctx);
- }
- if (Shift) {
- const MCExpr *ShiftExpr = MCConstantExpr::create(Shift, Ctx);
- Val = MCBinaryExpr::createShl(Val, ShiftExpr, Ctx);
- }
- return Val;
-}
-
-static const MCExpr *MaskShiftGet(const MCExpr *Val, uint32_t Mask,
- uint32_t Shift, MCContext &Ctx) {
- if (Shift) {
- const MCExpr *ShiftExpr = MCConstantExpr::create(Shift, Ctx);
- Val = MCBinaryExpr::createLShr(Val, ShiftExpr, Ctx);
- }
- if (Mask) {
- const MCExpr *MaskExpr = MCConstantExpr::create(Mask, Ctx);
- Val = MCBinaryExpr::createAnd(Val, MaskExpr, Ctx);
- }
- return Val;
-}
-
class PrintField {
public:
template <typename T, T AMDGPUMCKernelCodeT::*ptr,
@@ -305,10 +269,10 @@ static ArrayRef<PrintFx> getPrinterTable() {
const MCExpr *Value; \
if (PGMType == 0) { \
Value = \
- MaskShiftGet(C.compute_pgm_resource1_registers, Mask, Shift, Ctx); \
+ maskShiftGet(C.compute_pgm_resource1_registers, Mask, Shift, Ctx); \
} else { \
Value = \
- MaskShiftGet(C.compute_pgm_resource2_registers, Mask, Shift, Ctx); \
+ maskShiftGet(C.compute_pgm_resource2_registers, Mask, Shift, Ctx); \
} \
int64_t Val; \
if (Value->evaluateAsAbsolute(Val)) \
@@ -392,7 +356,7 @@ static ArrayRef<ParseFx> getParserTable() {
if (!parseExpr(MCParser, Value, Err)) \
return false; \
auto [Shift, Mask] = getShiftMask(Complement); \
- Value = MaskShiftSet(Value, Mask, Shift, Ctx); \
+ Value = maskShiftSet(Value, Mask, Shift, Ctx); \
const MCExpr *Compl = MCConstantExpr::create(Complement, Ctx); \
if (PGMType == 0) { \
C.compute_pgm_resource1_registers = MCBinaryExpr::createAnd( \
@@ -542,7 +506,7 @@ void AMDGPUMCKernelCodeT::EmitKernelCodeT(MCStreamer &OS, MCContext &Ctx) {
const MCExpr *CodeProps = MCConstantExpr::create(code_properties, Ctx);
CodeProps = MCBinaryExpr::createOr(
CodeProps,
- MaskShiftSet(is_dynamic_callstack,
+ maskShiftSet(is_dynamic_callstack,
(1 << AMD_CODE_PROPERTY_IS_DYNAMIC_CALLSTACK_WIDTH) - 1,
AMD_CODE_PROPERTY_IS_DYNAMIC_CALLSTACK_SHIFT, Ctx),
Ctx);
diff --git a/llvm/lib/Target/AMDGPU/Utils/SIDefinesUtils.h b/llvm/lib/Target/AMDGPU/Utils/SIDefinesUtils.h
new file mode 100644
index 0000000000000..6565618ac523e
--- /dev/null
+++ b/llvm/lib/Target/AMDGPU/Utils/SIDefinesUtils.h
@@ -0,0 +1,79 @@
+//===-- SIDefines.h - SI Helper Functions -----------------------*- C++ -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+/// \file - utility functions for the SIDefines and its common uses.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_LIB_TARGET_AMDGPU_UTILS_SIDEFINESUTILS_H
+#define LLVM_LIB_TARGET_AMDGPU_UTILS_SIDEFINESUTILS_H
+
+#include "llvm/MC/MCExpr.h"
+#include <utility>
+
+namespace llvm {
+class MCContext;
+namespace AMDGPU {
+
+/// Deduce the least significant bit aligned shift and mask values for a binary
+/// Complement \p Value (as they're defined in SIDefines.h as C_*) as a returned
+/// pair<shift, mask>. That is to say \p Value == ~(mask << shift)
+///
+/// For example, given C_00B848_FWD_PROGRESS (i.e., 0x7FFFFFFF) from
+/// SIDefines.h, this will return the pair as (31,1).
+constexpr inline std::pair<unsigned, unsigned> getShiftMask(unsigned Value) {
+ unsigned Shift = 0;
+ unsigned Mask = 0;
+
+ Mask = ~Value;
+ for (; !(Mask & 1); Shift++, Mask >>= 1) {
+ }
+
+ return std::make_pair(Shift, Mask);
+}
+
+/// Provided with the MCExpr * \p Val, uint32 \p Mask and \p Shift, will return
+/// the masked and left shifted, in said order of operations, MCExpr * created
+/// within the MCContext \p Ctx.
+///
+/// For example, given MCExpr *Val, Mask == 0xf, Shift == 6 the returned MCExpr
+/// * will be the equivalent of (Val & 0xf) << 6
+inline const MCExpr *maskShiftSet(const MCExpr *Val, uint32_t Mask,
+ uint32_t Shift, MCContext &Ctx) {
+ if (Mask) {
+ const MCExpr *MaskExpr = MCConstantExpr::create(Mask, Ctx);
+ Val = MCBinaryExpr::createAnd(Val, MaskExpr, Ctx);
+ }
+ if (Shift) {
+ const MCExpr *ShiftExpr = MCConstantExpr::create(Shift, Ctx);
+ Val = MCBinaryExpr::createShl(Val, ShiftExpr, Ctx);
+ }
+ return Val;
+}
+
+/// Provided with the MCExpr * \p Val, uint32 \p Mask and \p Shift, will return
+/// the right shifted and masked, in said order of operations, MCExpr * created
+/// within the MCContext \p Ctx.
+///
+/// For example, given MCExpr *Val, Mask == 0xf, Shift == 6 the returned MCExpr
+/// * will be the equivalent of (Val >> 6) & 0xf
+inline const MCExpr *maskShiftGet(const MCExpr *Val, uint32_t Mask,
+ uint32_t Shift, MCContext &Ctx) {
+ if (Shift) {
+ const MCExpr *ShiftExpr = MCConstantExpr::create(Shift, Ctx);
+ Val = MCBinaryExpr::createLShr(Val, ShiftExpr, Ctx);
+ }
+ if (Mask) {
+ const MCExpr *MaskExpr = MCConstantExpr::create(Mask, Ctx);
+ Val = MCBinaryExpr::createAnd(Val, MaskExpr, Ctx);
+ }
+ return Val;
+}
+
+} // end namespace AMDGPU
+} // end namespace llvm
+
+#endif // LLVM_LIB_TARGET_AMDGPU_UTILS_SIDEFINESUTILS_H
>From 82d7d3631433e3a6a32c487c0ae4d3f71de993fe Mon Sep 17 00:00:00 2001
From: Janek van Oirschot <janek.vanoirschot at amd.com>
Date: Fri, 31 May 2024 16:08:37 +0100
Subject: [PATCH 3/6] Feedback, decapitalize methods, stringref in unittest
where possible, remove redundant inline
---
.../AMDGPU/Utils/AMDGPUDelayedMCExpr.cpp | 4 +-
.../Target/AMDGPU/Utils/AMDGPUDelayedMCExpr.h | 4 +-
.../Target/AMDGPU/Utils/AMDGPUPALMetadata.cpp | 12 +-
llvm/lib/Target/AMDGPU/Utils/SIDefinesUtils.h | 2 +-
llvm/unittests/Target/AMDGPU/PALMetadata.cpp | 150 +++++++++---------
5 files changed, 86 insertions(+), 86 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/Utils/AMDGPUDelayedMCExpr.cpp b/llvm/lib/Target/AMDGPU/Utils/AMDGPUDelayedMCExpr.cpp
index 3955f557b9a25..a8da0cf8eb1ff 100644
--- a/llvm/lib/Target/AMDGPU/Utils/AMDGPUDelayedMCExpr.cpp
+++ b/llvm/lib/Target/AMDGPU/Utils/AMDGPUDelayedMCExpr.cpp
@@ -27,7 +27,7 @@ static msgpack::DocNode getNode(msgpack::DocNode DN, msgpack::Type Type,
}
}
-void DelayedMCExpr::AssignDocNode(msgpack::DocNode &DN, msgpack::Type Type,
+void DelayedMCExpr::assignDocNode(msgpack::DocNode &DN, msgpack::Type Type,
const MCExpr *Expr) {
MCValue Res;
if (Expr->evaluateAsRelocatable(Res, nullptr, nullptr)) {
@@ -40,7 +40,7 @@ void DelayedMCExpr::AssignDocNode(msgpack::DocNode &DN, msgpack::Type Type,
DelayedExprs.push_back(DelayedExpr{DN, Type, Expr});
}
-bool DelayedMCExpr::ResolveDelayedExpressions() {
+bool DelayedMCExpr::resolveDelayedExpressions() {
bool Success;
while (!DelayedExprs.empty()) {
diff --git a/llvm/lib/Target/AMDGPU/Utils/AMDGPUDelayedMCExpr.h b/llvm/lib/Target/AMDGPU/Utils/AMDGPUDelayedMCExpr.h
index c546660a0d996..729900efb403e 100644
--- a/llvm/lib/Target/AMDGPU/Utils/AMDGPUDelayedMCExpr.h
+++ b/llvm/lib/Target/AMDGPU/Utils/AMDGPUDelayedMCExpr.h
@@ -27,8 +27,8 @@ class DelayedMCExpr {
std::deque<DelayedExpr> DelayedExprs;
public:
- bool ResolveDelayedExpressions();
- void AssignDocNode(msgpack::DocNode &DN, msgpack::Type Type,
+ bool resolveDelayedExpressions();
+ void assignDocNode(msgpack::DocNode &DN, msgpack::Type Type,
const MCExpr *Expr);
void clear();
bool empty();
diff --git a/llvm/lib/Target/AMDGPU/Utils/AMDGPUPALMetadata.cpp b/llvm/lib/Target/AMDGPU/Utils/AMDGPUPALMetadata.cpp
index 4597dab142470..b3af32d33982a 100644
--- a/llvm/lib/Target/AMDGPU/Utils/AMDGPUPALMetadata.cpp
+++ b/llvm/lib/Target/AMDGPU/Utils/AMDGPUPALMetadata.cpp
@@ -224,7 +224,7 @@ void AMDGPUPALMetadata::setRegister(unsigned Reg, const MCExpr *Val,
REM[Reg] = Val;
(void)Unused;
}
- DelayedExprs.AssignDocNode(N, msgpack::Type::UInt, Val);
+ DelayedExprs.assignDocNode(N, msgpack::Type::UInt, Val);
}
// Set the entry point name for one shader.
@@ -350,7 +350,7 @@ void AMDGPUPALMetadata::setFunctionNumUsedVgprs(StringRef FnName,
void AMDGPUPALMetadata::setFunctionNumUsedVgprs(StringRef FnName,
const MCExpr *Val) {
auto Node = getShaderFunction(FnName);
- DelayedExprs.AssignDocNode(Node[".vgpr_count"], msgpack::Type::UInt, Val);
+ DelayedExprs.assignDocNode(Node[".vgpr_count"], msgpack::Type::UInt, Val);
}
// Set the number of used vgprs in the metadata.
@@ -363,7 +363,7 @@ void AMDGPUPALMetadata::setFunctionNumUsedSgprs(StringRef FnName,
void AMDGPUPALMetadata::setFunctionNumUsedSgprs(StringRef FnName,
const MCExpr *Val) {
auto Node = getShaderFunction(FnName);
- DelayedExprs.AssignDocNode(Node[".sgpr_count"], msgpack::Type::UInt, Val);
+ DelayedExprs.assignDocNode(Node[".sgpr_count"], msgpack::Type::UInt, Val);
}
// Set the hardware register bit in PAL metadata to enable wave32 on the
@@ -762,7 +762,7 @@ void AMDGPUPALMetadata::toString(std::string &String) {
String.clear();
if (!BlobType)
return;
- ResolvedAll = DelayedExprs.ResolveDelayedExpressions();
+ ResolvedAll = DelayedExprs.resolveDelayedExpressions();
raw_string_ostream Stream(String);
if (isLegacy()) {
if (MsgPackDoc.getRoot().getKind() == msgpack::Type::Nil)
@@ -812,7 +812,7 @@ void AMDGPUPALMetadata::toString(std::string &String) {
// a .note record of the specified AMD type. Returns an empty blob if
// there is no PAL metadata,
void AMDGPUPALMetadata::toBlob(unsigned Type, std::string &Blob) {
- ResolvedAll = DelayedExprs.ResolveDelayedExpressions();
+ ResolvedAll = DelayedExprs.resolveDelayedExpressions();
if (Type == ELF::NT_AMD_PAL_METADATA)
toLegacyBlob(Blob);
else if (Type)
@@ -1052,7 +1052,7 @@ void AMDGPUPALMetadata::setHwStage(unsigned CC, StringRef field, bool Val) {
void AMDGPUPALMetadata::setHwStage(unsigned CC, StringRef field,
msgpack::Type Type, const MCExpr *Val) {
- DelayedExprs.AssignDocNode(getHwStage(CC)[field], Type, Val);
+ DelayedExprs.assignDocNode(getHwStage(CC)[field], Type, Val);
}
void AMDGPUPALMetadata::setComputeRegisters(StringRef field, unsigned Val) {
diff --git a/llvm/lib/Target/AMDGPU/Utils/SIDefinesUtils.h b/llvm/lib/Target/AMDGPU/Utils/SIDefinesUtils.h
index 6565618ac523e..64d21de12c268 100644
--- a/llvm/lib/Target/AMDGPU/Utils/SIDefinesUtils.h
+++ b/llvm/lib/Target/AMDGPU/Utils/SIDefinesUtils.h
@@ -24,7 +24,7 @@ namespace AMDGPU {
///
/// For example, given C_00B848_FWD_PROGRESS (i.e., 0x7FFFFFFF) from
/// SIDefines.h, this will return the pair as (31,1).
-constexpr inline std::pair<unsigned, unsigned> getShiftMask(unsigned Value) {
+constexpr std::pair<unsigned, unsigned> getShiftMask(unsigned Value) {
unsigned Shift = 0;
unsigned Mask = 0;
diff --git a/llvm/unittests/Target/AMDGPU/PALMetadata.cpp b/llvm/unittests/Target/AMDGPU/PALMetadata.cpp
index f58a91c5df892..7d4f3836088ab 100644
--- a/llvm/unittests/Target/AMDGPU/PALMetadata.cpp
+++ b/llvm/unittests/Target/AMDGPU/PALMetadata.cpp
@@ -40,9 +40,9 @@ class PALMetadata : public testing::Test {
}
PALMetadata() {
- std::string Triple = "amdgcn--amdpal";
- std::string CPU = "gfx1010";
- std::string FS = "";
+ StringRef Triple = "amdgcn--amdpal";
+ StringRef CPU = "gfx1010";
+ StringRef FS = "";
std::string Error;
const Target *TheTarget = TargetRegistry::lookupTarget(Triple, Error);
@@ -67,18 +67,18 @@ class PALMetadata : public testing::Test {
};
TEST_F(PALMetadata, ResourceRegisterSetORsResolvableUnknown) {
- std::string yaml = "---\n"
- "amdpal.pipelines:\n"
- " - .hardware_stages:\n"
- " .es:\n"
- " .entry_point: Test\n"
- " .scratch_memory_size: 0\n"
- " .sgpr_count: 0x1\n"
- " .vgpr_count: 0x1\n"
- " .registers:\n"
- " \'0x2c4a (SPI_SHADER_PGM_RSRC1_VS)\': 0x2f0000\n"
- " \'0x2c4b (SPI_SHADER_PGM_RSRC2_VS)\': 0\n"
- "...\n";
+ StringRef yaml = "---\n"
+ "amdpal.pipelines:\n"
+ " - .hardware_stages:\n"
+ " .es:\n"
+ " .entry_point: Test\n"
+ " .scratch_memory_size: 0\n"
+ " .sgpr_count: 0x1\n"
+ " .vgpr_count: 0x1\n"
+ " .registers:\n"
+ " \'0x2c4a (SPI_SHADER_PGM_RSRC1_VS)\': 0x2f0000\n"
+ " \'0x2c4b (SPI_SHADER_PGM_RSRC2_VS)\': 0\n"
+ "...\n";
MCContext &MCCtx = MF->getContext();
auto CC = CallingConv::AMDGPU_VS;
@@ -100,18 +100,18 @@ TEST_F(PALMetadata, ResourceRegisterSetORsResolvableUnknown) {
}
TEST_F(PALMetadata, ResourceRegisterSetORsResolvableUnknowns) {
- std::string yaml = "---\n"
- "amdpal.pipelines:\n"
- " - .hardware_stages:\n"
- " .es:\n"
- " .entry_point: Test\n"
- " .scratch_memory_size: 0\n"
- " .sgpr_count: 0x1\n"
- " .vgpr_count: 0x1\n"
- " .registers:\n"
- " \'0x2c4a (SPI_SHADER_PGM_RSRC1_VS)\': 0x2f0000\n"
- " \'0x2c4b (SPI_SHADER_PGM_RSRC2_VS)\': 0\n"
- "...\n";
+ StringRef yaml = "---\n"
+ "amdpal.pipelines:\n"
+ " - .hardware_stages:\n"
+ " .es:\n"
+ " .entry_point: Test\n"
+ " .scratch_memory_size: 0\n"
+ " .sgpr_count: 0x1\n"
+ " .vgpr_count: 0x1\n"
+ " .registers:\n"
+ " \'0x2c4a (SPI_SHADER_PGM_RSRC1_VS)\': 0x2f0000\n"
+ " \'0x2c4b (SPI_SHADER_PGM_RSRC2_VS)\': 0\n"
+ "...\n";
MCContext &MCCtx = MF->getContext();
auto CC = CallingConv::AMDGPU_VS;
@@ -138,18 +138,18 @@ TEST_F(PALMetadata, ResourceRegisterSetORsResolvableUnknowns) {
}
TEST_F(PALMetadata, ResourceRegisterSetORsPreset) {
- std::string yaml = "---\n"
- "amdpal.pipelines:\n"
- " - .hardware_stages:\n"
- " .es:\n"
- " .entry_point: Test\n"
- " .scratch_memory_size: 0\n"
- " .sgpr_count: 0x1\n"
- " .vgpr_count: 0x1\n"
- " .registers:\n"
- " \'0x2c4a (SPI_SHADER_PGM_RSRC1_VS)\': 0x2f0000\n"
- " \'0x2c4b (SPI_SHADER_PGM_RSRC2_VS)\': 0x2a\n"
- "...\n";
+ StringRef yaml = "---\n"
+ "amdpal.pipelines:\n"
+ " - .hardware_stages:\n"
+ " .es:\n"
+ " .entry_point: Test\n"
+ " .scratch_memory_size: 0\n"
+ " .sgpr_count: 0x1\n"
+ " .vgpr_count: 0x1\n"
+ " .registers:\n"
+ " \'0x2c4a (SPI_SHADER_PGM_RSRC1_VS)\': 0x2f0000\n"
+ " \'0x2c4b (SPI_SHADER_PGM_RSRC2_VS)\': 0x2a\n"
+ "...\n";
MCContext &MCCtx = MF->getContext();
auto CC = CallingConv::AMDGPU_VS;
@@ -166,18 +166,18 @@ TEST_F(PALMetadata, ResourceRegisterSetORsPreset) {
}
TEST_F(PALMetadata, ResourceRegisterSetORs) {
- std::string yaml = "---\n"
- "amdpal.pipelines:\n"
- " - .hardware_stages:\n"
- " .es:\n"
- " .entry_point: Test\n"
- " .scratch_memory_size: 0\n"
- " .sgpr_count: 0x1\n"
- " .vgpr_count: 0x1\n"
- " .registers:\n"
- " \'0x2c4a (SPI_SHADER_PGM_RSRC1_VS)\': 0x2f0000\n"
- " \'0x2c4b (SPI_SHADER_PGM_RSRC2_VS)\': 0\n"
- "...\n";
+ StringRef yaml = "---\n"
+ "amdpal.pipelines:\n"
+ " - .hardware_stages:\n"
+ " .es:\n"
+ " .entry_point: Test\n"
+ " .scratch_memory_size: 0\n"
+ " .sgpr_count: 0x1\n"
+ " .vgpr_count: 0x1\n"
+ " .registers:\n"
+ " \'0x2c4a (SPI_SHADER_PGM_RSRC1_VS)\': 0x2f0000\n"
+ " \'0x2c4b (SPI_SHADER_PGM_RSRC2_VS)\': 0\n"
+ "...\n";
MCContext &MCCtx = MF->getContext();
auto CC = CallingConv::AMDGPU_VS;
@@ -195,18 +195,18 @@ TEST_F(PALMetadata, ResourceRegisterSetORs) {
}
TEST_F(PALMetadata, ResourceRegisterSetUnresolvedSym) {
- std::string yaml = "---\n"
- "amdpal.pipelines:\n"
- " - .hardware_stages:\n"
- " .es:\n"
- " .entry_point: Test\n"
- " .scratch_memory_size: 0\n"
- " .sgpr_count: 0x1\n"
- " .vgpr_count: 0x1\n"
- " .registers:\n"
- " \'0x2c4a (SPI_SHADER_PGM_RSRC1_VS)\': 0x2f0000\n"
- " \'0x2c4b (SPI_SHADER_PGM_RSRC2_VS)\': 0\n"
- "...\n";
+ StringRef yaml = "---\n"
+ "amdpal.pipelines:\n"
+ " - .hardware_stages:\n"
+ " .es:\n"
+ " .entry_point: Test\n"
+ " .scratch_memory_size: 0\n"
+ " .sgpr_count: 0x1\n"
+ " .vgpr_count: 0x1\n"
+ " .registers:\n"
+ " \'0x2c4a (SPI_SHADER_PGM_RSRC1_VS)\': 0x2f0000\n"
+ " \'0x2c4b (SPI_SHADER_PGM_RSRC2_VS)\': 0\n"
+ "...\n";
MCContext &MCCtx = MF->getContext();
auto CC = CallingConv::AMDGPU_VS;
@@ -221,18 +221,18 @@ TEST_F(PALMetadata, ResourceRegisterSetUnresolvedSym) {
}
TEST_F(PALMetadata, ResourceRegisterSetNoEmitUnresolved) {
- std::string yaml = "---\n"
- "amdpal.pipelines:\n"
- " - .hardware_stages:\n"
- " .es:\n"
- " .entry_point: Test\n"
- " .scratch_memory_size: 0\n"
- " .sgpr_count: 0x1\n"
- " .vgpr_count: 0x1\n"
- " .registers:\n"
- " \'0x2c4a (SPI_SHADER_PGM_RSRC1_VS)\': 0x2f0000\n"
- " \'0x2c4b (SPI_SHADER_PGM_RSRC2_VS)\': 0\n"
- "...\n";
+ StringRef yaml = "---\n"
+ "amdpal.pipelines:\n"
+ " - .hardware_stages:\n"
+ " .es:\n"
+ " .entry_point: Test\n"
+ " .scratch_memory_size: 0\n"
+ " .sgpr_count: 0x1\n"
+ " .vgpr_count: 0x1\n"
+ " .registers:\n"
+ " \'0x2c4a (SPI_SHADER_PGM_RSRC1_VS)\': 0x2f0000\n"
+ " \'0x2c4b (SPI_SHADER_PGM_RSRC2_VS)\': 0\n"
+ "...\n";
MCContext &MCCtx = MF->getContext();
auto CC = CallingConv::AMDGPU_VS;
>From 449f2ee1e8ad90bb0a9395c56fa562d1d45fdc70 Mon Sep 17 00:00:00 2001
From: Janek van Oirschot <janek.vanoirschot at amd.com>
Date: Fri, 7 Jun 2024 12:50:49 +0100
Subject: [PATCH 4/6] Feedback
---
llvm/lib/Target/AMDGPU/Utils/AMDGPUDelayedMCExpr.cpp | 7 ++-----
1 file changed, 2 insertions(+), 5 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/Utils/AMDGPUDelayedMCExpr.cpp b/llvm/lib/Target/AMDGPU/Utils/AMDGPUDelayedMCExpr.cpp
index a8da0cf8eb1ff..62325f5958849 100644
--- a/llvm/lib/Target/AMDGPU/Utils/AMDGPUDelayedMCExpr.cpp
+++ b/llvm/lib/Target/AMDGPU/Utils/AMDGPUDelayedMCExpr.cpp
@@ -41,15 +41,12 @@ void DelayedMCExpr::assignDocNode(msgpack::DocNode &DN, msgpack::Type Type,
}
bool DelayedMCExpr::resolveDelayedExpressions() {
- bool Success;
-
while (!DelayedExprs.empty()) {
DelayedExpr DE = DelayedExprs.front();
MCValue Res;
- Success = DE.Expr->evaluateAsRelocatable(Res, nullptr, nullptr);
- Success &= Res.isAbsolute();
- if (!Success)
+ if (!DE.Expr->evaluateAsRelocatable(Res, nullptr, nullptr) ||
+ !Res.isAbsolute())
return false;
DelayedExprs.pop_front();
>From 0bc7ebbb3fee5b527d043d476b74524141ae0e9f Mon Sep 17 00:00:00 2001
From: Janek van Oirschot <janek.vanoirschot at amd.com>
Date: Fri, 7 Jun 2024 19:21:43 +0100
Subject: [PATCH 5/6] Feedback
---
llvm/lib/Target/AMDGPU/Utils/AMDGPUPALMetadata.cpp | 7 ++++---
1 file changed, 4 insertions(+), 3 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/Utils/AMDGPUPALMetadata.cpp b/llvm/lib/Target/AMDGPU/Utils/AMDGPUPALMetadata.cpp
index b3af32d33982a..dd1a488b31a10 100644
--- a/llvm/lib/Target/AMDGPU/Utils/AMDGPUPALMetadata.cpp
+++ b/llvm/lib/Target/AMDGPU/Utils/AMDGPUPALMetadata.cpp
@@ -204,10 +204,11 @@ void AMDGPUPALMetadata::setRegister(unsigned Reg, const MCExpr *Val,
return;
}
auto &N = getRegisters()[MsgPackDoc.getNode(Reg)];
- bool RegSeenInREM = REM.find(Reg) != REM.end();
+ auto ExprIt = REM.find(Reg);
+ bool RegSeenInREM = ExprIt != REM.end();
if (RegSeenInREM) {
- Val = MCBinaryExpr::createOr(Val, REM[Reg], Ctx);
+ Val = MCBinaryExpr::createOr(Val, ExprIt->getSecond(), Ctx);
// This conditional may be redundant most of the time, but
// setRegister(unsigned, unsigned) could've been called while RegSeenInREM
// is true.
@@ -215,7 +216,7 @@ void AMDGPUPALMetadata::setRegister(unsigned Reg, const MCExpr *Val,
const MCExpr *NExpr = MCConstantExpr::create(N.getUInt(), Ctx);
Val = MCBinaryExpr::createOr(Val, NExpr, Ctx);
}
- REM[Reg] = Val;
+ ExprIt->getSecond() = Val;
} else if (N.getKind() == msgpack::Type::UInt) {
const MCExpr *NExpr = MCConstantExpr::create(N.getUInt(), Ctx);
Val = MCBinaryExpr::createOr(Val, NExpr, Ctx);
>From 73814b6b2d7c552f7b00a7dbcd5c7b1fd72f4bd5 Mon Sep 17 00:00:00 2001
From: Janek van Oirschot <janek.vanoirschot at amd.com>
Date: Fri, 7 Jun 2024 19:39:40 +0100
Subject: [PATCH 6/6] [AMDGPU] MCExpr-ify AMDGPU HSAMetadata
---
llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp | 47 ++++--
llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.h | 3 +-
.../AMDGPU/AMDGPUHSAMetadataStreamer.cpp | 34 ++--
.../Target/AMDGPU/AMDGPUHSAMetadataStreamer.h | 4 +
.../AMDGPU/AsmParser/AMDGPUAsmParser.cpp | 156 +++++++++++-------
.../MCTargetDesc/AMDGPUTargetStreamer.cpp | 54 +++---
.../MCTargetDesc/AMDGPUTargetStreamer.h | 15 +-
llvm/lib/Target/AMDGPU/SIProgramInfo.cpp | 39 -----
llvm/lib/Target/AMDGPU/SIProgramInfo.h | 4 -
.../AMDGPU/GlobalISel/flat-scratch-init.ll | 3 -
llvm/test/MC/AMDGPU/hsa-sym-expr-failure.s | 77 ---------
llvm/test/MC/AMDGPU/hsa-sym-exprs-gfx10.s | 48 +++---
llvm/test/MC/AMDGPU/hsa-sym-exprs-gfx11.s | 44 ++---
llvm/test/MC/AMDGPU/hsa-sym-exprs-gfx12.s | 42 ++---
llvm/test/MC/AMDGPU/hsa-sym-exprs-gfx7.s | 40 +++--
llvm/test/MC/AMDGPU/hsa-sym-exprs-gfx8.s | 40 +++--
llvm/test/MC/AMDGPU/hsa-sym-exprs-gfx90a.s | 46 +++---
llvm/test/MC/AMDGPU/hsa-tg-split.s | 2 +
llvm/test/MC/AMDGPU/hsa-v4.s | 1 +
.../MC/AMDGPU/hsa-v5-uses-dynamic-stack.s | 1 +
llvm/unittests/MC/AMDGPU/CMakeLists.txt | 3 -
.../MC/AMDGPU/SIProgramInfoMCExprs.cpp | 83 ----------
22 files changed, 346 insertions(+), 440 deletions(-)
delete mode 100644 llvm/unittests/MC/AMDGPU/SIProgramInfoMCExprs.cpp
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp b/llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp
index f0c7a5efb53a8..fdcd82a3528df 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp
@@ -249,14 +249,14 @@ void AMDGPUAsmPrinter::emitFunctionBodyEnd() {
getNameWithPrefix(KernelName, &MF->getFunction());
getTargetStreamer()->EmitAmdhsaKernelDescriptor(
STM, KernelName, getAmdhsaKernelDescriptor(*MF, CurrentProgramInfo),
- getMCExprValue(CurrentProgramInfo.NumVGPRsForWavesPerEU, Context),
- getMCExprValue(CurrentProgramInfo.NumSGPRsForWavesPerEU, Context) -
- IsaInfo::getNumExtraSGPRs(
- &STM, getMCExprValue(CurrentProgramInfo.VCCUsed, Context),
- getMCExprValue(CurrentProgramInfo.FlatUsed, Context),
- getTargetStreamer()->getTargetID()->isXnackOnOrAny()),
- getMCExprValue(CurrentProgramInfo.VCCUsed, Context),
- getMCExprValue(CurrentProgramInfo.FlatUsed, Context));
+ CurrentProgramInfo.NumVGPRsForWavesPerEU,
+ MCBinaryExpr::createSub(
+ CurrentProgramInfo.NumSGPRsForWavesPerEU,
+ AMDGPUVariadicMCExpr::createExtraSGPRs(
+ CurrentProgramInfo.VCCUsed, CurrentProgramInfo.FlatUsed,
+ getTargetStreamer()->getTargetID()->isXnackOnOrAny(), Context),
+ Context),
+ CurrentProgramInfo.VCCUsed, CurrentProgramInfo.FlatUsed);
Streamer.popSection();
}
@@ -401,9 +401,10 @@ void AMDGPUAsmPrinter::emitCommonFunctionComments(
false);
}
-uint16_t AMDGPUAsmPrinter::getAmdhsaKernelCodeProperties(
+const MCExpr *AMDGPUAsmPrinter::getAmdhsaKernelCodeProperties(
const MachineFunction &MF) const {
const SIMachineFunctionInfo &MFI = *MF.getInfo<SIMachineFunctionInfo>();
+ MCContext &Ctx = MF.getContext();
uint16_t KernelCodeProperties = 0;
const GCNUserSGPRUsageInfo &UserSGPRInfo = MFI.getUserSGPRInfo();
@@ -436,11 +437,19 @@ uint16_t AMDGPUAsmPrinter::getAmdhsaKernelCodeProperties(
amdhsa::KERNEL_CODE_PROPERTY_ENABLE_WAVEFRONT_SIZE32;
}
- if (getMCExprValue(CurrentProgramInfo.DynamicCallStack, MF.getContext()) &&
- CodeObjectVersion >= AMDGPU::AMDHSA_COV5)
- KernelCodeProperties |= amdhsa::KERNEL_CODE_PROPERTY_USES_DYNAMIC_STACK;
-
- return KernelCodeProperties;
+ // CurrentProgramInfo.DynamicCallStack is a MCExpr and could be
+ // un-evaluatable at this point so it cannot be conditionally checked here.
+ // Instead, we'll directly shift the possibly unknown MCExpr into its place
+ // and bitwise-or it into KernelCodeProperties.
+ const MCExpr *KernelCodePropExpr =
+ MCConstantExpr::create(KernelCodeProperties, Ctx);
+ const MCExpr *OrValue = MCConstantExpr::create(
+ amdhsa::KERNEL_CODE_PROPERTY_USES_DYNAMIC_STACK_SHIFT, Ctx);
+ OrValue = MCBinaryExpr::createShl(CurrentProgramInfo.DynamicCallStack,
+ OrValue, Ctx);
+ KernelCodePropExpr = MCBinaryExpr::createOr(KernelCodePropExpr, OrValue, Ctx);
+
+ return KernelCodePropExpr;
}
MCKernelDescriptor
@@ -463,11 +472,13 @@ AMDGPUAsmPrinter::getAmdhsaKernelDescriptor(const MachineFunction &MF,
KernelDescriptor.compute_pgm_rsrc1 = PI.getComputePGMRSrc1(STM, Ctx);
KernelDescriptor.compute_pgm_rsrc2 = PI.getComputePGMRSrc2(Ctx);
- KernelDescriptor.kernel_code_properties =
- MCConstantExpr::create(getAmdhsaKernelCodeProperties(MF), Ctx);
+ KernelDescriptor.kernel_code_properties = getAmdhsaKernelCodeProperties(MF);
- assert(STM.hasGFX90AInsts() ||
- getMCExprValue(CurrentProgramInfo.ComputePGMRSrc3GFX90A, Ctx) == 0);
+ int64_t PGRM_Rsrc3 = 1;
+ bool EvaluatableRsrc3 =
+ CurrentProgramInfo.ComputePGMRSrc3GFX90A->evaluateAsAbsolute(PGRM_Rsrc3);
+ assert(STM.hasGFX90AInsts() || !EvaluatableRsrc3 ||
+ static_cast<uint64_t>(PGRM_Rsrc3) == 0);
KernelDescriptor.compute_pgm_rsrc3 = CurrentProgramInfo.ComputePGMRSrc3GFX90A;
KernelDescriptor.kernarg_preload = MCConstantExpr::create(
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.h b/llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.h
index 87156f27fc6c5..12f6745fca7ee 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.h
+++ b/llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.h
@@ -69,8 +69,7 @@ class AMDGPUAsmPrinter final : public AsmPrinter {
const SIProgramInfo &CurrentProgramInfo,
bool isModuleEntryFunction, bool hasMAIInsts);
- uint16_t getAmdhsaKernelCodeProperties(
- const MachineFunction &MF) const;
+ const MCExpr *getAmdhsaKernelCodeProperties(const MachineFunction &MF) const;
AMDGPU::MCKernelDescriptor
getAmdhsaKernelDescriptor(const MachineFunction &MF,
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp b/llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp
index 7ab9ba2851332..efe47b2c3eed9 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp
@@ -464,16 +464,6 @@ MetadataStreamerMsgPackV4::getHSAKernelProps(const MachineFunction &MF,
const SIMachineFunctionInfo &MFI = *MF.getInfo<SIMachineFunctionInfo>();
const Function &F = MF.getFunction();
- auto GetMCExprValue = [&MF](const MCExpr *Value) {
- int64_t Val;
- if (!Value->evaluateAsAbsolute(Val)) {
- MCContext &Ctx = MF.getContext();
- Ctx.reportError(SMLoc(), "could not resolve expression when required.");
- Val = 0;
- }
- return static_cast<uint64_t>(Val);
- };
-
auto Kern = HSAMetadataDoc->getMapNode();
Align MaxKernArgAlign;
@@ -481,11 +471,12 @@ MetadataStreamerMsgPackV4::getHSAKernelProps(const MachineFunction &MF,
STM.getKernArgSegmentSize(F, MaxKernArgAlign));
Kern[".group_segment_fixed_size"] =
Kern.getDocument()->getNode(ProgramInfo.LDSSize);
- Kern[".private_segment_fixed_size"] =
- Kern.getDocument()->getNode(GetMCExprValue(ProgramInfo.ScratchSize));
+ DelayedExprs->assignDocNode(Kern[".private_segment_fixed_size"],
+ msgpack::Type::UInt, ProgramInfo.ScratchSize);
if (CodeObjectVersion >= AMDGPU::AMDHSA_COV5) {
- Kern[".uses_dynamic_stack"] = Kern.getDocument()->getNode(
- static_cast<bool>(GetMCExprValue(ProgramInfo.DynamicCallStack)));
+ DelayedExprs->assignDocNode(Kern[".uses_dynamic_stack"],
+ msgpack::Type::Boolean,
+ ProgramInfo.DynamicCallStack);
}
if (CodeObjectVersion >= AMDGPU::AMDHSA_COV5 && STM.supportsWGP())
@@ -497,15 +488,15 @@ MetadataStreamerMsgPackV4::getHSAKernelProps(const MachineFunction &MF,
Kern.getDocument()->getNode(std::max(Align(4), MaxKernArgAlign).value());
Kern[".wavefront_size"] =
Kern.getDocument()->getNode(STM.getWavefrontSize());
- Kern[".sgpr_count"] =
- Kern.getDocument()->getNode(GetMCExprValue(ProgramInfo.NumSGPR));
- Kern[".vgpr_count"] =
- Kern.getDocument()->getNode(GetMCExprValue(ProgramInfo.NumVGPR));
+ DelayedExprs->assignDocNode(Kern[".sgpr_count"], msgpack::Type::UInt,
+ ProgramInfo.NumSGPR);
+ DelayedExprs->assignDocNode(Kern[".vgpr_count"], msgpack::Type::UInt,
+ ProgramInfo.NumVGPR);
// Only add AGPR count to metadata for supported devices
if (STM.hasMAIInsts()) {
- Kern[".agpr_count"] =
- Kern.getDocument()->getNode(GetMCExprValue(ProgramInfo.NumAccVGPR));
+ DelayedExprs->assignDocNode(Kern[".agpr_count"], msgpack::Type::UInt,
+ ProgramInfo.NumAccVGPR);
}
Kern[".max_flat_workgroup_size"] =
@@ -527,6 +518,7 @@ MetadataStreamerMsgPackV4::getHSAKernelProps(const MachineFunction &MF,
}
bool MetadataStreamerMsgPackV4::emitTo(AMDGPUTargetStreamer &TargetStreamer) {
+ DelayedExprs->resolveDelayedExpressions();
return TargetStreamer.EmitHSAMetadata(*HSAMetadataDoc, true);
}
@@ -536,9 +528,11 @@ void MetadataStreamerMsgPackV4::begin(const Module &Mod,
emitTargetID(TargetID);
emitPrintf(Mod);
getRootMetadata("amdhsa.kernels") = HSAMetadataDoc->getArrayNode();
+ DelayedExprs->clear();
}
void MetadataStreamerMsgPackV4::end() {
+ DelayedExprs->resolveDelayedExpressions();
std::string HSAMetadataString;
raw_string_ostream StrOS(HSAMetadataString);
HSAMetadataDoc->toYAML(StrOS);
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.h b/llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.h
index 0e3bc63919f06..87de22dd0ca6a 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.h
+++ b/llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.h
@@ -15,6 +15,7 @@
#ifndef LLVM_LIB_TARGET_AMDGPU_MCTARGETDESC_AMDGPUHSAMETADATASTREAMER_H
#define LLVM_LIB_TARGET_AMDGPU_MCTARGETDESC_AMDGPUHSAMETADATASTREAMER_H
+#include "Utils/AMDGPUDelayedMCExpr.h"
#include "llvm/BinaryFormat/MsgPackDocument.h"
#include "llvm/Support/AMDGPUMetadata.h"
#include "llvm/Support/Alignment.h"
@@ -65,6 +66,9 @@ class MetadataStreamer {
class LLVM_EXTERNAL_VISIBILITY MetadataStreamerMsgPackV4
: public MetadataStreamer {
protected:
+ std::unique_ptr<DelayedMCExpr> DelayedExprs =
+ std::make_unique<DelayedMCExpr>();
+
std::unique_ptr<msgpack::Document> HSAMetadataDoc =
std::make_unique<msgpack::Document>();
diff --git a/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp b/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
index dcd4b22f4057a..a44d6ffe12fdf 100644
--- a/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
+++ b/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
@@ -1331,12 +1331,12 @@ class AMDGPUAsmParser : public MCTargetAsmParser {
/// \param SGPRRange [in] Token range, used for SGPR diagnostics.
/// \param VGPRBlocks [out] Result VGPR block count.
/// \param SGPRBlocks [out] Result SGPR block count.
- bool calculateGPRBlocks(const FeatureBitset &Features, bool VCCUsed,
- bool FlatScrUsed, bool XNACKUsed,
+ bool calculateGPRBlocks(const FeatureBitset &Features, const MCExpr *VCCUsed,
+ const MCExpr *FlatScrUsed, bool XNACKUsed,
std::optional<bool> EnableWavefrontSize32,
- unsigned NextFreeVGPR, SMRange VGPRRange,
- unsigned NextFreeSGPR, SMRange SGPRRange,
- unsigned &VGPRBlocks, unsigned &SGPRBlocks);
+ const MCExpr *NextFreeVGPR, SMRange VGPRRange,
+ const MCExpr *NextFreeSGPR, SMRange SGPRRange,
+ const MCExpr *&VGPRBlocks, const MCExpr *&SGPRBlocks);
bool ParseDirectiveAMDGCNTarget();
bool ParseDirectiveAMDHSACodeObjectVersion();
bool ParseDirectiveAMDHSAKernel();
@@ -5352,41 +5352,65 @@ bool AMDGPUAsmParser::OutOfRangeError(SMRange Range) {
}
bool AMDGPUAsmParser::calculateGPRBlocks(
- const FeatureBitset &Features, bool VCCUsed, bool FlatScrUsed,
- bool XNACKUsed, std::optional<bool> EnableWavefrontSize32,
- unsigned NextFreeVGPR, SMRange VGPRRange, unsigned NextFreeSGPR,
- SMRange SGPRRange, unsigned &VGPRBlocks, unsigned &SGPRBlocks) {
+ const FeatureBitset &Features, const MCExpr *VCCUsed,
+ const MCExpr *FlatScrUsed, bool XNACKUsed,
+ std::optional<bool> EnableWavefrontSize32, const MCExpr *NextFreeVGPR,
+ SMRange VGPRRange, const MCExpr *NextFreeSGPR, SMRange SGPRRange,
+ const MCExpr *&VGPRBlocks, const MCExpr *&SGPRBlocks) {
// TODO(scott.linder): These calculations are duplicated from
// AMDGPUAsmPrinter::getSIProgramInfo and could be unified.
IsaVersion Version = getIsaVersion(getSTI().getCPU());
+ MCContext &Ctx = getContext();
- unsigned NumVGPRs = NextFreeVGPR;
- unsigned NumSGPRs = NextFreeSGPR;
+ const MCExpr *NumSGPRs = NextFreeSGPR;
+ int64_t evaluatedSGPRs;
if (Version.Major >= 10)
- NumSGPRs = 0;
+ NumSGPRs = MCConstantExpr::create(0, Ctx);
else {
unsigned MaxAddressableNumSGPRs =
IsaInfo::getAddressableNumSGPRs(&getSTI());
- if (Version.Major >= 8 && !Features.test(FeatureSGPRInitBug) &&
- NumSGPRs > MaxAddressableNumSGPRs)
+ if (NumSGPRs->evaluateAsAbsolute(evaluatedSGPRs) && Version.Major >= 8 &&
+ !Features.test(FeatureSGPRInitBug) &&
+ static_cast<uint64_t>(evaluatedSGPRs) > MaxAddressableNumSGPRs)
return OutOfRangeError(SGPRRange);
- NumSGPRs +=
- IsaInfo::getNumExtraSGPRs(&getSTI(), VCCUsed, FlatScrUsed, XNACKUsed);
+ const MCExpr *ExtraSGPRs = AMDGPUVariadicMCExpr::createExtraSGPRs(
+ VCCUsed, FlatScrUsed, XNACKUsed, Ctx);
+ NumSGPRs = MCBinaryExpr::createAdd(NumSGPRs, ExtraSGPRs, Ctx);
- if ((Version.Major <= 7 || Features.test(FeatureSGPRInitBug)) &&
- NumSGPRs > MaxAddressableNumSGPRs)
+ if (NumSGPRs->evaluateAsAbsolute(evaluatedSGPRs) &&
+ (Version.Major <= 7 || Features.test(FeatureSGPRInitBug)) &&
+ static_cast<uint64_t>(evaluatedSGPRs) > MaxAddressableNumSGPRs)
return OutOfRangeError(SGPRRange);
if (Features.test(FeatureSGPRInitBug))
- NumSGPRs = IsaInfo::FIXED_NUM_SGPRS_FOR_INIT_BUG;
- }
+ NumSGPRs =
+ MCConstantExpr::create(IsaInfo::FIXED_NUM_SGPRS_FOR_INIT_BUG, Ctx);
+ }
+
+ // The MCExpr equivalent of getNumSGPRBlocks/getNumVGPRBlocks:
+ // (alignTo(max(1u, NumGPR), GPREncodingGranule) / GPREncodingGranule) - 1
+ auto GetNumGPRBlocks = [&Ctx](const MCExpr *NumGPR,
+ unsigned Granule) -> const MCExpr * {
+ const MCExpr *OneConst = MCConstantExpr::create(1ul, Ctx);
+ const MCExpr *GranuleConst = MCConstantExpr::create(Granule, Ctx);
+ const MCExpr *MaxNumGPR =
+ AMDGPUVariadicMCExpr::createMax({NumGPR, OneConst}, Ctx);
+ const MCExpr *AlignToGPR =
+ AMDGPUVariadicMCExpr::createAlignTo(MaxNumGPR, GranuleConst, Ctx);
+ const MCExpr *DivGPR =
+ MCBinaryExpr::createDiv(AlignToGPR, GranuleConst, Ctx);
+ const MCExpr *SubGPR = MCBinaryExpr::createSub(DivGPR, OneConst, Ctx);
+ return SubGPR;
+ };
- VGPRBlocks = IsaInfo::getEncodedNumVGPRBlocks(&getSTI(), NumVGPRs,
- EnableWavefrontSize32);
- SGPRBlocks = IsaInfo::getNumSGPRBlocks(&getSTI(), NumSGPRs);
+ VGPRBlocks = GetNumGPRBlocks(
+ NextFreeVGPR,
+ IsaInfo::getVGPREncodingGranule(&getSTI(), EnableWavefrontSize32));
+ SGPRBlocks =
+ GetNumGPRBlocks(NumSGPRs, IsaInfo::getSGPREncodingGranule(&getSTI()));
return false;
}
@@ -5410,14 +5434,17 @@ bool AMDGPUAsmParser::ParseDirectiveAMDHSAKernel() {
IsaVersion IVersion = getIsaVersion(getSTI().getCPU());
+ const MCExpr *ZeroExpr = MCConstantExpr::create(0, getContext());
+ const MCExpr *OneExpr = MCConstantExpr::create(1, getContext());
+
SMRange VGPRRange;
- uint64_t NextFreeVGPR = 0;
- uint64_t AccumOffset = 0;
+ const MCExpr *NextFreeVGPR = ZeroExpr;
+ const MCExpr *AccumOffset = MCConstantExpr::create(0, getContext());
uint64_t SharedVGPRCount = 0;
uint64_t PreloadLength = 0;
uint64_t PreloadOffset = 0;
SMRange SGPRRange;
- uint64_t NextFreeSGPR = 0;
+ const MCExpr *NextFreeSGPR = ZeroExpr;
// Count the number of user SGPRs implied from the enabled feature bits.
unsigned ImpliedUserSGPRCount = 0;
@@ -5425,8 +5452,8 @@ bool AMDGPUAsmParser::ParseDirectiveAMDHSAKernel() {
// Track if the asm explicitly contains the directive for the user SGPR
// count.
std::optional<unsigned> ExplicitUserSGPRCount;
- bool ReserveVCC = true;
- bool ReserveFlatScr = true;
+ const MCExpr *ReserveVCC = OneExpr;
+ const MCExpr *ReserveFlatScr = OneExpr;
std::optional<bool> EnableWavefrontSize32;
while (true) {
@@ -5620,34 +5647,29 @@ bool AMDGPUAsmParser::ParseDirectiveAMDHSAKernel() {
COMPUTE_PGM_RSRC2_ENABLE_VGPR_WORKITEM_ID, ExprVal,
ValRange);
} else if (ID == ".amdhsa_next_free_vgpr") {
- EXPR_RESOLVE_OR_ERROR(EvaluatableExpr);
VGPRRange = ValRange;
- NextFreeVGPR = Val;
+ NextFreeVGPR = ExprVal;
} else if (ID == ".amdhsa_next_free_sgpr") {
- EXPR_RESOLVE_OR_ERROR(EvaluatableExpr);
SGPRRange = ValRange;
- NextFreeSGPR = Val;
+ NextFreeSGPR = ExprVal;
} else if (ID == ".amdhsa_accum_offset") {
if (!isGFX90A())
return Error(IDRange.Start, "directive requires gfx90a+", IDRange);
- EXPR_RESOLVE_OR_ERROR(EvaluatableExpr);
- AccumOffset = Val;
+ AccumOffset = ExprVal;
} else if (ID == ".amdhsa_reserve_vcc") {
- EXPR_RESOLVE_OR_ERROR(EvaluatableExpr);
- if (!isUInt<1>(Val))
+ if (EvaluatableExpr && !isUInt<1>(Val))
return OutOfRangeError(ValRange);
- ReserveVCC = Val;
+ ReserveVCC = ExprVal;
} else if (ID == ".amdhsa_reserve_flat_scratch") {
- EXPR_RESOLVE_OR_ERROR(EvaluatableExpr);
if (IVersion.Major < 7)
return Error(IDRange.Start, "directive requires gfx7+", IDRange);
if (hasArchitectedFlatScratch())
return Error(IDRange.Start,
"directive is not supported with architected flat scratch",
IDRange);
- if (!isUInt<1>(Val))
+ if (EvaluatableExpr && !isUInt<1>(Val))
return OutOfRangeError(ValRange);
- ReserveFlatScr = Val;
+ ReserveFlatScr = ExprVal;
} else if (ID == ".amdhsa_reserve_xnack_mask") {
if (IVersion.Major < 8)
return Error(IDRange.Start, "directive requires gfx8+", IDRange);
@@ -5771,8 +5793,8 @@ bool AMDGPUAsmParser::ParseDirectiveAMDHSAKernel() {
if (!Seen.contains(".amdhsa_next_free_sgpr"))
return TokError(".amdhsa_next_free_sgpr directive is required");
- unsigned VGPRBlocks;
- unsigned SGPRBlocks;
+ const MCExpr *VGPRBlocks;
+ const MCExpr *SGPRBlocks;
if (calculateGPRBlocks(getFeatureBits(), ReserveVCC, ReserveFlatScr,
getTargetStreamer().getTargetID()->isXnackOnOrAny(),
EnableWavefrontSize32, NextFreeVGPR,
@@ -5780,19 +5802,26 @@ bool AMDGPUAsmParser::ParseDirectiveAMDHSAKernel() {
SGPRBlocks))
return true;
- if (!isUInt<COMPUTE_PGM_RSRC1_GRANULATED_WORKITEM_VGPR_COUNT_WIDTH>(
- VGPRBlocks))
+ int64_t EvaluatedVGPRBlocks;
+ bool VGPRBlocksEvaluatable =
+ VGPRBlocks->evaluateAsAbsolute(EvaluatedVGPRBlocks);
+ if (VGPRBlocksEvaluatable &&
+ !isUInt<COMPUTE_PGM_RSRC1_GRANULATED_WORKITEM_VGPR_COUNT_WIDTH>(
+ static_cast<uint64_t>(EvaluatedVGPRBlocks))) {
return OutOfRangeError(VGPRRange);
+ }
AMDGPU::MCKernelDescriptor::bits_set(
- KD.compute_pgm_rsrc1, MCConstantExpr::create(VGPRBlocks, getContext()),
+ KD.compute_pgm_rsrc1, VGPRBlocks,
COMPUTE_PGM_RSRC1_GRANULATED_WORKITEM_VGPR_COUNT_SHIFT,
COMPUTE_PGM_RSRC1_GRANULATED_WORKITEM_VGPR_COUNT, getContext());
- if (!isUInt<COMPUTE_PGM_RSRC1_GRANULATED_WAVEFRONT_SGPR_COUNT_WIDTH>(
- SGPRBlocks))
+ int64_t EvaluatedSGPRBlocks;
+ if (SGPRBlocks->evaluateAsAbsolute(EvaluatedSGPRBlocks) &&
+ !isUInt<COMPUTE_PGM_RSRC1_GRANULATED_WAVEFRONT_SGPR_COUNT_WIDTH>(
+ static_cast<uint64_t>(EvaluatedSGPRBlocks)))
return OutOfRangeError(SGPRRange);
AMDGPU::MCKernelDescriptor::bits_set(
- KD.compute_pgm_rsrc1, MCConstantExpr::create(SGPRBlocks, getContext()),
+ KD.compute_pgm_rsrc1, SGPRBlocks,
COMPUTE_PGM_RSRC1_GRANULATED_WAVEFRONT_SGPR_COUNT_SHIFT,
COMPUTE_PGM_RSRC1_GRANULATED_WAVEFRONT_SGPR_COUNT, getContext());
@@ -5822,16 +5851,28 @@ bool AMDGPUAsmParser::ParseDirectiveAMDHSAKernel() {
if (isGFX90A()) {
if (!Seen.contains(".amdhsa_accum_offset"))
return TokError(".amdhsa_accum_offset directive is required");
- if (AccumOffset < 4 || AccumOffset > 256 || (AccumOffset & 3))
+ int64_t EvaluatedAccum;
+ bool AccumEvaluatable = AccumOffset->evaluateAsAbsolute(EvaluatedAccum);
+ uint64_t UEvaluatedAccum = EvaluatedAccum;
+ if (AccumEvaluatable &&
+ (UEvaluatedAccum < 4 || UEvaluatedAccum > 256 || (UEvaluatedAccum & 3)))
return TokError("accum_offset should be in range [4..256] in "
"increments of 4");
- if (AccumOffset > alignTo(std::max((uint64_t)1, NextFreeVGPR), 4))
+
+ int64_t EvaluatedNumVGPR;
+ if (NextFreeVGPR->evaluateAsAbsolute(EvaluatedNumVGPR) &&
+ AccumEvaluatable &&
+ UEvaluatedAccum >
+ alignTo(std::max((uint64_t)1, (uint64_t)EvaluatedNumVGPR), 4))
return TokError("accum_offset exceeds total VGPR allocation");
- MCKernelDescriptor::bits_set(
- KD.compute_pgm_rsrc3,
- MCConstantExpr::create(AccumOffset / 4 - 1, getContext()),
- COMPUTE_PGM_RSRC3_GFX90A_ACCUM_OFFSET_SHIFT,
- COMPUTE_PGM_RSRC3_GFX90A_ACCUM_OFFSET, getContext());
+ const MCExpr *AdjustedAccum = MCBinaryExpr::createSub(
+ MCBinaryExpr::createDiv(
+ AccumOffset, MCConstantExpr::create(4, getContext()), getContext()),
+ MCConstantExpr::create(1, getContext()), getContext());
+ MCKernelDescriptor::bits_set(KD.compute_pgm_rsrc3, AdjustedAccum,
+ COMPUTE_PGM_RSRC3_GFX90A_ACCUM_OFFSET_SHIFT,
+ COMPUTE_PGM_RSRC3_GFX90A_ACCUM_OFFSET,
+ getContext());
}
if (IVersion.Major >= 10 && IVersion.Major < 12) {
@@ -5840,7 +5881,10 @@ bool AMDGPUAsmParser::ParseDirectiveAMDHSAKernel() {
return TokError("shared_vgpr_count directive not valid on "
"wavefront size 32");
}
- if (SharedVGPRCount * 2 + VGPRBlocks > 63) {
+
+ if (VGPRBlocksEvaluatable &&
+ (SharedVGPRCount * 2 + static_cast<uint64_t>(EvaluatedVGPRBlocks) >
+ 63)) {
return TokError("shared_vgpr_count*2 + "
"compute_pgm_rsrc1.GRANULATED_WORKITEM_VGPR_COUNT cannot "
"exceed 63\n");
diff --git a/llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.cpp b/llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.cpp
index 00e64e3419ba0..52db390f41d77 100644
--- a/llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.cpp
+++ b/llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.cpp
@@ -315,8 +315,9 @@ bool AMDGPUTargetAsmStreamer::EmitCodeEnd(const MCSubtargetInfo &STI) {
void AMDGPUTargetAsmStreamer::EmitAmdhsaKernelDescriptor(
const MCSubtargetInfo &STI, StringRef KernelName,
- const MCKernelDescriptor &KD, uint64_t NextVGPR, uint64_t NextSGPR,
- bool ReserveVCC, bool ReserveFlatScr) {
+ const MCKernelDescriptor &KD, const MCExpr *NextVGPR,
+ const MCExpr *NextSGPR, const MCExpr *ReserveVCC,
+ const MCExpr *ReserveFlatScr) {
IsaVersion IVersion = getIsaVersion(STI.getCPU());
const MCAsmInfo *MAI = getContext().getAsmInfo();
@@ -335,16 +336,25 @@ void AMDGPUTargetAsmStreamer::EmitAmdhsaKernelDescriptor(
OS << '\n';
};
+ auto EmitMCExpr = [&](const MCExpr *Value) {
+ int64_t evaluatableValue;
+ if (Value->evaluateAsAbsolute(evaluatableValue)) {
+ OS << static_cast<uint64_t>(evaluatableValue);
+ } else {
+ Value->print(OS, MAI);
+ }
+ };
+
OS << "\t\t.amdhsa_group_segment_fixed_size ";
- KD.group_segment_fixed_size->print(OS, MAI);
+ EmitMCExpr(KD.group_segment_fixed_size);
OS << '\n';
OS << "\t\t.amdhsa_private_segment_fixed_size ";
- KD.private_segment_fixed_size->print(OS, MAI);
+ EmitMCExpr(KD.private_segment_fixed_size);
OS << '\n';
OS << "\t\t.amdhsa_kernarg_size ";
- KD.kernarg_size->print(OS, MAI);
+ EmitMCExpr(KD.kernarg_size);
OS << '\n';
PrintField(
@@ -429,8 +439,13 @@ void AMDGPUTargetAsmStreamer::EmitAmdhsaKernelDescriptor(
".amdhsa_system_vgpr_workitem_id");
// These directives are required.
- OS << "\t\t.amdhsa_next_free_vgpr " << NextVGPR << '\n';
- OS << "\t\t.amdhsa_next_free_sgpr " << NextSGPR << '\n';
+ OS << "\t\t.amdhsa_next_free_vgpr ";
+ EmitMCExpr(NextVGPR);
+ OS << '\n';
+
+ OS << "\t\t.amdhsa_next_free_sgpr ";
+ EmitMCExpr(NextSGPR);
+ OS << '\n';
if (AMDGPU::isGFX90A(STI)) {
// MCExpr equivalent of taking the (accum_offset + 1) * 4.
@@ -443,19 +458,19 @@ void AMDGPUTargetAsmStreamer::EmitAmdhsaKernelDescriptor(
accum_bits = MCBinaryExpr::createMul(
accum_bits, MCConstantExpr::create(4, getContext()), getContext());
OS << "\t\t.amdhsa_accum_offset ";
- int64_t IVal;
- if (accum_bits->evaluateAsAbsolute(IVal)) {
- OS << static_cast<uint64_t>(IVal);
- } else {
- accum_bits->print(OS, MAI);
- }
+ EmitMCExpr(accum_bits);
OS << '\n';
}
- if (!ReserveVCC)
- OS << "\t\t.amdhsa_reserve_vcc " << ReserveVCC << '\n';
- if (IVersion.Major >= 7 && !ReserveFlatScr && !hasArchitectedFlatScratch(STI))
- OS << "\t\t.amdhsa_reserve_flat_scratch " << ReserveFlatScr << '\n';
+ OS << "\t\t.amdhsa_reserve_vcc ";
+ EmitMCExpr(ReserveVCC);
+ OS << '\n';
+
+ if (IVersion.Major >= 7 && !hasArchitectedFlatScratch(STI)) {
+ OS << "\t\t.amdhsa_reserve_flat_scratch ";
+ EmitMCExpr(ReserveFlatScr);
+ OS << '\n';
+ }
switch (CodeObjectVersion) {
default:
@@ -908,8 +923,9 @@ bool AMDGPUTargetELFStreamer::EmitCodeEnd(const MCSubtargetInfo &STI) {
void AMDGPUTargetELFStreamer::EmitAmdhsaKernelDescriptor(
const MCSubtargetInfo &STI, StringRef KernelName,
- const MCKernelDescriptor &KernelDescriptor, uint64_t NextVGPR,
- uint64_t NextSGPR, bool ReserveVCC, bool ReserveFlatScr) {
+ const MCKernelDescriptor &KernelDescriptor, const MCExpr *NextVGPR,
+ const MCExpr *NextSGPR, const MCExpr *ReserveVCC,
+ const MCExpr *ReserveFlatScr) {
auto &Streamer = getStreamer();
auto &Context = Streamer.getContext();
diff --git a/llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.h b/llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.h
index e5c90060cb5d0..bf1538c71d154 100644
--- a/llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.h
+++ b/llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.h
@@ -94,8 +94,9 @@ class AMDGPUTargetStreamer : public MCTargetStreamer {
virtual void
EmitAmdhsaKernelDescriptor(const MCSubtargetInfo &STI, StringRef KernelName,
const AMDGPU::MCKernelDescriptor &KernelDescriptor,
- uint64_t NextVGPR, uint64_t NextSGPR,
- bool ReserveVCC, bool ReserveFlatScr) {}
+ const MCExpr *NextVGPR, const MCExpr *NextSGPR,
+ const MCExpr *ReserveVCC,
+ const MCExpr *ReserveFlatScr) {}
static StringRef getArchNameFromElfMach(unsigned ElfMach);
static unsigned getElfMach(StringRef GPU);
@@ -151,8 +152,9 @@ class AMDGPUTargetAsmStreamer final : public AMDGPUTargetStreamer {
void
EmitAmdhsaKernelDescriptor(const MCSubtargetInfo &STI, StringRef KernelName,
const AMDGPU::MCKernelDescriptor &KernelDescriptor,
- uint64_t NextVGPR, uint64_t NextSGPR,
- bool ReserveVCC, bool ReserveFlatScr) override;
+ const MCExpr *NextVGPR, const MCExpr *NextSGPR,
+ const MCExpr *ReserveVCC,
+ const MCExpr *ReserveFlatScr) override;
};
class AMDGPUTargetELFStreamer final : public AMDGPUTargetStreamer {
@@ -207,8 +209,9 @@ class AMDGPUTargetELFStreamer final : public AMDGPUTargetStreamer {
void
EmitAmdhsaKernelDescriptor(const MCSubtargetInfo &STI, StringRef KernelName,
const AMDGPU::MCKernelDescriptor &KernelDescriptor,
- uint64_t NextVGPR, uint64_t NextSGPR,
- bool ReserveVCC, bool ReserveFlatScr) override;
+ const MCExpr *NextVGPR, const MCExpr *NextSGPR,
+ const MCExpr *ReserveVCC,
+ const MCExpr *ReserveFlatScr) override;
};
}
#endif
diff --git a/llvm/lib/Target/AMDGPU/SIProgramInfo.cpp b/llvm/lib/Target/AMDGPU/SIProgramInfo.cpp
index 0d40816cdd4b8..212edff097837 100644
--- a/llvm/lib/Target/AMDGPU/SIProgramInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/SIProgramInfo.cpp
@@ -161,45 +161,6 @@ static const MCExpr *MaskShift(const MCExpr *Val, uint32_t Mask, uint32_t Shift,
return Val;
}
-uint64_t SIProgramInfo::getComputePGMRSrc1(const GCNSubtarget &ST) const {
- int64_t VBlocks, SBlocks;
- VGPRBlocks->evaluateAsAbsolute(VBlocks);
- SGPRBlocks->evaluateAsAbsolute(SBlocks);
-
- uint64_t Reg = S_00B848_VGPRS(static_cast<uint64_t>(VBlocks)) |
- S_00B848_SGPRS(static_cast<uint64_t>(SBlocks)) |
- getComputePGMRSrc1Reg(*this, ST);
-
- return Reg;
-}
-
-uint64_t SIProgramInfo::getPGMRSrc1(CallingConv::ID CC,
- const GCNSubtarget &ST) const {
- if (AMDGPU::isCompute(CC)) {
- return getComputePGMRSrc1(ST);
- }
- int64_t VBlocks, SBlocks;
- VGPRBlocks->evaluateAsAbsolute(VBlocks);
- SGPRBlocks->evaluateAsAbsolute(SBlocks);
-
- return getPGMRSrc1Reg(*this, CC, ST) |
- S_00B848_VGPRS(static_cast<uint64_t>(VBlocks)) |
- S_00B848_SGPRS(static_cast<uint64_t>(SBlocks));
-}
-
-uint64_t SIProgramInfo::getComputePGMRSrc2() const {
- int64_t ScratchEn;
- ScratchEnable->evaluateAsAbsolute(ScratchEn);
- return ScratchEn | getComputePGMRSrc2Reg(*this);
-}
-
-uint64_t SIProgramInfo::getPGMRSrc2(CallingConv::ID CC) const {
- if (AMDGPU::isCompute(CC))
- return getComputePGMRSrc2();
-
- return 0;
-}
-
const MCExpr *SIProgramInfo::getComputePGMRSrc1(const GCNSubtarget &ST,
MCContext &Ctx) const {
uint64_t Reg = getComputePGMRSrc1Reg(*this, ST);
diff --git a/llvm/lib/Target/AMDGPU/SIProgramInfo.h b/llvm/lib/Target/AMDGPU/SIProgramInfo.h
index e66e5a194c8b5..c358a2d9db10b 100644
--- a/llvm/lib/Target/AMDGPU/SIProgramInfo.h
+++ b/llvm/lib/Target/AMDGPU/SIProgramInfo.h
@@ -98,16 +98,12 @@ struct LLVM_EXTERNAL_VISIBILITY SIProgramInfo {
void reset(const MachineFunction &MF);
/// Compute the value of the ComputePGMRsrc1 register.
- uint64_t getComputePGMRSrc1(const GCNSubtarget &ST) const;
- uint64_t getPGMRSrc1(CallingConv::ID CC, const GCNSubtarget &ST) const;
const MCExpr *getComputePGMRSrc1(const GCNSubtarget &ST,
MCContext &Ctx) const;
const MCExpr *getPGMRSrc1(CallingConv::ID CC, const GCNSubtarget &ST,
MCContext &Ctx) const;
/// Compute the value of the ComputePGMRsrc2 register.
- uint64_t getComputePGMRSrc2() const;
- uint64_t getPGMRSrc2(CallingConv::ID CC) const;
const MCExpr *getComputePGMRSrc2(MCContext &Ctx) const;
const MCExpr *getPGMRSrc2(CallingConv::ID CC, MCContext &Ctx) const;
};
diff --git a/llvm/test/CodeGen/AMDGPU/GlobalISel/flat-scratch-init.ll b/llvm/test/CodeGen/AMDGPU/GlobalISel/flat-scratch-init.ll
index a8aa6c780b86a..7cd99fcfd5e74 100644
--- a/llvm/test/CodeGen/AMDGPU/GlobalISel/flat-scratch-init.ll
+++ b/llvm/test/CodeGen/AMDGPU/GlobalISel/flat-scratch-init.ll
@@ -16,7 +16,6 @@
; RW-FLAT-NOT: .amdhsa_enable_private_segment
; RO-FLAT-NOT: .amdhsa_system_sgpr_private_segment_wavefront_offset
; RO-FLAT: .amdhsa_enable_private_segment 1
-; GCN-NOT: .amdhsa_reserve_flat_scratch
; GCN: COMPUTE_PGM_RSRC2:SCRATCH_EN: 1
; RW-FLAT: COMPUTE_PGM_RSRC2:USER_SGPR: 6
; RO-FLAT: COMPUTE_PGM_RSRC2:USER_SGPR: 0
@@ -41,7 +40,6 @@ define amdgpu_kernel void @stack_object_addrspacecast_in_kernel_no_calls() {
; RO-FLAT-NOT: .amdhsa_system_sgpr_private_segment_wavefront_offset
; RO-FLAT: .amdhsa_enable_private_segment 1
; RW-FLAT: .amdhsa_reserve_flat_scratch 0
-; RO-FLAT-NOT: .amdhsa_reserve_flat_scratch
; GCN: COMPUTE_PGM_RSRC2:SCRATCH_EN: 1
; RW-FLAT: COMPUTE_PGM_RSRC2:USER_SGPR: 6
; RO-FLAT: COMPUTE_PGM_RSRC2:USER_SGPR: 0
@@ -62,7 +60,6 @@ define amdgpu_kernel void @stack_object_in_kernel_no_calls() {
; RO-FLAT-NOT: .amdhsa_system_sgpr_private_segment_wavefront_offset
; RO-FLAT: .amdhsa_enable_private_segment 0
; RW-FLAT: .amdhsa_reserve_flat_scratch 0
-; RO-FLAT-NOT: .amdhsa_reserve_flat_scratch 0
; GCN: COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; RW-FLAT: COMPUTE_PGM_RSRC2:USER_SGPR: 4
; RO-FLAT: COMPUTE_PGM_RSRC2:USER_SGPR: 0
diff --git a/llvm/test/MC/AMDGPU/hsa-sym-expr-failure.s b/llvm/test/MC/AMDGPU/hsa-sym-expr-failure.s
index fab3e893352b2..4515a86a09154 100644
--- a/llvm/test/MC/AMDGPU/hsa-sym-expr-failure.s
+++ b/llvm/test/MC/AMDGPU/hsa-sym-expr-failure.s
@@ -182,83 +182,6 @@ wavefront_size32:
// ASM: error: directive should have resolvable expression
// ASM-NEXT: .amdhsa_wavefront_size32
-.p2align 8
-.type next_free_vgpr, at function
-next_free_vgpr:
- s_endpgm
-
-.p2align 6
-.amdhsa_kernel next_free_vgpr
- .amdhsa_next_free_vgpr defined_boolean
- .amdhsa_next_free_sgpr 0
- .amdhsa_accum_offset 4
-.end_amdhsa_kernel
-
-// ASM: error: directive should have resolvable expression
-// ASM-NEXT: .amdhsa_next_free_vgpr
-
-.p2align 8
-.type next_free_sgpr, at function
-next_free_sgpr:
- s_endpgm
-
-.p2align 6
-.amdhsa_kernel next_free_sgpr
- .amdhsa_next_free_vgpr 0
- .amdhsa_next_free_sgpr defined_boolean
- .amdhsa_accum_offset 4
-.end_amdhsa_kernel
-
-// ASM: error: directive should have resolvable expression
-// ASM-NEXT: .amdhsa_next_free_sgpr
-
-.p2align 8
-.type accum_offset, at function
-accum_offset:
- s_endpgm
-
-.p2align 6
-.amdhsa_kernel accum_offset
- .amdhsa_next_free_vgpr 0
- .amdhsa_next_free_sgpr 0
- .amdhsa_accum_offset defined_boolean
-.end_amdhsa_kernel
-
-// ASM: error: directive should have resolvable expression
-// ASM-NEXT: .amdhsa_accum_offset
-
-.p2align 8
-.type reserve_vcc, at function
-reserve_vcc:
- s_endpgm
-
-.p2align 6
-.amdhsa_kernel reserve_vcc
- .amdhsa_next_free_vgpr 0
- .amdhsa_next_free_sgpr 0
- .amdhsa_accum_offset 4
- .amdhsa_reserve_vcc defined_boolean
-.end_amdhsa_kernel
-
-// ASM: error: directive should have resolvable expression
-// ASM-NEXT: .amdhsa_reserve_vcc
-
-.p2align 8
-.type reserve_flat_scratch, at function
-reserve_flat_scratch:
- s_endpgm
-
-.p2align 6
-.amdhsa_kernel reserve_flat_scratch
- .amdhsa_next_free_vgpr 0
- .amdhsa_next_free_sgpr 0
- .amdhsa_accum_offset 4
- .amdhsa_reserve_flat_scratch defined_boolean
-.end_amdhsa_kernel
-
-// ASM: error: directive should have resolvable expression
-// ASM-NEXT: .amdhsa_reserve_flat_scratch
-
.p2align 8
.type shared_vgpr_count, at function
shared_vgpr_count:
diff --git a/llvm/test/MC/AMDGPU/hsa-sym-exprs-gfx10.s b/llvm/test/MC/AMDGPU/hsa-sym-exprs-gfx10.s
index 95af59c413ae6..af4cb1a008f9d 100644
--- a/llvm/test/MC/AMDGPU/hsa-sym-exprs-gfx10.s
+++ b/llvm/test/MC/AMDGPU/hsa-sym-exprs-gfx10.s
@@ -10,12 +10,12 @@
// OBJDUMP-NEXT: 0000 2b000000 2c000000 00000000 00000000
// OBJDUMP-NEXT: 0010 00000000 00000000 00000000 00000000
// OBJDUMP-NEXT: 0020 00000000 00000000 00000000 00000000
-// OBJDUMP-NEXT: 0030 00f0afe4 801f007f 000c0000 00000000
+// OBJDUMP-NEXT: 0030 05f0afe4 801f007f 000c0000 00000000
// expr_defined
// OBJDUMP-NEXT: 0040 2a000000 2b000000 00000000 00000000
// OBJDUMP-NEXT: 0050 00000000 00000000 00000000 00000000
// OBJDUMP-NEXT: 0060 00000000 00000000 00000000 00000000
-// OBJDUMP-NEXT: 0070 00f0afe4 801f007f 000c0000 00000000
+// OBJDUMP-NEXT: 0070 05f0afe4 801f007f 000c0000 00000000
.text
// ASM: .text
@@ -61,8 +61,10 @@ expr_defined:
.amdhsa_exception_fp_ieee_inexact defined_boolean
.amdhsa_exception_int_div_zero defined_boolean
.amdhsa_uses_dynamic_stack defined_boolean
- .amdhsa_next_free_vgpr 0
- .amdhsa_next_free_sgpr 0
+ .amdhsa_next_free_vgpr defined_value+4
+ .amdhsa_next_free_sgpr defined_value+5
+ .amdhsa_reserve_vcc defined_boolean
+ .amdhsa_reserve_flat_scratch defined_boolean
.end_amdhsa_kernel
.set defined_value, 41
@@ -94,8 +96,10 @@ expr_defined:
.amdhsa_exception_fp_ieee_inexact defined_boolean
.amdhsa_exception_int_div_zero defined_boolean
.amdhsa_uses_dynamic_stack defined_boolean
- .amdhsa_next_free_vgpr 0
- .amdhsa_next_free_sgpr 0
+ .amdhsa_next_free_vgpr defined_value+3
+ .amdhsa_next_free_sgpr defined_value+4
+ .amdhsa_reserve_vcc defined_boolean
+ .amdhsa_reserve_flat_scratch defined_boolean
.end_amdhsa_kernel
// ASM: .amdhsa_kernel expr_defined_later
@@ -117,19 +121,21 @@ expr_defined:
// ASM-NEXT: .amdhsa_system_sgpr_workgroup_id_z (((((((((((((((((((((((((((((0&(~128))|(1<<7))&(~6144))|(defined_2_bits<<11))&(~128))|(defined_boolean<<7))&(~256))|(defined_boolean<<8))&(~512))|(defined_boolean<<9))&(~1024))|(defined_boolean<<10))&(~16777216))|(defined_boolean<<24))&(~33554432))|(defined_boolean<<25))&(~67108864))|(defined_boolean<<26))&(~134217728))|(defined_boolean<<27))&(~268435456))|(defined_boolean<<28))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~62))|(0<<1))&512)>>9
// ASM-NEXT: .amdhsa_system_sgpr_workgroup_info (((((((((((((((((((((((((((((0&(~128))|(1<<7))&(~6144))|(defined_2_bits<<11))&(~128))|(defined_boolean<<7))&(~256))|(defined_boolean<<8))&(~512))|(defined_boolean<<9))&(~1024))|(defined_boolean<<10))&(~16777216))|(defined_boolean<<24))&(~33554432))|(defined_boolean<<25))&(~67108864))|(defined_boolean<<26))&(~134217728))|(defined_boolean<<27))&(~268435456))|(defined_boolean<<28))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~62))|(0<<1))&1024)>>10
// ASM-NEXT: .amdhsa_system_vgpr_workitem_id (((((((((((((((((((((((((((((0&(~128))|(1<<7))&(~6144))|(defined_2_bits<<11))&(~128))|(defined_boolean<<7))&(~256))|(defined_boolean<<8))&(~512))|(defined_boolean<<9))&(~1024))|(defined_boolean<<10))&(~16777216))|(defined_boolean<<24))&(~33554432))|(defined_boolean<<25))&(~67108864))|(defined_boolean<<26))&(~134217728))|(defined_boolean<<27))&(~268435456))|(defined_boolean<<28))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~62))|(0<<1))&6144)>>11
-// ASM-NEXT: .amdhsa_next_free_vgpr 0
-// ASM-NEXT: .amdhsa_next_free_sgpr 0
+// ASM-NEXT: .amdhsa_next_free_vgpr defined_value+4
+// ASM-NEXT: .amdhsa_next_free_sgpr defined_value+5
+// ASM-NEXT: .amdhsa_reserve_vcc defined_boolean
+// ASM-NEXT: .amdhsa_reserve_flat_scratch defined_boolean
// ASM-NEXT: .amdhsa_reserve_xnack_mask 1
-// ASM-NEXT: .amdhsa_float_round_mode_32 (((((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~63))|(0<<0))&(~960))|(0<<6))&12288)>>12
-// ASM-NEXT: .amdhsa_float_round_mode_16_64 (((((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~63))|(0<<0))&(~960))|(0<<6))&49152)>>14
-// ASM-NEXT: .amdhsa_float_denorm_mode_32 (((((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~63))|(0<<0))&(~960))|(0<<6))&196608)>>16
-// ASM-NEXT: .amdhsa_float_denorm_mode_16_64 (((((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~63))|(0<<0))&(~960))|(0<<6))&786432)>>18
-// ASM-NEXT: .amdhsa_dx10_clamp (((((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~63))|(0<<0))&(~960))|(0<<6))&2097152)>>21
-// ASM-NEXT: .amdhsa_ieee_mode (((((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~63))|(0<<0))&(~960))|(0<<6))&8388608)>>23
-// ASM-NEXT: .amdhsa_fp16_overflow (((((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~63))|(0<<0))&(~960))|(0<<6))&67108864)>>26
-// ASM-NEXT: .amdhsa_workgroup_processor_mode (((((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~63))|(0<<0))&(~960))|(0<<6))&536870912)>>29
-// ASM-NEXT: .amdhsa_memory_ordered (((((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~63))|(0<<0))&(~960))|(0<<6))&1073741824)>>30
-// ASM-NEXT: .amdhsa_forward_progress (((((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~63))|(0<<0))&(~960))|(0<<6))&2147483648)>>31
+// ASM-NEXT: .amdhsa_float_round_mode_32 (((((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~63))|((((alignto(max(defined_value+4, 1), 8))/8)-1)<<0))&(~960))|((((alignto(max(0, 1), 8))/8)-1)<<6))&12288)>>12
+// ASM-NEXT: .amdhsa_float_round_mode_16_64 (((((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~63))|((((alignto(max(defined_value+4, 1), 8))/8)-1)<<0))&(~960))|((((alignto(max(0, 1), 8))/8)-1)<<6))&49152)>>14
+// ASM-NEXT: .amdhsa_float_denorm_mode_32 (((((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~63))|((((alignto(max(defined_value+4, 1), 8))/8)-1)<<0))&(~960))|((((alignto(max(0, 1), 8))/8)-1)<<6))&196608)>>16
+// ASM-NEXT: .amdhsa_float_denorm_mode_16_64 (((((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~63))|((((alignto(max(defined_value+4, 1), 8))/8)-1)<<0))&(~960))|((((alignto(max(0, 1), 8))/8)-1)<<6))&786432)>>18
+// ASM-NEXT: .amdhsa_dx10_clamp (((((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~63))|((((alignto(max(defined_value+4, 1), 8))/8)-1)<<0))&(~960))|((((alignto(max(0, 1), 8))/8)-1)<<6))&2097152)>>21
+// ASM-NEXT: .amdhsa_ieee_mode (((((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~63))|((((alignto(max(defined_value+4, 1), 8))/8)-1)<<0))&(~960))|((((alignto(max(0, 1), 8))/8)-1)<<6))&8388608)>>23
+// ASM-NEXT: .amdhsa_fp16_overflow (((((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~63))|((((alignto(max(defined_value+4, 1), 8))/8)-1)<<0))&(~960))|((((alignto(max(0, 1), 8))/8)-1)<<6))&67108864)>>26
+// ASM-NEXT: .amdhsa_workgroup_processor_mode (((((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~63))|((((alignto(max(defined_value+4, 1), 8))/8)-1)<<0))&(~960))|((((alignto(max(0, 1), 8))/8)-1)<<6))&536870912)>>29
+// ASM-NEXT: .amdhsa_memory_ordered (((((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~63))|((((alignto(max(defined_value+4, 1), 8))/8)-1)<<0))&(~960))|((((alignto(max(0, 1), 8))/8)-1)<<6))&1073741824)>>30
+// ASM-NEXT: .amdhsa_forward_progress (((((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~63))|((((alignto(max(defined_value+4, 1), 8))/8)-1)<<0))&(~960))|((((alignto(max(0, 1), 8))/8)-1)<<6))&2147483648)>>31
// ASM-NEXT: .amdhsa_shared_vgpr_count 0
// ASM-NEXT: .amdhsa_exception_fp_ieee_invalid_op (((((((((((((((((((((((((((((0&(~128))|(1<<7))&(~6144))|(defined_2_bits<<11))&(~128))|(defined_boolean<<7))&(~256))|(defined_boolean<<8))&(~512))|(defined_boolean<<9))&(~1024))|(defined_boolean<<10))&(~16777216))|(defined_boolean<<24))&(~33554432))|(defined_boolean<<25))&(~67108864))|(defined_boolean<<26))&(~134217728))|(defined_boolean<<27))&(~268435456))|(defined_boolean<<28))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~62))|(0<<1))&16777216)>>24
// ASM-NEXT: .amdhsa_exception_fp_denorm_src (((((((((((((((((((((((((((((0&(~128))|(1<<7))&(~6144))|(defined_2_bits<<11))&(~128))|(defined_boolean<<7))&(~256))|(defined_boolean<<8))&(~512))|(defined_boolean<<9))&(~1024))|(defined_boolean<<10))&(~16777216))|(defined_boolean<<24))&(~33554432))|(defined_boolean<<25))&(~67108864))|(defined_boolean<<26))&(~134217728))|(defined_boolean<<27))&(~268435456))|(defined_boolean<<28))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~62))|(0<<1))&33554432)>>25
@@ -166,8 +172,10 @@ expr_defined:
// ASM-NEXT: .amdhsa_system_sgpr_workgroup_id_z 1
// ASM-NEXT: .amdhsa_system_sgpr_workgroup_info 1
// ASM-NEXT: .amdhsa_system_vgpr_workitem_id 3
-// ASM-NEXT: .amdhsa_next_free_vgpr 0
-// ASM-NEXT: .amdhsa_next_free_sgpr 0
+// ASM-NEXT: .amdhsa_next_free_vgpr 44
+// ASM-NEXT: .amdhsa_next_free_sgpr 45
+// ASM-NEXT: .amdhsa_reserve_vcc 1
+// ASM-NEXT: .amdhsa_reserve_flat_scratch 1
// ASM-NEXT: .amdhsa_reserve_xnack_mask 1
// ASM-NEXT: .amdhsa_float_round_mode_32 3
// ASM-NEXT: .amdhsa_float_round_mode_16_64 3
diff --git a/llvm/test/MC/AMDGPU/hsa-sym-exprs-gfx11.s b/llvm/test/MC/AMDGPU/hsa-sym-exprs-gfx11.s
index e1107fb69ba41..b6e4ddde3d7f9 100644
--- a/llvm/test/MC/AMDGPU/hsa-sym-exprs-gfx11.s
+++ b/llvm/test/MC/AMDGPU/hsa-sym-exprs-gfx11.s
@@ -10,12 +10,12 @@
// OBJDUMP-NEXT: 0000 2b000000 2c000000 00000000 00000000
// OBJDUMP-NEXT: 0010 00000000 00000000 00000000 00000000
// OBJDUMP-NEXT: 0020 00000000 00000000 00000000 00000000
-// OBJDUMP-NEXT: 0030 00f0afe4 811f007f 000c0000 00000000
+// OBJDUMP-NEXT: 0030 05f0afe4 811f007f 000c0000 00000000
// expr_defined
// OBJDUMP-NEXT: 0040 2a000000 2b000000 00000000 00000000
// OBJDUMP-NEXT: 0050 00000000 00000000 00000000 00000000
// OBJDUMP-NEXT: 0060 00000000 00000000 00000000 00000000
-// OBJDUMP-NEXT: 0070 00f0afe4 811f007f 000c0000 00000000
+// OBJDUMP-NEXT: 0070 05f0afe4 811f007f 000c0000 00000000
.text
// ASM: .text
@@ -62,8 +62,9 @@ expr_defined:
.amdhsa_exception_int_div_zero defined_boolean
.amdhsa_enable_private_segment defined_boolean
.amdhsa_uses_dynamic_stack defined_boolean
- .amdhsa_next_free_vgpr 0
- .amdhsa_next_free_sgpr 0
+ .amdhsa_next_free_vgpr defined_value+4
+ .amdhsa_next_free_sgpr defined_value+5
+ .amdhsa_reserve_vcc defined_boolean
.end_amdhsa_kernel
.set defined_value, 41
@@ -96,8 +97,9 @@ expr_defined:
.amdhsa_exception_int_div_zero defined_boolean
.amdhsa_enable_private_segment defined_boolean
.amdhsa_uses_dynamic_stack defined_boolean
- .amdhsa_next_free_vgpr 0
- .amdhsa_next_free_sgpr 0
+ .amdhsa_next_free_vgpr defined_value+3
+ .amdhsa_next_free_sgpr defined_value+4
+ .amdhsa_reserve_vcc defined_boolean
.end_amdhsa_kernel
// ASM: .amdhsa_kernel expr_defined_later
@@ -117,18 +119,19 @@ expr_defined:
// ASM-NEXT: .amdhsa_system_sgpr_workgroup_id_z (((((((((((((((((((((((((((((((0&(~128))|(1<<7))&(~6144))|(defined_2_bits<<11))&(~128))|(defined_boolean<<7))&(~256))|(defined_boolean<<8))&(~512))|(defined_boolean<<9))&(~1024))|(defined_boolean<<10))&(~16777216))|(defined_boolean<<24))&(~33554432))|(defined_boolean<<25))&(~67108864))|(defined_boolean<<26))&(~134217728))|(defined_boolean<<27))&(~268435456))|(defined_boolean<<28))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~1))|(defined_boolean<<0))&(~62))|(0<<1))&512)>>9
// ASM-NEXT: .amdhsa_system_sgpr_workgroup_info (((((((((((((((((((((((((((((((0&(~128))|(1<<7))&(~6144))|(defined_2_bits<<11))&(~128))|(defined_boolean<<7))&(~256))|(defined_boolean<<8))&(~512))|(defined_boolean<<9))&(~1024))|(defined_boolean<<10))&(~16777216))|(defined_boolean<<24))&(~33554432))|(defined_boolean<<25))&(~67108864))|(defined_boolean<<26))&(~134217728))|(defined_boolean<<27))&(~268435456))|(defined_boolean<<28))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~1))|(defined_boolean<<0))&(~62))|(0<<1))&1024)>>10
// ASM-NEXT: .amdhsa_system_vgpr_workitem_id (((((((((((((((((((((((((((((((0&(~128))|(1<<7))&(~6144))|(defined_2_bits<<11))&(~128))|(defined_boolean<<7))&(~256))|(defined_boolean<<8))&(~512))|(defined_boolean<<9))&(~1024))|(defined_boolean<<10))&(~16777216))|(defined_boolean<<24))&(~33554432))|(defined_boolean<<25))&(~67108864))|(defined_boolean<<26))&(~134217728))|(defined_boolean<<27))&(~268435456))|(defined_boolean<<28))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~1))|(defined_boolean<<0))&(~62))|(0<<1))&6144)>>11
-// ASM-NEXT: .amdhsa_next_free_vgpr 0
-// ASM-NEXT: .amdhsa_next_free_sgpr 0
-// ASM-NEXT: .amdhsa_float_round_mode_32 (((((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~63))|(0<<0))&(~960))|(0<<6))&12288)>>12
-// ASM-NEXT: .amdhsa_float_round_mode_16_64 (((((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~63))|(0<<0))&(~960))|(0<<6))&49152)>>14
-// ASM-NEXT: .amdhsa_float_denorm_mode_32 (((((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~63))|(0<<0))&(~960))|(0<<6))&196608)>>16
-// ASM-NEXT: .amdhsa_float_denorm_mode_16_64 (((((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~63))|(0<<0))&(~960))|(0<<6))&786432)>>18
-// ASM-NEXT: .amdhsa_dx10_clamp (((((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~63))|(0<<0))&(~960))|(0<<6))&2097152)>>21
-// ASM-NEXT: .amdhsa_ieee_mode (((((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~63))|(0<<0))&(~960))|(0<<6))&8388608)>>23
-// ASM-NEXT: .amdhsa_fp16_overflow (((((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~63))|(0<<0))&(~960))|(0<<6))&67108864)>>26
-// ASM-NEXT: .amdhsa_workgroup_processor_mode (((((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~63))|(0<<0))&(~960))|(0<<6))&536870912)>>29
-// ASM-NEXT: .amdhsa_memory_ordered (((((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~63))|(0<<0))&(~960))|(0<<6))&1073741824)>>30
-// ASM-NEXT: .amdhsa_forward_progress (((((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~63))|(0<<0))&(~960))|(0<<6))&2147483648)>>31
+// ASM-NEXT: .amdhsa_next_free_vgpr defined_value+4
+// ASM-NEXT: .amdhsa_next_free_sgpr defined_value+5
+// ASM-NEXT: .amdhsa_reserve_vcc defined_boolean
+// ASM-NEXT: .amdhsa_float_round_mode_32 (((((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~63))|((((alignto(max(defined_value+4, 1), 8))/8)-1)<<0))&(~960))|((((alignto(max(0, 1), 8))/8)-1)<<6))&12288)>>12
+// ASM-NEXT: .amdhsa_float_round_mode_16_64 (((((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~63))|((((alignto(max(defined_value+4, 1), 8))/8)-1)<<0))&(~960))|((((alignto(max(0, 1), 8))/8)-1)<<6))&49152)>>14
+// ASM-NEXT: .amdhsa_float_denorm_mode_32 (((((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~63))|((((alignto(max(defined_value+4, 1), 8))/8)-1)<<0))&(~960))|((((alignto(max(0, 1), 8))/8)-1)<<6))&196608)>>16
+// ASM-NEXT: .amdhsa_float_denorm_mode_16_64 (((((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~63))|((((alignto(max(defined_value+4, 1), 8))/8)-1)<<0))&(~960))|((((alignto(max(0, 1), 8))/8)-1)<<6))&786432)>>18
+// ASM-NEXT: .amdhsa_dx10_clamp (((((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~63))|((((alignto(max(defined_value+4, 1), 8))/8)-1)<<0))&(~960))|((((alignto(max(0, 1), 8))/8)-1)<<6))&2097152)>>21
+// ASM-NEXT: .amdhsa_ieee_mode (((((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~63))|((((alignto(max(defined_value+4, 1), 8))/8)-1)<<0))&(~960))|((((alignto(max(0, 1), 8))/8)-1)<<6))&8388608)>>23
+// ASM-NEXT: .amdhsa_fp16_overflow (((((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~63))|((((alignto(max(defined_value+4, 1), 8))/8)-1)<<0))&(~960))|((((alignto(max(0, 1), 8))/8)-1)<<6))&67108864)>>26
+// ASM-NEXT: .amdhsa_workgroup_processor_mode (((((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~63))|((((alignto(max(defined_value+4, 1), 8))/8)-1)<<0))&(~960))|((((alignto(max(0, 1), 8))/8)-1)<<6))&536870912)>>29
+// ASM-NEXT: .amdhsa_memory_ordered (((((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~63))|((((alignto(max(defined_value+4, 1), 8))/8)-1)<<0))&(~960))|((((alignto(max(0, 1), 8))/8)-1)<<6))&1073741824)>>30
+// ASM-NEXT: .amdhsa_forward_progress (((((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~63))|((((alignto(max(defined_value+4, 1), 8))/8)-1)<<0))&(~960))|((((alignto(max(0, 1), 8))/8)-1)<<6))&2147483648)>>31
// ASM-NEXT: .amdhsa_shared_vgpr_count 0
// ASM-NEXT: .amdhsa_exception_fp_ieee_invalid_op (((((((((((((((((((((((((((((((0&(~128))|(1<<7))&(~6144))|(defined_2_bits<<11))&(~128))|(defined_boolean<<7))&(~256))|(defined_boolean<<8))&(~512))|(defined_boolean<<9))&(~1024))|(defined_boolean<<10))&(~16777216))|(defined_boolean<<24))&(~33554432))|(defined_boolean<<25))&(~67108864))|(defined_boolean<<26))&(~134217728))|(defined_boolean<<27))&(~268435456))|(defined_boolean<<28))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~1))|(defined_boolean<<0))&(~62))|(0<<1))&16777216)>>24
// ASM-NEXT: .amdhsa_exception_fp_denorm_src (((((((((((((((((((((((((((((((0&(~128))|(1<<7))&(~6144))|(defined_2_bits<<11))&(~128))|(defined_boolean<<7))&(~256))|(defined_boolean<<8))&(~512))|(defined_boolean<<9))&(~1024))|(defined_boolean<<10))&(~16777216))|(defined_boolean<<24))&(~33554432))|(defined_boolean<<25))&(~67108864))|(defined_boolean<<26))&(~134217728))|(defined_boolean<<27))&(~268435456))|(defined_boolean<<28))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~1))|(defined_boolean<<0))&(~62))|(0<<1))&33554432)>>25
@@ -163,8 +166,9 @@ expr_defined:
// ASM-NEXT: .amdhsa_system_sgpr_workgroup_id_z 1
// ASM-NEXT: .amdhsa_system_sgpr_workgroup_info 1
// ASM-NEXT: .amdhsa_system_vgpr_workitem_id 3
-// ASM-NEXT: .amdhsa_next_free_vgpr 0
-// ASM-NEXT: .amdhsa_next_free_sgpr 0
+// ASM-NEXT: .amdhsa_next_free_vgpr 44
+// ASM-NEXT: .amdhsa_next_free_sgpr 45
+// ASM-NEXT: .amdhsa_reserve_vcc 1
// ASM-NEXT: .amdhsa_float_round_mode_32 3
// ASM-NEXT: .amdhsa_float_round_mode_16_64 3
// ASM-NEXT: .amdhsa_float_denorm_mode_32 3
diff --git a/llvm/test/MC/AMDGPU/hsa-sym-exprs-gfx12.s b/llvm/test/MC/AMDGPU/hsa-sym-exprs-gfx12.s
index 449616d35186b..a80000dc44dac 100644
--- a/llvm/test/MC/AMDGPU/hsa-sym-exprs-gfx12.s
+++ b/llvm/test/MC/AMDGPU/hsa-sym-exprs-gfx12.s
@@ -10,12 +10,12 @@
// OBJDUMP-NEXT: 0000 2b000000 2c000000 00000000 00000000
// OBJDUMP-NEXT: 0010 00000000 00000000 00000000 00000000
// OBJDUMP-NEXT: 0020 00000000 00000000 00000000 00000000
-// OBJDUMP-NEXT: 0030 00f02fe4 811f007f 000c0000 00000000
+// OBJDUMP-NEXT: 0030 05f02fe4 811f007f 000c0000 00000000
// expr_defined
// OBJDUMP-NEXT: 0040 2a000000 2b000000 00000000 00000000
// OBJDUMP-NEXT: 0050 00000000 00000000 00000000 00000000
// OBJDUMP-NEXT: 0060 00000000 00000000 00000000 00000000
-// OBJDUMP-NEXT: 0070 00f02fe4 811f007f 000c0000 00000000
+// OBJDUMP-NEXT: 0070 05f02fe4 811f007f 000c0000 00000000
.text
// ASM: .text
@@ -63,8 +63,9 @@ expr_defined:
.amdhsa_round_robin_scheduling defined_boolean
.amdhsa_enable_private_segment defined_boolean
.amdhsa_uses_dynamic_stack defined_boolean
- .amdhsa_next_free_vgpr 0
- .amdhsa_next_free_sgpr 0
+ .amdhsa_next_free_vgpr defined_value+4
+ .amdhsa_next_free_sgpr defined_value+5
+ .amdhsa_reserve_vcc defined_boolean
.end_amdhsa_kernel
.set defined_value, 41
@@ -98,8 +99,9 @@ expr_defined:
.amdhsa_round_robin_scheduling defined_boolean
.amdhsa_enable_private_segment defined_boolean
.amdhsa_uses_dynamic_stack defined_boolean
- .amdhsa_next_free_vgpr 0
- .amdhsa_next_free_sgpr 0
+ .amdhsa_next_free_vgpr defined_value+3
+ .amdhsa_next_free_sgpr defined_value+4
+ .amdhsa_reserve_vcc defined_boolean
.end_amdhsa_kernel
// ASM: .amdhsa_kernel expr_defined_later
@@ -119,17 +121,18 @@ expr_defined:
// ASM-NEXT: .amdhsa_system_sgpr_workgroup_id_z (((((((((((((((((((((((((((((((0&(~128))|(1<<7))&(~6144))|(defined_2_bits<<11))&(~128))|(defined_boolean<<7))&(~256))|(defined_boolean<<8))&(~512))|(defined_boolean<<9))&(~1024))|(defined_boolean<<10))&(~16777216))|(defined_boolean<<24))&(~33554432))|(defined_boolean<<25))&(~67108864))|(defined_boolean<<26))&(~134217728))|(defined_boolean<<27))&(~268435456))|(defined_boolean<<28))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~1))|(defined_boolean<<0))&(~62))|(0<<1))&512)>>9
// ASM-NEXT: .amdhsa_system_sgpr_workgroup_info (((((((((((((((((((((((((((((((0&(~128))|(1<<7))&(~6144))|(defined_2_bits<<11))&(~128))|(defined_boolean<<7))&(~256))|(defined_boolean<<8))&(~512))|(defined_boolean<<9))&(~1024))|(defined_boolean<<10))&(~16777216))|(defined_boolean<<24))&(~33554432))|(defined_boolean<<25))&(~67108864))|(defined_boolean<<26))&(~134217728))|(defined_boolean<<27))&(~268435456))|(defined_boolean<<28))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~1))|(defined_boolean<<0))&(~62))|(0<<1))&1024)>>10
// ASM-NEXT: .amdhsa_system_vgpr_workitem_id (((((((((((((((((((((((((((((((0&(~128))|(1<<7))&(~6144))|(defined_2_bits<<11))&(~128))|(defined_boolean<<7))&(~256))|(defined_boolean<<8))&(~512))|(defined_boolean<<9))&(~1024))|(defined_boolean<<10))&(~16777216))|(defined_boolean<<24))&(~33554432))|(defined_boolean<<25))&(~67108864))|(defined_boolean<<26))&(~134217728))|(defined_boolean<<27))&(~268435456))|(defined_boolean<<28))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~1))|(defined_boolean<<0))&(~62))|(0<<1))&6144)>>11
-// ASM-NEXT: .amdhsa_next_free_vgpr 0
-// ASM-NEXT: .amdhsa_next_free_sgpr 0
-// ASM-NEXT: .amdhsa_float_round_mode_32 (((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~2097152))|(defined_boolean<<21))&(~63))|(0<<0))&(~960))|(0<<6))&12288)>>12
-// ASM-NEXT: .amdhsa_float_round_mode_16_64 (((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~2097152))|(defined_boolean<<21))&(~63))|(0<<0))&(~960))|(0<<6))&49152)>>14
-// ASM-NEXT: .amdhsa_float_denorm_mode_32 (((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~2097152))|(defined_boolean<<21))&(~63))|(0<<0))&(~960))|(0<<6))&196608)>>16
-// ASM-NEXT: .amdhsa_float_denorm_mode_16_64 (((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~2097152))|(defined_boolean<<21))&(~63))|(0<<0))&(~960))|(0<<6))&786432)>>18
-// ASM-NEXT: .amdhsa_fp16_overflow (((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~2097152))|(defined_boolean<<21))&(~63))|(0<<0))&(~960))|(0<<6))&67108864)>>26
-// ASM-NEXT: .amdhsa_workgroup_processor_mode (((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~2097152))|(defined_boolean<<21))&(~63))|(0<<0))&(~960))|(0<<6))&536870912)>>29
-// ASM-NEXT: .amdhsa_memory_ordered (((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~2097152))|(defined_boolean<<21))&(~63))|(0<<0))&(~960))|(0<<6))&1073741824)>>30
-// ASM-NEXT: .amdhsa_forward_progress (((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~2097152))|(defined_boolean<<21))&(~63))|(0<<0))&(~960))|(0<<6))&2147483648)>>31
-// ASM-NEXT: .amdhsa_round_robin_scheduling (((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~2097152))|(defined_boolean<<21))&(~63))|(0<<0))&(~960))|(0<<6))&2097152)>>21
+// ASM-NEXT: .amdhsa_next_free_vgpr defined_value+4
+// ASM-NEXT: .amdhsa_next_free_sgpr defined_value+5
+// ASM-NEXT: .amdhsa_reserve_vcc defined_boolean
+// ASM-NEXT: .amdhsa_float_round_mode_32 (((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~2097152))|(defined_boolean<<21))&(~63))|((((alignto(max(defined_value+4, 1), 8))/8)-1)<<0))&(~960))|((((alignto(max(0, 1), 8))/8)-1)<<6))&12288)>>12
+// ASM-NEXT: .amdhsa_float_round_mode_16_64 (((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~2097152))|(defined_boolean<<21))&(~63))|((((alignto(max(defined_value+4, 1), 8))/8)-1)<<0))&(~960))|((((alignto(max(0, 1), 8))/8)-1)<<6))&49152)>>14
+// ASM-NEXT: .amdhsa_float_denorm_mode_32 (((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~2097152))|(defined_boolean<<21))&(~63))|((((alignto(max(defined_value+4, 1), 8))/8)-1)<<0))&(~960))|((((alignto(max(0, 1), 8))/8)-1)<<6))&196608)>>16
+// ASM-NEXT: .amdhsa_float_denorm_mode_16_64 (((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~2097152))|(defined_boolean<<21))&(~63))|((((alignto(max(defined_value+4, 1), 8))/8)-1)<<0))&(~960))|((((alignto(max(0, 1), 8))/8)-1)<<6))&786432)>>18
+// ASM-NEXT: .amdhsa_fp16_overflow (((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~2097152))|(defined_boolean<<21))&(~63))|((((alignto(max(defined_value+4, 1), 8))/8)-1)<<0))&(~960))|((((alignto(max(0, 1), 8))/8)-1)<<6))&67108864)>>26
+// ASM-NEXT: .amdhsa_workgroup_processor_mode (((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~2097152))|(defined_boolean<<21))&(~63))|((((alignto(max(defined_value+4, 1), 8))/8)-1)<<0))&(~960))|((((alignto(max(0, 1), 8))/8)-1)<<6))&536870912)>>29
+// ASM-NEXT: .amdhsa_memory_ordered (((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~2097152))|(defined_boolean<<21))&(~63))|((((alignto(max(defined_value+4, 1), 8))/8)-1)<<0))&(~960))|((((alignto(max(0, 1), 8))/8)-1)<<6))&1073741824)>>30
+// ASM-NEXT: .amdhsa_forward_progress (((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~2097152))|(defined_boolean<<21))&(~63))|((((alignto(max(defined_value+4, 1), 8))/8)-1)<<0))&(~960))|((((alignto(max(0, 1), 8))/8)-1)<<6))&2147483648)>>31
+// ASM-NEXT: .amdhsa_round_robin_scheduling (((((((((((((((((((((((((((((0&(~786432))|(3<<18))&(~536870912))|(1<<29))&(~1073741824))|(1<<30))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~67108864))|(defined_boolean<<26))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~2147483648))|(defined_boolean<<31))&(~2097152))|(defined_boolean<<21))&(~63))|((((alignto(max(defined_value+4, 1), 8))/8)-1)<<0))&(~960))|((((alignto(max(0, 1), 8))/8)-1)<<6))&2097152)>>21
// ASM-NEXT: .amdhsa_exception_fp_ieee_invalid_op (((((((((((((((((((((((((((((((0&(~128))|(1<<7))&(~6144))|(defined_2_bits<<11))&(~128))|(defined_boolean<<7))&(~256))|(defined_boolean<<8))&(~512))|(defined_boolean<<9))&(~1024))|(defined_boolean<<10))&(~16777216))|(defined_boolean<<24))&(~33554432))|(defined_boolean<<25))&(~67108864))|(defined_boolean<<26))&(~134217728))|(defined_boolean<<27))&(~268435456))|(defined_boolean<<28))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~1))|(defined_boolean<<0))&(~62))|(0<<1))&16777216)>>24
// ASM-NEXT: .amdhsa_exception_fp_denorm_src (((((((((((((((((((((((((((((((0&(~128))|(1<<7))&(~6144))|(defined_2_bits<<11))&(~128))|(defined_boolean<<7))&(~256))|(defined_boolean<<8))&(~512))|(defined_boolean<<9))&(~1024))|(defined_boolean<<10))&(~16777216))|(defined_boolean<<24))&(~33554432))|(defined_boolean<<25))&(~67108864))|(defined_boolean<<26))&(~134217728))|(defined_boolean<<27))&(~268435456))|(defined_boolean<<28))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~1))|(defined_boolean<<0))&(~62))|(0<<1))&33554432)>>25
// ASM-NEXT: .amdhsa_exception_fp_ieee_div_zero (((((((((((((((((((((((((((((((0&(~128))|(1<<7))&(~6144))|(defined_2_bits<<11))&(~128))|(defined_boolean<<7))&(~256))|(defined_boolean<<8))&(~512))|(defined_boolean<<9))&(~1024))|(defined_boolean<<10))&(~16777216))|(defined_boolean<<24))&(~33554432))|(defined_boolean<<25))&(~67108864))|(defined_boolean<<26))&(~134217728))|(defined_boolean<<27))&(~268435456))|(defined_boolean<<28))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~1))|(defined_boolean<<0))&(~62))|(0<<1))&67108864)>>26
@@ -163,8 +166,9 @@ expr_defined:
// ASM-NEXT: .amdhsa_system_sgpr_workgroup_id_z 1
// ASM-NEXT: .amdhsa_system_sgpr_workgroup_info 1
// ASM-NEXT: .amdhsa_system_vgpr_workitem_id 3
-// ASM-NEXT: .amdhsa_next_free_vgpr 0
-// ASM-NEXT: .amdhsa_next_free_sgpr 0
+// ASM-NEXT: .amdhsa_next_free_vgpr 44
+// ASM-NEXT: .amdhsa_next_free_sgpr 45
+// ASM-NEXT: .amdhsa_reserve_vcc 1
// ASM-NEXT: .amdhsa_float_round_mode_32 3
// ASM-NEXT: .amdhsa_float_round_mode_16_64 3
// ASM-NEXT: .amdhsa_float_denorm_mode_32 3
diff --git a/llvm/test/MC/AMDGPU/hsa-sym-exprs-gfx7.s b/llvm/test/MC/AMDGPU/hsa-sym-exprs-gfx7.s
index c7e05441b45ff..7ab2e2b28a0e6 100644
--- a/llvm/test/MC/AMDGPU/hsa-sym-exprs-gfx7.s
+++ b/llvm/test/MC/AMDGPU/hsa-sym-exprs-gfx7.s
@@ -10,12 +10,12 @@
// OBJDUMP-NEXT: 0000 2b000000 2c000000 00000000 00000000
// OBJDUMP-NEXT: 0010 00000000 00000000 00000000 00000000
// OBJDUMP-NEXT: 0020 00000000 00000000 00000000 00000000
-// OBJDUMP-NEXT: 0030 00f0af00 801f007f 00080000 00000000
+// OBJDUMP-NEXT: 0030 8bf1af00 801f007f 00080000 00000000
// expr_defined
// OBJDUMP-NEXT: 0040 2a000000 2b000000 00000000 00000000
// OBJDUMP-NEXT: 0050 00000000 00000000 00000000 00000000
// OBJDUMP-NEXT: 0060 00000000 00000000 00000000 00000000
-// OBJDUMP-NEXT: 0070 00f0af00 801f007f 00080000 00000000
+// OBJDUMP-NEXT: 0070 8af1af00 801f007f 00080000 00000000
.text
// ASM: .text
@@ -57,8 +57,10 @@ expr_defined:
.amdhsa_exception_fp_ieee_inexact defined_boolean
.amdhsa_exception_int_div_zero defined_boolean
.amdhsa_uses_dynamic_stack defined_boolean
- .amdhsa_next_free_vgpr 0
- .amdhsa_next_free_sgpr 0
+ .amdhsa_next_free_vgpr defined_value+4
+ .amdhsa_next_free_sgpr defined_value+5
+ .amdhsa_reserve_vcc defined_boolean
+ .amdhsa_reserve_flat_scratch defined_boolean
.end_amdhsa_kernel
.set defined_value, 41
@@ -86,8 +88,10 @@ expr_defined:
.amdhsa_exception_fp_ieee_inexact defined_boolean
.amdhsa_exception_int_div_zero defined_boolean
.amdhsa_uses_dynamic_stack defined_boolean
- .amdhsa_next_free_vgpr 0
- .amdhsa_next_free_sgpr 0
+ .amdhsa_next_free_vgpr defined_value+3
+ .amdhsa_next_free_sgpr defined_value+4
+ .amdhsa_reserve_vcc defined_boolean
+ .amdhsa_reserve_flat_scratch defined_boolean
.end_amdhsa_kernel
// ASM: .amdhsa_kernel expr_defined_later
@@ -108,14 +112,16 @@ expr_defined:
// ASM-NEXT: .amdhsa_system_sgpr_workgroup_id_z (((((((((((((((((((((((((((((0&(~128))|(1<<7))&(~6144))|(defined_2_bits<<11))&(~128))|(defined_boolean<<7))&(~256))|(defined_boolean<<8))&(~512))|(defined_boolean<<9))&(~1024))|(defined_boolean<<10))&(~16777216))|(defined_boolean<<24))&(~33554432))|(defined_boolean<<25))&(~67108864))|(defined_boolean<<26))&(~134217728))|(defined_boolean<<27))&(~268435456))|(defined_boolean<<28))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~62))|(0<<1))&512)>>9
// ASM-NEXT: .amdhsa_system_sgpr_workgroup_info (((((((((((((((((((((((((((((0&(~128))|(1<<7))&(~6144))|(defined_2_bits<<11))&(~128))|(defined_boolean<<7))&(~256))|(defined_boolean<<8))&(~512))|(defined_boolean<<9))&(~1024))|(defined_boolean<<10))&(~16777216))|(defined_boolean<<24))&(~33554432))|(defined_boolean<<25))&(~67108864))|(defined_boolean<<26))&(~134217728))|(defined_boolean<<27))&(~268435456))|(defined_boolean<<28))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~62))|(0<<1))&1024)>>10
// ASM-NEXT: .amdhsa_system_vgpr_workitem_id (((((((((((((((((((((((((((((0&(~128))|(1<<7))&(~6144))|(defined_2_bits<<11))&(~128))|(defined_boolean<<7))&(~256))|(defined_boolean<<8))&(~512))|(defined_boolean<<9))&(~1024))|(defined_boolean<<10))&(~16777216))|(defined_boolean<<24))&(~33554432))|(defined_boolean<<25))&(~67108864))|(defined_boolean<<26))&(~134217728))|(defined_boolean<<27))&(~268435456))|(defined_boolean<<28))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~62))|(0<<1))&6144)>>11
-// ASM-NEXT: .amdhsa_next_free_vgpr 0
-// ASM-NEXT: .amdhsa_next_free_sgpr 0
-// ASM-NEXT: .amdhsa_float_round_mode_32 (((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~63))|(0<<0))&(~960))|(0<<6))&12288)>>12
-// ASM-NEXT: .amdhsa_float_round_mode_16_64 (((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~63))|(0<<0))&(~960))|(0<<6))&49152)>>14
-// ASM-NEXT: .amdhsa_float_denorm_mode_32 (((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~63))|(0<<0))&(~960))|(0<<6))&196608)>>16
-// ASM-NEXT: .amdhsa_float_denorm_mode_16_64 (((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~63))|(0<<0))&(~960))|(0<<6))&786432)>>18
-// ASM-NEXT: .amdhsa_dx10_clamp (((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~63))|(0<<0))&(~960))|(0<<6))&2097152)>>21
-// ASM-NEXT: .amdhsa_ieee_mode (((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~63))|(0<<0))&(~960))|(0<<6))&8388608)>>23
+// ASM-NEXT: .amdhsa_next_free_vgpr defined_value+4
+// ASM-NEXT: .amdhsa_next_free_sgpr defined_value+5
+// ASM-NEXT: .amdhsa_reserve_vcc defined_boolean
+// ASM-NEXT: .amdhsa_reserve_flat_scratch defined_boolean
+// ASM-NEXT: .amdhsa_float_round_mode_32 (((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~63))|((((alignto(max(defined_value+4, 1), 4))/4)-1)<<0))&(~960))|((((alignto(max((defined_value+5)+(extrasgprs(defined_boolean, defined_boolean, 0)), 1), 8))/8)-1)<<6))&12288)>>12
+// ASM-NEXT: .amdhsa_float_round_mode_16_64 (((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~63))|((((alignto(max(defined_value+4, 1), 4))/4)-1)<<0))&(~960))|((((alignto(max((defined_value+5)+(extrasgprs(defined_boolean, defined_boolean, 0)), 1), 8))/8)-1)<<6))&49152)>>14
+// ASM-NEXT: .amdhsa_float_denorm_mode_32 (((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~63))|((((alignto(max(defined_value+4, 1), 4))/4)-1)<<0))&(~960))|((((alignto(max((defined_value+5)+(extrasgprs(defined_boolean, defined_boolean, 0)), 1), 8))/8)-1)<<6))&196608)>>16
+// ASM-NEXT: .amdhsa_float_denorm_mode_16_64 (((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~63))|((((alignto(max(defined_value+4, 1), 4))/4)-1)<<0))&(~960))|((((alignto(max((defined_value+5)+(extrasgprs(defined_boolean, defined_boolean, 0)), 1), 8))/8)-1)<<6))&786432)>>18
+// ASM-NEXT: .amdhsa_dx10_clamp (((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~63))|((((alignto(max(defined_value+4, 1), 4))/4)-1)<<0))&(~960))|((((alignto(max((defined_value+5)+(extrasgprs(defined_boolean, defined_boolean, 0)), 1), 8))/8)-1)<<6))&2097152)>>21
+// ASM-NEXT: .amdhsa_ieee_mode (((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~63))|((((alignto(max(defined_value+4, 1), 4))/4)-1)<<0))&(~960))|((((alignto(max((defined_value+5)+(extrasgprs(defined_boolean, defined_boolean, 0)), 1), 8))/8)-1)<<6))&8388608)>>23
// ASM-NEXT: .amdhsa_exception_fp_ieee_invalid_op (((((((((((((((((((((((((((((0&(~128))|(1<<7))&(~6144))|(defined_2_bits<<11))&(~128))|(defined_boolean<<7))&(~256))|(defined_boolean<<8))&(~512))|(defined_boolean<<9))&(~1024))|(defined_boolean<<10))&(~16777216))|(defined_boolean<<24))&(~33554432))|(defined_boolean<<25))&(~67108864))|(defined_boolean<<26))&(~134217728))|(defined_boolean<<27))&(~268435456))|(defined_boolean<<28))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~62))|(0<<1))&16777216)>>24
// ASM-NEXT: .amdhsa_exception_fp_denorm_src (((((((((((((((((((((((((((((0&(~128))|(1<<7))&(~6144))|(defined_2_bits<<11))&(~128))|(defined_boolean<<7))&(~256))|(defined_boolean<<8))&(~512))|(defined_boolean<<9))&(~1024))|(defined_boolean<<10))&(~16777216))|(defined_boolean<<24))&(~33554432))|(defined_boolean<<25))&(~67108864))|(defined_boolean<<26))&(~134217728))|(defined_boolean<<27))&(~268435456))|(defined_boolean<<28))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~62))|(0<<1))&33554432)>>25
// ASM-NEXT: .amdhsa_exception_fp_ieee_div_zero (((((((((((((((((((((((((((((0&(~128))|(1<<7))&(~6144))|(defined_2_bits<<11))&(~128))|(defined_boolean<<7))&(~256))|(defined_boolean<<8))&(~512))|(defined_boolean<<9))&(~1024))|(defined_boolean<<10))&(~16777216))|(defined_boolean<<24))&(~33554432))|(defined_boolean<<25))&(~67108864))|(defined_boolean<<26))&(~134217728))|(defined_boolean<<27))&(~268435456))|(defined_boolean<<28))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~62))|(0<<1))&67108864)>>26
@@ -150,8 +156,10 @@ expr_defined:
// ASM-NEXT: .amdhsa_system_sgpr_workgroup_id_z 1
// ASM-NEXT: .amdhsa_system_sgpr_workgroup_info 1
// ASM-NEXT: .amdhsa_system_vgpr_workitem_id 3
-// ASM-NEXT: .amdhsa_next_free_vgpr 0
-// ASM-NEXT: .amdhsa_next_free_sgpr 0
+// ASM-NEXT: .amdhsa_next_free_vgpr 44
+// ASM-NEXT: .amdhsa_next_free_sgpr 45
+// ASM-NEXT: .amdhsa_reserve_vcc 1
+// ASM-NEXT: .amdhsa_reserve_flat_scratch 1
// ASM-NEXT: .amdhsa_float_round_mode_32 3
// ASM-NEXT: .amdhsa_float_round_mode_16_64 3
// ASM-NEXT: .amdhsa_float_denorm_mode_32 3
diff --git a/llvm/test/MC/AMDGPU/hsa-sym-exprs-gfx8.s b/llvm/test/MC/AMDGPU/hsa-sym-exprs-gfx8.s
index 49a5015987a65..caccde7ba0e3b 100644
--- a/llvm/test/MC/AMDGPU/hsa-sym-exprs-gfx8.s
+++ b/llvm/test/MC/AMDGPU/hsa-sym-exprs-gfx8.s
@@ -11,12 +11,12 @@
// OBJDUMP-NEXT: 0000 2b000000 2c000000 00000000 00000000
// OBJDUMP-NEXT: 0010 00000000 00000000 00000000 00000000
// OBJDUMP-NEXT: 0020 00000000 00000000 00000000 00000000
-// OBJDUMP-NEXT: 0030 00f0af00 801f007f 00080000 00000000
+// OBJDUMP-NEXT: 0030 8bf1af00 801f007f 00080000 00000000
// expr_defined
// OBJDUMP-NEXT: 0040 2a000000 2b000000 00000000 00000000
// OBJDUMP-NEXT: 0050 00000000 00000000 00000000 00000000
// OBJDUMP-NEXT: 0060 00000000 00000000 00000000 00000000
-// OBJDUMP-NEXT: 0070 00f0af00 801f007f 00080000 00000000
+// OBJDUMP-NEXT: 0070 8af1af00 801f007f 00080000 00000000
.text
// ASM: .text
@@ -58,8 +58,10 @@ expr_defined:
.amdhsa_exception_fp_ieee_inexact defined_boolean
.amdhsa_exception_int_div_zero defined_boolean
.amdhsa_uses_dynamic_stack defined_boolean
- .amdhsa_next_free_vgpr 0
- .amdhsa_next_free_sgpr 0
+ .amdhsa_next_free_vgpr defined_value+4
+ .amdhsa_next_free_sgpr defined_value+5
+ .amdhsa_reserve_vcc defined_boolean
+ .amdhsa_reserve_flat_scratch defined_boolean
.end_amdhsa_kernel
.set defined_value, 41
@@ -87,8 +89,10 @@ expr_defined:
.amdhsa_exception_fp_ieee_inexact defined_boolean
.amdhsa_exception_int_div_zero defined_boolean
.amdhsa_uses_dynamic_stack defined_boolean
- .amdhsa_next_free_vgpr 0
- .amdhsa_next_free_sgpr 0
+ .amdhsa_next_free_vgpr defined_value+3
+ .amdhsa_next_free_sgpr defined_value+4
+ .amdhsa_reserve_vcc defined_boolean
+ .amdhsa_reserve_flat_scratch defined_boolean
.end_amdhsa_kernel
// ASM: .amdhsa_kernel expr_defined_later
@@ -109,15 +113,17 @@ expr_defined:
// ASM-NEXT: .amdhsa_system_sgpr_workgroup_id_z (((((((((((((((((((((((((((((0&(~128))|(1<<7))&(~6144))|(defined_2_bits<<11))&(~128))|(defined_boolean<<7))&(~256))|(defined_boolean<<8))&(~512))|(defined_boolean<<9))&(~1024))|(defined_boolean<<10))&(~16777216))|(defined_boolean<<24))&(~33554432))|(defined_boolean<<25))&(~67108864))|(defined_boolean<<26))&(~134217728))|(defined_boolean<<27))&(~268435456))|(defined_boolean<<28))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~62))|(0<<1))&512)>>9
// ASM-NEXT: .amdhsa_system_sgpr_workgroup_info (((((((((((((((((((((((((((((0&(~128))|(1<<7))&(~6144))|(defined_2_bits<<11))&(~128))|(defined_boolean<<7))&(~256))|(defined_boolean<<8))&(~512))|(defined_boolean<<9))&(~1024))|(defined_boolean<<10))&(~16777216))|(defined_boolean<<24))&(~33554432))|(defined_boolean<<25))&(~67108864))|(defined_boolean<<26))&(~134217728))|(defined_boolean<<27))&(~268435456))|(defined_boolean<<28))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~62))|(0<<1))&1024)>>10
// ASM-NEXT: .amdhsa_system_vgpr_workitem_id (((((((((((((((((((((((((((((0&(~128))|(1<<7))&(~6144))|(defined_2_bits<<11))&(~128))|(defined_boolean<<7))&(~256))|(defined_boolean<<8))&(~512))|(defined_boolean<<9))&(~1024))|(defined_boolean<<10))&(~16777216))|(defined_boolean<<24))&(~33554432))|(defined_boolean<<25))&(~67108864))|(defined_boolean<<26))&(~134217728))|(defined_boolean<<27))&(~268435456))|(defined_boolean<<28))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~62))|(0<<1))&6144)>>11
-// ASM-NEXT: .amdhsa_next_free_vgpr 0
-// ASM-NEXT: .amdhsa_next_free_sgpr 0
+// ASM-NEXT: .amdhsa_next_free_vgpr defined_value+4
+// ASM-NEXT: .amdhsa_next_free_sgpr defined_value+5
+// ASM-NEXT: .amdhsa_reserve_vcc defined_boolean
+// ASM-NEXT: .amdhsa_reserve_flat_scratch defined_boolean
// ASM-NEXT: .amdhsa_reserve_xnack_mask 1
-// ASM-NEXT: .amdhsa_float_round_mode_32 (((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~63))|(0<<0))&(~960))|(0<<6))&12288)>>12
-// ASM-NEXT: .amdhsa_float_round_mode_16_64 (((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~63))|(0<<0))&(~960))|(0<<6))&49152)>>14
-// ASM-NEXT: .amdhsa_float_denorm_mode_32 (((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~63))|(0<<0))&(~960))|(0<<6))&196608)>>16
-// ASM-NEXT: .amdhsa_float_denorm_mode_16_64 (((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~63))|(0<<0))&(~960))|(0<<6))&786432)>>18
-// ASM-NEXT: .amdhsa_dx10_clamp (((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~63))|(0<<0))&(~960))|(0<<6))&2097152)>>21
-// ASM-NEXT: .amdhsa_ieee_mode (((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~63))|(0<<0))&(~960))|(0<<6))&8388608)>>23
+// ASM-NEXT: .amdhsa_float_round_mode_32 (((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~63))|((((alignto(max(defined_value+4, 1), 4))/4)-1)<<0))&(~960))|((((alignto(max((defined_value+5)+(extrasgprs(defined_boolean, defined_boolean, 1)), 1), 8))/8)-1)<<6))&12288)>>12
+// ASM-NEXT: .amdhsa_float_round_mode_16_64 (((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~63))|((((alignto(max(defined_value+4, 1), 4))/4)-1)<<0))&(~960))|((((alignto(max((defined_value+5)+(extrasgprs(defined_boolean, defined_boolean, 1)), 1), 8))/8)-1)<<6))&49152)>>14
+// ASM-NEXT: .amdhsa_float_denorm_mode_32 (((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~63))|((((alignto(max(defined_value+4, 1), 4))/4)-1)<<0))&(~960))|((((alignto(max((defined_value+5)+(extrasgprs(defined_boolean, defined_boolean, 1)), 1), 8))/8)-1)<<6))&196608)>>16
+// ASM-NEXT: .amdhsa_float_denorm_mode_16_64 (((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~63))|((((alignto(max(defined_value+4, 1), 4))/4)-1)<<0))&(~960))|((((alignto(max((defined_value+5)+(extrasgprs(defined_boolean, defined_boolean, 1)), 1), 8))/8)-1)<<6))&786432)>>18
+// ASM-NEXT: .amdhsa_dx10_clamp (((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~63))|((((alignto(max(defined_value+4, 1), 4))/4)-1)<<0))&(~960))|((((alignto(max((defined_value+5)+(extrasgprs(defined_boolean, defined_boolean, 1)), 1), 8))/8)-1)<<6))&2097152)>>21
+// ASM-NEXT: .amdhsa_ieee_mode (((((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~12288))|(defined_2_bits<<12))&(~49152))|(defined_2_bits<<14))&(~196608))|(defined_2_bits<<16))&(~786432))|(defined_2_bits<<18))&(~63))|((((alignto(max(defined_value+4, 1), 4))/4)-1)<<0))&(~960))|((((alignto(max((defined_value+5)+(extrasgprs(defined_boolean, defined_boolean, 1)), 1), 8))/8)-1)<<6))&8388608)>>23
// ASM-NEXT: .amdhsa_exception_fp_ieee_invalid_op (((((((((((((((((((((((((((((0&(~128))|(1<<7))&(~6144))|(defined_2_bits<<11))&(~128))|(defined_boolean<<7))&(~256))|(defined_boolean<<8))&(~512))|(defined_boolean<<9))&(~1024))|(defined_boolean<<10))&(~16777216))|(defined_boolean<<24))&(~33554432))|(defined_boolean<<25))&(~67108864))|(defined_boolean<<26))&(~134217728))|(defined_boolean<<27))&(~268435456))|(defined_boolean<<28))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~62))|(0<<1))&16777216)>>24
// ASM-NEXT: .amdhsa_exception_fp_denorm_src (((((((((((((((((((((((((((((0&(~128))|(1<<7))&(~6144))|(defined_2_bits<<11))&(~128))|(defined_boolean<<7))&(~256))|(defined_boolean<<8))&(~512))|(defined_boolean<<9))&(~1024))|(defined_boolean<<10))&(~16777216))|(defined_boolean<<24))&(~33554432))|(defined_boolean<<25))&(~67108864))|(defined_boolean<<26))&(~134217728))|(defined_boolean<<27))&(~268435456))|(defined_boolean<<28))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~62))|(0<<1))&33554432)>>25
// ASM-NEXT: .amdhsa_exception_fp_ieee_div_zero (((((((((((((((((((((((((((((0&(~128))|(1<<7))&(~6144))|(defined_2_bits<<11))&(~128))|(defined_boolean<<7))&(~256))|(defined_boolean<<8))&(~512))|(defined_boolean<<9))&(~1024))|(defined_boolean<<10))&(~16777216))|(defined_boolean<<24))&(~33554432))|(defined_boolean<<25))&(~67108864))|(defined_boolean<<26))&(~134217728))|(defined_boolean<<27))&(~268435456))|(defined_boolean<<28))&(~536870912))|(defined_boolean<<29))&(~1073741824))|(defined_boolean<<30))&(~62))|(0<<1))&67108864)>>26
@@ -152,8 +158,10 @@ expr_defined:
// ASM-NEXT: .amdhsa_system_sgpr_workgroup_id_z 1
// ASM-NEXT: .amdhsa_system_sgpr_workgroup_info 1
// ASM-NEXT: .amdhsa_system_vgpr_workitem_id 3
-// ASM-NEXT: .amdhsa_next_free_vgpr 0
-// ASM-NEXT: .amdhsa_next_free_sgpr 0
+// ASM-NEXT: .amdhsa_next_free_vgpr 44
+// ASM-NEXT: .amdhsa_next_free_sgpr 45
+// ASM-NEXT: .amdhsa_reserve_vcc 1
+// ASM-NEXT: .amdhsa_reserve_flat_scratch 1
// ASM-NEXT: .amdhsa_reserve_xnack_mask 1
// ASM-NEXT: .amdhsa_float_round_mode_32 3
// ASM-NEXT: .amdhsa_float_round_mode_16_64 3
diff --git a/llvm/test/MC/AMDGPU/hsa-sym-exprs-gfx90a.s b/llvm/test/MC/AMDGPU/hsa-sym-exprs-gfx90a.s
index b7f89239160fc..216ae4c42a3d9 100644
--- a/llvm/test/MC/AMDGPU/hsa-sym-exprs-gfx90a.s
+++ b/llvm/test/MC/AMDGPU/hsa-sym-exprs-gfx90a.s
@@ -10,12 +10,12 @@
// OBJDUMP-NEXT: 0000 00000000 00000000 00000000 00000000
// OBJDUMP-NEXT: 0010 00000000 00000000 00000000 00000000
// OBJDUMP-NEXT: 0020 00000000 00000000 00000000 00000100
-// OBJDUMP-NEXT: 0030 0000ac04 81000000 00000000 00000000
+// OBJDUMP-NEXT: 0030 4000ac04 81000000 00000000 00000000
// expr_defined
// OBJDUMP-NEXT: 0040 00000000 00000000 00000000 00000000
// OBJDUMP-NEXT: 0050 00000000 00000000 00000000 00000000
// OBJDUMP-NEXT: 0060 00000000 00000000 00000000 00000100
-// OBJDUMP-NEXT: 0070 0000ac04 81000000 00000000 00000000
+// OBJDUMP-NEXT: 0070 4000ac04 81000000 00000000 00000000
.text
// ASM: .text
@@ -43,9 +43,11 @@ expr_defined:
.amdhsa_ieee_mode defined_boolean
.amdhsa_fp16_overflow defined_boolean
.amdhsa_tg_split defined_boolean
- .amdhsa_next_free_vgpr 0
- .amdhsa_next_free_sgpr 0
+ .amdhsa_next_free_vgpr defined_boolean+1
+ .amdhsa_next_free_sgpr defined_boolean+2
.amdhsa_accum_offset 4
+ .amdhsa_reserve_vcc defined_boolean
+ .amdhsa_reserve_flat_scratch defined_boolean
.end_amdhsa_kernel
.set defined_boolean, 1
@@ -57,9 +59,11 @@ expr_defined:
.amdhsa_ieee_mode defined_boolean
.amdhsa_fp16_overflow defined_boolean
.amdhsa_tg_split defined_boolean
- .amdhsa_next_free_vgpr 0
- .amdhsa_next_free_sgpr 0
+ .amdhsa_next_free_vgpr defined_boolean+1
+ .amdhsa_next_free_sgpr defined_boolean+2
.amdhsa_accum_offset 4
+ .amdhsa_reserve_vcc defined_boolean
+ .amdhsa_reserve_flat_scratch defined_boolean
.end_amdhsa_kernel
// ASM: .amdhsa_kernel expr_defined_later
@@ -82,18 +86,20 @@ expr_defined:
// ASM-NEXT: .amdhsa_system_sgpr_workgroup_id_z (((((((0&(~128))|(1<<7))&(~1))|(defined_boolean<<0))&(~62))|(0<<1))&512)>>9
// ASM-NEXT: .amdhsa_system_sgpr_workgroup_info (((((((0&(~128))|(1<<7))&(~1))|(defined_boolean<<0))&(~62))|(0<<1))&1024)>>10
// ASM-NEXT: .amdhsa_system_vgpr_workitem_id (((((((0&(~128))|(1<<7))&(~1))|(defined_boolean<<0))&(~62))|(0<<1))&6144)>>11
-// ASM-NEXT: .amdhsa_next_free_vgpr 0
-// ASM-NEXT: .amdhsa_next_free_sgpr 0
-// ASM-NEXT: .amdhsa_accum_offset (((((((0&(~65536))|(defined_boolean<<16))&(~63))|(0<<0))&63)>>0)+1)*4
+// ASM-NEXT: .amdhsa_next_free_vgpr defined_boolean+1
+// ASM-NEXT: .amdhsa_next_free_sgpr defined_boolean+2
+// ASM-NEXT: .amdhsa_accum_offset (((((((0&(~65536))|(defined_boolean<<16))&(~63))|(((4/4)-1)<<0))&63)>>0)+1)*4
+// ASM-NEXT: .amdhsa_reserve_vcc defined_boolean
+// ASM-NEXT: .amdhsa_reserve_flat_scratch defined_boolean
// ASM-NEXT: .amdhsa_reserve_xnack_mask 1
-// ASM-NEXT: .amdhsa_float_round_mode_32 (((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~2097152))|(defined_boolean<<21))&(~8388608))|(defined_boolean<<23))&(~67108864))|(defined_boolean<<26))&(~63))|(0<<0))&(~960))|(0<<6))&12288)>>12
-// ASM-NEXT: .amdhsa_float_round_mode_16_64 (((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~2097152))|(defined_boolean<<21))&(~8388608))|(defined_boolean<<23))&(~67108864))|(defined_boolean<<26))&(~63))|(0<<0))&(~960))|(0<<6))&49152)>>14
-// ASM-NEXT: .amdhsa_float_denorm_mode_32 (((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~2097152))|(defined_boolean<<21))&(~8388608))|(defined_boolean<<23))&(~67108864))|(defined_boolean<<26))&(~63))|(0<<0))&(~960))|(0<<6))&196608)>>16
-// ASM-NEXT: .amdhsa_float_denorm_mode_16_64 (((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~2097152))|(defined_boolean<<21))&(~8388608))|(defined_boolean<<23))&(~67108864))|(defined_boolean<<26))&(~63))|(0<<0))&(~960))|(0<<6))&786432)>>18
-// ASM-NEXT: .amdhsa_dx10_clamp (((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~2097152))|(defined_boolean<<21))&(~8388608))|(defined_boolean<<23))&(~67108864))|(defined_boolean<<26))&(~63))|(0<<0))&(~960))|(0<<6))&2097152)>>21
-// ASM-NEXT: .amdhsa_ieee_mode (((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~2097152))|(defined_boolean<<21))&(~8388608))|(defined_boolean<<23))&(~67108864))|(defined_boolean<<26))&(~63))|(0<<0))&(~960))|(0<<6))&8388608)>>23
-// ASM-NEXT: .amdhsa_fp16_overflow (((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~2097152))|(defined_boolean<<21))&(~8388608))|(defined_boolean<<23))&(~67108864))|(defined_boolean<<26))&(~63))|(0<<0))&(~960))|(0<<6))&67108864)>>26
-// ASM-NEXT: .amdhsa_tg_split (((((0&(~65536))|(defined_boolean<<16))&(~63))|(0<<0))&65536)>>16
+// ASM-NEXT: .amdhsa_float_round_mode_32 (((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~2097152))|(defined_boolean<<21))&(~8388608))|(defined_boolean<<23))&(~67108864))|(defined_boolean<<26))&(~63))|((((alignto(max(defined_boolean+1, 1), 8))/8)-1)<<0))&(~960))|((((alignto(max((defined_boolean+2)+(extrasgprs(defined_boolean, defined_boolean, 1)), 1), 8))/8)-1)<<6))&12288)>>12
+// ASM-NEXT: .amdhsa_float_round_mode_16_64 (((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~2097152))|(defined_boolean<<21))&(~8388608))|(defined_boolean<<23))&(~67108864))|(defined_boolean<<26))&(~63))|((((alignto(max(defined_boolean+1, 1), 8))/8)-1)<<0))&(~960))|((((alignto(max((defined_boolean+2)+(extrasgprs(defined_boolean, defined_boolean, 1)), 1), 8))/8)-1)<<6))&49152)>>14
+// ASM-NEXT: .amdhsa_float_denorm_mode_32 (((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~2097152))|(defined_boolean<<21))&(~8388608))|(defined_boolean<<23))&(~67108864))|(defined_boolean<<26))&(~63))|((((alignto(max(defined_boolean+1, 1), 8))/8)-1)<<0))&(~960))|((((alignto(max((defined_boolean+2)+(extrasgprs(defined_boolean, defined_boolean, 1)), 1), 8))/8)-1)<<6))&196608)>>16
+// ASM-NEXT: .amdhsa_float_denorm_mode_16_64 (((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~2097152))|(defined_boolean<<21))&(~8388608))|(defined_boolean<<23))&(~67108864))|(defined_boolean<<26))&(~63))|((((alignto(max(defined_boolean+1, 1), 8))/8)-1)<<0))&(~960))|((((alignto(max((defined_boolean+2)+(extrasgprs(defined_boolean, defined_boolean, 1)), 1), 8))/8)-1)<<6))&786432)>>18
+// ASM-NEXT: .amdhsa_dx10_clamp (((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~2097152))|(defined_boolean<<21))&(~8388608))|(defined_boolean<<23))&(~67108864))|(defined_boolean<<26))&(~63))|((((alignto(max(defined_boolean+1, 1), 8))/8)-1)<<0))&(~960))|((((alignto(max((defined_boolean+2)+(extrasgprs(defined_boolean, defined_boolean, 1)), 1), 8))/8)-1)<<6))&2097152)>>21
+// ASM-NEXT: .amdhsa_ieee_mode (((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~2097152))|(defined_boolean<<21))&(~8388608))|(defined_boolean<<23))&(~67108864))|(defined_boolean<<26))&(~63))|((((alignto(max(defined_boolean+1, 1), 8))/8)-1)<<0))&(~960))|((((alignto(max((defined_boolean+2)+(extrasgprs(defined_boolean, defined_boolean, 1)), 1), 8))/8)-1)<<6))&8388608)>>23
+// ASM-NEXT: .amdhsa_fp16_overflow (((((((((((((((((0&(~786432))|(3<<18))&(~2097152))|(1<<21))&(~8388608))|(1<<23))&(~2097152))|(defined_boolean<<21))&(~8388608))|(defined_boolean<<23))&(~67108864))|(defined_boolean<<26))&(~63))|((((alignto(max(defined_boolean+1, 1), 8))/8)-1)<<0))&(~960))|((((alignto(max((defined_boolean+2)+(extrasgprs(defined_boolean, defined_boolean, 1)), 1), 8))/8)-1)<<6))&67108864)>>26
+// ASM-NEXT: .amdhsa_tg_split (((((0&(~65536))|(defined_boolean<<16))&(~63))|(((4/4)-1)<<0))&65536)>>16
// ASM-NEXT: .amdhsa_exception_fp_ieee_invalid_op (((((((0&(~128))|(1<<7))&(~1))|(defined_boolean<<0))&(~62))|(0<<1))&16777216)>>24
// ASM-NEXT: .amdhsa_exception_fp_denorm_src (((((((0&(~128))|(1<<7))&(~1))|(defined_boolean<<0))&(~62))|(0<<1))&33554432)>>25
// ASM-NEXT: .amdhsa_exception_fp_ieee_div_zero (((((((0&(~128))|(1<<7))&(~1))|(defined_boolean<<0))&(~62))|(0<<1))&67108864)>>26
@@ -126,9 +132,11 @@ expr_defined:
// ASM-NEXT: .amdhsa_system_sgpr_workgroup_id_z 0
// ASM-NEXT: .amdhsa_system_sgpr_workgroup_info 0
// ASM-NEXT: .amdhsa_system_vgpr_workitem_id 0
-// ASM-NEXT: .amdhsa_next_free_vgpr 0
-// ASM-NEXT: .amdhsa_next_free_sgpr 0
+// ASM-NEXT: .amdhsa_next_free_vgpr 2
+// ASM-NEXT: .amdhsa_next_free_sgpr 3
// ASM-NEXT: .amdhsa_accum_offset 4
+// ASM-NEXT: .amdhsa_reserve_vcc 1
+// ASM-NEXT: .amdhsa_reserve_flat_scratch 1
// ASM-NEXT: .amdhsa_reserve_xnack_mask 1
// ASM-NEXT: .amdhsa_float_round_mode_32 0
// ASM-NEXT: .amdhsa_float_round_mode_16_64 0
diff --git a/llvm/test/MC/AMDGPU/hsa-tg-split.s b/llvm/test/MC/AMDGPU/hsa-tg-split.s
index 5a4d3e2c279c5..78235e7125810 100644
--- a/llvm/test/MC/AMDGPU/hsa-tg-split.s
+++ b/llvm/test/MC/AMDGPU/hsa-tg-split.s
@@ -55,6 +55,8 @@ minimal:
// ASM-NEXT: .amdhsa_next_free_vgpr 0
// ASM-NEXT: .amdhsa_next_free_sgpr 0
// ASM-NEXT: .amdhsa_accum_offset 4
+// ASM-NEXT: .amdhsa_reserve_vcc 1
+// ASM-NEXT: .amdhsa_reserve_flat_scratch 1
// ASM-NEXT: .amdhsa_reserve_xnack_mask 1
// ASM-NEXT: .amdhsa_float_round_mode_32 0
// ASM-NEXT: .amdhsa_float_round_mode_16_64 0
diff --git a/llvm/test/MC/AMDGPU/hsa-v4.s b/llvm/test/MC/AMDGPU/hsa-v4.s
index e19dba0f5fd0f..5aa7a58578dec 100644
--- a/llvm/test/MC/AMDGPU/hsa-v4.s
+++ b/llvm/test/MC/AMDGPU/hsa-v4.s
@@ -193,6 +193,7 @@ disabled_user_sgpr:
// ASM: .amdhsa_next_free_vgpr 0
// ASM-NEXT: .amdhsa_next_free_sgpr 27
// ASM-NEXT: .amdhsa_reserve_vcc 0
+// ASM-NEXT: .amdhsa_reserve_flat_scratch 1
// ASM-NEXT: .amdhsa_reserve_xnack_mask 1
// ASM: .amdhsa_float_denorm_mode_16_64 0
// ASM-NEXT: .amdhsa_dx10_clamp 0
diff --git a/llvm/test/MC/AMDGPU/hsa-v5-uses-dynamic-stack.s b/llvm/test/MC/AMDGPU/hsa-v5-uses-dynamic-stack.s
index 4c8849e8540ba..ec3d5bae22d17 100644
--- a/llvm/test/MC/AMDGPU/hsa-v5-uses-dynamic-stack.s
+++ b/llvm/test/MC/AMDGPU/hsa-v5-uses-dynamic-stack.s
@@ -200,6 +200,7 @@ disabled_user_sgpr:
// ASM: .amdhsa_next_free_vgpr 0
// ASM-NEXT: .amdhsa_next_free_sgpr 27
// ASM-NEXT: .amdhsa_reserve_vcc 0
+// ASM-NEXT: .amdhsa_reserve_flat_scratch 1
// ASM-NEXT: .amdhsa_reserve_xnack_mask 1
// ASM: .amdhsa_float_denorm_mode_16_64 0
// ASM-NEXT: .amdhsa_dx10_clamp 0
diff --git a/llvm/unittests/MC/AMDGPU/CMakeLists.txt b/llvm/unittests/MC/AMDGPU/CMakeLists.txt
index be8ff572e6f7d..0a399772e019e 100644
--- a/llvm/unittests/MC/AMDGPU/CMakeLists.txt
+++ b/llvm/unittests/MC/AMDGPU/CMakeLists.txt
@@ -7,8 +7,6 @@ set(LLVM_LINK_COMPONENTS
AMDGPUCodeGen
AMDGPUDesc
AMDGPUInfo
- CodeGen
- Core
MC
Support
TargetParser
@@ -16,5 +14,4 @@ set(LLVM_LINK_COMPONENTS
add_llvm_unittest(AMDGPUMCTests
DwarfRegMappings.cpp
- SIProgramInfoMCExprs.cpp
)
diff --git a/llvm/unittests/MC/AMDGPU/SIProgramInfoMCExprs.cpp b/llvm/unittests/MC/AMDGPU/SIProgramInfoMCExprs.cpp
deleted file mode 100644
index 57828a728931d..0000000000000
--- a/llvm/unittests/MC/AMDGPU/SIProgramInfoMCExprs.cpp
+++ /dev/null
@@ -1,83 +0,0 @@
-//===- llvm/unittests/MC/AMDGPU/SIProgramInfoMCExprs.cpp ------------------===//
-//
-// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
-// See https://llvm.org/LICENSE.txt for license information.
-// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-//
-//===----------------------------------------------------------------------===//
-
-#include "AMDGPUHSAMetadataStreamer.h"
-#include "SIProgramInfo.h"
-#include "llvm/CodeGen/MachineFunction.h"
-#include "llvm/CodeGen/MachineModuleInfo.h"
-#include "llvm/MC/MCContext.h"
-#include "llvm/MC/MCExpr.h"
-#include "llvm/MC/MCStreamer.h"
-#include "llvm/MC/MCSymbol.h"
-#include "llvm/MC/MCTargetOptions.h"
-#include "llvm/MC/TargetRegistry.h"
-#include "llvm/Support/TargetSelect.h"
-#include "llvm/Target/TargetMachine.h"
-#include "gtest/gtest.h"
-
-using namespace llvm;
-
-class SIProgramInfoMCExprsTest : public testing::Test {
-protected:
- std::unique_ptr<LLVMTargetMachine> TM;
- std::unique_ptr<LLVMContext> Ctx;
- std::unique_ptr<MachineModuleInfo> MMI;
- std::unique_ptr<MachineFunction> MF;
- std::unique_ptr<Module> M;
-
- SIProgramInfo PI;
-
- static void SetUpTestSuite() {
- LLVMInitializeAMDGPUTargetInfo();
- LLVMInitializeAMDGPUTarget();
- LLVMInitializeAMDGPUTargetMC();
- }
-
- SIProgramInfoMCExprsTest() {
- std::string Triple = "amdgcn-amd-amdhsa";
- std::string CPU = "gfx1010";
- std::string FS = "";
-
- std::string Error;
- const Target *TheTarget = TargetRegistry::lookupTarget(Triple, Error);
- TargetOptions Options;
-
- TM.reset(static_cast<LLVMTargetMachine *>(TheTarget->createTargetMachine(
- Triple, CPU, FS, Options, std::nullopt, std::nullopt)));
-
- Ctx = std::make_unique<LLVMContext>();
- M = std::make_unique<Module>("Module", *Ctx);
- M->setDataLayout(TM->createDataLayout());
- auto *FType = FunctionType::get(Type::getVoidTy(*Ctx), false);
- auto *F = Function::Create(FType, GlobalValue::ExternalLinkage, "Test", *M);
- MMI = std::make_unique<MachineModuleInfo>(TM.get());
-
- auto *ST = TM->getSubtargetImpl(*F);
-
- MF = std::make_unique<MachineFunction>(*F, *TM, *ST, 1, *MMI);
- MF->initTargetMachineFunctionInfo(*ST);
- PI.reset(*MF.get());
- }
-};
-
-TEST_F(SIProgramInfoMCExprsTest, TestDeathHSAKernelEmit) {
- MCContext &Ctx = MF->getContext();
- MCSymbol *Sym = Ctx.getOrCreateSymbol("Unknown");
- PI.ScratchSize = MCSymbolRefExpr::create(Sym, Ctx);
-
- auto &Func = MF->getFunction();
- Func.setCallingConv(CallingConv::AMDGPU_KERNEL);
- AMDGPU::HSAMD::MetadataStreamerMsgPackV4 MD;
-
- testing::internal::CaptureStderr();
- MD.emitKernel(*MF, PI);
- std::string err = testing::internal::GetCapturedStderr();
- EXPECT_EQ(
- err, "<unknown>:0: error: could not resolve expression when required.\n");
- EXPECT_TRUE(Ctx.hadError());
-}
More information about the llvm-commits
mailing list