Index: clang/test/Misc/backend-resource-limit-diagnostics.hip
===================================================================
--- /dev/null
+++ clang/test/Misc/backend-resource-limit-diagnostics.hip
@@ -0,0 +1,19 @@
+// RUN: not %clang_cc1 -debug-info-kind=standalone -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
+
+#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: <stdin>:[[@LINE-5]]:0: local memory (480000) exceeds limit (65536) in function '_Z12use_huge_ldsIiEvv'
+ template
+ __global__ void use_huge_lds<int>();
+
+// CHECK: error: <stdin>:[[@LINE-9]]:0: local memory (960000) exceeds limit (65536) in function '_Z12use_huge_ldsIdEvv'
+ template
+ __global__ void use_huge_lds<double>();
Index: clang/test/Misc/backend-resource-limit-diagnostics.cl
===================================================================
--- clang/test/Misc/backend-resource-limit-diagnostics.cl
+++ clang/test/Misc/backend-resource-limit-diagnostics.cl
@@ -1,9 +1,8 @@
 // REQUIRES: amdgpu-registered-target
-// RUN: not %clang_cc1 -emit-codegen-only -triple=amdgcn-- %s 2>&1 | FileCheck %s
+// RUN: not %clang_cc1 -debug-info-kind=standalone -x cl -emit-codegen-only -triple=amdgcn-- < %s 2>&1 | FileCheck %s
 
-// CHECK: error: local memory (480000) exceeds limit (32768) in function 'use_huge_lds'
-kernel void use_huge_lds()
-{
+// CHECK: error: <stdin>:[[@LINE+1]]:0: local memory (480000) exceeds limit (32768) in function 'use_huge_lds'
+kernel void use_huge_lds() {
     volatile local int huge[120000];
     huge[0] = 2;
 }
Index: clang/test/Misc/backend-stack-frame-diagnostics-fallback.cpp
===================================================================
--- clang/test/Misc/backend-stack-frame-diagnostics-fallback.cpp
+++ clang/test/Misc/backend-stack-frame-diagnostics-fallback.cpp
@@ -13,7 +13,7 @@
   };
 
   // CHECK: warning: stack frame size ([[#]]) exceeds limit ([[#]]) in 'frameSizeThunkWarning::B::f()'
-  // CHECK: warning: stack frame size ([[#]]) exceeds limit ([[#]]) in function '_ZTv0_n12_N21frameSizeThunkWarning1B1fEv'
+  // CHECK: warning: <unknown>:0:0: stack frame size ([[#]]) exceeds limit ([[#]]) in function '_ZTv0_n12_N21frameSizeThunkWarning1B1fEv'
   void B::f() {
     volatile int x = 0; // Ensure there is stack usage.
   }
Index: llvm/include/llvm/IR/DiagnosticInfo.h
===================================================================
--- llvm/include/llvm/IR/DiagnosticInfo.h
+++ llvm/include/llvm/IR/DiagnosticInfo.h
@@ -181,62 +181,6 @@
   }
 };
 
