| Issue |
131517
|
| Summary |
Clang crashes for large structs in CUDA
|
| Labels |
|
| Assignees |
|
| Reporter |
AdUhTkJm
|
Consider this `big.cu`:
```cpp
struct image {
int x[1920 * 1080];
};
__device__ image f() {
return {};
}
```
With `clang++ big.cu`, it crashes with the following stack trace:
```
clang-20: /home/aduhtkjm/llvm/llvm-project/llvm/include/llvm/CodeGen/SelectionDAGNodes.h:1166: llvm::SDNode::SDNode(unsigned int, unsigned int, DebugLoc, SDVTList): Assertion `NumValues == VTs.NumVTs && "NumValues wasn't wide enough for its operands!"' failed.
PLEASE submit a bug report to https://github.com/llvm/llvm-project/issues/ and include the crash backtrace, preprocessed source, and associated run script.
Stack dump:
0. Program arguments: /home/aduhtkjm/llvm/llvm-project/build/bin/clang-20 -cc1 -triple nvptx64-nvidia-cuda -aux-triple x86_64-unknown-linux-gnu -S -dumpdir a- -disable-free -clear-ast-before-backend -main-file-name big.cu -mrelocation-model static -mframe-pointer=all -fno-rounding-math -no-integrated-as -aux-target-cpu x86-64 -fcuda-is-device -mllvm -enable-memcpyopt-without-libcalls -fno-threadsafe-statics -fcuda-allow-variadic-functions -mlink-builtin-bitcode /usr/local/cuda/nvvm/libdevice/libdevice.10.bc -target-sdk-version=12.8 -target-cpu sm_52 -target-feature +ptx87 -debugger-tuning=gdb -fno-dwarf-directory-asm -fdebug-compilation-dir=/home/aduhtkjm/llvm/tamper -resource-dir /home/aduhtkjm/llvm/llvm-project/build/lib/clang/20 -internal-isystem /home/aduhtkjm/llvm/llvm-project/build/lib/clang/20/include/cuda_wrappers -include __clang_cuda_runtime_wrapper.h -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/13/../../../../include/c++/13 -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/13/../../../../include/x86_64-linux-gnu/c++/13 -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/13/../../../../include/c++/13/backward -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/13/../../../../include/c++/13 -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/13/../../../../include/x86_64-linux-gnu/c++/13 -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/13/../../../../include/c++/13/backward -internal-isystem /home/aduhtkjm/llvm/llvm-project/build/lib/clang/20/include -internal-isystem /usr/local/include -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/13/../../../../x86_64-linux-gnu/include -internal-externc-isystem /usr/include/x86_64-linux-gnu -internal-externc-isystem /include -internal-externc-isystem /usr/include -internal-isystem /usr/local/cuda/include -internal-isystem /home/aduhtkjm/llvm/llvm-project/build/lib/clang/20/include -internal-isystem /usr/local/include -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/13/../../../../x86_64-linux-gnu/include -internal-externc-isystem /usr/include/x86_64-linux-gnu -internal-externc-isystem /include -internal-externc-isystem /usr/include -fdeprecated-macro -fno-autolink -ferror-limit 19 --offload-new-driver -fgnuc-version=4.2.1 -fskip-odr-check-in-gmf -fcxx-exceptions -fexceptions -fcolor-diagnostics -cuid=51d3a9fb8176529a -D__GCC_HAVE_DWARF2_CFI_ASM=1 -o /tmp/big-sm_52-6f00ea.s -x cuda big.cu
1. <eof> parser at end of file
2. Code generation
3. Running pass 'Function Pass Manager' on module 'big.cu'.
4. Running pass 'NVPTX DAG->DAG Pattern Instruction Selection' on function '@_Z1f3big'
#0 0x0000561ce4ab2988 llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) (/home/aduhtkjm/llvm/llvm-project/build/bin/clang-20+0x4fb8988)
#1 0x0000561ce4ab043e llvm::sys::RunSignalHandlers() (/home/aduhtkjm/llvm/llvm-project/build/bin/clang-20+0x4fb643e)
#2 0x0000561ce4ab3018 SignalHandler(int) Signals.cpp:0:0
#3 0x00007f3f59671520 (/lib/x86_64-linux-gnu/libc.so.6+0x42520)
#4 0x00007f3f596c59fc __pthread_kill_implementation ./nptl/pthread_kill.c:44:76
#5 0x00007f3f596c59fc __pthread_kill_internal ./nptl/pthread_kill.c:78:10
#6 0x00007f3f596c59fc pthread_kill ./nptl/pthread_kill.c:89:10
#7 0x00007f3f59671476 gsignal ./signal/../sysdeps/posix/raise.c:27:6
#8 0x00007f3f596577f3 abort ./stdlib/abort.c:81:7
#9 0x00007f3f5965771b _nl_load_domain ./intl/loadmsgcat.c:1177:9
#10 0x00007f3f59668e96 (/lib/x86_64-linux-gnu/libc.so.6+0x39e96)
#11 0x0000561ce3a11722 (/home/aduhtkjm/llvm/llvm-project/build/bin/clang-20+0x3f17722)
#12 0x0000561ce6580767 llvm::SDNode* llvm::SelectionDAG::newSDNode<llvm::SDNode, unsigned int&, unsigned int, llvm::DebugLoc const&, llvm::SDVTList&>(unsigned int&, unsigned int&&, llvm::DebugLoc const&, llvm::SDVTList&) SelectionDAG.cpp:0:0
#13 0x0000561ce65a6cdb llvm::SelectionDAG::getNode(unsigned int, llvm::SDLoc const&, llvm::SDVTList, llvm::ArrayRef<llvm::SDValue>, llvm::SDNodeFlags) (/home/aduhtkjm/llvm/llvm-project/build/bin/clang-20+0x6aaccdb)
#14 0x0000561ce650ee35 llvm::SelectionDAGBuilder::visitLoad(llvm::LoadInst const&) (/home/aduhtkjm/llvm/llvm-project/build/bin/clang-20+0x6a14e35)
#15 0x0000561ce65065f4 llvm::SelectionDAGBuilder::visit(llvm::Instruction const&) (/home/aduhtkjm/llvm/llvm-project/build/bin/clang-20+0x6a0c5f4)
#16 0x0000561ce65c9215 llvm::SelectionDAGISel::SelectBasicBlock(llvm::ilist_iterator_w_bits<llvm::ilist_detail::node_options<llvm::Instruction, true, false, void, true, llvm::BasicBlock>, false, true>, llvm::ilist_iterator_w_bits<llvm::ilist_detail::node_options<llvm::Instruction, true, false, void, true, llvm::BasicBlock>, false, true>, bool&) (/home/aduhtkjm/llvm/llvm-project/build/bin/clang-20+0x6acf215)
#17 0x0000561ce65c7ecd llvm::SelectionDAGISel::SelectAllBasicBlocks(llvm::Function const&) (/home/aduhtkjm/llvm/llvm-project/build/bin/clang-20+0x6acdecd)
#18 0x0000561ce65c5890 llvm::SelectionDAGISel::runOnMachineFunction(llvm::MachineFunction&) (/home/aduhtkjm/llvm/llvm-project/build/bin/clang-20+0x6acb890)
#19 0x0000561ce3b58cc5 llvm::NVPTXDAGToDAGISel::runOnMachineFunction(llvm::MachineFunction&) NVPTXISelDAGToDAG.cpp:0:0
#20 0x0000561ce65c30f0 llvm::SelectionDAGISelLegacy::runOnMachineFunction(llvm::MachineFunction&) (/home/aduhtkjm/llvm/llvm-project/build/bin/clang-20+0x6ac90f0)
#21 0x0000561ce3f80457 llvm::MachineFunctionPass::runOnFunction(llvm::Function&) (/home/aduhtkjm/llvm/llvm-project/build/bin/clang-20+0x4486457)
#22 0x0000561ce4565c95 llvm::FPPassManager::runOnFunction(llvm::Function&) (/home/aduhtkjm/llvm/llvm-project/build/bin/clang-20+0x4a6bc95)
#23 0x0000561ce456e0f2 llvm::FPPassManager::runOnModule(llvm::Module&) (/home/aduhtkjm/llvm/llvm-project/build/bin/clang-20+0x4a740f2)
#24 0x0000561ce4566706 llvm::legacy::PassManagerImpl::run(llvm::Module&) (/home/aduhtkjm/llvm/llvm-project/build/bin/clang-20+0x4a6c706)
#25 0x0000561ce52eb2e3 clang::emitBackendOutput(clang::CompilerInstance&, clang::CodeGenOptions&, llvm::StringRef, llvm::Module*, clang::BackendAction, llvm::IntrusiveRefCntPtr<llvm::vfs::FileSystem>, std::unique_ptr<llvm::raw_pwrite_stream, std::default_delete<llvm::raw_pwrite_stream>>, clang::BackendConsumer*) (/home/aduhtkjm/llvm/llvm-project/build/bin/clang-20+0x57f12e3)
#26 0x0000561ce5314410 clang::BackendConsumer::HandleTranslationUnit(clang::ASTContext&) (/home/aduhtkjm/llvm/llvm-project/build/bin/clang-20+0x581a410)
#27 0x0000561ce833ed99 clang::ParseAST(clang::Sema&, bool, bool) (/home/aduhtkjm/llvm/llvm-project/build/bin/clang-20+0x8844d99)
#28 0x0000561ce579b5c6 clang::FrontendAction::Execute() (/home/aduhtkjm/llvm/llvm-project/build/bin/clang-20+0x5ca15c6)
#29 0x0000561ce570788d clang::CompilerInstance::ExecuteAction(clang::FrontendAction&) (/home/aduhtkjm/llvm/llvm-project/build/bin/clang-20+0x5c0d88d)
#30 0x0000561ce587c5da clang::ExecuteCompilerInvocation(clang::CompilerInstance*) (/home/aduhtkjm/llvm/llvm-project/build/bin/clang-20+0x5d825da)
#31 0x0000561ce3718625 cc1_main(llvm::ArrayRef<char const*>, char const*, void*) (/home/aduhtkjm/llvm/llvm-project/build/bin/clang-20+0x3c1e625)
#32 0x0000561ce371472f ExecuteCC1Tool(llvm::SmallVectorImpl<char const*>&, llvm::ToolContext const&) driver.cpp:0:0
#33 0x0000561ce37137ff clang_main(int, char**, llvm::ToolContext const&) (/home/aduhtkjm/llvm/llvm-project/build/bin/clang-20+0x3c197ff)
#34 0x0000561ce3723977 main (/home/aduhtkjm/llvm/llvm-project/build/bin/clang-20+0x3c29977)
#35 0x00007f3f59658d90 __libc_start_call_main ./csu/../sysdeps/nptl/libc_start_call_main.h:58:16
#36 0x00007f3f59658e40 call_init ./csu/../csu/libc-start.c:128:20
#37 0x00007f3f59658e40 __libc_start_main ./csu/../csu/libc-start.c:379:5
#38 0x0000561ce37122a5 _start (/home/aduhtkjm/llvm/llvm-project/build/bin/clang-20+0x3c182a5)
clang++: error: unable to execute command: Aborted (core dumped)
clang++: error: clang frontend command failed due to signal (use -v to see invocation)
clang version 20.0.0git (https://github.com/llvm/clangir.git 40b5b6288618ceab8cdd69402407b1a8e86aeb84)
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /home/aduhtkjm/llvm/llvm-project/build/bin
Build config: +assertions
```
The threshold of crashing seem to be 65536. Is this intentional?
![]()
_______________________________________________
llvm-bugs mailing list
llvm-bugs@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-bugs
