[LLVM 17.0.0 RC4][OpenMP][AMDGPU offload] Stack protector makes the linker crash

[admnd]

When invoking LLVM 17.0.0 RC4 with OpenMP/AMDGPU offload and the stack protector being active, the linker fails with:

Assertion `(i >= FTy->getNumParams() || FTy->getParamType(i) == Args[i]->getType()) && "Calling a function with a bad signature!"' failed.

Command line and full stack dump :

$ /usr/lib/llvm/17/bin/clang --rocm-device-lib-path=/usr/lib/amdgcn/bitcode -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa -x c openmp_offload.c

clang-linker-wrapper: /var/tmp/portage/sys-devel/llvm-17.0.0_rc4/work/llvm/lib/IR/Instructions.cpp:652: void llvm::CallInst::init(llvm::FunctionType*, llvm::Value*, llvm::ArrayRef<llvm::Value*>, llvm::ArrayRef<llvm::OperandBundleDefT<llvm::Value*> >, const llvm::Twine&): Assertion `(i >= FTy->getNumParams() || FTy->getParamType(i) == Args[i]->getType()) && "Calling a function with a bad signature!"' failed.
PLEASE submit a bug report to https://github.com/llvm/llvm-project/issues/ and include the crash backtrace.
Stack dump:
0.      Program arguments: /usr/lib/llvm/17/bin/clang-linker-wrapper --host-triple=x86_64-pc-linux-gnu --linker-path=/usr/bin/x86_64-pc-linux-gnu-ld.bfd -- -pie --hash-style=gnu --eh-frame-hdr -m elf_x86_64 -dynamic-linker /lib64/ld-linux-x86-64.so.2 -o a.out /usr/lib/gcc/x86_64-pc-linux-gnu/13/../../../../lib64/Scrt1.o /usr/lib/gcc/x86_64-pc-linux-gnu/13/../../../../lib64/crti.o /usr/lib/gcc/x86_64-pc-linux-gnu/13/crtbeginS.o -L/usr/lib/gcc/x86_64-pc-linux-gnu/13 -L/usr/lib/gcc/x86_64-pc-linux-gnu/13/../../../../lib64 -L/lib/../lib64 -L/usr/lib/../lib64 -L/usr/lib/gcc/x86_64-pc-linux-gnu/13/../../../../x86_64-pc-linux-gnu/lib -L/lib -L/usr/lib -z relro /tmp/openmp_offload-c3b2da.o -lomp -lomptarget -lomptarget.devicertl -L/usr/lib/llvm/17/lib64 -lgcc --as-needed -lgcc_s --no-as-needed -lpthread -lc -lgcc --as-needed -lgcc_s --no-as-needed /usr/lib/gcc/x86_64-pc-linux-gnu/13/crtendS.o /usr/lib/gcc/x86_64-pc-linux-gnu/13/../../../../lib64/crtn.o
1.      Running pass 'CallGraph Pass Manager' on module 'ld-temp.o'.
2.      Running pass 'Insert stack protectors' on function '@__omp_offloading_43_17b92b_main_l12_omp_outlined'
 #0 0x00007fd91f47eaee llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0xf15aee)
 #1 0x00007fd91f435a8b llvm::sys::RunSignalHandlers() (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0xecca8b)
 #2 0x00007fd91f435c37 (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0xeccc37)
 #3 0x00007fd91dfd44a0 (/lib64/libc.so.6+0x3a4a0)
 #4 0x00007fd91e0258ec __pthread_kill_implementation /usr/src/debug/sys-libs/glibc-2.38-r1/glibc-2.38/nptl/pthread_kill.c:44:76
 #5 0x00007fd91dfd43f6 gsignal /usr/src/debug/sys-libs/glibc-2.38-r1/glibc-2.38/signal/../sysdeps/posix/raise.c:27:6
 #6 0x00007fd91dfbc8f5 abort /usr/src/debug/sys-libs/glibc-2.38-r1/glibc-2.38/stdlib/abort.c:81:7
 #7 0x00007fd91dfbc819 _nl_load_domain /usr/src/debug/sys-libs/glibc-2.38-r1/glibc-2.38/intl/loadmsgcat.c:1177:9
 #8 0x00007fd91dfccad6 (/lib64/libc.so.6+0x32ad6)
 #9 0x00007fd91f616428 llvm::CallInst::init(llvm::FunctionType*, llvm::Value*, llvm::ArrayRef<llvm::Value*>, llvm::ArrayRef<llvm::OperandBundleDefT<llvm::Value*>>, llvm::Twine const&) (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x10ad428)
