Author: Matt Arsenault Date: 2022-10-28T21:42:57-07:00 New Revision: 0ebd4638af1f71788ca55f521ed8e1ed8cab518d
URL: https://github.com/llvm/llvm-project/commit/0ebd4638af1f71788ca55f521ed8e1ed8cab518d DIFF: https://github.com/llvm/llvm-project/commit/0ebd4638af1f71788ca55f521ed8e1ed8cab518d.diff LOG: clang: Improve errors for DiagnosticInfoResourceLimit 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. 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. Added: Modified: 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 Removed: ################################################################################ diff --git a/clang/include/clang/Basic/DiagnosticFrontendKinds.td b/clang/include/clang/Basic/DiagnosticFrontendKinds.td index 07ecbe332f8b3..1bd31ca2fc446 100644 --- a/clang/include/clang/Basic/DiagnosticFrontendKinds.td +++ b/clang/include/clang/Basic/DiagnosticFrontendKinds.td @@ -35,6 +35,11 @@ def note_fe_backend_frame_larger_than: Note<"%0">, BackendInfo; 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; diff --git a/clang/lib/CodeGen/CodeGenAction.cpp b/clang/lib/CodeGen/CodeGenAction.cpp index 12c6b3f49c432..52d0417a4fa6b 100644 --- a/clang/lib/CodeGen/CodeGenAction.cpp +++ b/clang/lib/CodeGen/CodeGenAction.cpp @@ -435,6 +435,11 @@ namespace clang { /// \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 @@ BackendConsumer::StackSizeDiagHandler(const llvm::DiagnosticInfoStackSize &D) { 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 @@ void BackendConsumer::DiagnosticHandlerImpl(const DiagnosticInfo &DI) { 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; diff --git a/clang/test/Misc/backend-resource-limit-diagnostics.cl b/clang/test/Misc/backend-resource-limit-diagnostics.cl index 347bafc55acb3..993d0ebaa2152 100644 --- a/clang/test/Misc/backend-resource-limit-diagnostics.cl +++ b/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; diff --git a/clang/test/Misc/backend-resource-limit-diagnostics.hip b/clang/test/Misc/backend-resource-limit-diagnostics.hip index 3cdd3ac86ce71..4479f067be3c3 100644 --- a/clang/test/Misc/backend-resource-limit-diagnostics.hip +++ b/clang/test/Misc/backend-resource-limit-diagnostics.hip @@ -10,10 +10,10 @@ __global__ void use_huge_lds() { huge[0] = 2; } -// CHECK: error: <stdin>:[[#@LINE-5]]:0: local memory (480000) exceeds limit (65536) in function '_Z12use_huge_ldsIiEvv' +// CHECK: <stdin>:[[@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: <stdin>:[[@LINE-9]]:17: error: local memory (960000) exceeds limit (65536) in 'void use_huge_lds<double>()' template __global__ void use_huge_lds<double>(); _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits