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; def err_fe_backend_plugin: Error<"%0">, BackendInfo; + +def warn_fe_backend_resource_limit: Warning<"%0 (%1) exceeds limit (%2) in '%3'">, BackendInfo, InGroup; +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; def note_fe_backend_plugin: Note<"%0">, BackendInfo; 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(DI))) + return; + ComputeDiagID(Severity, backend_resource_limit, DiagID); + break; case DK_Linker: ComputeDiagID(Severity, linking_module, DiagID); break; 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: :[[@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()' template __global__ void use_huge_lds(); -// CHECK: error: :[[@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()' template __global__ void use_huge_lds(); 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: :[[@LINE+1]]:0: local memory (480000) exceeds limit (32768) in function 'use_huge_lds' +// CHECK: :[[@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;