<table border="1" cellspacing="0" cellpadding="8">
<tr>
<th>Issue</th>
<td>
<a href=https://github.com/llvm/llvm-project/issues/131517>131517</a>
</td>
</tr>
<tr>
<th>Summary</th>
<td>
Clang crashes for large structs in CUDA
</td>
</tr>
<tr>
<th>Labels</th>
<td>
</td>
</tr>
<tr>
<th>Assignees</th>
<td>
</td>
</tr>
<tr>
<th>Reporter</th>
<td>
AdUhTkJm
</td>
</tr>
</table>
<pre>
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?
</pre>
<img width="1" height="1" alt="" src="http://email.email.llvm.org/o/eJzsWt1y3LaSfhrqBoUpEvy_0AU14_HxlpOoLMVna29YINDk4IgkuAAoyW-_BZCjIUey4uRMsrtVcdkWBQLdX_8B3U1QrUXTA1x78Y0X767oaA5SXRf818P9w390V5Xk3663steCg0LmIDTyEr8SzYaNXuJ7YeH5hX1wf9kweH6hjRqZQaKjDSAvvfH8AiGERG_QsxffBDnxkUcKFPiZ78U7L7QzvPT44BdlyeFRMCjLmUjtkcwj-YKYAjOq3g2kCwoLLJ5f_FOYg4XLWto3HrnxyA06QSdbJAxiiuoDaPRk55oDoFq2rXwSfYO0oewBGUUZnMnp-YWjiYnVAPLI_iA78Mie8vFgHv7VeWTfto_HH3hQ8l_AzGlU9KwdOZwGtpLDR-g9sr-DFpgRst8VH3-WHPTm4IVFECSJZeWmh4UXFnc7-3b1TLKxd-bkVtlWwPPfd1CNzWfJ7PPd7uv9Z6GNR3JLudAalOVrNfbz2H2l7QgaeeHOC3fo673e2MF7jTySeCRBHiGnWU9U9x5JDXoSHBD0cmwOqJYKCaORHEDRnmuPBB4hHklRTUULfOP5xe3nD8XdB6THqhMGUVSNDVIwSGWQkehgzKCtiGTvkX0jzGGsNkx-X71C6xG0R_aI9lZop2VnVmdnVFH2MBmUbNGgYFCSgdbAkZajmobtSqq1ZIIa4EiNPdJMicFYvHfOJ_jYDZNL-Bvk_twq2SjaIaqasYPe6N_pF9UoWm5_CusDR-dCmLEAYaPE0ALqHwfznES4fxRcUMxGThGm4_Px_XOWlEmEx_6hl089bkU_PuOmHxG-Q9hC5kIhihHmQtOqBVwrAIRZC1Rhqg2uoJYKsFUS9Bzhjooe16IF3NMO5tBBuFPQSkatr-BOcmhtpBjBEO5qRTvAgxS9AeWFO9q2CNe9xEqOPRd9gztqDgj3EtspjbIqxlTPclDVgMFsGK0sOIkQrq2UWGg87QgId1ZxCEPvJOigY8M3ORhs41eOBreiYrRt9cTWHBRQrmkNeMKojySpjXL8SJWgXDBcj70LOm05iP4BW4MY0eNKGCY5WFuOWlnDSUZba6ORU4_s-8fJnKKaEC6fN4G_qRg6yqX5A34EpYXsvXAXkE2GliLrrozJy0gN1IzK8r0ZzHOWIsxt7DagsBl70TdeuGt4NYnJn6iqMRcKmJHqG6a6Q7h2CzCT3SDayVpcWKO855aGdgMohBVMAWGX_BE_bkV19GOP7K0nO5foaYuF_qYNdJegutpJrUHKJ0WHAZS2_KbgL0s3v3Sv1dgb0cFx2ubwNq7Z0o5bw5hH9nNsvcSUR_ZB6JH9ZvP6vwWk6dBxk_9cTm-s-uuYLzlZa1H28EQV_1u3f7VuLxdP7wo174A_MvUPyv_G3Nfc4Nk-sNdcv2-691f_fhY_oqb5oHh3xd-G-8sNV3MYFDCXgHSUKTkdpXQ00iYACNeglFS4FTYzDXKEsazrVlKOe3jCXIlHe1DWTT-yxbEebcgmQLjWD2LAkivMDsAesOhx09U2-Xh-xvDMYJizjXr1C5OtVJgL2vRST-kKGwX3wl0c8JDmdZUFaRKTnCK8K8uP2235j-Lrh3L3z-LLnpTb_aeyuPvJZhcISyu06QaXVTbYpRc4qX0f6EYj_IxcAjnXQ34RzJmsF25B1l74AQ1UaVCIGmSzQVkjmwp6fkHmmbZmQQ30oFx24flFOL_5MvY2RUED1bZeSPdzdoVu7cBPtKcNKFsIyB51ko-tzXPSGQpJbaIdfYfUz19v7_8T7YqP2As_7IqP6JYaa2r0qZ_qTsvmpYaaeRyzO0vAi_zyv4I6rERjf_Ntmh76yH_2fd-Pk4BBRCuSZ9mi1tLf9PRwq0RvXA1wP1UR2WmSok-l1EYB7VyJtJ2Krhy54vXfrAQ8cuM_R3WV5ZmthGfYwRlsPwrhLdhfxv5OND1t_0F73oLSx3r6gsiSKIQTMnKGLPSDDK0weCSb9TMN6w0bbFHlu38TlXCmktZhHedJGsSudZC9bFBvbDe2ANhouUkmYCS2EI-wojVBFuc1Q2U5THVC-SDathTd0IKt4ZxXI7ur9YOx--Zy2oZ5YRFFXlikyUw8_hHi8wb1Ltk0sxX_UQnJW2SXa96lleVLWum5QqM0QY12BnBkpsfjbq6_aQ6DracHqcWzR_aKCg2OMEm9sDiKnq3pxmlah4hWtox3dAyf7OWGJlyBFXRenr9aHlSo7NvSbrcll7YQRVNe5KS0w51uGJ1oBUFqwVgj26Dw19SSDPLk93lNmEOeTF5jCa7CLKRBkBJyudgJ6yBNCTmxW8VOEmd-mqSvGz9k1QxaNI2mkR6ejh2i7eu1552h5K1m0WnZsW2EmOz1cfqS6rGTlHjhh1dtqDeIJ69o_A4WdstYyHu2b1gNhmsN0oTx6l1tNWC-00FbQvgReMvRQin67QvUZyZwPTOnqO0rs-5b2uiLbs0JpYzx6uRe0Vo5PkAYf0c5N5YsqOnFo9DCfJaUr049O2CP3oViLgo-iCCMT-DjM_BJXEc_DH4FfJkv_EnYfRbX0Ql7ssbOchJ8T_Gf7qBdjt5QLdhNK9nDSgbRCm1KYWwGJlX5VFbC6JW3TTM4GCpmgr3kUMop41xNXSjEeqZRo9soatpq9_AoBV--OK1coJu8-mWNm3vu6f-fUFdStpd3DFaTYOHU6ZljpMD4DzpG0bYnOfTKOV6S7j_JuxkHxk9CZGdCxFnu_6YQaux_6X-i7CB6OOJdCfHq3cWlqLLcP0mRr476Ks4YW8aoqz52xcd7eQkxHDVL5Ujx9VFG_DO1hn79nlo_Q0PZt_8zys39-qRcss6j6syP4mVicwbH1ooLQd6U4M-CHkVZEsXpCfq6nont7r30i_2tRXssbP-3QNOkYvlpXyHhGWjwa_IDoH9yFfnaU-ahSwNOI78-Zb4kOgOcpH6yANwunHsB_VM3nOLwL4GdsNQ_1QdklZTEBCoCIZqacw4GdMLcTB_VfhnNMNpEZPF6677QgLJHGe0ZzNnlaob7OvzLfP69yj6NEn1jE83V8FH64ozaDKV4OTSXp6lRoxaP8AXqbW9ujVodto_1HJF70cKd66zN56Q2fHoz9uK_RyiHs5WKPpXDkxIGymOLZLGIQ03H1p74LRj4jYXhh5nna5G2stdjB8rJfEGLx2kdEAhPFl-lcnEYRFHgv4fHDU6Nj3tFez19j_u1F2euUNzdb2Vv4Pni53WcBTQKFpvxKunIwhB4ni9FuKVKQ3F3vwZ4Bx2d_W9KjY4_Lwg1y6KI5_kJ6iq1iNO8ilmyhLpXVmcvHu3GPjwDGw1cutEVMxrEbBH7-Rqbn2YZR-8F9xLdMQSz7wtzaTdgPs-yU9q27nzGWcpiTpf4Z6QnMR7nT_-_tYddOAJ5RmJOT7jXuUQaZAmJEWNB2VGxPgQWVTg7UHVMh4vjLrIae6kYLgs_ZAEk5HQ0r_ujYRpEKanRUdfb4N7F1EKIu4627Vf3eX867t6Q5exUsDTmzWRVAkxfLl6nmmF4hilM63pyhaNS556IZe00VPwYx0uqMU_r-qTGaA2ZhHmaItctvCBPkufpKRU8b_PGGc99VJatqFipDVWmZLRty5eeJdPj637q3K59a5G77xW7FnAys0zOWULkI7dA9MKcc5keLWnsSE89UmIpkqOt07coroR4C_-blMM0t4BnwtmZExFCYzSRvKgfZITO4bS43eeFBXJf7ezD6K4IISMRTHGFmOw62nN3262SygC3kJhU4G50AX-foBtH9bw_H6nNl9kQHx2vuZfukWzUgPCjGwNAYrFxvjBB83dDRPyNv_Eb4XT027fe3GKhNnZB5FdxlZAsS4KMAa0yxnmSRz6J_LQKaAZZQqHK5j7UvbtdZKX53m0xO8l9Q0DubpedOjX_fddroW0LfCfUv3G_zS9ca85uEbVoJkI39Hj9UJ9dtLw_ADIHBfogW_cV0t3kc_czATqr3gpQEsdhskGf9HRFVVgLWWK09cL9Fb8OeR7m9AqugzQiQUiyxL86XIdVkrAwztKojjnjLK79qgJOK5IFjAbkSlwTn8R-GCRBGKWRv_HjKoyA1xDWPElT7kU-dFS0GyvuRqrmyt1AvA7CIA7Sq5ZW0Orj3Vp17ZRSjY32Ir8V2ujTOiNMC9db5xXHO6m1VKi1BkNTj8sKhra_7oqrUbXXf_xy5Izu8Zr8TwAAAP__X5y5fQ">