DiagnosticInfo: Report function location for resource limits

[arsenm]

Summary

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.

[arsenm]

Update clang test

[paulkirth]

This is basically LGTM from my perspective. I'd suggest getting some feedback from someone who works more on this part of Clang before landing, but I think it fixes the inconsistency nicely.

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

2

Does this need // REQUIRES: amdgpu-registered-target, like the test in backend-resourse-limit-diagnostics.cl?

3

Is this planned for this patch? or for future refactoring in multiple tests?

[arsenm]

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

3

I'm not sure. I currently cannot comprehend why this flag exists. It's only used in a handful of tests, and its effect seems to be to just create an edge case to crash codegen when there's no target streamer

[arsenm]

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

3

A little more digging suggests this is a feature that is only properly implemented for 5/24 targets. Still don't really see the point

[aeubanks]

the reason clang tries to handle diagnostics is to be able to actually print the source code as context which is nice. but of course we can't do that without clang so there's a default handler

If clang prints it, the source location is first followed by "warning". The backend diagnostic, as printed by other tools puts "warning" first.

It looks like the backend diagnostic doesn't even print the location?

llvm/include/llvm/IR/DiagnosticInfo.h

356

any reason for moving this code? it makes it hard to see the diff

llvm/lib/IR/DiagnosticInfo.cpp

76–78

we should guard this with isLocationAvailable(), the <unknown>:0:0 isn't very useful

[paulkirth]

I think you may also need to rebase, both the X86/warn-stack.ll and ARM/warn-stack.ll were recently updated to reflect changes in the backend diagnostics.

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

3

I think it's ideal for testing diagnostics in the backend, where you need codegen to occur but the actual output is unnecessary. IIUC this flag allows you to avoid using a non-portable way to throw away the output instead of redirecting to /dev/null.

[arsenm]

llvm/include/llvm/IR/DiagnosticInfo.h

356

Moving after DiagnosticInfoWithLocationBase since it's now the base class

llvm/lib/IR/DiagnosticInfo.cpp

76–78

Everywhere else here prints the unknown. I'd rather keep it consistent and remove them all together later

[arsenm]

In D136954#3892864, @aeubanks wrote:

If clang prints it, the source location is first followed by "warning". The backend diagnostic, as printed by other tools puts "warning" first.

It looks like the backend diagnostic doesn't even print the location?

That's the point of this patch, to fix it for this case. There are still other remarks which don't

[MaskRay]

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

3

FIXME: looks like an unused check prefix. Drop :

14

[[@LINE-5]] uses deprecated syntax. Use [[#@LINE-5]]

[MaskRay]

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

2

-emit-codegen-only is preferred to -S -o /dev/null

[MaskRay]

LGTM.

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

2

Ignore me. I assume that the object MC backend isn't available. -S -o /dev/null looks fine.

10

2-space indentation is more common

15

Unintended indentation

[arsenm]

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

2

Fixed this in D136983, The null streamer is broken in 19/24 in tree targets

[arsenm]

c62745e1678e26185dc93ea30f858c91c03323f9

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

+// 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>();

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

 // 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;
 }

llvm/include/llvm/IR/DiagnosticInfo.h

public:
   /// \see DiagnosticInfo::print.
   void print(DiagnosticPrinter &DP) const override;
 
   static bool classof(const DiagnosticInfo *DI) {
     return DI->getKind() == DK_InlineAsm;
   }
 };
 