#10 0x00007fd92307822d (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x4b0f22d)
#11 0x00007fd91fbef7f0 llvm::StackProtector::InsertStackProtectors() (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x16867f0)
#12 0x00007fd91fbf0354 llvm::StackProtector::runOnFunction(llvm::Function&) (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x1687354)
#13 0x00007fd91f63fa0b llvm::FPPassManager::runOnFunction(llvm::Function&) (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x10d6a0b)
#14 0x00007fd920c9c862 (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x2733862)
#15 0x00007fd91f63edca llvm::legacy::PassManagerImpl::run(llvm::Module&) (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x10d5dca)
#16 0x00007fd920f65fbd (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x29fcfbd)
#17 0x00007fd920f670b8 llvm::lto::backend(llvm::lto::Config const&, std::function<llvm::Expected<std::unique_ptr<llvm::CachedFileStream, std::default_delete<llvm::CachedFileStream>>> (unsigned int, llvm::Twine const&)>, unsigned int, llvm::Module&, llvm::ModuleSummaryIndex&) (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x29fe0b8)
#18 0x00007fd920f67fb5 llvm::lto::LTO::runRegularLTO(std::function<llvm::Expected<std::unique_ptr<llvm::CachedFileStream, std::default_delete<llvm::CachedFileStream>>> (unsigned int, llvm::Twine const&)>) (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x29fefb5)
#19 0x00007fd920f685fa llvm::lto::LTO::run(std::function<llvm::Expected<std::unique_ptr<llvm::CachedFileStream, std::default_delete<llvm::CachedFileStream>>> (unsigned int, llvm::Twine const&)>, std::function<llvm::Expected<std::function<llvm::Expected<std::unique_ptr<llvm::CachedFileStream, std::default_delete<llvm::CachedFileStream>>> (unsigned int, llvm::Twine const&)>> (unsigned int, llvm::StringRef, llvm::Twine const&)>) (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x29ff5fa)
#20 0x0000557e27d8dc07 (/usr/lib/llvm/17/bin/clang-linker-wrapper+0x34c07)
#21 0x0000557e27d801a3 (/usr/lib/llvm/17/bin/clang-linker-wrapper+0x271a3)
#22 0x0000557e27d723eb (/usr/lib/llvm/17/bin/clang-linker-wrapper+0x193eb)
#23 0x0000557e27d67f29 (/usr/lib/llvm/17/bin/clang-linker-wrapper+0xef29)
#24 0x00007fd91dfbe31e __libc_start_call_main /usr/src/debug/sys-libs/glibc-2.38-r1/glibc-2.38/csu/../sysdeps/nptl/libc_start_call_main.h:74:3
#25 0x00007fd91dfbe3d9 call_init /usr/src/debug/sys-libs/glibc-2.38-r1/glibc-2.38/csu/../csu/libc-start.c:128:20
#26 0x00007fd91dfbe3d9 __libc_start_main /usr/src/debug/sys-libs/glibc-2.38-r1/glibc-2.38/csu/../csu/libc-start.c:347:5
#27 0x0000557e27d68e35 (/usr/lib/llvm/17/bin/clang-linker-wrapper+0xfe35)
 #0 0x00007fd91f47eaee llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0xf15aee)
 #1 0x00007fd91f435a8b llvm::sys::RunSignalHandlers() (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0xecca8b)
 #2 0x00007fd91f435c37 (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0xeccc37)
 #3 0x00007fd91dfd44a0 (/lib64/libc.so.6+0x3a4a0)
 #4 0x00007fd91e0258ec __pthread_kill_implementation /usr/src/debug/sys-libs/glibc-2.38-r1/glibc-2.38/nptl/pthread_kill.c:44:76
 #5 0x00007fd91dfd43f6 gsignal /usr/src/debug/sys-libs/glibc-2.38-r1/glibc-2.38/signal/../sysdeps/posix/raise.c:27:6
 #6 0x00007fd91dfbc8f5 abort /usr/src/debug/sys-libs/glibc-2.38-r1/glibc-2.38/stdlib/abort.c:81:7
 #7 0x00007fd91dfbc819 _nl_load_domain /usr/src/debug/sys-libs/glibc-2.38-r1/glibc-2.38/intl/loadmsgcat.c:1177:9
 #8 0x00007fd91dfccad6 (/lib64/libc.so.6+0x32ad6)
 #9 0x00007fd91f616428 llvm::CallInst::init(llvm::FunctionType*, llvm::Value*, llvm::ArrayRef<llvm::Value*>, llvm::ArrayRef<llvm::OperandBundleDefT<llvm::Value*>>, llvm::Twine const&) (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x10ad428)
