DiagnosticInfo: Report function location for resource limits

We have some odd redundancy where clang specially handles
the stack size case. If clang prints it, the source location is first
followed by "warning". The backend diagnostic, as printed by other tools
puts "warning" first.

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'
+// CHECK: error: <stdin>:[[@LINE+1]]:0: local memory (480000) exceeds limit (32768) in function 'use_huge_lds'
-kernel void use_huge_lds()
+kernel void use_huge_lds() {
-{
     volatile local int huge[120000];
     huge[0] = 2;
 }

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

@@ -1,5 +1,5 @@
 // REQUIRES: amdgpu-registered-target
-// RUN: not %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx906 -x hip -fcuda-is-device -emit-codegen-only %s 2>&1 | FileCheck %s
+// RUN: not %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx906 -debug-info-kind=standalone -x hip -fcuda-is-device -emit-codegen-only < %s 2>&1 | FileCheck -DFILE=%s %s
 
 #define __global__ __attribute__((global))
 #define __shared__ __attribute__((shared))
@@ -10,10 +10,10 @@ __global__ void use_huge_lds() {
   huge[0] = 2;
 }
 
-// CHECK: error: local memory (480000) exceeds limit (65536) in function '_Z12use_huge_ldsIiEvv'
+// 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: local memory (960000) exceeds limit (65536) in function '_Z12use_huge_ldsIdEvv'
+// CHECK: error: <stdin>:[[#@LINE-9]]:0: local memory (960000) exceeds limit (65536) in function '_Z12use_huge_ldsIdEvv'
 template
 __global__ void use_huge_lds<double>();

clang/test/Misc/backend-stack-frame-diagnostics-fallback.cpp

@@ -13,7 +13,7 @@ namespace frameSizeThunkWarning {
   };
 
   // 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.
   }

llvm/include/llvm/IR/DiagnosticInfo.h

@@ -181,62 +181,6 @@ public:
   }
 };
 
-/// 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 @@ private:
   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 {

llvm/lib/IR/DiagnosticInfo.cpp

@@ -65,9 +65,17 @@ void DiagnosticInfoInlineAsm::print(DiagnosticPrinter &DP) const {
     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 ("
+  DP << getLocationStr() << ": " << getResourceName() << " ("
-     << getResourceLimit() << ") in function '" << getFunction() << '\'';
+     << getResourceSize() << ") exceeds limit (" << getResourceLimit()
+     << ") in function '" << getFunction() << '\'';
 }
 
 void DiagnosticInfoDebugMetadataVersion::print(DiagnosticPrinter &DP) const {

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 @@ define amdgpu_kernel void @use_too_many_sgprs_tahiti() #0 {
   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 @@ define amdgpu_kernel void @use_too_many_sgprs_bonaire() #1 {
   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 @@ define amdgpu_kernel void @use_too_many_sgprs_bonaire_flat_scr() #1 {
   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 @@ define amdgpu_kernel void @use_too_many_sgprs_iceland() #2 {
   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]}" ()

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 @@ entry:
   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:

llvm/test/CodeGen/ARM/warn-stack.ll

@@ -11,7 +11,7 @@ entry:
   ret void
 }
 
-; CHECK: warning: stack frame size ([[STCK:[0-9]+]]) exceeds limit (80) in function 'warn'
+; CHECK: warning: <unknown>:0:0: stack frame size ([[STCK:[0-9]+]]) exceeds limit (80) in function 'warn'
 ; CHECK: {{[0-9]+}}/[[STCK]] ({{.*}}%) spills, {{[0-9]+}}/[[STCK]] ({{.*}}%) variables
 define i32 @warn() nounwind ssp "frame-pointer"="all" "warn-stack-size"="80" {
 entry:

llvm/test/CodeGen/X86/warn-stack.ll

@@ -11,7 +11,7 @@ entry:
   ret void
 }
 
-; CHECK: warning: stack frame size ([[STCK:[0-9]+]]) exceeds limit (80) in function 'warn'
+; CHECK: warning: <unknown>:0:0: stack frame size ([[STCK:[0-9]+]]) exceeds limit (80) in function 'warn'
 ; CHECK: {{[0-9]+}}/[[STCK]] ({{.*}}%) spills, {{[0-9]+}}/[[STCK]] ({{.*}}%) variables
 define void @warn() nounwind ssp "warn-stack-size"="80" {
 entry:
@@ -25,7 +25,7 @@ entry:
 ; combined stack size of the machine stack and unsafe stack will exceed the
 ; warning threshold
 
-; CHECK: warning: stack frame size ([[STCK:[0-9]+]]) exceeds limit (80) in function 'warn_safestack'
+; CHECK: warning: <unknown>:0:0: stack frame size ([[STCK:[0-9]+]]) exceeds limit (80) in function 'warn_safestack'
 ; CHECK: {{[0-9]+}}/[[STCK]] ({{.*}}%) spills, {{[0-9]+}}/[[STCK]] ({{.*}}%) variables, {{[0-9]+}}/[[STCK]] ({{.*}}%) unsafe stack
 define i32 @warn_safestack() nounwind ssp safestack "warn-stack-size"="80" {
 entry: