Index: clang/test/Misc/backend-resource-limit-diagnostics.hip
===================================================================
--- clang/test/Misc/backend-resource-limit-diagnostics.hip
+++ clang/test/Misc/backend-resource-limit-diagnostics.hip
@@ -1,5 +1,4 @@
-// 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))
Index: llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp
===================================================================
--- llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp
+++ llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp
@@ -88,6 +88,8 @@
 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 @@
   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 @@
 
   // 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 @@
 
     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;
   }
 
Index: llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUMCTargetDesc.cpp
===================================================================
--- llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUMCTargetDesc.cpp
+++ llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUMCTargetDesc.cpp
@@ -105,6 +105,10 @@
   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 @@
                                             createAMDGPUAsmTargetStreamer);
   TargetRegistry::RegisterObjectTargetStreamer(
       getTheGCNTarget(), createAMDGPUObjectTargetStreamer);
+  TargetRegistry::RegisterNullTargetStreamer(getTheGCNTarget(),
+                                             createAMDGPUNullTargetStreamer);
 }
Index: llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.h
===================================================================
--- llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.h
+++ llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.h
@@ -45,25 +45,25 @@
 
   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);
@@ -78,18 +78,18 @@
   /// 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);