Index: clang/test/Misc/backend-resource-limit-diagnostics.hip =================================================================== --- /dev/null +++ clang/test/Misc/backend-resource-limit-diagnostics.hip @@ -0,0 +1,19 @@ +// 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 +// FIXME: Use -emit-codegen-only + +#define __global__ __attribute__((global)) +#define __shared__ __attribute__((shared)) + +template +__global__ void use_huge_lds() { + volatile __shared__ T huge[120000]; + huge[0] = 2; +} + +// CHECK: error: :[[@LINE-5]]:0: local memory (480000) exceeds limit (65536) in function '_Z12use_huge_ldsIiEvv' + template + __global__ void use_huge_lds(); + +// CHECK: error: :[[@LINE-9]]:0: local memory (960000) exceeds limit (65536) in function '_Z12use_huge_ldsIdEvv' + 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,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: :[[@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; } Index: llvm/include/llvm/IR/DiagnosticInfo.h =================================================================== --- llvm/include/llvm/IR/DiagnosticInfo.h +++ llvm/include/llvm/IR/DiagnosticInfo.h @@ -181,62 +181,6 @@ } }; -/// 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 @@ 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 { Index: llvm/lib/IR/DiagnosticInfo.cpp =================================================================== --- llvm/lib/IR/DiagnosticInfo.cpp +++ llvm/lib/IR/DiagnosticInfo.cpp @@ -65,9 +65,17 @@ 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 { Index: llvm/test/CodeGen/AMDGPU/exceed-max-sgprs.ll =================================================================== --- llvm/test/CodeGen/AMDGPU/exceed-max-sgprs.ll +++ 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: :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 @@ ret void } -; ERROR: error: scalar registers (106) exceeds limit (104) in function 'use_too_many_sgprs_bonaire' +; ERROR: error: :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 @@ ret void } -; ERROR: error: scalar registers (108) exceeds limit (104) in function 'use_too_many_sgprs_bonaire_flat_scr' +; ERROR: error: :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 @@ ret void } -; ERROR: error: scalar registers (98) exceeds limit (96) in function 'use_too_many_sgprs_iceland' +; ERROR: error: :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 @@ ret void } -; ERROR: error: addressable scalar registers (103) exceeds limit (102) in function 'use_too_many_sgprs_fiji' +; ERROR: error: :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]}" () Index: llvm/test/CodeGen/AMDGPU/stack-size-overflow.ll =================================================================== --- llvm/test/CodeGen/AMDGPU/stack-size-overflow.ll +++ 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: :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 @@ ret void } -; ERROR: error: stack frame size (262117) exceeds limit (262112) in function 'stack_size_limit_wave32' +; ERROR: error: :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: Index: llvm/test/CodeGen/ARM/warn-stack.ll =================================================================== --- llvm/test/CodeGen/ARM/warn-stack.ll +++ llvm/test/CodeGen/ARM/warn-stack.ll @@ -12,7 +12,7 @@ ret void } -; CHECK: warning: stack frame size (92) exceeds limit (80) in function 'warn' +; CHECK: warning: :0:0: stack frame size (92) exceeds limit (80) in function 'warn' define void @warn() nounwind ssp "frame-pointer"="all" "warn-stack-size"="80" { entry: %buffer = alloca [80 x i8], align 1 Index: llvm/test/CodeGen/X86/warn-stack.ll =================================================================== --- llvm/test/CodeGen/X86/warn-stack.ll +++ llvm/test/CodeGen/X86/warn-stack.ll @@ -11,7 +11,7 @@ ret void } -; CHECK: warning: stack frame size (88) exceeds limit (80) in function 'warn' +; CHECK: warning: :0:0: stack frame size (88) exceeds limit (80) in function 'warn' define void @warn() nounwind ssp "warn-stack-size"="80" { entry: %buffer = alloca [80 x i8], align 1 @@ -24,7 +24,7 @@ ; combined stack size of the machine stack and unsafe stack will exceed the ; warning threshold -; CHECK: warning: stack frame size (120) exceeds limit (80) in function 'warn_safestack' +; CHECK: warning: :0:0: stack frame size (120) exceeds limit (80) in function 'warn_safestack' define void @warn_safestack() nounwind ssp safestack "warn-stack-size"="80" { entry: %buffer = alloca [80 x i8], align 1