-/// 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 {
 private:
   /// The module that is concerned by this debug metadata version diagnostic.
   const Module &M;
   /// The actual metadata version.
   unsigned MetadataVersion;

 private:
   /// Function where this diagnostic is triggered.
   const Function &Fn;
 
   /// Debug location where this diagnostic is triggered.
   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 {
 public:
   /// Used to set IsVerbose via the stream interface.
   struct setIsVerbose {};
 
   /// When an instance of this is inserted into the stream, the arguments

llvm/lib/IR/DiagnosticInfo.cpp

 }
 
 void DiagnosticInfoInlineAsm::print(DiagnosticPrinter &DP) const {
   DP << getMsgStr();
   if (getLocCookie())
     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 {
   DP << "ignoring debug info with an invalid version (" << getMetadataVersion()
      << ") in " << getModule();
 }
 
 void DiagnosticInfoIgnoringInvalidDebugMetadata::print(

llvm/test/CodeGen/AMDGPU/exceed-max-sgprs.ll

 ; 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]}" ()
   call void asm sideeffect "", "~{s[16:23]}" ()
   call void asm sideeffect "", "~{s[24:31]}" ()
   call void asm sideeffect "", "~{s[32:39]}" ()
   call void asm sideeffect "", "~{s[40:47]}" ()
   call void asm sideeffect "", "~{s[48:55]}" ()
   call void asm sideeffect "", "~{s[56:63]}" ()
   call void asm sideeffect "", "~{s[64:71]}" ()
   call void asm sideeffect "", "~{s[72:79]}" ()
   call void asm sideeffect "", "~{s[80:87]}" ()
   call void asm sideeffect "", "~{s[88:95]}" ()
   call void asm sideeffect "", "~{s[96:103]}" ()
   call void asm sideeffect "", "~{vcc}" ()
   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]}" ()
   call void asm sideeffect "", "~{s[16:23]}" ()
   call void asm sideeffect "", "~{s[24:31]}" ()
   call void asm sideeffect "", "~{s[32:39]}" ()
   call void asm sideeffect "", "~{s[40:47]}" ()
   call void asm sideeffect "", "~{s[48:55]}" ()
   call void asm sideeffect "", "~{s[56:63]}" ()
   call void asm sideeffect "", "~{s[64:71]}" ()
   call void asm sideeffect "", "~{s[72:79]}" ()
   call void asm sideeffect "", "~{s[80:87]}" ()
   call void asm sideeffect "", "~{s[88:95]}" ()
   call void asm sideeffect "", "~{s[96:103]}" ()
   call void asm sideeffect "", "~{vcc}" ()
   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]}" ()
   call void asm sideeffect "", "~{s[16:23]}" ()
   call void asm sideeffect "", "~{s[24:31]}" ()
   call void asm sideeffect "", "~{s[32:39]}" ()
   call void asm sideeffect "", "~{s[40:47]}" ()
   call void asm sideeffect "", "~{s[48:55]}" ()
   call void asm sideeffect "", "~{s[56:63]}" ()
   call void asm sideeffect "", "~{s[64:71]}" ()
   call void asm sideeffect "", "~{s[72:79]}" ()
   call void asm sideeffect "", "~{s[80:87]}" ()
   call void asm sideeffect "", "~{s[88:95]}" ()
   call void asm sideeffect "", "~{s[96:103]}" ()
   call void asm sideeffect "", "~{vcc}" ()
   call void asm sideeffect "", "~{flat_scratch}" ()
   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]}" ()
   call void asm sideeffect "", "~{s[8:15]}" ()
   call void asm sideeffect "", "~{s[16:23]}" ()
   call void asm sideeffect "", "~{s[24:31]}" ()
   call void asm sideeffect "", "~{s[32:39]}" ()
   call void asm sideeffect "", "~{s[40:47]}" ()
   call void asm sideeffect "", "~{s[48:55]}" ()
   call void asm sideeffect "", "~{s[56:63]}" ()
   call void asm sideeffect "", "~{s[64:71]}" ()
   call void asm sideeffect "", "~{s[72:79]}" ()
   call void asm sideeffect "", "~{s[80:87]}" ()
   call void asm sideeffect "", "~{s[88:95]}" ()
   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]}" ()
   call void asm sideeffect "", "~{s[16:23]}" ()
   call void asm sideeffect "", "~{s[24:31]}" ()
   call void asm sideeffect "", "~{s[32:39]}" ()
   call void asm sideeffect "", "~{s[40:47]}" ()
   call void asm sideeffect "", "~{s[48:55]}" ()

llvm/test/CodeGen/AMDGPU/stack-size-overflow.ll

 ; RUN: not llc -march=amdgcn < %s 2>&1 | FileCheck -check-prefix=ERROR %s
 ; RUN: not llc -march=amdgcn < %s | FileCheck -check-prefix=GCN %s
 
 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:
   %alloca = alloca [131057 x i8], align 1, addrspace(5)
   %alloca.bc = bitcast [131057 x i8] addrspace(5)* %alloca to i8 addrspace(5)*
   call void @llvm.memset.p5i8.i32(i8 addrspace(5)* %alloca.bc, i8 9, i32 131057, i32 1, i1 true)
   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:
   %alloca = alloca [262113 x i8], align 1, addrspace(5)
   %alloca.bc = bitcast [262113 x i8] addrspace(5)* %alloca to i8 addrspace(5)*
   call void @llvm.memset.p5i8.i32(i8 addrspace(5)* %alloca.bc, i8 9, i32 262113, i32 1, i1 true)
   ret void
 }

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

 ; RUN: llc -mtriple thumbv7-apple-ios3.0.0 < %s 2>&1 >/dev/null | FileCheck %s
 ; Check the internal option that warns when the stack frame size exceeds the
 ; given amount.
 ; <rdar://13987214>
 
 ; CHECK-NOT: nowarn
 define void @nowarn() nounwind ssp "frame-pointer"="all" "warn-stack-size"="80" {
 entry:
   %buffer = alloca [12 x i8], align 1
   %arraydecay = getelementptr inbounds [12 x i8], [12 x i8]* %buffer, i64 0, i64 0
   call void @doit(i8* %arraydecay) nounwind
   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
   %arraydecay = getelementptr inbounds [80 x i8], [80 x i8]* %buffer, i64 0, i64 0
   call void @doit(i8* %arraydecay) nounwind
   ret void
 }
 
 declare void @doit(i8*)

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

 ; RUN: llc -mtriple x86_64-apple-macosx10.8.0 < %s 2>&1 >/dev/null | FileCheck %s
 ; Check the internal option that warns when the stack frame size exceeds the
 ; given amount.
 ; <rdar://13987214>
 
 ; CHECK-NOT: nowarn
 define void @nowarn() nounwind ssp "warn-stack-size"="80" {
 entry:
   %buffer = alloca [12 x i8], align 1
   call void @doit(ptr %buffer) nounwind
   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
   call void @doit(ptr %buffer) nounwind
   ret void
 }
 
 ; Ensure that warn-stack-size also considers the size of the unsafe stack.
 ; With safestack enabled the machine stack size is well below 80, but the
 ; 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
   call void @doit(ptr %buffer) nounwind
   ret void
 }
 declare void @doit(ptr)