-/// Diagnostic information for stack size etc. reporting.
-/// This is basically a function and a size.
-class DiagnosticInfoResourceLimit : public DiagnosticInfo {
-private:
-  /// The function that is concerned by this resource limit diagnostic.
-  const Function &Fn;
-
-  /// Description of the resource type (e.g. stack size)
-  const char *ResourceName;
-
-  /// The computed size usage
-  uint64_t ResourceSize;
-
-  // Threshould passed
-  uint64_t ResourceLimit;
-
-public:
-  /// \p The function that is concerned by this stack size diagnostic.
-  /// \p The computed stack size.
-  DiagnosticInfoResourceLimit(const Function &Fn, const char *ResourceName,
-                              uint64_t ResourceSize, uint64_t ResourceLimit,
-                              DiagnosticSeverity Severity = DS_Warning,
-                              DiagnosticKind Kind = DK_ResourceLimit)
-      : DiagnosticInfo(Kind, Severity), Fn(Fn), ResourceName(ResourceName),
-        ResourceSize(ResourceSize), ResourceLimit(ResourceLimit) {}
-
-  const Function &getFunction() const { return Fn; }
-  const char *getResourceName() const { return ResourceName; }
-  uint64_t getResourceSize() const { return ResourceSize; }
-  uint64_t getResourceLimit() const { return ResourceLimit; }
-
-  /// \see DiagnosticInfo::print.
-  void print(DiagnosticPrinter &DP) const override;
-
-  static bool classof(const DiagnosticInfo *DI) {
-    return DI->getKind() == DK_ResourceLimit || DI->getKind() == DK_StackSize;
-  }
-};
-
-class DiagnosticInfoStackSize : public DiagnosticInfoResourceLimit {
-  void anchor() override;
-public:
-  DiagnosticInfoStackSize(const Function &Fn, uint64_t StackSize,
-                          uint64_t StackLimit,
-                          DiagnosticSeverity Severity = DS_Warning)
-      : DiagnosticInfoResourceLimit(Fn, "stack frame size", StackSize,
-                                    StackLimit, Severity, DK_StackSize) {}
-
-  uint64_t getStackSize() const { return getResourceSize(); }
-  uint64_t getStackLimit() const { return getResourceLimit(); }
-
-  static bool classof(const DiagnosticInfo *DI) {
-    return DI->getKind() == DK_StackSize;
-  }
-};
-
 /// Diagnostic information for debug metadata version reporting.
 /// This is basically a module and a version.
 class DiagnosticInfoDebugMetadataVersion : public DiagnosticInfo {
@@ -409,6 +353,61 @@
   DiagnosticLocation Loc;
 };
 
+/// Diagnostic information for stack size etc. reporting.
+/// This is basically a function and a size.
+class DiagnosticInfoResourceLimit : public DiagnosticInfoWithLocationBase {
+private:
+  /// The function that is concerned by this resource limit diagnostic.
+  const Function &Fn;
+
+  /// Description of the resource type (e.g. stack size)
+  const char *ResourceName;
+
+  /// The computed size usage
+  uint64_t ResourceSize;
+
+  // Threshould passed
+  uint64_t ResourceLimit;
+
+public:
+  /// \p The function that is concerned by this stack size diagnostic.
+  /// \p The computed stack size.
+  DiagnosticInfoResourceLimit(const Function &Fn, const char *ResourceName,
+                              uint64_t ResourceSize, uint64_t ResourceLimit,
+                              DiagnosticSeverity Severity = DS_Warning,
+                              DiagnosticKind Kind = DK_ResourceLimit);
+
+  const Function &getFunction() const { return Fn; }
+  const char *getResourceName() const { return ResourceName; }
+  uint64_t getResourceSize() const { return ResourceSize; }
+  uint64_t getResourceLimit() const { return ResourceLimit; }
+
+  /// \see DiagnosticInfo::print.
+  void print(DiagnosticPrinter &DP) const override;
+
+  static bool classof(const DiagnosticInfo *DI) {
+    return DI->getKind() == DK_ResourceLimit || DI->getKind() == DK_StackSize;
+  }
+};
+
+class DiagnosticInfoStackSize : public DiagnosticInfoResourceLimit {
+  void anchor() override;
+
+public:
+  DiagnosticInfoStackSize(const Function &Fn, uint64_t StackSize,
+                          uint64_t StackLimit,
+                          DiagnosticSeverity Severity = DS_Warning)
+      : DiagnosticInfoResourceLimit(Fn, "stack frame size", StackSize,
+                                    StackLimit, Severity, DK_StackSize) {}
+
+  uint64_t getStackSize() const { return getResourceSize(); }
+  uint64_t getStackLimit() const { return getResourceLimit(); }
+
+  static bool classof(const DiagnosticInfo *DI) {
+    return DI->getKind() == DK_StackSize;
+  }
+};
+
 /// Common features for diagnostics dealing with optimization remarks
 /// that are used by both IR and MIR passes.
 class DiagnosticInfoOptimizationBase : public DiagnosticInfoWithLocationBase {
Index: llvm/lib/IR/DiagnosticInfo.cpp
===================================================================
--- llvm/lib/IR/DiagnosticInfo.cpp
+++ llvm/lib/IR/DiagnosticInfo.cpp
@@ -65,9 +65,17 @@
     DP << " at line " << getLocCookie();
 }
 