#10 0x00007fd92307822d (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x4b0f22d)
#11 0x00007fd91fbef7f0 llvm::StackProtector::InsertStackProtectors() (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x16867f0)
#12 0x00007fd91fbf0354 llvm::StackProtector::runOnFunction(llvm::Function&) (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x1687354)
#13 0x00007fd91f63fa0b llvm::FPPassManager::runOnFunction(llvm::Function&) (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x10d6a0b)
#14 0x00007fd920c9c862 (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x2733862)
#15 0x00007fd91f63edca llvm::legacy::PassManagerImpl::run(llvm::Module&) (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x10d5dca)
#16 0x00007fd920f65fbd (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x29fcfbd)
#17 0x00007fd920f670b8 llvm::lto::backend(llvm::lto::Config const&, std::function<llvm::Expected<std::unique_ptr<llvm::CachedFileStream, std::default_delete<llvm::CachedFileStream>>> (unsigned int, llvm::Twine const&)>, unsigned int, llvm::Module&, llvm::ModuleSummaryIndex&) (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x29fe0b8)
#18 0x00007fd920f67fb5 llvm::lto::LTO::runRegularLTO(std::function<llvm::Expected<std::unique_ptr<llvm::CachedFileStream, std::default_delete<llvm::CachedFileStream>>> (unsigned int, llvm::Twine const&)>) (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x29fefb5)
#19 0x00007fd920f685fa llvm::lto::LTO::run(std::function<llvm::Expected<std::unique_ptr<llvm::CachedFileStream, std::default_delete<llvm::CachedFileStream>>> (unsigned int, llvm::Twine const&)>, std::function<llvm::Expected<std::function<llvm::Expected<std::unique_ptr<llvm::CachedFileStream, std::default_delete<llvm::CachedFileStream>>> (unsigned int, llvm::Twine const&)>> (unsigned int, llvm::StringRef, llvm::Twine const&)>) (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x29ff5fa)
#20 0x0000557e27d8dc07 (/usr/lib/llvm/17/bin/clang-linker-wrapper+0x34c07)
#21 0x0000557e27d801a3 (/usr/lib/llvm/17/bin/clang-linker-wrapper+0x271a3)
#22 0x0000557e27d723eb (/usr/lib/llvm/17/bin/clang-linker-wrapper+0x193eb)
#23 0x0000557e27d67f29 (/usr/lib/llvm/17/bin/clang-linker-wrapper+0xef29)
#24 0x00007fd91dfbe31e __libc_start_call_main /usr/src/debug/sys-libs/glibc-2.38-r1/glibc-2.38/csu/../sysdeps/nptl/libc_start_call_main.h:74:3
#25 0x00007fd91dfbe3d9 call_init /usr/src/debug/sys-libs/glibc-2.38-r1/glibc-2.38/csu/../csu/libc-start.c:128:20
#26 0x00007fd91dfbe3d9 __libc_start_main /usr/src/debug/sys-libs/glibc-2.38-r1/glibc-2.38/csu/../csu/libc-start.c:347:5
#27 0x0000557e27d68e35 (/usr/lib/llvm/17/bin/clang-linker-wrapper+0xfe35)
clang: error: unable to execute command: Aborted

When the stack protector is disabled (-fno-stack-protector) the compilation succeeds and the compiled program runs normally :

$ /usr/lib/llvm/17/bin/clang -fno-stack-protector  --rocm-device-lib-path=/usr/lib/amdgcn/bitcode -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa -x c openmp_offload.c
$ ./a.out 
number of devices: 1
ran on GPU 0: yes, 160 teams, limit of 256 threads

openmp_offload.c :

#include <stdio.h>
#include <omp.h>
int main(int argc, char const * argv[]) {
    int ndev;
    int nteams;
    int thread_lmt;
    int on_host;
 
    ndev = omp_get_num_devices();
    printf("number of devices: %d\n", ndev);
    for (int i = 0; i < omp_get_num_devices(); ++i) {
        #pragma omp target teams \
                           device(i) \
                           map(from:on_host) map(from:nteams) \
                           map(from:thread_lmt)
        #pragma omp parallel
        #pragma omp master
        if (0 == omp_get_team_num()) {
            on_host = omp_is_initial_device();
            nteams = omp_get_num_teams();
            thread_lmt = omp_get_thread_limit();
        }
        printf("ran on GPU %d: %s, %d teams, "
               "limit of %d threads\n",
               i, on_host ? "no" : "yes", nteams, thread_lmt);
    }
    return 0;
}

[llvmbot]

@llvm/issue-subscribers-openmp

When invoking LLVM 17.0.0 RC4 with OpenMP/AMDGPU offload **and** the stack protector being active, the linker fails with: ``` Assertion `(i >= FTy->getNumParams() || FTy->getParamType(i) == Args[i]->getType()) && "Calling a function with a bad signature!"' failed. ``` Command line and full stack dump : ``` $ /usr/lib/llvm/17/bin/clang --rocm-device-lib-path=/usr/lib/amdgcn/bitcode -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa -x c openmp_offload.c

