[clang] 7d568cd - AMDGPU: Register a null MC streamer for -emit-codegen-only
Matt Arsenault via cfe-commits
cfe-commits at lists.llvm.org
Fri Oct 28 16:39:16 PDT 2022
Author: Matt Arsenault
Date: 2022-10-28T16:39:09-07:00
New Revision: 7d568cdc9d68ec295f1e141c8bcad344873f51cd
URL: https://github.com/llvm/llvm-project/commit/7d568cdc9d68ec295f1e141c8bcad344873f51cd
DIFF: https://github.com/llvm/llvm-project/commit/7d568cdc9d68ec295f1e141c8bcad344873f51cd.diff
LOG: AMDGPU: Register a null MC streamer for -emit-codegen-only
For some reason null is a valid MC target, used from clang with
-emit-codegen-only. Previously the target streamer was null,
which was inconsistently null checked resulting in crashes
if using amdhsa.
Added:
clang/test/Misc/backend-resource-limit-diagnostics.hip
Modified:
llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp
llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUMCTargetDesc.cpp
llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.h
Removed:
################################################################################
diff --git a/clang/test/Misc/backend-resource-limit-diagnostics.hip b/clang/test/Misc/backend-resource-limit-diagnostics.hip
new file mode 100644
index 0000000000000..ac31dced434a6
--- /dev/null
+++ b/clang/test/Misc/backend-resource-limit-diagnostics.hip
@@ -0,0 +1,18 @@
+// RUN: not %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx906 -x hip -fcuda-is-device -emit-codegen-only %s 2>&1 | FileCheck %s
+
+#define __global__ __attribute__((global))
+#define __shared__ __attribute__((shared))
+
+template<typename T>
+__global__ void use_huge_lds() {
+ volatile __shared__ T huge[120000];
+ huge[0] = 2;
+}
+
+// CHECK: error: local memory (480000) exceeds limit (65536) in function '_Z12use_huge_ldsIiEvv'
+template
+__global__ void use_huge_lds<int>();
+
+// CHECK: error: local memory (960000) exceeds limit (65536) in function '_Z12use_huge_ldsIdEvv'
+template
+__global__ void use_huge_lds<double>();
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp b/llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp
index be08b7f721051..fd930fcaa6434 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp
@@ -88,6 +88,8 @@ extern "C" void LLVM_EXTERNAL_VISIBILITY LLVMInitializeAMDGPUAsmPrinter() {
AMDGPUAsmPrinter::AMDGPUAsmPrinter(TargetMachine &TM,
std::unique_ptr<MCStreamer> Streamer)
: AsmPrinter(TM, std::move(Streamer)) {
+ assert(OutStreamer && "AsmPrinter constructed without streamer");
+
if (TM.getTargetTriple().getOS() == Triple::AMDHSA) {
if (isHsaAbiVersion2(getGlobalSTI())) {
HSAMetadataStream.reset(new HSAMD::MetadataStreamerYamlV2());
@@ -158,10 +160,6 @@ void AMDGPUAsmPrinter::emitEndOfAsmFile(Module &M) {
if (!IsTargetStreamerInitialized)
initTargetStreamer(M);
- // Following code requires TargetStreamer to be present.
- if (!getTargetStreamer())
- return;
-
if (TM.getTargetTriple().getOS() != Triple::AMDHSA ||
isHsaAbiVersion2(getGlobalSTI()))
getTargetStreamer()->EmitISAVersion();
@@ -197,7 +195,7 @@ void AMDGPUAsmPrinter::emitFunctionBodyStart() {
// TODO: Which one is called first, emitStartOfAsmFile or
// emitFunctionBodyStart?
- if (getTargetStreamer() && !getTargetStreamer()->getTargetID())
+ if (!getTargetStreamer()->getTargetID())
initializeTargetID(*F.getParent());
const auto &FunctionTargetID = STM.getTargetID();
@@ -338,8 +336,8 @@ void AMDGPUAsmPrinter::emitGlobalVariable(const GlobalVariable *GV) {
emitVisibility(GVSym, GV->getVisibility(), !GV->isDeclaration());
emitLinkage(GV, GVSym);
- if (auto TS = getTargetStreamer())
- TS->emitAMDGPULDS(GVSym, Size, Alignment);
+ auto TS = getTargetStreamer();
+ TS->emitAMDGPULDS(GVSym, Size, Alignment);
return;
}
diff --git a/llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUMCTargetDesc.cpp b/llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUMCTargetDesc.cpp
index fba4b1a3db661..2aa0572811b3c 100644
--- a/llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUMCTargetDesc.cpp
+++ b/llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUMCTargetDesc.cpp
@@ -105,6 +105,10 @@ static MCTargetStreamer * createAMDGPUObjectTargetStreamer(
return new AMDGPUTargetELFStreamer(S, STI);
}
+static MCTargetStreamer *createAMDGPUNullTargetStreamer(MCStreamer &S) {
+ return new AMDGPUTargetStreamer(S);
+}
+
static MCStreamer *createMCStreamer(const Triple &T, MCContext &Context,
std::unique_ptr<MCAsmBackend> &&MAB,
std::unique_ptr<MCObjectWriter> &&OW,
@@ -172,4 +176,6 @@ extern "C" LLVM_EXTERNAL_VISIBILITY void LLVMInitializeAMDGPUTargetMC() {
createAMDGPUAsmTargetStreamer);
TargetRegistry::RegisterObjectTargetStreamer(
getTheGCNTarget(), createAMDGPUObjectTargetStreamer);
+ TargetRegistry::RegisterNullTargetStreamer(getTheGCNTarget(),
+ createAMDGPUNullTargetStreamer);
}
diff --git a/llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.h b/llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.h
index a857fd00a8555..2bb9c4a6396b2 100644
--- a/llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.h
+++ b/llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.h
@@ -45,25 +45,25 @@ class AMDGPUTargetStreamer : public MCTargetStreamer {
AMDGPUPALMetadata *getPALMetadata() { return &PALMetadata; }
- virtual void EmitDirectiveAMDGCNTarget() = 0;
+ virtual void EmitDirectiveAMDGCNTarget(){};
virtual void EmitDirectiveHSACodeObjectVersion(uint32_t Major,
- uint32_t Minor) = 0;
+ uint32_t Minor){};
virtual void EmitDirectiveHSACodeObjectISAV2(uint32_t Major, uint32_t Minor,
uint32_t Stepping,
StringRef VendorName,
- StringRef ArchName) = 0;
+ StringRef ArchName){};
- virtual void EmitAMDKernelCodeT(const amd_kernel_code_t &Header) = 0;
+ virtual void EmitAMDKernelCodeT(const amd_kernel_code_t &Header){};
- virtual void EmitAMDGPUSymbolType(StringRef SymbolName, unsigned Type) = 0;
+ virtual void EmitAMDGPUSymbolType(StringRef SymbolName, unsigned Type){};
- virtual void emitAMDGPULDS(MCSymbol *Symbol, unsigned Size,
- Align Alignment) = 0;
+ virtual void emitAMDGPULDS(MCSymbol *Symbol, unsigned Size, Align Alignment) {
+ }
/// \returns True on success, false on failure.
- virtual bool EmitISAVersion() = 0;
+ virtual bool EmitISAVersion() { return true; }
/// \returns True on success, false on failure.
virtual bool EmitHSAMetadataV2(StringRef HSAMetadataString);
@@ -78,18 +78,22 @@ class AMDGPUTargetStreamer : public MCTargetStreamer {
/// the \p HSAMetadata structure is updated with the correct types.
///
/// \returns True on success, false on failure.
- virtual bool EmitHSAMetadata(msgpack::Document &HSAMetadata, bool Strict) = 0;
+ virtual bool EmitHSAMetadata(msgpack::Document &HSAMetadata, bool Strict) {
+ return true;
+ }
/// \returns True on success, false on failure.
- virtual bool EmitHSAMetadata(const AMDGPU::HSAMD::Metadata &HSAMetadata) = 0;
+ virtual bool EmitHSAMetadata(const AMDGPU::HSAMD::Metadata &HSAMetadata) {
+ return true;
+ }
/// \returns True on success, false on failure.
- virtual bool EmitCodeEnd(const MCSubtargetInfo &STI) = 0;
+ virtual bool EmitCodeEnd(const MCSubtargetInfo &STI) { return true; }
virtual void EmitAmdhsaKernelDescriptor(
const MCSubtargetInfo &STI, StringRef KernelName,
const amdhsa::kernel_descriptor_t &KernelDescriptor, uint64_t NextVGPR,
- uint64_t NextSGPR, bool ReserveVCC, bool ReserveFlatScr) = 0;
+ uint64_t NextSGPR, bool ReserveVCC, bool ReserveFlatScr){};
static StringRef getArchNameFromElfMach(unsigned ElfMach);
static unsigned getElfMach(StringRef GPU);
More information about the cfe-commits
mailing list