+DiagnosticInfoResourceLimit::DiagnosticInfoResourceLimit(
+    const Function &Fn, const char *ResourceName, uint64_t ResourceSize,
+    uint64_t ResourceLimit, DiagnosticSeverity Severity, DiagnosticKind Kind)
+    : DiagnosticInfoWithLocationBase(Kind, Severity, Fn, Fn.getSubprogram()),
+      Fn(Fn), ResourceName(ResourceName), ResourceSize(ResourceSize),
+      ResourceLimit(ResourceLimit) {}
+
 void DiagnosticInfoResourceLimit::print(DiagnosticPrinter &DP) const {
-  DP << getResourceName() << " (" << getResourceSize() << ") exceeds limit ("
-     << getResourceLimit() << ") in function '" << getFunction() << '\'';
+  DP << getLocationStr() << ": " << getResourceName() << " ("
+     << getResourceSize() << ") exceeds limit (" << getResourceLimit()
+     << ") in function '" << getFunction() << '\'';
 }
 
 void DiagnosticInfoDebugMetadataVersion::print(DiagnosticPrinter &DP) const {
Index: llvm/test/CodeGen/AMDGPU/exceed-max-sgprs.ll
===================================================================
--- llvm/test/CodeGen/AMDGPU/exceed-max-sgprs.ll
+++ llvm/test/CodeGen/AMDGPU/exceed-max-sgprs.ll
@@ -1,6 +1,6 @@
 ; RUN: not llc -march=amdgcn -verify-machineinstrs < %s 2>&1 | FileCheck -check-prefix=ERROR %s
 
-; ERROR: error: scalar registers (106) exceeds limit (104) in function 'use_too_many_sgprs_tahiti'
+; ERROR: error: <unknown>:0:0: scalar registers (106) exceeds limit (104) in function 'use_too_many_sgprs_tahiti'
 define amdgpu_kernel void @use_too_many_sgprs_tahiti() #0 {
   call void asm sideeffect "", "~{s[0:7]}" ()
   call void asm sideeffect "", "~{s[8:15]}" ()
@@ -19,7 +19,7 @@
   ret void
 }
 
-; ERROR: error: scalar registers (106) exceeds limit (104) in function 'use_too_many_sgprs_bonaire'
+; ERROR: error: <unknown>:0:0: scalar registers (106) exceeds limit (104) in function 'use_too_many_sgprs_bonaire'
 define amdgpu_kernel void @use_too_many_sgprs_bonaire() #1 {
   call void asm sideeffect "", "~{s[0:7]}" ()
   call void asm sideeffect "", "~{s[8:15]}" ()
@@ -38,7 +38,7 @@
   ret void
 }
 
-; ERROR: error: scalar registers (108) exceeds limit (104) in function 'use_too_many_sgprs_bonaire_flat_scr'
+; ERROR: error: <unknown>:0:0: scalar registers (108) exceeds limit (104) in function 'use_too_many_sgprs_bonaire_flat_scr'
 define amdgpu_kernel void @use_too_many_sgprs_bonaire_flat_scr() #1 {
   call void asm sideeffect "", "~{s[0:7]}" ()
   call void asm sideeffect "", "~{s[8:15]}" ()
@@ -58,7 +58,7 @@
   ret void
 }
 