clang-linker-wrapper: /var/tmp/portage/sys-devel/llvm-17.0.0_rc4/work/llvm/lib/IR/Instructions.cpp:652: void llvm::CallInst::init(llvm::FunctionType*, llvm::Value*, llvm::ArrayRef<llvm::Value*>, llvm::ArrayRef<llvm::OperandBundleDefT<llvm::Value*> >, const llvm::Twine&): Assertion `(i >= FTy->getNumParams() || FTy->getParamType(i) == Args[i]->getType()) && "Calling a function with a bad signature!"' failed.
PLEASE submit a bug report to https://github.com/llvm/llvm-project/issues/ and include the crash backtrace.
Stack dump:
0. Program arguments: /usr/lib/llvm/17/bin/clang-linker-wrapper --host-triple=x86_64-pc-linux-gnu --linker-path=/usr/bin/x86_64-pc-linux-gnu-ld.bfd -- -pie --hash-style=gnu --eh-frame-hdr -m elf_x86_64 -dynamic-linker /lib64/ld-linux-x86-64.so.2 -o a.out /usr/lib/gcc/x86_64-pc-linux-gnu/13/../../../../lib64/Scrt1.o /usr/lib/gcc/x86_64-pc-linux-gnu/13/../../../../lib64/crti.o /usr/lib/gcc/x86_64-pc-linux-gnu/13/crtbeginS.o -L/usr/lib/gcc/x86_64-pc-linux-gnu/13 -L/usr/lib/gcc/x86_64-pc-linux-gnu/13/../../../../lib64 -L/lib/../lib64 -L/usr/lib/../lib64 -L/usr/lib/gcc/x86_64-pc-linux-gnu/13/../../../../x86_64-pc-linux-gnu/lib -L/lib -L/usr/lib -z relro /tmp/openmp_offload-c3b2da.o -lomp -lomptarget -lomptarget.devicertl -L/usr/lib/llvm/17/lib64 -lgcc --as-needed -lgcc_s --no-as-needed -lpthread -lc -lgcc --as-needed -lgcc_s --no-as-needed /usr/lib/gcc/x86_64-pc-linux-gnu/13/crtendS.o /usr/lib/gcc/x86_64-pc-linux-gnu/13/../../../../lib64/crtn.o

1.  Running pass 'CallGraph Pass Manager' on module 'ld-temp.o'.
   

2.  Running pass 'Insert stack protectors' on function '@<!-- -->__omp_offloading_43_17b92b_main_l12_omp_outlined'
   

#0 0x00007fd91f47eaee llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0xf15aee)
#1 0x00007fd91f435a8b llvm::sys::RunSignalHandlers() (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0xecca8b)
#2 0x00007fd91f435c37 (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0xeccc37)
#3 0x00007fd91dfd44a0 (/lib64/libc.so.6+0x3a4a0)
#4 0x00007fd91e0258ec __pthread_kill_implementation /usr/src/debug/sys-libs/glibc-2.38-r1/glibc-2.38/nptl/pthread_kill.c:44:76
#5 0x00007fd91dfd43f6 gsignal /usr/src/debug/sys-libs/glibc-2.38-r1/glibc-2.38/signal/../sysdeps/posix/raise.c:27:6
#6 0x00007fd91dfbc8f5 abort /usr/src/debug/sys-libs/glibc-2.38-r1/glibc-2.38/stdlib/abort.c:81:7
#7 0x00007fd91dfbc819 _nl_load_domain /usr/src/debug/sys-libs/glibc-2.38-r1/glibc-2.38/intl/loadmsgcat.c:1177:9
#8 0x00007fd91dfccad6 (/lib64/libc.so.6+0x32ad6)
#9 0x00007fd91f616428 llvm::CallInst::init(llvm::FunctionType*, llvm::Value*, llvm::ArrayRef<llvm::Value*>, llvm::ArrayRef<llvm::OperandBundleDefT<llvm::Value*>>, llvm::Twine const&) (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x10ad428)
#10 0x00007fd92307822d (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x4b0f22d)
#11 0x00007fd91fbef7f0 llvm::StackProtector::InsertStackProtectors() (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x16867f0)
#12 0x00007fd91fbf0354 llvm::StackProtector::runOnFunction(llvm::Function&) (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x1687354)
#13 0x00007fd91f63fa0b llvm::FPPassManager::runOnFunction(llvm::Function&) (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x10d6a0b)
#14 0x00007fd920c9c862 (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x2733862)
#15 0x00007fd91f63edca llvm::legacy::PassManagerImpl::run(llvm::Module&) (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x10d5dca)
#16 0x00007fd920f65fbd (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x29fcfbd)
#17 0x00007fd920f670b8 llvm::lto::backend(llvm::lto::Config const&, std::function<llvm::Expected<std::unique_ptr<llvm::CachedFileStream, std::default_delete<llvm::CachedFileStream>>> (unsigned int, llvm::Twine const&)>, unsigned int, llvm::Module&, llvm::ModuleSummaryIndex&) (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x29fe0b8)
#18 0x00007fd920f67fb5 llvm::lto::LTO::runRegularLTO(std::function<llvm::Expected<std::unique_ptr<llvm::CachedFileStream, std::default_delete<llvm::CachedFileStream>>> (unsigned int, llvm::Twine const&)>) (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x29fefb5)
#19 0x00007fd920f685fa llvm::lto::LTO::run(std::function<llvm::Expected<std::unique_ptr<llvm::CachedFileStream, std::default_delete<llvm::CachedFileStream>>> (unsigned int, llvm::Twine const&)>, std::function<llvm::Expected<std::function<llvm::Expected<std::unique_ptr<llvm::CachedFileStream, std::default_delete<llvm::CachedFileStream>>> (unsigned int, llvm::Twine const&)>> (unsigned int, llvm::StringRef, llvm::Twine const&)>) (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x29ff5fa)
#20 0x0000557e27d8dc07 (/usr/lib/llvm/17/bin/clang-linker-wrapper+0x34c07)
#21 0x0000557e27d801a3 (/usr/lib/llvm/17/bin/clang-linker-wrapper+0x271a3)
#22 0x0000557e27d723eb (/usr/lib/llvm/17/bin/clang-linker-wrapper+0x193eb)
#23 0x0000557e27d67f29 (/usr/lib/llvm/17/bin/clang-linker-wrapper+0xef29)
#24 0x00007fd91dfbe31e __libc_start_call_main /usr/src/debug/sys-libs/glibc-2.38-r1/glibc-2.38/csu/../sysdeps/nptl/libc_start_call_main.h:74:3
#25 0x00007fd91dfbe3d9 call_init /usr/src/debug/sys-libs/glibc-2.38-r1/glibc-2.38/csu/../csu/libc-start.c:128:20
#26 0x00007fd91dfbe3d9 __libc_start_main /usr/src/debug/sys-libs/glibc-2.38-r1/glibc-2.38/csu/../csu/libc-start.c:347:5
#27 0x0000557e27d68e35 (/usr/lib/llvm/17/bin/clang-linker-wrapper+0xfe35)
#0 0x00007fd91f47eaee llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0xf15aee)
#1 0x00007fd91f435a8b llvm::sys::RunSignalHandlers() (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0xecca8b)
#2 0x00007fd91f435c37 (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0xeccc37)
#3 0x00007fd91dfd44a0 (/lib64/libc.so.6+0x3a4a0)
#4 0x00007fd91e0258ec __pthread_kill_implementation /usr/src/debug/sys-libs/glibc-2.38-r1/glibc-2.38/nptl/pthread_kill.c:44:76
#5 0x00007fd91dfd43f6 gsignal /usr/src/debug/sys-libs/glibc-2.38-r1/glibc-2.38/signal/../sysdeps/posix/raise.c:27:6
#6 0x00007fd91dfbc8f5 abort /usr/src/debug/sys-libs/glibc-2.38-r1/glibc-2.38/stdlib/abort.c:81:7
#7 0x00007fd91dfbc819 _nl_load_domain /usr/src/debug/sys-libs/glibc-2.38-r1/glibc-2.38/intl/loadmsgcat.c:1177:9
#8 0x00007fd91dfccad6 (/lib64/libc.so.6+0x32ad6)
#9 0x00007fd91f616428 llvm::CallInst::init(llvm::FunctionType*, llvm::Value*, llvm::ArrayRef<llvm::Value*>, llvm::ArrayRef<llvm::OperandBundleDefT<llvm::Value*>>, llvm::Twine const&) (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x10ad428)
#10 0x00007fd92307822d (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x4b0f22d)
#11 0x00007fd91fbef7f0 llvm::StackProtector::InsertStackProtectors() (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x16867f0)
#12 0x00007fd91fbf0354 llvm::StackProtector::runOnFunction(llvm::Function&) (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x1687354)
#13 0x00007fd91f63fa0b llvm::FPPassManager::runOnFunction(llvm::Function&) (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x10d6a0b)
#14 0x00007fd920c9c862 (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x2733862)
#15 0x00007fd91f63edca llvm::legacy::PassManagerImpl::run(llvm::Module&) (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x10d5dca)
#16 0x00007fd920f65fbd (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x29fcfbd)
#17 0x00007fd920f670b8 llvm::lto::backend(llvm::lto::Config const&, std::function<llvm::Expected<std::unique_ptr<llvm::CachedFileStream, std::default_delete<llvm::CachedFileStream>>> (unsigned int, llvm::Twine const&)>, unsigned int, llvm::Module&, llvm::ModuleSummaryIndex&) (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x29fe0b8)
#18 0x00007fd920f67fb5 llvm::lto::LTO::runRegularLTO(std::function<llvm::Expected<std::unique_ptr<llvm::CachedFileStream, std::default_delete<llvm::CachedFileStream>>> (unsigned int, llvm::Twine const&)>) (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x29fefb5)
#19 0x00007fd920f685fa llvm::lto::LTO::run(std::function<llvm::Expected<std::unique_ptr<llvm::CachedFileStream, std::default_delete<llvm::CachedFileStream>>> (unsigned int, llvm::Twine const&)>, std::function<llvm::Expected<std::function<llvm::Expected<std::unique_ptr<llvm::CachedFileStream, std::default_delete<llvm::CachedFileStream>>> (unsigned int, llvm::Twine const&)>> (unsigned int, llvm::StringRef, llvm::Twine const&)>) (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x29ff5fa)
#20 0x0000557e27d8dc07 (/usr/lib/llvm/17/bin/clang-linker-wrapper+0x34c07)
#21 0x0000557e27d801a3 (/usr/lib/llvm/17/bin/clang-linker-wrapper+0x271a3)
#22 0x0000557e27d723eb (/usr/lib/llvm/17/bin/clang-linker-wrapper+0x193eb)
#23 0x0000557e27d67f29 (/usr/lib/llvm/17/bin/clang-linker-wrapper+0xef29)
#24 0x00007fd91dfbe31e __libc_start_call_main /usr/src/debug/sys-libs/glibc-2.38-r1/glibc-2.38/csu/../sysdeps/nptl/libc_start_call_main.h:74:3
#25 0x00007fd91dfbe3d9 call_init /usr/src/debug/sys-libs/glibc-2.38-r1/glibc-2.38/csu/../csu/libc-start.c:128:20
#26 0x00007fd91dfbe3d9 __libc_start_main /usr/src/debug/sys-libs/glibc-2.38-r1/glibc-2.38/csu/../csu/libc-start.c:347:5
#27 0x0000557e27d68e35 (/usr/lib/llvm/17/bin/clang-linker-wrapper+0xfe35)
clang: error: unable to execute command: Aborted

When the stack protector is disabled (`-fno-stack-protector`) the compilation succeeds and the compiled program runs normally :

$ /usr/lib/llvm/17/bin/clang -fno-stack-protector --rocm-device-lib-path=/usr/lib/amdgcn/bitcode -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa -x c openmp_offload.c
$ ./a.out
number of devices: 1
ran on GPU 0: yes, 160 teams, limit of 256 threads

**openmp_offload.c** :

#include <stdio.h>
#include <omp.h>
int main(int argc, char const * argv[]) {
int ndev;
int nteams;
int thread_lmt;
int on_host;

ndev = omp_get_num_devices();
printf("number of devices: %d\n", ndev);
for (int i = 0; i < omp_get_num_devices(); ++i) {
    #pragma omp target teams \
                       device(i) \
                       map(from:on_host) map(from:nteams) \
                       map(from:thread_lmt)
    #pragma omp parallel
    #pragma omp master
    if (0 == omp_get_team_num()) {
        on_host = omp_is_initial_device();
        nteams = omp_get_num_teams();
        thread_lmt = omp_get_thread_limit();
    }
    printf("ran on GPU %d: %s, %d teams, "
           "limit of %d threads\n",
           i, on_host ? "no" : "yes", nteams, thread_lmt);
}
return 0;

}

</details>

[llvmbot]

@llvm/issue-subscribers-backend-amdgpu

When invoking LLVM 17.0.0 RC4 with OpenMP/AMDGPU offload **and** the stack protector being active, the linker fails with: ``` Assertion `(i >= FTy->getNumParams() || FTy->getParamType(i) == Args[i]->getType()) && "Calling a function with a bad signature!"' failed. ``` Command line and full stack dump : ``` $ /usr/lib/llvm/17/bin/clang --rocm-device-lib-path=/usr/lib/amdgcn/bitcode -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa -x c openmp_offload.c

clang-linker-wrapper: /var/tmp/portage/sys-devel/llvm-17.0.0_rc4/work/llvm/lib/IR/Instructions.cpp:652: void llvm::CallInst::init(llvm::FunctionType*, llvm::Value*, llvm::ArrayRef<llvm::Value*>, llvm::ArrayRef<llvm::OperandBundleDefT<llvm::Value*> >, const llvm::Twine&): Assertion `(i >= FTy->getNumParams() || FTy->getParamType(i) == Args[i]->getType()) && "Calling a function with a bad signature!"' failed.
PLEASE submit a bug report to https://github.com/llvm/llvm-project/issues/ and include the crash backtrace.
Stack dump:
0. Program arguments: /usr/lib/llvm/17/bin/clang-linker-wrapper --host-triple=x86_64-pc-linux-gnu --linker-path=/usr/bin/x86_64-pc-linux-gnu-ld.bfd -- -pie --hash-style=gnu --eh-frame-hdr -m elf_x86_64 -dynamic-linker /lib64/ld-linux-x86-64.so.2 -o a.out /usr/lib/gcc/x86_64-pc-linux-gnu/13/../../../../lib64/Scrt1.o /usr/lib/gcc/x86_64-pc-linux-gnu/13/../../../../lib64/crti.o /usr/lib/gcc/x86_64-pc-linux-gnu/13/crtbeginS.o -L/usr/lib/gcc/x86_64-pc-linux-gnu/13 -L/usr/lib/gcc/x86_64-pc-linux-gnu/13/../../../../lib64 -L/lib/../lib64 -L/usr/lib/../lib64 -L/usr/lib/gcc/x86_64-pc-linux-gnu/13/../../../../x86_64-pc-linux-gnu/lib -L/lib -L/usr/lib -z relro /tmp/openmp_offload-c3b2da.o -lomp -lomptarget -lomptarget.devicertl -L/usr/lib/llvm/17/lib64 -lgcc --as-needed -lgcc_s --no-as-needed -lpthread -lc -lgcc --as-needed -lgcc_s --no-as-needed /usr/lib/gcc/x86_64-pc-linux-gnu/13/crtendS.o /usr/lib/gcc/x86_64-pc-linux-gnu/13/../../../../lib64/crtn.o

