arsenm created this revision.
arsenm added reviewers: nickdesaulniers, yaxunl, aaron.ballman, qcolombet,
aeubanks, olista01, dnovillo, echristo, MaskRay.
Herald added subscribers: kosarev, StephenFan, tpr.
Herald added a project: All.
arsenm requested review of this revision.
Herald added a subscriber: wdng.
Print source location info and demangle the name, compared
to the default behavior.
Several observations:
1. Specially handling this seems to give source locations
without enabling debug info, and also gives columns compared
to the backend diagnostic. This seems like a bug?
2. We're duplicating diagnostic effort in DiagnosticInfo
and clang. This feels wrong, but clang can demangle and I guess
have better debug info available? Should clang really have any of this
code? For the purposes of this diagnostic, the important piece
is just reading the source location out of the llvm::Function.
3. lld is not duplicating the same effort as clang with LTO, and just
directly printing the DiagnosticInfo as-is. e.g.
$ clang -fgpu-rdc
lld: error: local memory (480000) exceeds limit (65536) in function
'_Z12use_huge_ldsIiEvv'
lld: error: local memory (960000) exceeds limit (65536) in function
'_Z12use_huge_ldsIdEvv'
$ clang -fno-gpu-rdc
backend-resource-limit-diagnostics.hip:8:17: error: local memory
(480000) exceeds limit (65536) in 'void use_huge_lds<int>()'
__global__ void use_huge_lds() {
^
backend-resource-limit-diagnostics.hip:8:17: error: local memory
(960000) exceeds limit (65536) in 'void use_huge_lds<double>()'
2 errors generated when compiling for gfx90a.
4. Backend errors are not observed with -save-temps and -fno-gpu-rdc or -flto,
and the compile incorrectly succeeds.
5. The backend version prints error: <location info>; clang prints <location
info>: error:
6. -emit-codegen-only is totally broken for AMDGPU. MC
gets a null target streamer. I do not understand why this
is a thing. This just creates a horrible edge case.
Just work around this by emitting actual code instead of blocking
this patch.
https://reviews.llvm.org/D136959
Files:
clang/include/clang/Basic/DiagnosticFrontendKinds.td
clang/lib/CodeGen/CodeGenAction.cpp
clang/test/Misc/backend-resource-limit-diagnostics.cl
clang/test/Misc/backend-resource-limit-diagnostics.hip
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,4 +1,4 @@
-// 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
+// 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
#define __global__ __attribute__((global))
@@ -10,10 +10,10 @@
huge[0] = 2;
}
-// CHECK: error: <stdin>:[[@LINE-5]]:0: local memory (480000) exceeds limit (65536) in function '_Z12use_huge_ldsIiEvv'
+// CHECK: :[[@LINE-5]]:17: error: local memory (480000) exceeds limit (65536) in 'void use_huge_lds<int>()'
template
__global__ void use_huge_lds<int>();
-// CHECK: error: <stdin>:[[@LINE-9]]:0: local memory (960000) exceeds limit (65536) in function '_Z12use_huge_ldsIdEvv'
+// CHECK: :[[@LINE-9]]:17: error: local memory (960000) exceeds limit (65536) in 'void use_huge_lds<double>()'
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,7 +1,7 @@
// REQUIRES: amdgpu-registered-target
-// RUN: not %clang_cc1 -debug-info-kind=standalone -x cl -emit-codegen-only -triple=amdgcn-- < %s 2>&1 | FileCheck %s
+// RUN: not %clang_cc1 -x cl -emit-codegen-only -triple=amdgcn-- < %s 2>&1 | FileCheck %s
-// CHECK: error: <stdin>:[[@LINE+1]]:0: local memory (480000) exceeds limit (32768) in function 'use_huge_lds'
+// CHECK: <stdin>:[[@LINE+1]]:13: error: local memory (480000) exceeds limit (32768) in 'use_huge_lds'
kernel void use_huge_lds() {
volatile local int huge[120000];
huge[0] = 2;
Index: clang/lib/CodeGen/CodeGenAction.cpp
===================================================================
--- clang/lib/CodeGen/CodeGenAction.cpp
+++ clang/lib/CodeGen/CodeGenAction.cpp
@@ -435,6 +435,11 @@
/// \return True if the diagnostic has been successfully reported, false
/// otherwise.
bool StackSizeDiagHandler(const llvm::DiagnosticInfoStackSize &D);
+ /// Specialized handler for ResourceLimit diagnostic.
+ /// \return True if the diagnostic has been successfully reported, false
+ /// otherwise.
+ bool ResourceLimitDiagHandler(const llvm::DiagnosticInfoResourceLimit &D);
+
/// Specialized handler for unsupported backend feature diagnostic.
void UnsupportedDiagHandler(const llvm::DiagnosticInfoUnsupported &D);
/// Specialized handlers for optimization remarks.
@@ -631,6 +636,20 @@
return true;
}
+bool BackendConsumer::ResourceLimitDiagHandler(
+ const llvm::DiagnosticInfoResourceLimit &D) {
+ auto Loc = getFunctionSourceLocation(D.getFunction());
+ if (!Loc)
+ return false;
+ unsigned DiagID = diag::err_fe_backend_resource_limit;
+ ComputeDiagID(D.getSeverity(), backend_resource_limit, DiagID);
+
+ Diags.Report(*Loc, DiagID)
+ << D.getResourceName() << D.getResourceSize() << D.getResourceLimit()
+ << llvm::demangle(D.getFunction().getName().str());
+ return true;
+}
+
const FullSourceLoc BackendConsumer::getBestLocationFromDebugLoc(
const llvm::DiagnosticInfoWithLocationBase &D, bool &BadDebugInfo,
StringRef &Filename, unsigned &Line, unsigned &Column) const {
@@ -874,6 +893,11 @@
return;
ComputeDiagID(Severity, backend_frame_larger_than, DiagID);
break;
+ case llvm::DK_ResourceLimit:
+ if (ResourceLimitDiagHandler(cast<DiagnosticInfoResourceLimit>(DI)))
+ return;
+ ComputeDiagID(Severity, backend_resource_limit, DiagID);
+ break;
case DK_Linker:
ComputeDiagID(Severity, linking_module, DiagID);
break;
Index: clang/include/clang/Basic/DiagnosticFrontendKinds.td
===================================================================
--- clang/include/clang/Basic/DiagnosticFrontendKinds.td
+++ clang/include/clang/Basic/DiagnosticFrontendKinds.td
@@ -35,6 +35,12 @@
def warn_fe_backend_plugin: Warning<"%0">, BackendInfo, InGroup<BackendPlugin>;
def err_fe_backend_plugin: Error<"%0">, BackendInfo;
+
+def warn_fe_backend_resource_limit: Warning<"%0 (%1) exceeds limit (%2) in '%3'">, BackendInfo, InGroup<BackendPlugin>;
+def err_fe_backend_resource_limit: Error<"%0 (%1) exceeds limit (%2) in '%3'">, BackendInfo;
+def note_fe_backend_resource_limit: Note<"%0 (%1) exceeds limit (%2) in '%3'">, BackendInfo;
+
+
def remark_fe_backend_plugin: Remark<"%0">, BackendInfo, InGroup<RemarkBackendPlugin>;
def note_fe_backend_plugin: Note<"%0">, BackendInfo;
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