-; ERROR: error: scalar registers (98) exceeds limit (96) in function 'use_too_many_sgprs_iceland'
+; ERROR: error: <unknown>:0:0: scalar registers (98) exceeds limit (96) in function 'use_too_many_sgprs_iceland'
 define amdgpu_kernel void @use_too_many_sgprs_iceland() #2 {
   call void asm sideeffect "", "~{vcc}" ()
   call void asm sideeffect "", "~{s[0:7]}" ()
@@ -76,7 +76,7 @@
   ret void
 }
 
-; ERROR: error: addressable scalar registers (103) exceeds limit (102) in function 'use_too_many_sgprs_fiji'
+; ERROR: error: <unknown>:0:0: addressable scalar registers (103) exceeds limit (102) in function 'use_too_many_sgprs_fiji'
 define amdgpu_kernel void @use_too_many_sgprs_fiji() #3 {
   call void asm sideeffect "", "~{s[0:7]}" ()
   call void asm sideeffect "", "~{s[8:15]}" ()
Index: llvm/test/CodeGen/AMDGPU/stack-size-overflow.ll
===================================================================
--- llvm/test/CodeGen/AMDGPU/stack-size-overflow.ll
+++ llvm/test/CodeGen/AMDGPU/stack-size-overflow.ll
@@ -3,7 +3,7 @@
 
 declare void @llvm.memset.p5i8.i32(i8 addrspace(5)* nocapture, i8, i32, i32, i1) #1
 
-; ERROR: error: stack frame size (131061) exceeds limit (131056) in function 'stack_size_limit_wave64'
+; ERROR: error: <unknown>:0:0: stack frame size (131061) exceeds limit (131056) in function 'stack_size_limit_wave64'
 ; GCN: ; ScratchSize: 131061
 define amdgpu_kernel void @stack_size_limit_wave64() #0 {
 entry:
@@ -13,7 +13,7 @@
   ret void
 }
 
-; ERROR: error: stack frame size (262117) exceeds limit (262112) in function 'stack_size_limit_wave32'
+; ERROR: error: <unknown>:0:0: stack frame size (262117) exceeds limit (262112) in function 'stack_size_limit_wave32'
 ; GCN: ; ScratchSize: 262117
 define amdgpu_kernel void @stack_size_limit_wave32() #1 {
 entry:
Index: llvm/test/CodeGen/ARM/warn-stack.ll
===================================================================
--- llvm/test/CodeGen/ARM/warn-stack.ll
+++ llvm/test/CodeGen/ARM/warn-stack.ll
@@ -12,7 +12,7 @@
   ret void
 }
 
-; CHECK: warning: stack frame size (92) exceeds limit (80) in function 'warn'
+; CHECK: warning: <unknown>:0:0: stack frame size (92) exceeds limit (80) in function 'warn'
 define void @warn() nounwind ssp "frame-pointer"="all" "warn-stack-size"="80" {
 entry:
   %buffer = alloca [80 x i8], align 1
Index: llvm/test/CodeGen/X86/warn-stack.ll
===================================================================
--- llvm/test/CodeGen/X86/warn-stack.ll
+++ llvm/test/CodeGen/X86/warn-stack.ll
@@ -11,7 +11,7 @@
   ret void
 }
 
-; CHECK: warning: stack frame size (88) exceeds limit (80) in function 'warn'
+; CHECK: warning: <unknown>:0:0: stack frame size (88) exceeds limit (80) in function 'warn'
 define void @warn() nounwind ssp "warn-stack-size"="80" {
 entry:
   %buffer = alloca [80 x i8], align 1
@@ -24,7 +24,7 @@
 ; combined stack size of the machine stack and unsafe stack will exceed the
 ; warning threshold
 
-; CHECK: warning: stack frame size (120) exceeds limit (80) in function 'warn_safestack'
+; CHECK: warning: <unknown>:0:0: stack frame size (120) exceeds limit (80) in function 'warn_safestack'
 define void @warn_safestack() nounwind ssp safestack "warn-stack-size"="80" {
 entry:
   %buffer = alloca [80 x i8], align 1