AMDGPU: Register a null MC streamer for -emit-codegen-only

[arsenm]

Summary

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.

[MaskRay]

LGTM with some nits

clang/test/Misc/backend-resource-limit-diagnostics.hip

1

The layering is strange. This patch ideally should land before the clang patch.

llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.h

66

{ return true; }

Use clang-format

[arsenm]

7d568cdc9d68ec295f1e141c8bcad344873f51cd

clang/test/Misc/backend-resource-limit-diagnostics.hip

-// RUN: not %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx906 -x hip -fcuda-is-device -S -o /dev/null %s 2>&1 | FileCheck %s
-// FIXME: Use -emit-codegen-only
+// 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;

llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp

TargetRegistry::RegisterAsmPrinter(getTheAMDGPUTarget(),
                                      llvm::createR600AsmPrinterPass);
   TargetRegistry::RegisterAsmPrinter(getTheGCNTarget(),
                                      createAMDGPUAsmPrinterPass);
 }
 
 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());
     } else if (isHsaAbiVersion3(getGlobalSTI())) {
       HSAMetadataStream.reset(new HSAMD::MetadataStreamerMsgPackV3());
     } else if (isHsaAbiVersion5(getGlobalSTI())) {
       HSAMetadataStream.reset(new HSAMD::MetadataStreamerMsgPackV5());
     } else {
getTargetStreamer()->EmitDirectiveHSACodeObjectISAV2(
       Version.Major, Version.Minor, Version.Stepping, "AMD", "AMDGPU");
 }
 
 void AMDGPUAsmPrinter::emitEndOfAsmFile(Module &M) {
   // Init target streamer if it has not yet happened
   if (!IsTargetStreamerInitialized)
     initTargetStreamer(M);
 
-  // Following code requires TargetStreamer to be present.
-  if (!getTargetStreamer())
-    return;
-
   if (TM.getTargetTriple().getOS() != Triple::AMDHSA ||
       isHsaAbiVersion2(getGlobalSTI()))
     getTargetStreamer()->EmitISAVersion();
 
   // Emit HSA Metadata (NT_AMD_AMDGPU_HSA_METADATA).
   // Emit HSA Metadata (NT_AMD_HSA_METADATA).
   if (TM.getTargetTriple().getOS() == Triple::AMDHSA) {
     HSAMetadataStream->end();

 
 void AMDGPUAsmPrinter::emitFunctionBodyStart() {
   const SIMachineFunctionInfo &MFI = *MF->getInfo<SIMachineFunctionInfo>();
   const GCNSubtarget &STM = MF->getSubtarget<GCNSubtarget>();
   const Function &F = MF->getFunction();
 
   // TODO: Which one is called first, emitStartOfAsmFile or
   // emitFunctionBodyStart?
-  if (getTargetStreamer() && !getTargetStreamer()->getTargetID())
+  if (!getTargetStreamer()->getTargetID())
     initializeTargetID(*F.getParent());
 
   const auto &FunctionTargetID = STM.getTargetID();
   // Make sure function's xnack settings are compatible with module's
   // xnack settings.
   if (FunctionTargetID.isXnackSupported() &&
       FunctionTargetID.getXnackSetting() != IsaInfo::TargetIDSetting::Any &&
       FunctionTargetID.getXnackSetting() != getTargetStreamer()->getTargetID()->getXnackSetting()) {
if (GVSym->isDefined() || GVSym->isVariable())
                          "' is already defined");
 
     const DataLayout &DL = GV->getParent()->getDataLayout();
     uint64_t Size = DL.getTypeAllocSize(GV->getValueType());
     Align Alignment = GV->getAlign().value_or(Align(4));
 
     emitVisibility(GVSym, GV->getVisibility(), !GV->isDeclaration());
     emitLinkage(GV, GVSym);
-    if (auto TS = getTargetStreamer())
-      TS->emitAMDGPULDS(GVSym, Size, Alignment);
+    auto TS = getTargetStreamer();
     return;
   }
 
   AsmPrinter::emitGlobalVariable(GV);
 }
 
 bool AMDGPUAsmPrinter::doFinalization(Module &M) {
   // Pad with s_code_end to help tools and guard against instruction prefetch

llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUMCTargetDesc.cpp

 }
 
 static MCTargetStreamer * createAMDGPUObjectTargetStreamer(
                                                    MCStreamer &S,
                                                    const MCSubtargetInfo &STI) {
   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,
                                     std::unique_ptr<MCCodeEmitter> &&Emitter,
                                     bool RelaxAll) {
   return createAMDGPUELFStreamer(T, Context, std::move(MAB), std::move(OW),
                                  std::move(Emitter), RelaxAll);
 }
extern "C" LLVM_EXTERNAL_VISIBILITY void LLVMInitializeAMDGPUTargetMC() {
   // GCN specific registration
   TargetRegistry::RegisterMCCodeEmitter(getTheGCNTarget(),
                                         createSIMCCodeEmitter);
 
   TargetRegistry::RegisterAsmTargetStreamer(getTheGCNTarget(),
                                             createAMDGPUAsmTargetStreamer);
   TargetRegistry::RegisterObjectTargetStreamer(
       getTheGCNTarget(), createAMDGPUObjectTargetStreamer);
+  TargetRegistry::RegisterNullTargetStreamer(getTheGCNTarget(),
+                                             createAMDGPUNullTargetStreamer);
 }

llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.h

protected:
 
   MCContext &getContext() const { return Streamer.getContext(); }
 
 public:
   AMDGPUTargetStreamer(MCStreamer &S) : MCTargetStreamer(S) {}
 
   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;
+                             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);
 
   /// \returns True on success, false on failure.
   virtual bool EmitHSAMetadataV3(StringRef HSAMetadataString);
 
   /// Emit HSA Metadata
   ///
   /// When \p Strict is true, known metadata elements must already be
   /// well-typed. When \p Strict is false, known types are inferred and
   /// 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);
 
   const Optional<AMDGPU::IsaInfo::AMDGPUTargetID> &getTargetID() const {
     return TargetID;
   }
   Optional<AMDGPU::IsaInfo::AMDGPUTargetID> &getTargetID() {