1.  Running pass 'CallGraph Pass Manager' on module 'ld-temp.o'.
   

2.  Running pass 'Insert stack protectors' on function '@<!-- -->__omp_offloading_43_17b92b_main_l12_omp_outlined'
   

#0 0x00007fd91f47eaee llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0xf15aee)
#1 0x00007fd91f435a8b llvm::sys::RunSignalHandlers() (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0xecca8b)
#2 0x00007fd91f435c37 (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0xeccc37)
#3 0x00007fd91dfd44a0 (/lib64/libc.so.6+0x3a4a0)
#4 0x00007fd91e0258ec __pthread_kill_implementation /usr/src/debug/sys-libs/glibc-2.38-r1/glibc-2.38/nptl/pthread_kill.c:44:76
#5 0x00007fd91dfd43f6 gsignal /usr/src/debug/sys-libs/glibc-2.38-r1/glibc-2.38/signal/../sysdeps/posix/raise.c:27:6
#6 0x00007fd91dfbc8f5 abort /usr/src/debug/sys-libs/glibc-2.38-r1/glibc-2.38/stdlib/abort.c:81:7
#7 0x00007fd91dfbc819 _nl_load_domain /usr/src/debug/sys-libs/glibc-2.38-r1/glibc-2.38/intl/loadmsgcat.c:1177:9
#8 0x00007fd91dfccad6 (/lib64/libc.so.6+0x32ad6)
#9 0x00007fd91f616428 llvm::CallInst::init(llvm::FunctionType*, llvm::Value*, llvm::ArrayRef<llvm::Value*>, llvm::ArrayRef<llvm::OperandBundleDefT<llvm::Value*>>, llvm::Twine const&) (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x10ad428)
#10 0x00007fd92307822d (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x4b0f22d)
#11 0x00007fd91fbef7f0 llvm::StackProtector::InsertStackProtectors() (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x16867f0)
#12 0x00007fd91fbf0354 llvm::StackProtector::runOnFunction(llvm::Function&) (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x1687354)
#13 0x00007fd91f63fa0b llvm::FPPassManager::runOnFunction(llvm::Function&) (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x10d6a0b)
#14 0x00007fd920c9c862 (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x2733862)
#15 0x00007fd91f63edca llvm::legacy::PassManagerImpl::run(llvm::Module&) (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x10d5dca)
#16 0x00007fd920f65fbd (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x29fcfbd)
#17 0x00007fd920f670b8 llvm::lto::backend(llvm::lto::Config const&, std::function<llvm::Expected<std::unique_ptr<llvm::CachedFileStream, std::default_delete<llvm::CachedFileStream>>> (unsigned int, llvm::Twine const&)>, unsigned int, llvm::Module&, llvm::ModuleSummaryIndex&) (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x29fe0b8)
#18 0x00007fd920f67fb5 llvm::lto::LTO::runRegularLTO(std::function<llvm::Expected<std::unique_ptr<llvm::CachedFileStream, std::default_delete<llvm::CachedFileStream>>> (unsigned int, llvm::Twine const&)>) (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x29fefb5)
#19 0x00007fd920f685fa llvm::lto::LTO::run(std::function<llvm::Expected<std::unique_ptr<llvm::CachedFileStream, std::default_delete<llvm::CachedFileStream>>> (unsigned int, llvm::Twine const&)>, std::function<llvm::Expected<std::function<llvm::Expected<std::unique_ptr<llvm::CachedFileStream, std::default_delete<llvm::CachedFileStream>>> (unsigned int, llvm::Twine const&)>> (unsigned int, llvm::StringRef, llvm::Twine const&)>) (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x29ff5fa)
#20 0x0000557e27d8dc07 (/usr/lib/llvm/17/bin/clang-linker-wrapper+0x34c07)
#21 0x0000557e27d801a3 (/usr/lib/llvm/17/bin/clang-linker-wrapper+0x271a3)
#22 0x0000557e27d723eb (/usr/lib/llvm/17/bin/clang-linker-wrapper+0x193eb)
#23 0x0000557e27d67f29 (/usr/lib/llvm/17/bin/clang-linker-wrapper+0xef29)
#24 0x00007fd91dfbe31e __libc_start_call_main /usr/src/debug/sys-libs/glibc-2.38-r1/glibc-2.38/csu/../sysdeps/nptl/libc_start_call_main.h:74:3
#25 0x00007fd91dfbe3d9 call_init /usr/src/debug/sys-libs/glibc-2.38-r1/glibc-2.38/csu/../csu/libc-start.c:128:20
#26 0x00007fd91dfbe3d9 __libc_start_main /usr/src/debug/sys-libs/glibc-2.38-r1/glibc-2.38/csu/../csu/libc-start.c:347:5
#27 0x0000557e27d68e35 (/usr/lib/llvm/17/bin/clang-linker-wrapper+0xfe35)
#0 0x00007fd91f47eaee llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0xf15aee)
#1 0x00007fd91f435a8b llvm::sys::RunSignalHandlers() (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0xecca8b)
#2 0x00007fd91f435c37 (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0xeccc37)
#3 0x00007fd91dfd44a0 (/lib64/libc.so.6+0x3a4a0)
#4 0x00007fd91e0258ec __pthread_kill_implementation /usr/src/debug/sys-libs/glibc-2.38-r1/glibc-2.38/nptl/pthread_kill.c:44:76
#5 0x00007fd91dfd43f6 gsignal /usr/src/debug/sys-libs/glibc-2.38-r1/glibc-2.38/signal/../sysdeps/posix/raise.c:27:6
#6 0x00007fd91dfbc8f5 abort /usr/src/debug/sys-libs/glibc-2.38-r1/glibc-2.38/stdlib/abort.c:81:7
#7 0x00007fd91dfbc819 _nl_load_domain /usr/src/debug/sys-libs/glibc-2.38-r1/glibc-2.38/intl/loadmsgcat.c:1177:9
#8 0x00007fd91dfccad6 (/lib64/libc.so.6+0x32ad6)
#9 0x00007fd91f616428 llvm::CallInst::init(llvm::FunctionType*, llvm::Value*, llvm::ArrayRef<llvm::Value*>, llvm::ArrayRef<llvm::OperandBundleDefT<llvm::Value*>>, llvm::Twine const&) (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x10ad428)
#10 0x00007fd92307822d (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x4b0f22d)
#11 0x00007fd91fbef7f0 llvm::StackProtector::InsertStackProtectors() (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x16867f0)
#12 0x00007fd91fbf0354 llvm::StackProtector::runOnFunction(llvm::Function&) (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x1687354)
#13 0x00007fd91f63fa0b llvm::FPPassManager::runOnFunction(llvm::Function&) (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x10d6a0b)
#14 0x00007fd920c9c862 (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x2733862)
#15 0x00007fd91f63edca llvm::legacy::PassManagerImpl::run(llvm::Module&) (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x10d5dca)
#16 0x00007fd920f65fbd (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x29fcfbd)
#17 0x00007fd920f670b8 llvm::lto::backend(llvm::lto::Config const&, std::function<llvm::Expected<std::unique_ptr<llvm::CachedFileStream, std::default_delete<llvm::CachedFileStream>>> (unsigned int, llvm::Twine const&)>, unsigned int, llvm::Module&, llvm::ModuleSummaryIndex&) (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x29fe0b8)
#18 0x00007fd920f67fb5 llvm::lto::LTO::runRegularLTO(std::function<llvm::Expected<std::unique_ptr<llvm::CachedFileStream, std::default_delete<llvm::CachedFileStream>>> (unsigned int, llvm::Twine const&)>) (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x29fefb5)
#19 0x00007fd920f685fa llvm::lto::LTO::run(std::function<llvm::Expected<std::unique_ptr<llvm::CachedFileStream, std::default_delete<llvm::CachedFileStream>>> (unsigned int, llvm::Twine const&)>, std::function<llvm::Expected<std::function<llvm::Expected<std::unique_ptr<llvm::CachedFileStream, std::default_delete<llvm::CachedFileStream>>> (unsigned int, llvm::Twine const&)>> (unsigned int, llvm::StringRef, llvm::Twine const&)>) (/usr/lib/llvm/17/bin/../lib64/libLLVM-17.so+0x29ff5fa)
#20 0x0000557e27d8dc07 (/usr/lib/llvm/17/bin/clang-linker-wrapper+0x34c07)
#21 0x0000557e27d801a3 (/usr/lib/llvm/17/bin/clang-linker-wrapper+0x271a3)
#22 0x0000557e27d723eb (/usr/lib/llvm/17/bin/clang-linker-wrapper+0x193eb)
#23 0x0000557e27d67f29 (/usr/lib/llvm/17/bin/clang-linker-wrapper+0xef29)
#24 0x00007fd91dfbe31e __libc_start_call_main /usr/src/debug/sys-libs/glibc-2.38-r1/glibc-2.38/csu/../sysdeps/nptl/libc_start_call_main.h:74:3
#25 0x00007fd91dfbe3d9 call_init /usr/src/debug/sys-libs/glibc-2.38-r1/glibc-2.38/csu/../csu/libc-start.c:128:20
#26 0x00007fd91dfbe3d9 __libc_start_main /usr/src/debug/sys-libs/glibc-2.38-r1/glibc-2.38/csu/../csu/libc-start.c:347:5
#27 0x0000557e27d68e35 (/usr/lib/llvm/17/bin/clang-linker-wrapper+0xfe35)
clang: error: unable to execute command: Aborted

