[clang] c62745e - DiagnosticInfo: Report function location for resource limits
Matt Arsenault via cfe-commits
cfe-commits at lists.llvm.org
Fri Oct 28 21:43:08 PDT 2022
Author: Matt Arsenault
Date: 2022-10-28T21:42:57-07:00
New Revision: c62745e1678e26185dc93ea30f858c91c03323f9
URL: https://github.com/llvm/llvm-project/commit/c62745e1678e26185dc93ea30f858c91c03323f9
DIFF: https://github.com/llvm/llvm-project/commit/c62745e1678e26185dc93ea30f858c91c03323f9.diff
LOG: 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.
Added:
Modified:
clang/test/Misc/backend-resource-limit-diagnostics.cl
clang/test/Misc/backend-resource-limit-diagnostics.hip
clang/test/Misc/backend-stack-frame-diagnostics-fallback.cpp
llvm/include/llvm/IR/DiagnosticInfo.h
llvm/lib/IR/DiagnosticInfo.cpp
llvm/test/CodeGen/AMDGPU/exceed-max-sgprs.ll
llvm/test/CodeGen/AMDGPU/stack-size-overflow.ll
llvm/test/CodeGen/ARM/warn-stack.ll
llvm/test/CodeGen/X86/warn-stack.ll
Removed:
################################################################################
diff --git a/clang/test/Misc/backend-resource-limit-diagnostics.cl b/clang/test/Misc/backend-resource-limit-diagnostics.cl
index fd37e0bd0ef86..347bafc55acb3 100644
--- a/clang/test/Misc/backend-resource-limit-diagnostics.cl
+++ b/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;
}
diff --git a/clang/test/Misc/backend-resource-limit-diagnostics.hip b/clang/test/Misc/backend-resource-limit-diagnostics.hip
index b9a61e5aadad6..3cdd3ac86ce71 100644
--- a/clang/test/Misc/backend-resource-limit-diagnostics.hip
+++ b/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>();
diff --git a/clang/test/Misc/backend-stack-frame-diagnostics-fallback.cpp b/clang/test/Misc/backend-stack-frame-diagnostics-fallback.cpp
index 0eca44f49f27f..17456a90da275 100644
--- a/clang/test/Misc/backend-stack-frame-diagnostics-fallback.cpp
+++ b/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.
}
diff --git a/llvm/include/llvm/IR/DiagnosticInfo.h b/llvm/include/llvm/IR/DiagnosticInfo.h
index da37801b6d192..0a1c4abdd56af 100644
--- a/llvm/include/llvm/IR/DiagnosticInfo.h
+++ b/llvm/include/llvm/IR/DiagnosticInfo.h
@@ -181,62 +181,6 @@ class DiagnosticInfoInlineAsm : public DiagnosticInfo {
}
};
-/// 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 @@ class DiagnosticInfoWithLocationBase : public DiagnosticInfo {
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 {
diff --git a/llvm/lib/IR/DiagnosticInfo.cpp b/llvm/lib/IR/DiagnosticInfo.cpp
index 50fe6829ad86f..aed38b6927bde 100644
--- a/llvm/lib/IR/DiagnosticInfo.cpp
+++ b/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 ("
- << getResourceLimit() << ") in function '" << getFunction() << '\'';
+ DP << getLocationStr() << ": " << getResourceName() << " ("
+ << getResourceSize() << ") exceeds limit (" << getResourceLimit()
+ << ") in function '" << getFunction() << '\'';
}
void DiagnosticInfoDebugMetadataVersion::print(DiagnosticPrinter &DP) const {
diff --git a/llvm/test/CodeGen/AMDGPU/exceed-max-sgprs.ll b/llvm/test/CodeGen/AMDGPU/exceed-max-sgprs.ll
index ad2c6debda7e9..aea3c11a00f37 100644
--- a/llvm/test/CodeGen/AMDGPU/exceed-max-sgprs.ll
+++ b/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]}" ()
diff --git a/llvm/test/CodeGen/AMDGPU/stack-size-overflow.ll b/llvm/test/CodeGen/AMDGPU/stack-size-overflow.ll
index 976c8809aff23..9060fcd1ecc02 100644
--- a/llvm/test/CodeGen/AMDGPU/stack-size-overflow.ll
+++ b/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:
diff --git a/llvm/test/CodeGen/ARM/warn-stack.ll b/llvm/test/CodeGen/ARM/warn-stack.ll
index 4f23ec8db5cba..511597c37286c 100644
--- a/llvm/test/CodeGen/ARM/warn-stack.ll
+++ b/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:
diff --git a/llvm/test/CodeGen/X86/warn-stack.ll b/llvm/test/CodeGen/X86/warn-stack.ll
index 29c5216cf39ff..6bed7a0ad5477 100644
--- a/llvm/test/CodeGen/X86/warn-stack.ll
+++ b/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:
More information about the cfe-commits
mailing list