Clang 11.0.0 CUDA compiler bug: tuple conversion constructor (fixed in 11.0.1)

[havogt]

The following program crashes Clang 11.0.0 (in the Cray provided version, however -fno-cray, -fno-cray-gpu doesn't fix it):

#include <gridtools/common/tuple.hpp>
__global__ void test_kernel() {
    gridtools::tuple<int, double> tpl{'a', 'b'};
}

Seems to be fixed in Clang 11.0.1 (tested on my notebook), but could also be related to the Cray wrapper or daint specific setup.

Error

/opt/cray/pe/craype/2.7.3/bin/CC -I/scratch/snx3000/vogtha/gridtools/include -isystem /project/c14/install/daint/boost/boost_1_67_0/include  --cuda-gpu-arch=sm_60 -xcuda --cuda-path=/usr/local/cuda-11.0/bin/.. -std=c++14 test_tuple.cu -c
clang-11: /b/worker/clang-release-craystable-x86/build/src/llvm/lib/CodeGen/SelectionDAG/LegalizeDAG.cpp:2427: llvm::SDValue {anonymous}::SelectionDAGLegalize::ExpandLegalINT_TO_FP(llvm::SDNode*, llvm::SDValue&): Assertion `!isSigned && "Legalize cannot Expand SINT_TO_FP for i64 yet"' failed.
PLEASE submit a bug report to Cray and include the crash backtrace, preprocessed source, and associated run script.
Stack dump:
0.      Program arguments: /opt/cray/pe/cce/11.0.0/cce-clang/x86_64/bin/clang-11 -cc1 -triple nvptx64-nvidia-cuda -mllvm -cray-omp-opt-fork-call -mllvm -cray-omp-parallel-opt-call -fcray-omp-target-debug -fcray-gpu -fcray-openmp-rename-outlined-funcs -mllvm -cray-openmp-rename-outlined -flocal-restrict -mllvm -cray-enhanced-asm=1 -fenhanced-asm=1 -mllvm -cray-enhanced-ir=1 -fenhanced-ir=1 -fomp-local-offload-table -ffortran -aux-triple x86_64-unknown-linux-gnu -S -disable-free -main-file-name test_tuple.cu -mrelocation-model static -mframe-pointer=all -fno-rounding-math -fno-verbose-asm -no-integrated-as -aux-target-cpu x86-64 -fcuda-is-device -mlink-builtin-bitcode /usr/local/cuda-11.0/bin/../nvvm/libdevice/libdevice.10.bc -target-feature +ptx70 -target-sdk-version=11.0 -target-cpu sm_60 -fno-split-dwarf-inlining -debugger-tuning=gdb -resource-dir /opt/cray/pe/cce/11.0.0/cce-clang/x86_64/lib/clang/11.0.0 -internal-isystem /opt/cray/pe/cce/11.0.0/cce-clang/x86_64/lib/clang/11.0.0/include/cuda_wrappers -internal-isystem /usr/local/cuda-11.0/bin/../include -include __clang_cuda_runtime_wrapper.h -isystem /project/c14/install/daint/boost/boost_1_67_0/include -isystem /opt/cray/pe/cce/11.0.0/cce-clang/x86_64/lib/clang/11.0.0/include -isystem /opt/cray/pe/cce/11.0.0/cce/x86_64/include/craylibs -D __CRAYXC -D __CRAY_HASWELL -D __CRAYXT_COMPUTE_LINUX_TARGET -I /scratch/snx3000/vogtha/gridtools/include -I /opt/cray/pe/mpt/7.7.16/gni/mpich-crayclang/9.0/include -I /opt/cray/pe/libsci/20.09.1/CRAYCLANG/9.0/x86_64/include -I /usr/local/cuda-11.0/include -I /usr/local/cuda-11.0/nvvm/include -I /opt/cray/rca/2.2.20-7.0.2.1_2.55__g8e3fb5b.ari/include -I /opt/cray/alps/6.6.59-7.0.2.1_3.39__g872a8d62.ari/include -I /opt/cray/xpmem/2.2.20-7.0.2.1_2.45__g87eb960.ari/include -I /opt/cray/gni-headers/5.0.12.0-7.0.2.1_2.10__g3b1768f.ari/include -I /opt/cray/dmapp/7.1.1-7.0.2.1_2.52__g38cf134.ari/include -I /opt/cray/pe/pmi/5.0.17/include -I /opt/cray/ugni/6.0.14.0-7.0.2.1_3.46__ge78e5b0.ari/include -I /opt/cray/udreg/2.3.2-7.0.2.1_2.27__g8175d3d.ari/include -I /opt/cray/pe/atp/3.8.1/include -I /opt/cray/wlm_detect/1.3.3-7.0.2.1_2.10__g7109084.ari/include -I /opt/cray/krca/2.2.7-7.0.2.1_2.46__ge897ee1.ari/include -I /opt/cray-hss-devel/9.0.0/include -I/apps/dom/UES/xalt/xalt2/software/xalt/2.8.10/include -internal-isystem /opt/gcc/8.1.0/snos/lib/gcc/x86_64-suse-linux/8.1.0/../../../../include/g++ -internal-isystem /opt/gcc/8.1.0/snos/lib/gcc/x86_64-suse-linux/8.1.0/../../../../include/g++/x86_64-suse-linux -internal-isystem /opt/gcc/8.1.0/snos/lib/gcc/x86_64-suse-linux/8.1.0/../../../../include/g++/backward -internal-isystem /opt/gcc/8.1.0/snos/lib/gcc/x86_64-suse-linux/8.1.0/../../../../include/g++ -internal-isystem /opt/gcc/8.1.0/snos/lib/gcc/x86_64-suse-linux/8.1.0/../../../../include/g++/x86_64-suse-linux -internal-isystem /opt/gcc/8.1.0/snos/lib/gcc/x86_64-suse-linux/8.1.0/../../../../include/g++/backward -internal-isystem /usr/local/include -internal-isystem /opt/cray/pe/cce/11.0.0/cce-clang/x86_64/lib/clang/11.0.0/include -internal-externc-isystem /include -internal-externc-isystem /usr/include -internal-isystem /usr/local/include -internal-isystem /opt/cray/pe/cce/11.0.0/cce-clang/x86_64/lib/clang/11.0.0/include -internal-externc-isystem /include -internal-externc-isystem /usr/include -std=c++14 -fdeprecated-macro -fno-dwarf-directory-asm -fno-autolink -fdebug-compilation-dir /scratch/snx3000/vogtha/gridtools/build2/tests/unit_tests/common -ferror-limit 19 -fgnuc-version=4.2.1 -fcxx-exceptions -fexceptions -fcolor-diagnostics -o /tmp/test_tuple-35f123.s -x cuda test_tuple.cu 
1.      <eof> parser at end of file
2.      Code generation
3.      Running pass 'Function Pass Manager' on module 'test_tuple.cu'.
4.      Running pass 'NVPTX DAG->DAG Pattern Instruction Selection' on function '@_Z11test_kernelv'
 #0 0x0000000002def2b9 llvm::sys::PrintStackTrace(llvm::raw_ostream&) (/opt/cray/pe/cce/11.0.0/cce-clang/x86_64/bin/clang-11+0x2def2b9)
 #1 0x0000000002ded003 llvm::sys::RunSignalHandlers() (/opt/cray/pe/cce/11.0.0/cce-clang/x86_64/bin/clang-11+0x2ded003)
 #2 0x0000000002ded137 SignalHandler(int) (/opt/cray/pe/cce/11.0.0/cce-clang/x86_64/bin/clang-11+0x2ded137)
 #3 0x00002ada60d072cf (/apps/dom/UES/xalt/xalt2/software/xalt/2.8.10/lib64/libpthread.so.0+0x132cf)
 #4 0x00002ada61e4751f raise (/lib64/libc.so.6+0x3951f)
 #5 0x00002ada61e48b00 abort (/lib64/libc.so.6+0x3ab00)
 #6 0x00002ada61e3fb19 __assert_fail_base (/lib64/libc.so.6+0x31b19)
 #7 0x00002ada61e3fb91 __assert_fail (/lib64/libc.so.6+0x31b91)
 #8 0x0000000003aeb3eb (anonymous namespace)::SelectionDAGLegalize::ExpandNode(llvm::SDNode*) (/opt/cray/pe/cce/11.0.0/cce-clang/x86_64/bin/clang-11+0x3aeb3eb)
 #9 0x0000000003ad9d7f (anonymous namespace)::SelectionDAGLegalize::LegalizeOp(llvm::SDNode*) (/opt/cray/pe/cce/11.0.0/cce-clang/x86_64/bin/clang-11+0x3ad9d7f)
#10 0x0000000003adf2f3 llvm::SelectionDAG::Legalize() (/opt/cray/pe/cce/11.0.0/cce-clang/x86_64/bin/clang-11+0x3adf2f3)
#11 0x0000000003b8c4f4 llvm::SelectionDAGISel::CodeGenAndEmitDAG() (/opt/cray/pe/cce/11.0.0/cce-clang/x86_64/bin/clang-11+0x3b8c4f4)
#12 0x0000000003b905fc llvm::SelectionDAGISel::SelectAllBasicBlocks(llvm::Function const&) (/opt/cray/pe/cce/11.0.0/cce-clang/x86_64/bin/clang-11+0x3b905fc)
#13 0x0000000003b91f56 llvm::SelectionDAGISel::runOnMachineFunction(llvm::MachineFunction&) (/opt/cray/pe/cce/11.0.0/cce-clang/x86_64/bin/clang-11+0x3b91f56)
#14 0x0000000002284a3d llvm::MachineFunctionPass::runOnFunction(llvm::Function&) (/opt/cray/pe/cce/11.0.0/cce-clang/x86_64/bin/clang-11+0x2284a3d)
#15 0x000000000267ce9e llvm::FPPassManager::runOnFunction(llvm::Function&) (/opt/cray/pe/cce/11.0.0/cce-clang/x86_64/bin/clang-11+0x267ce9e)
#16 0x000000000267d978 llvm::FPPassManager::runOnModule(llvm::Module&) (/opt/cray/pe/cce/11.0.0/cce-clang/x86_64/bin/clang-11+0x267d978)
#17 0x000000000267c82a llvm::legacy::PassManagerImpl::run(llvm::Module&) (/opt/cray/pe/cce/11.0.0/cce-clang/x86_64/bin/clang-11+0x267c82a)
#18 0x000000000307d1d8 (anonymous namespace)::EmitAssemblyHelper::EmitAssembly(clang::BackendAction, std::unique_ptr<llvm::raw_pwrite_stream, std::default_delete<llvm::raw_pwrite_stream> >) (/opt/cray/pe/cce/11.0.0/cce-clang/x86_64/bin/clang-11+0x307d1d8)
#19 0x000000000307dd4d clang::EmitBackendOutput(clang::DiagnosticsEngine&, clang::HeaderSearchOptions const&, clang::CodeGenOptions const&, clang::TargetOptions const&, clang::LangOptions const&, clang::FrontendOptions const&, llvm::DataLayout const&, llvm::Module*, clang::BackendAction, std::unique_ptr<llvm::raw_pwrite_stream, std::default_delete<llvm::raw_pwrite_stream> >) (/opt/cray/pe/cce/11.0.0/cce-clang/x86_64/bin/clang-11+0x307dd4d)
#20 0x0000000003c97fa6 clang::BackendConsumer::HandleTranslationUnit(clang::ASTContext&) (/opt/cray/pe/cce/11.0.0/cce-clang/x86_64/bin/clang-11+0x3c97fa6)
#21 0x0000000004740d68 clang::ParseAST(clang::Sema&, bool, bool) (/opt/cray/pe/cce/11.0.0/cce-clang/x86_64/bin/clang-11+0x4740d68)
#22 0x0000000003c96897 clang::CodeGenAction::ExecuteAction() (/opt/cray/pe/cce/11.0.0/cce-clang/x86_64/bin/clang-11+0x3c96897)
#23 0x000000000367b7d0 clang::FrontendAction::Execute() (/opt/cray/pe/cce/11.0.0/cce-clang/x86_64/bin/clang-11+0x367b7d0)
#24 0x0000000003633dd8 clang::CompilerInstance::ExecuteAction(clang::FrontendAction&) (/opt/cray/pe/cce/11.0.0/cce-clang/x86_64/bin/clang-11+0x3633dd8)
#25 0x0000000003737f00 clang::ExecuteCompilerInvocation(clang::CompilerInstance*) (/opt/cray/pe/cce/11.0.0/cce-clang/x86_64/bin/clang-11+0x3737f00)
#26 0x0000000000cbd69c cc1_main(llvm::ArrayRef<char const*>, char const*, void*) (/opt/cray/pe/cce/11.0.0/cce-clang/x86_64/bin/clang-11+0xcbd69c)
#27 0x0000000000cbbe26 ExecuteCC1Tool(llvm::SmallVectorImpl<char const*>&) (/opt/cray/pe/cce/11.0.0/cce-clang/x86_64/bin/clang-11+0xcbbe26)
#28 0x0000000000c0c02c main (/opt/cray/pe/cce/11.0.0/cce-clang/x86_64/bin/clang-11+0xc0c02c)
#29 0x00002ada61e32349 __libc_start_main (/lib64/libc.so.6+0x24349)
#30 0x0000000000cb8c88 _start /home/abuild/rpmbuild/BUILD/glibc-2.19/csu/../sysdeps/x86_64/start.S:122:0
clang-11: error: unable to execute command: Aborted (core dumped)
clang-11: error: clang frontend command failed due to signal (use -v to see invocation)
Cray clang version 11.0.0  (1563bb12fb3eabb03515d4cc4aeaedd757aea56a)
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /opt/cray/pe/cce/11.0.0/cce-clang/x86_64/share/../bin
clang-11: note: diagnostic msg: 
********************
PLEASE ATTACH THE FOLLOWING FILES TO THE BUG REPORT:
Preprocessed source(s) and associated run script(s) are located at:
clang-11: note: diagnostic msg: /tmp/test_tuple-51e055.cu
clang-11: note: diagnostic msg: /tmp/test_tuple-4ae26f.cu
clang-11: note: diagnostic msg: /tmp/test_tuple-51e055.sh
clang-11: note: diagnostic msg: 
********************

[havogt]

This particular test is disabled in #1613 for Clang 11.0.0.

[havogt]

Let's keep this issue open until daint has a newer version with a fix.

[havogt]

Closing as the CUDA version is probably not in use anymore.