When the stack protector is disabled (`-fno-stack-protector`) the compilation succeeds and the compiled program runs normally :

$ /usr/lib/llvm/17/bin/clang -fno-stack-protector --rocm-device-lib-path=/usr/lib/amdgcn/bitcode -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa -x c openmp_offload.c
$ ./a.out
number of devices: 1
ran on GPU 0: yes, 160 teams, limit of 256 threads

**openmp_offload.c** :

#include <stdio.h>
#include <omp.h>
int main(int argc, char const * argv[]) {
int ndev;
int nteams;
int thread_lmt;
int on_host;

ndev = omp_get_num_devices();
printf("number of devices: %d\n", ndev);
for (int i = 0; i < omp_get_num_devices(); ++i) {
    #pragma omp target teams \
                       device(i) \
                       map(from:on_host) map(from:nteams) \
                       map(from:thread_lmt)
    #pragma omp parallel
    #pragma omp master
    if (0 == omp_get_team_num()) {
        on_host = omp_is_initial_device();
        nteams = omp_get_num_teams();
        thread_lmt = omp_get_thread_limit();
    }
    printf("ran on GPU %d: %s, %d teams, "
           "limit of %d threads\n",
           i, on_host ? "no" : "yes", nteams, thread_lmt);
}
return 0;

}

