clang: Improve errors for DiagnosticInfoResourceLimit

[arsenm]

Summary

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.

[aeubanks]

see https://reviews.llvm.org/D110665 for how we get line/col from a function. basically, we have a map from function name hashes to source location available in clang

if we don't have clang source information for functions available, there's not much more we can do other than rely on debug info, which is what the default LLVM handlers use when possible. that's why lld doesn't have its own handler

Backend errors are not observed with -save-temps and -fno-gpu-rdc or -flto, and the compile incorrectly succeeds.

not sure why this would be happening

clang/include/clang/Basic/DiagnosticFrontendKinds.td

42

ultra nit: extra blank line

[arsenm]

0ebd4638af1f71788ca55f521ed8e1ed8cab518d

clang/include/clang/Basic/DiagnosticFrontendKinds.td

def warn_fe_frame_larger_than : Warning<"stack frame size (%0) exceeds limit (%1) in '%2'">,
     BackendInfo, InGroup<BackendFrameLargerThan>;
 def warn_fe_backend_frame_larger_than: Warning<"%0">,
     BackendInfo, InGroup<BackendFrameLargerThan>;
 def err_fe_backend_frame_larger_than: Error<"%0">, BackendInfo;
 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;
 
 def warn_fe_override_module : Warning<
     "overriding the module target triple with %0">,
     InGroup<DiagGroup<"override-module">>;
 def warn_fe_backend_unsupported_fp_rounding : Warning<
     "overriding currently unsupported rounding mode on this target">,

clang/lib/CodeGen/CodeGenAction.cpp

public:
     /// otherwise.
     bool InlineAsmDiagHandler(const llvm::DiagnosticInfoInlineAsm &D);
     /// Specialized handler for diagnostics reported using SMDiagnostic.
     void SrcMgrDiagHandler(const llvm::DiagnosticInfoSrcMgr &D);
     /// Specialized handler for StackSize diagnostic.
     /// \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.
     /// Note that these handlers only accept remarks and they always handle
     /// them.
     void EmitOptimizationMessage(const llvm::DiagnosticInfoOptimizationBase &D,
                                  unsigned DiagID);
     void
BackendConsumer::StackSizeDiagHandler(const llvm::DiagnosticInfoStackSize &D) {
   // FIXME: Shouldn't need to truncate to uint32_t
   Diags.Report(*Loc, diag::warn_fe_frame_larger_than)
       << static_cast<uint32_t>(D.getStackSize())
       << static_cast<uint32_t>(D.getStackLimit())
       << llvm::demangle(D.getFunction().getName().str());
   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 {
   SourceManager &SourceMgr = Context->getSourceManager();
   FileManager &FileMgr = SourceMgr.getFileManager();
   SourceLocation DILoc;
 
   if (D.isLocationAvailable()) {
void BackendConsumer::DiagnosticHandlerImpl(const DiagnosticInfo &DI) {
   case llvm::DK_SrcMgr:
     SrcMgrDiagHandler(cast<DiagnosticInfoSrcMgr>(DI));
     return;
   case llvm::DK_StackSize:
     if (StackSizeDiagHandler(cast<DiagnosticInfoStackSize>(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;
   case llvm::DK_OptimizationRemark:
     // Optimization remarks are always handled completely by this
     // handler. There is no generic way of emitting them.
     OptimizationRemarkHandler(cast<OptimizationRemark>(DI));
     return;

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

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

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