mirrors/llvm-project
0
0
Fork 0
mirror of https://github.com/llvm/llvm-project.git synced 2025-04-26 16:56:06 +00:00
Code Issues Projects Releases Wiki Activity
llvm-project/clang/test/Misc/backend-resource-limit-diagnostics.cl

Ignoring revisions in .git-blame-ignore-revs. Click here to bypass and see the normal blame view.

9 lines
317 B
Common Lisp
Raw Normal View History

Update for DiagnosticInfoStackSize changes llvm-svn: 273178
2016-06-20 18:13:09 +00:00
// REQUIRES: amdgpu-registered-target
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.
2022-10-27 16:28:46 -07:00
// RUN: not %clang_cc1 -x cl -emit-codegen-only -triple=amdgcn-- < %s 2>&1 | FileCheck %s
Update for DiagnosticInfoStackSize changes llvm-svn: 273178
2016-06-20 18:13:09 +00:00
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.
2022-10-27 16:28:46 -07:00
// CHECK: <stdin>:[[@LINE+1]]:13: error: local memory (480000) exceeds limit (32768) in 'use_huge_lds'
DiagnosticInfo: Report function location for resource limits We have some odd redundancy where clang specially handles the stack size case. If clang prints it, the source location is first followed by "warning". The backend diagnostic, as printed by other tools puts "warning" first.
2022-10-27 16:29:26 -07:00
kernel void use_huge_lds() {
Update for DiagnosticInfoStackSize changes llvm-svn: 273178
2016-06-20 18:13:09 +00:00
volatile local int huge[120000];
huge[0] = 2;
}
Reference in New Issue Copy Permalink