</details>

[shiltian]

Does GPU even support stack protector?

[jhuber6]

Yeah, I'm unsure why we're getting the stack protector instantiated here. Maybe it's a system default or something from the configuration file? Worst case scenario you can do -Xarch_device -fno-stack-protector to only change this for the GPU compilation. I don't think we handle any of this stuff on the GPU.

[shiltian]

Can we close this one for now?

[admnd]

The bug is still here (clang/llvm 17.0.3) and as the stack protector is enabled by default the issue will still be triggered. So it should be fixed in a way or another (turn off the stack protector and put warning in that case, etc) and/or explicitly document the fact that GPU offloading requires the stack protector to be disabled depending on how you fix/work the issue around.

Closing the bug will not magically fix the issue ;) You guys know the software more than me, there is probably a quick-win acceptable solution. I know this issue is not probably on your high priority list, but please leave it open as long as the problem persists.

[shiltian]

@jhuber6 Probably we need something in the driver to check if the stack protector is enabled, don't pass it to device compilation? Or even the worst "fix" is to emit an error saying this is not supported on offloading?

[jhuber6]

@jhuber6 Probably we need something in the driver to check if the stack protector is enabled, don't pass it to device compilation? Or even the worst "fix" is to emit an error saying this is not supported on offloading?

I'd need to look into it. The solution is to either make the clang driver ignore the option, or to make the clang compiler not generate the code for the stack protectors on architectures where it doesn't make sense.