search

llvm / llvm-project

The LLVM Project is a collection of modular and reusable compiler and toolchain technologies.
http://llvm.org
Other
28.26k stars 11.67k forks source link

[OpenMP] Dynamic schedule compilation fail #70069

Open doru1004 opened 11 months ago

doru1004 commented 11 months ago

After this commit: https://github.com/llvm/llvm-project/commit/d3921e467005733daa8e63df4b553052c8ec72c1

The following example fails to compile:

#include <omp.h>
#include <stdio.h>
#include <stdlib.h>

#define MAX_N 25000

void reset_input(double *a, double *a_h, double *b, double *c) {
  for(int i = 0 ; i < MAX_N ; i++) {
    a[i] = a_h[i] = i;
    b[i] = i*2;
    c[i] = i-3;
  }
}

int main(int argc, char *argv[]) {
  double * a = (double *) malloc(MAX_N * sizeof(double));
  double * a_h = (double *) malloc(MAX_N * sizeof(double));
  double * d = (double *) malloc(MAX_N * sizeof(double));
  double * d_h = (double *) malloc(MAX_N * sizeof(double));
  double * b = (double *) malloc(MAX_N * sizeof(double));
  double * c = (double *) malloc(MAX_N * sizeof(double));

#pragma omp target enter data map(to:a[:MAX_N],b[:MAX_N],c[:MAX_N],d[:MAX_N])
  int n = 32;
    reset_input(a, a_h, b, c);

#pragma omp target update to(a[:n],b[:n],c[:n])

#pragma omp target teams distribute parallel for schedule(dynamic)
      for (int i = 0; i < n; ++i) {
        a[i] += b[i] + c[i];
      }

      for (int i = 0; i < n; ++i)
        a_h[i] += b[i] + c[i];

#pragma omp target update from(a[:n])

    for (int i = 0; i < n; ++i) {
      if (a_h[i] != a[i]) {
        printf("Error at n = %d, i = %d: host = %lf, device = %lf\n", n, i, a_h[i], a[i]);
        return 1;
      }
    }
  printf("Succeeded\n");

  #pragma omp target exit data map(delete:a[:MAX_N],b[:MAX_N],c[:MAX_N],d[:MAX_N])

  return 0;
}

Compiled with:

clang++ -std=c++11 -fopenmp-targets=amdgcn-amd-amdhsa -fopenmp -O3 -fno-exceptions test.cpp -o test

The compilation error is:

 "/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper" --opt-level=O3 --host-triple=x86_64-unknown-linux-gnu --wrapper-verbose --save-temps --linker-path=/home/dobercea/rocm/trunk_1.0/bin/ld.lld -- -pie -z relro --hash-style=gnu --eh-frame-hdr -m elf_x86_64 -dynamic-linker /lib64/ld-linux-x86-64.so.2 -o test /lib/x86_64-linux-gnu/Scrt1.o /lib/x86_64-linux-gnu/crti.o /usr/lib/gcc/x86_64-linux-gnu/9/crtbeginS.o -L/usr/lib/gcc/x86_64-linux-gnu/9 -L/usr/lib/gcc/x86_64-linux-gnu/9/../../../../lib64 -L/lib/x86_64-linux-gnu -L/lib/../lib64 -L/usr/lib/x86_64-linux-gnu -L/usr/lib/../lib64 -L/lib -L/usr/lib test-host-x86_64-unknown-linux-gnu.o -lstdc++ -lm -lomp -lomptarget -lomptarget.devicertl -L/home/dobercea/rocm/trunk_1.0/lib -lgcc_s -lgcc -lpthread -lc -lgcc_s -lgcc /usr/lib/gcc/x86_64-linux-gnu/9/crtendS.o /lib/x86_64-linux-gnu/crtn.o
LLVM ERROR: Cannot select: t17: i32 = GlobalAddress<ptr addrspace(5) @_ZL12ThreadDSTPtr> 0
In function: __omp_offloading_fd00_28c3b6a_main_l34
PLEASE submit a bug report to https://github.com/llvm/llvm-project/issues/ and include the crash backtrace.
Stack dump:
0.      Program arguments: /home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper --opt-level=O3 --host-triple=x86_64-unknown-linux-gnu --wrapper-verbose --save-temps --linker-path=/home/dobercea/rocm/trunk_1.0/bin/ld.lld -- -pie -z relro --hash-style=gnu --eh-frame-hdr -m elf_x86_64 -dynamic-linker /lib64/ld-linux-x86-64.so.2 -o test /lib/x86_64-linux-gnu/Scrt1.o /lib/x86_64-linux-gnu/crti.o /usr/lib/gcc/x86_64-linux-gnu/9/crtbeginS.o -L/usr/lib/gcc/x86_64-linux-gnu/9 -L/usr/lib/gcc/x86_64-linux-gnu/9/../../../../lib64 -L/lib/x86_64-linux-gnu -L/lib/../lib64 -L/usr/lib/x86_64-linux-gnu -L/usr/lib/../lib64 -L/lib -L/usr/lib test-host-x86_64-unknown-linux-gnu.o -lstdc++ -lm -lomp -lomptarget -lomptarget.devicertl -L/home/dobercea/rocm/trunk_1.0/lib -lgcc_s -lgcc -lpthread -lc -lgcc_s -lgcc /usr/lib/gcc/x86_64-linux-gnu/9/crtendS.o /lib/x86_64-linux-gnu/crtn.o
1.      Running pass 'CallGraph Pass Manager' on module 'ld-temp.o'.
2.      Running pass 'AMDGPU DAG->DAG Pattern Instruction Selection' on function '@__omp_offloading_fd00_28c3b6a_main_l34'
 #0 0x0000559b76f12fe4 PrintStackTraceSignalHandler(void*) Signals.cpp:0:0
 #1 0x0000559b76f10814 SignalHandler(int) Signals.cpp:0:0
 #2 0x00007f74d189c420 __restore_rt (/lib/x86_64-linux-gnu/libpthread.so.0+0x14420)
 #3 0x00007f74d133900b raise /build/glibc-SzIz7B/glibc-2.31/signal/../sysdeps/unix/sysv/linux/raise.c:51:1
 #4 0x00007f74d1318859 abort /build/glibc-SzIz7B/glibc-2.31/stdlib/abort.c:81:7
 #5 0x0000559b75d54bd8 llvm::ConvertUTF8toUTF32(unsigned char const**, unsigned char const*, unsigned int**, unsigned int*, llvm::ConversionFlags) (.cold) ConvertUTF.cpp:0:0
 #6 0x0000559b7789ef5d llvm::SelectionDAGISel::CannotYetSelect(llvm::SDNode*) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x1dccf5d)
 #7 0x0000559b778a169a llvm::SelectionDAGISel::SelectCodeCommon(llvm::SDNode*, unsigned char const*, unsigned int) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x1dcf69a)
 #8 0x0000559b764bf877 AMDGPUDAGToDAGISel::Select(llvm::SDNode*) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x9ed877)
 #9 0x0000559b7789be90 llvm::SelectionDAGISel::DoInstructionSelection() (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x1dc9e90)
#10 0x0000559b778a9095 llvm::SelectionDAGISel::CodeGenAndEmitDAG() (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x1dd7095)
#11 0x0000559b778ac438 llvm::SelectionDAGISel::SelectAllBasicBlocks(llvm::Function const&) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x1dda438)
#12 0x0000559b778ae102 llvm::SelectionDAGISel::runOnMachineFunction(llvm::MachineFunction&) (.part.0) SelectionDAGISel.cpp:0:0
#13 0x0000559b764c8ad9 AMDGPUDAGToDAGISel::runOnMachineFunction(llvm::MachineFunction&) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x9f6ad9)
#14 0x0000559b7701dde1 llvm::MachineFunctionPass::runOnFunction(llvm::Function&) (.part.0) MachineFunctionPass.cpp:0:0
#15 0x0000559b7688f4c1 llvm::FPPassManager::runOnFunction(llvm::Function&) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0xdbd4c1)
#16 0x0000559b76b22cf7 (anonymous namespace)::CGPassManager::runOnModule(llvm::Module&) CallGraphSCCPass.cpp:0:0
#17 0x0000559b7688ff92 llvm::legacy::PassManagerImpl::run(llvm::Module&) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0xdbdf92)
#18 0x0000559b774d5b15 codegen(llvm::lto::Config const&, llvm::TargetMachine*, 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 const&) LTOBackend.cpp:0:0
#19 0x0000559b774d60ed 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&) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x1a040ed)
#20 0x0000559b774cc6fc llvm::lto::LTO::runRegularLTO(std::function<llvm::Expected<std::unique_ptr<llvm::CachedFileStream, std::default_delete<llvm::CachedFileStream>>> (unsigned int, llvm::Twine const&)>) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x19fa6fc)
#21 0x0000559b774ccda8 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&)>) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x19fada8)
#22 0x0000559b75e06f3d (anonymous namespace)::linkBitcodeFiles(llvm::SmallVectorImpl<llvm::object::OffloadFile>&, llvm::SmallVectorImpl<llvm::StringRef>&, llvm::opt::ArgList const&) (.constprop.0) ClangLinkerWrapper.cpp:0:0
#23 0x0000559b75e0dc6a llvm::Error (anonymous namespace)::linkAndWrapDeviceFiles(llvm::SmallVectorImpl<llvm::object::OffloadFile>&, llvm::opt::InputArgList const&, char**, int)::'lambda'(auto&)::operator()<llvm::SmallVector<llvm::object::OffloadFile, 3u>>(auto&) const ClangLinkerWrapper.cpp:0:0
#24 0x0000559b75e14215 (anonymous namespace)::linkAndWrapDeviceFiles(llvm::SmallVectorImpl<llvm::object::OffloadFile>&, llvm::opt::InputArgList const&, char**, int) ClangLinkerWrapper.cpp:0:0
#25 0x0000559b75d5a746 main (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x288746)
#26 0x00007f74d131a083 __libc_start_main /build/glibc-SzIz7B/glibc-2.31/csu/../csu/libc-start.c:342:3
#27 0x0000559b75df66de _start (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x3246de)
 #0 0x0000559b76f12fe4 PrintStackTraceSignalHandler(void*) Signals.cpp:0:0
 #1 0x0000559b76f10814 SignalHandler(int) Signals.cpp:0:0
 #2 0x00007f74d189c420 __restore_rt (/lib/x86_64-linux-gnu/libpthread.so.0+0x14420)
 #3 0x00007f74d133900b raise /build/glibc-SzIz7B/glibc-2.31/signal/../sysdeps/unix/sysv/linux/raise.c:51:1
 #4 0x00007f74d1318859 abort /build/glibc-SzIz7B/glibc-2.31/stdlib/abort.c:81:7
 #5 0x0000559b75d54bd8 llvm::ConvertUTF8toUTF32(unsigned char const**, unsigned char const*, unsigned int**, unsigned int*, llvm::ConversionFlags) (.cold) ConvertUTF.cpp:0:0
 #6 0x0000559b7789ef5d llvm::SelectionDAGISel::CannotYetSelect(llvm::SDNode*) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x1dccf5d)
 #7 0x0000559b778a169a llvm::SelectionDAGISel::SelectCodeCommon(llvm::SDNode*, unsigned char const*, unsigned int) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x1dcf69a)
 #8 0x0000559b764bf877 AMDGPUDAGToDAGISel::Select(llvm::SDNode*) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x9ed877)
 #9 0x0000559b7789be90 llvm::SelectionDAGISel::DoInstructionSelection() (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x1dc9e90)
#10 0x0000559b778a9095 llvm::SelectionDAGISel::CodeGenAndEmitDAG() (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x1dd7095)
#11 0x0000559b778ac438 llvm::SelectionDAGISel::SelectAllBasicBlocks(llvm::Function const&) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x1dda438)
#12 0x0000559b778ae102 llvm::SelectionDAGISel::runOnMachineFunction(llvm::MachineFunction&) (.part.0) SelectionDAGISel.cpp:0:0
#13 0x0000559b764c8ad9 AMDGPUDAGToDAGISel::runOnMachineFunction(llvm::MachineFunction&) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x9f6ad9)
#14 0x0000559b7701dde1 llvm::MachineFunctionPass::runOnFunction(llvm::Function&) (.part.0) MachineFunctionPass.cpp:0:0
#15 0x0000559b7688f4c1 llvm::FPPassManager::runOnFunction(llvm::Function&) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0xdbd4c1)
#16 0x0000559b76b22cf7 (anonymous namespace)::CGPassManager::runOnModule(llvm::Module&) CallGraphSCCPass.cpp:0:0
#17 0x0000559b7688ff92 llvm::legacy::PassManagerImpl::run(llvm::Module&) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0xdbdf92)
#18 0x0000559b774d5b15 codegen(llvm::lto::Config const&, llvm::TargetMachine*, 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 const&) LTOBackend.cpp:0:0
#19 0x0000559b774d60ed 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&) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x1a040ed)
#20 0x0000559b774cc6fc llvm::lto::LTO::runRegularLTO(std::function<llvm::Expected<std::unique_ptr<llvm::CachedFileStream, std::default_delete<llvm::CachedFileStream>>> (unsigned int, llvm::Twine const&)>) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x19fa6fc)
#21 0x0000559b774ccda8 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&)>) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x19fada8)
#22 0x0000559b75e06f3d (anonymous namespace)::linkBitcodeFiles(llvm::SmallVectorImpl<llvm::object::OffloadFile>&, llvm::SmallVectorImpl<llvm::StringRef>&, llvm::opt::ArgList const&) (.constprop.0) ClangLinkerWrapper.cpp:0:0
#23 0x0000559b75e0dc6a llvm::Error (anonymous namespace)::linkAndWrapDeviceFiles(llvm::SmallVectorImpl<llvm::object::OffloadFile>&, llvm::opt::InputArgList const&, char**, int)::'lambda'(auto&)::operator()<llvm::SmallVector<llvm::object::OffloadFile, 3u>>(auto&) const ClangLinkerWrapper.cpp:0:0
#24 0x0000559b75e14215 (anonymous namespace)::linkAndWrapDeviceFiles(llvm::SmallVectorImpl<llvm::object::OffloadFile>&, llvm::opt::InputArgList const&, char**, int) ClangLinkerWrapper.cpp:0:0
#25 0x0000559b75d5a746 main (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x288746)
#26 0x00007f74d131a083 __libc_start_main /build/glibc-SzIz7B/glibc-2.31/csu/../csu/libc-start.c:342:3
#27 0x0000559b75df66de _start (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x3246de)
llvmbot commented 11 months ago

@llvm/issue-subscribers-backend-amdgpu

Author: Gheorghe-Teodor Bercea (doru1004)

After this commit: https://github.com/llvm/llvm-project/commit/d3921e467005733daa8e63df4b553052c8ec72c1 The following example fails to compile: ``` #include <omp.h> #include <stdio.h> #include <stdlib.h> #define MAX_N 25000 void reset_input(double *a, double *a_h, double *b, double *c) { for(int i = 0 ; i < MAX_N ; i++) { a[i] = a_h[i] = i; b[i] = i*2; c[i] = i-3; } } int main(int argc, char *argv[]) { double * a = (double *) malloc(MAX_N * sizeof(double)); double * a_h = (double *) malloc(MAX_N * sizeof(double)); double * d = (double *) malloc(MAX_N * sizeof(double)); double * d_h = (double *) malloc(MAX_N * sizeof(double)); double * b = (double *) malloc(MAX_N * sizeof(double)); double * c = (double *) malloc(MAX_N * sizeof(double)); #pragma omp target enter data map(to:a[:MAX_N],b[:MAX_N],c[:MAX_N],d[:MAX_N]) int n = 32; reset_input(a, a_h, b, c); #pragma omp target update to(a[:n],b[:n],c[:n]) #pragma omp target teams distribute parallel for schedule(dynamic) for (int i = 0; i < n; ++i) { a[i] += b[i] + c[i]; } for (int i = 0; i < n; ++i) a_h[i] += b[i] + c[i]; #pragma omp target update from(a[:n]) for (int i = 0; i < n; ++i) { if (a_h[i] != a[i]) { printf("Error at n = %d, i = %d: host = %lf, device = %lf\n", n, i, a_h[i], a[i]); return 1; } } printf("Succeeded\n"); #pragma omp target exit data map(delete:a[:MAX_N],b[:MAX_N],c[:MAX_N],d[:MAX_N]) return 0; } ``` Compiled with: ``` clang++ -std=c++11 -fopenmp-targets=amdgcn-amd-amdhsa -fopenmp -O3 -fno-exceptions test.cpp -o test ``` The compilation error is: ``` "/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper" --opt-level=O3 --host-triple=x86_64-unknown-linux-gnu --wrapper-verbose --save-temps --linker-path=/home/dobercea/rocm/trunk_1.0/bin/ld.lld -- -pie -z relro --hash-style=gnu --eh-frame-hdr -m elf_x86_64 -dynamic-linker /lib64/ld-linux-x86-64.so.2 -o test /lib/x86_64-linux-gnu/Scrt1.o /lib/x86_64-linux-gnu/crti.o /usr/lib/gcc/x86_64-linux-gnu/9/crtbeginS.o -L/usr/lib/gcc/x86_64-linux-gnu/9 -L/usr/lib/gcc/x86_64-linux-gnu/9/../../../../lib64 -L/lib/x86_64-linux-gnu -L/lib/../lib64 -L/usr/lib/x86_64-linux-gnu -L/usr/lib/../lib64 -L/lib -L/usr/lib test-host-x86_64-unknown-linux-gnu.o -lstdc++ -lm -lomp -lomptarget -lomptarget.devicertl -L/home/dobercea/rocm/trunk_1.0/lib -lgcc_s -lgcc -lpthread -lc -lgcc_s -lgcc /usr/lib/gcc/x86_64-linux-gnu/9/crtendS.o /lib/x86_64-linux-gnu/crtn.o LLVM ERROR: Cannot select: t17: i32 = GlobalAddress<ptr addrspace(5) @_ZL12ThreadDSTPtr> 0 In function: __omp_offloading_fd00_28c3b6a_main_l34 PLEASE submit a bug report to https://github.com/llvm/llvm-project/issues/ and include the crash backtrace. Stack dump: 0. Program arguments: /home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper --opt-level=O3 --host-triple=x86_64-unknown-linux-gnu --wrapper-verbose --save-temps --linker-path=/home/dobercea/rocm/trunk_1.0/bin/ld.lld -- -pie -z relro --hash-style=gnu --eh-frame-hdr -m elf_x86_64 -dynamic-linker /lib64/ld-linux-x86-64.so.2 -o test /lib/x86_64-linux-gnu/Scrt1.o /lib/x86_64-linux-gnu/crti.o /usr/lib/gcc/x86_64-linux-gnu/9/crtbeginS.o -L/usr/lib/gcc/x86_64-linux-gnu/9 -L/usr/lib/gcc/x86_64-linux-gnu/9/../../../../lib64 -L/lib/x86_64-linux-gnu -L/lib/../lib64 -L/usr/lib/x86_64-linux-gnu -L/usr/lib/../lib64 -L/lib -L/usr/lib test-host-x86_64-unknown-linux-gnu.o -lstdc++ -lm -lomp -lomptarget -lomptarget.devicertl -L/home/dobercea/rocm/trunk_1.0/lib -lgcc_s -lgcc -lpthread -lc -lgcc_s -lgcc /usr/lib/gcc/x86_64-linux-gnu/9/crtendS.o /lib/x86_64-linux-gnu/crtn.o 1. Running pass 'CallGraph Pass Manager' on module 'ld-temp.o'. 2. Running pass 'AMDGPU DAG->DAG Pattern Instruction Selection' on function '@__omp_offloading_fd00_28c3b6a_main_l34' #0 0x0000559b76f12fe4 PrintStackTraceSignalHandler(void*) Signals.cpp:0:0 #1 0x0000559b76f10814 SignalHandler(int) Signals.cpp:0:0 #2 0x00007f74d189c420 __restore_rt (/lib/x86_64-linux-gnu/libpthread.so.0+0x14420) #3 0x00007f74d133900b raise /build/glibc-SzIz7B/glibc-2.31/signal/../sysdeps/unix/sysv/linux/raise.c:51:1 #4 0x00007f74d1318859 abort /build/glibc-SzIz7B/glibc-2.31/stdlib/abort.c:81:7 #5 0x0000559b75d54bd8 llvm::ConvertUTF8toUTF32(unsigned char const**, unsigned char const*, unsigned int**, unsigned int*, llvm::ConversionFlags) (.cold) ConvertUTF.cpp:0:0 #6 0x0000559b7789ef5d llvm::SelectionDAGISel::CannotYetSelect(llvm::SDNode*) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x1dccf5d) #7 0x0000559b778a169a llvm::SelectionDAGISel::SelectCodeCommon(llvm::SDNode*, unsigned char const*, unsigned int) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x1dcf69a) #8 0x0000559b764bf877 AMDGPUDAGToDAGISel::Select(llvm::SDNode*) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x9ed877) #9 0x0000559b7789be90 llvm::SelectionDAGISel::DoInstructionSelection() (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x1dc9e90) #10 0x0000559b778a9095 llvm::SelectionDAGISel::CodeGenAndEmitDAG() (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x1dd7095) #11 0x0000559b778ac438 llvm::SelectionDAGISel::SelectAllBasicBlocks(llvm::Function const&) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x1dda438) #12 0x0000559b778ae102 llvm::SelectionDAGISel::runOnMachineFunction(llvm::MachineFunction&) (.part.0) SelectionDAGISel.cpp:0:0 #13 0x0000559b764c8ad9 AMDGPUDAGToDAGISel::runOnMachineFunction(llvm::MachineFunction&) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x9f6ad9) #14 0x0000559b7701dde1 llvm::MachineFunctionPass::runOnFunction(llvm::Function&) (.part.0) MachineFunctionPass.cpp:0:0 #15 0x0000559b7688f4c1 llvm::FPPassManager::runOnFunction(llvm::Function&) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0xdbd4c1) #16 0x0000559b76b22cf7 (anonymous namespace)::CGPassManager::runOnModule(llvm::Module&) CallGraphSCCPass.cpp:0:0 #17 0x0000559b7688ff92 llvm::legacy::PassManagerImpl::run(llvm::Module&) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0xdbdf92) #18 0x0000559b774d5b15 codegen(llvm::lto::Config const&, llvm::TargetMachine*, 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 const&) LTOBackend.cpp:0:0 #19 0x0000559b774d60ed 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&) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x1a040ed) #20 0x0000559b774cc6fc llvm::lto::LTO::runRegularLTO(std::function<llvm::Expected<std::unique_ptr<llvm::CachedFileStream, std::default_delete<llvm::CachedFileStream>>> (unsigned int, llvm::Twine const&)>) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x19fa6fc) #21 0x0000559b774ccda8 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&)>) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x19fada8) #22 0x0000559b75e06f3d (anonymous namespace)::linkBitcodeFiles(llvm::SmallVectorImpl<llvm::object::OffloadFile>&, llvm::SmallVectorImpl<llvm::StringRef>&, llvm::opt::ArgList const&) (.constprop.0) ClangLinkerWrapper.cpp:0:0 #23 0x0000559b75e0dc6a llvm::Error (anonymous namespace)::linkAndWrapDeviceFiles(llvm::SmallVectorImpl<llvm::object::OffloadFile>&, llvm::opt::InputArgList const&, char**, int)::'lambda'(auto&)::operator()<llvm::SmallVector<llvm::object::OffloadFile, 3u>>(auto&) const ClangLinkerWrapper.cpp:0:0 #24 0x0000559b75e14215 (anonymous namespace)::linkAndWrapDeviceFiles(llvm::SmallVectorImpl<llvm::object::OffloadFile>&, llvm::opt::InputArgList const&, char**, int) ClangLinkerWrapper.cpp:0:0 #25 0x0000559b75d5a746 main (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x288746) #26 0x00007f74d131a083 __libc_start_main /build/glibc-SzIz7B/glibc-2.31/csu/../csu/libc-start.c:342:3 #27 0x0000559b75df66de _start (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x3246de) #0 0x0000559b76f12fe4 PrintStackTraceSignalHandler(void*) Signals.cpp:0:0 #1 0x0000559b76f10814 SignalHandler(int) Signals.cpp:0:0 #2 0x00007f74d189c420 __restore_rt (/lib/x86_64-linux-gnu/libpthread.so.0+0x14420) #3 0x00007f74d133900b raise /build/glibc-SzIz7B/glibc-2.31/signal/../sysdeps/unix/sysv/linux/raise.c:51:1 #4 0x00007f74d1318859 abort /build/glibc-SzIz7B/glibc-2.31/stdlib/abort.c:81:7 #5 0x0000559b75d54bd8 llvm::ConvertUTF8toUTF32(unsigned char const**, unsigned char const*, unsigned int**, unsigned int*, llvm::ConversionFlags) (.cold) ConvertUTF.cpp:0:0 #6 0x0000559b7789ef5d llvm::SelectionDAGISel::CannotYetSelect(llvm::SDNode*) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x1dccf5d) #7 0x0000559b778a169a llvm::SelectionDAGISel::SelectCodeCommon(llvm::SDNode*, unsigned char const*, unsigned int) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x1dcf69a) #8 0x0000559b764bf877 AMDGPUDAGToDAGISel::Select(llvm::SDNode*) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x9ed877) #9 0x0000559b7789be90 llvm::SelectionDAGISel::DoInstructionSelection() (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x1dc9e90) #10 0x0000559b778a9095 llvm::SelectionDAGISel::CodeGenAndEmitDAG() (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x1dd7095) #11 0x0000559b778ac438 llvm::SelectionDAGISel::SelectAllBasicBlocks(llvm::Function const&) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x1dda438) #12 0x0000559b778ae102 llvm::SelectionDAGISel::runOnMachineFunction(llvm::MachineFunction&) (.part.0) SelectionDAGISel.cpp:0:0 #13 0x0000559b764c8ad9 AMDGPUDAGToDAGISel::runOnMachineFunction(llvm::MachineFunction&) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x9f6ad9) #14 0x0000559b7701dde1 llvm::MachineFunctionPass::runOnFunction(llvm::Function&) (.part.0) MachineFunctionPass.cpp:0:0 #15 0x0000559b7688f4c1 llvm::FPPassManager::runOnFunction(llvm::Function&) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0xdbd4c1) #16 0x0000559b76b22cf7 (anonymous namespace)::CGPassManager::runOnModule(llvm::Module&) CallGraphSCCPass.cpp:0:0 #17 0x0000559b7688ff92 llvm::legacy::PassManagerImpl::run(llvm::Module&) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0xdbdf92) #18 0x0000559b774d5b15 codegen(llvm::lto::Config const&, llvm::TargetMachine*, 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 const&) LTOBackend.cpp:0:0 #19 0x0000559b774d60ed 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&) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x1a040ed) #20 0x0000559b774cc6fc llvm::lto::LTO::runRegularLTO(std::function<llvm::Expected<std::unique_ptr<llvm::CachedFileStream, std::default_delete<llvm::CachedFileStream>>> (unsigned int, llvm::Twine const&)>) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x19fa6fc) #21 0x0000559b774ccda8 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&)>) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x19fada8) #22 0x0000559b75e06f3d (anonymous namespace)::linkBitcodeFiles(llvm::SmallVectorImpl<llvm::object::OffloadFile>&, llvm::SmallVectorImpl<llvm::StringRef>&, llvm::opt::ArgList const&) (.constprop.0) ClangLinkerWrapper.cpp:0:0 #23 0x0000559b75e0dc6a llvm::Error (anonymous namespace)::linkAndWrapDeviceFiles(llvm::SmallVectorImpl<llvm::object::OffloadFile>&, llvm::opt::InputArgList const&, char**, int)::'lambda'(auto&)::operator()<llvm::SmallVector<llvm::object::OffloadFile, 3u>>(auto&) const ClangLinkerWrapper.cpp:0:0 #24 0x0000559b75e14215 (anonymous namespace)::linkAndWrapDeviceFiles(llvm::SmallVectorImpl<llvm::object::OffloadFile>&, llvm::opt::InputArgList const&, char**, int) ClangLinkerWrapper.cpp:0:0 #25 0x0000559b75d5a746 main (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x288746) #26 0x00007f74d131a083 __libc_start_main /build/glibc-SzIz7B/glibc-2.31/csu/../csu/libc-start.c:342:3 #27 0x0000559b75df66de _start (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x3246de) ```
llvmbot commented 11 months ago

@llvm/issue-subscribers-openmp

Author: Gheorghe-Teodor Bercea (doru1004)

After this commit: https://github.com/llvm/llvm-project/commit/d3921e467005733daa8e63df4b553052c8ec72c1 The following example fails to compile: ``` #include <omp.h> #include <stdio.h> #include <stdlib.h> #define MAX_N 25000 void reset_input(double *a, double *a_h, double *b, double *c) { for(int i = 0 ; i < MAX_N ; i++) { a[i] = a_h[i] = i; b[i] = i*2; c[i] = i-3; } } int main(int argc, char *argv[]) { double * a = (double *) malloc(MAX_N * sizeof(double)); double * a_h = (double *) malloc(MAX_N * sizeof(double)); double * d = (double *) malloc(MAX_N * sizeof(double)); double * d_h = (double *) malloc(MAX_N * sizeof(double)); double * b = (double *) malloc(MAX_N * sizeof(double)); double * c = (double *) malloc(MAX_N * sizeof(double)); #pragma omp target enter data map(to:a[:MAX_N],b[:MAX_N],c[:MAX_N],d[:MAX_N]) int n = 32; reset_input(a, a_h, b, c); #pragma omp target update to(a[:n],b[:n],c[:n]) #pragma omp target teams distribute parallel for schedule(dynamic) for (int i = 0; i < n; ++i) { a[i] += b[i] + c[i]; } for (int i = 0; i < n; ++i) a_h[i] += b[i] + c[i]; #pragma omp target update from(a[:n]) for (int i = 0; i < n; ++i) { if (a_h[i] != a[i]) { printf("Error at n = %d, i = %d: host = %lf, device = %lf\n", n, i, a_h[i], a[i]); return 1; } } printf("Succeeded\n"); #pragma omp target exit data map(delete:a[:MAX_N],b[:MAX_N],c[:MAX_N],d[:MAX_N]) return 0; } ``` Compiled with: ``` clang++ -std=c++11 -fopenmp-targets=amdgcn-amd-amdhsa -fopenmp -O3 -fno-exceptions test.cpp -o test ``` The compilation error is: ``` "/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper" --opt-level=O3 --host-triple=x86_64-unknown-linux-gnu --wrapper-verbose --save-temps --linker-path=/home/dobercea/rocm/trunk_1.0/bin/ld.lld -- -pie -z relro --hash-style=gnu --eh-frame-hdr -m elf_x86_64 -dynamic-linker /lib64/ld-linux-x86-64.so.2 -o test /lib/x86_64-linux-gnu/Scrt1.o /lib/x86_64-linux-gnu/crti.o /usr/lib/gcc/x86_64-linux-gnu/9/crtbeginS.o -L/usr/lib/gcc/x86_64-linux-gnu/9 -L/usr/lib/gcc/x86_64-linux-gnu/9/../../../../lib64 -L/lib/x86_64-linux-gnu -L/lib/../lib64 -L/usr/lib/x86_64-linux-gnu -L/usr/lib/../lib64 -L/lib -L/usr/lib test-host-x86_64-unknown-linux-gnu.o -lstdc++ -lm -lomp -lomptarget -lomptarget.devicertl -L/home/dobercea/rocm/trunk_1.0/lib -lgcc_s -lgcc -lpthread -lc -lgcc_s -lgcc /usr/lib/gcc/x86_64-linux-gnu/9/crtendS.o /lib/x86_64-linux-gnu/crtn.o LLVM ERROR: Cannot select: t17: i32 = GlobalAddress<ptr addrspace(5) @_ZL12ThreadDSTPtr> 0 In function: __omp_offloading_fd00_28c3b6a_main_l34 PLEASE submit a bug report to https://github.com/llvm/llvm-project/issues/ and include the crash backtrace. Stack dump: 0. Program arguments: /home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper --opt-level=O3 --host-triple=x86_64-unknown-linux-gnu --wrapper-verbose --save-temps --linker-path=/home/dobercea/rocm/trunk_1.0/bin/ld.lld -- -pie -z relro --hash-style=gnu --eh-frame-hdr -m elf_x86_64 -dynamic-linker /lib64/ld-linux-x86-64.so.2 -o test /lib/x86_64-linux-gnu/Scrt1.o /lib/x86_64-linux-gnu/crti.o /usr/lib/gcc/x86_64-linux-gnu/9/crtbeginS.o -L/usr/lib/gcc/x86_64-linux-gnu/9 -L/usr/lib/gcc/x86_64-linux-gnu/9/../../../../lib64 -L/lib/x86_64-linux-gnu -L/lib/../lib64 -L/usr/lib/x86_64-linux-gnu -L/usr/lib/../lib64 -L/lib -L/usr/lib test-host-x86_64-unknown-linux-gnu.o -lstdc++ -lm -lomp -lomptarget -lomptarget.devicertl -L/home/dobercea/rocm/trunk_1.0/lib -lgcc_s -lgcc -lpthread -lc -lgcc_s -lgcc /usr/lib/gcc/x86_64-linux-gnu/9/crtendS.o /lib/x86_64-linux-gnu/crtn.o 1. Running pass 'CallGraph Pass Manager' on module 'ld-temp.o'. 2. Running pass 'AMDGPU DAG->DAG Pattern Instruction Selection' on function '@__omp_offloading_fd00_28c3b6a_main_l34' #0 0x0000559b76f12fe4 PrintStackTraceSignalHandler(void*) Signals.cpp:0:0 #1 0x0000559b76f10814 SignalHandler(int) Signals.cpp:0:0 #2 0x00007f74d189c420 __restore_rt (/lib/x86_64-linux-gnu/libpthread.so.0+0x14420) #3 0x00007f74d133900b raise /build/glibc-SzIz7B/glibc-2.31/signal/../sysdeps/unix/sysv/linux/raise.c:51:1 #4 0x00007f74d1318859 abort /build/glibc-SzIz7B/glibc-2.31/stdlib/abort.c:81:7 #5 0x0000559b75d54bd8 llvm::ConvertUTF8toUTF32(unsigned char const**, unsigned char const*, unsigned int**, unsigned int*, llvm::ConversionFlags) (.cold) ConvertUTF.cpp:0:0 #6 0x0000559b7789ef5d llvm::SelectionDAGISel::CannotYetSelect(llvm::SDNode*) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x1dccf5d) #7 0x0000559b778a169a llvm::SelectionDAGISel::SelectCodeCommon(llvm::SDNode*, unsigned char const*, unsigned int) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x1dcf69a) #8 0x0000559b764bf877 AMDGPUDAGToDAGISel::Select(llvm::SDNode*) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x9ed877) #9 0x0000559b7789be90 llvm::SelectionDAGISel::DoInstructionSelection() (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x1dc9e90) #10 0x0000559b778a9095 llvm::SelectionDAGISel::CodeGenAndEmitDAG() (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x1dd7095) #11 0x0000559b778ac438 llvm::SelectionDAGISel::SelectAllBasicBlocks(llvm::Function const&) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x1dda438) #12 0x0000559b778ae102 llvm::SelectionDAGISel::runOnMachineFunction(llvm::MachineFunction&) (.part.0) SelectionDAGISel.cpp:0:0 #13 0x0000559b764c8ad9 AMDGPUDAGToDAGISel::runOnMachineFunction(llvm::MachineFunction&) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x9f6ad9) #14 0x0000559b7701dde1 llvm::MachineFunctionPass::runOnFunction(llvm::Function&) (.part.0) MachineFunctionPass.cpp:0:0 #15 0x0000559b7688f4c1 llvm::FPPassManager::runOnFunction(llvm::Function&) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0xdbd4c1) #16 0x0000559b76b22cf7 (anonymous namespace)::CGPassManager::runOnModule(llvm::Module&) CallGraphSCCPass.cpp:0:0 #17 0x0000559b7688ff92 llvm::legacy::PassManagerImpl::run(llvm::Module&) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0xdbdf92) #18 0x0000559b774d5b15 codegen(llvm::lto::Config const&, llvm::TargetMachine*, 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 const&) LTOBackend.cpp:0:0 #19 0x0000559b774d60ed 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&) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x1a040ed) #20 0x0000559b774cc6fc llvm::lto::LTO::runRegularLTO(std::function<llvm::Expected<std::unique_ptr<llvm::CachedFileStream, std::default_delete<llvm::CachedFileStream>>> (unsigned int, llvm::Twine const&)>) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x19fa6fc) #21 0x0000559b774ccda8 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&)>) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x19fada8) #22 0x0000559b75e06f3d (anonymous namespace)::linkBitcodeFiles(llvm::SmallVectorImpl<llvm::object::OffloadFile>&, llvm::SmallVectorImpl<llvm::StringRef>&, llvm::opt::ArgList const&) (.constprop.0) ClangLinkerWrapper.cpp:0:0 #23 0x0000559b75e0dc6a llvm::Error (anonymous namespace)::linkAndWrapDeviceFiles(llvm::SmallVectorImpl<llvm::object::OffloadFile>&, llvm::opt::InputArgList const&, char**, int)::'lambda'(auto&)::operator()<llvm::SmallVector<llvm::object::OffloadFile, 3u>>(auto&) const ClangLinkerWrapper.cpp:0:0 #24 0x0000559b75e14215 (anonymous namespace)::linkAndWrapDeviceFiles(llvm::SmallVectorImpl<llvm::object::OffloadFile>&, llvm::opt::InputArgList const&, char**, int) ClangLinkerWrapper.cpp:0:0 #25 0x0000559b75d5a746 main (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x288746) #26 0x00007f74d131a083 __libc_start_main /build/glibc-SzIz7B/glibc-2.31/csu/../csu/libc-start.c:342:3 #27 0x0000559b75df66de _start (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x3246de) #0 0x0000559b76f12fe4 PrintStackTraceSignalHandler(void*) Signals.cpp:0:0 #1 0x0000559b76f10814 SignalHandler(int) Signals.cpp:0:0 #2 0x00007f74d189c420 __restore_rt (/lib/x86_64-linux-gnu/libpthread.so.0+0x14420) #3 0x00007f74d133900b raise /build/glibc-SzIz7B/glibc-2.31/signal/../sysdeps/unix/sysv/linux/raise.c:51:1 #4 0x00007f74d1318859 abort /build/glibc-SzIz7B/glibc-2.31/stdlib/abort.c:81:7 #5 0x0000559b75d54bd8 llvm::ConvertUTF8toUTF32(unsigned char const**, unsigned char const*, unsigned int**, unsigned int*, llvm::ConversionFlags) (.cold) ConvertUTF.cpp:0:0 #6 0x0000559b7789ef5d llvm::SelectionDAGISel::CannotYetSelect(llvm::SDNode*) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x1dccf5d) #7 0x0000559b778a169a llvm::SelectionDAGISel::SelectCodeCommon(llvm::SDNode*, unsigned char const*, unsigned int) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x1dcf69a) #8 0x0000559b764bf877 AMDGPUDAGToDAGISel::Select(llvm::SDNode*) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x9ed877) #9 0x0000559b7789be90 llvm::SelectionDAGISel::DoInstructionSelection() (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x1dc9e90) #10 0x0000559b778a9095 llvm::SelectionDAGISel::CodeGenAndEmitDAG() (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x1dd7095) #11 0x0000559b778ac438 llvm::SelectionDAGISel::SelectAllBasicBlocks(llvm::Function const&) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x1dda438) #12 0x0000559b778ae102 llvm::SelectionDAGISel::runOnMachineFunction(llvm::MachineFunction&) (.part.0) SelectionDAGISel.cpp:0:0 #13 0x0000559b764c8ad9 AMDGPUDAGToDAGISel::runOnMachineFunction(llvm::MachineFunction&) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x9f6ad9) #14 0x0000559b7701dde1 llvm::MachineFunctionPass::runOnFunction(llvm::Function&) (.part.0) MachineFunctionPass.cpp:0:0 #15 0x0000559b7688f4c1 llvm::FPPassManager::runOnFunction(llvm::Function&) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0xdbd4c1) #16 0x0000559b76b22cf7 (anonymous namespace)::CGPassManager::runOnModule(llvm::Module&) CallGraphSCCPass.cpp:0:0 #17 0x0000559b7688ff92 llvm::legacy::PassManagerImpl::run(llvm::Module&) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0xdbdf92) #18 0x0000559b774d5b15 codegen(llvm::lto::Config const&, llvm::TargetMachine*, 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 const&) LTOBackend.cpp:0:0 #19 0x0000559b774d60ed 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&) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x1a040ed) #20 0x0000559b774cc6fc llvm::lto::LTO::runRegularLTO(std::function<llvm::Expected<std::unique_ptr<llvm::CachedFileStream, std::default_delete<llvm::CachedFileStream>>> (unsigned int, llvm::Twine const&)>) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x19fa6fc) #21 0x0000559b774ccda8 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&)>) (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x19fada8) #22 0x0000559b75e06f3d (anonymous namespace)::linkBitcodeFiles(llvm::SmallVectorImpl<llvm::object::OffloadFile>&, llvm::SmallVectorImpl<llvm::StringRef>&, llvm::opt::ArgList const&) (.constprop.0) ClangLinkerWrapper.cpp:0:0 #23 0x0000559b75e0dc6a llvm::Error (anonymous namespace)::linkAndWrapDeviceFiles(llvm::SmallVectorImpl<llvm::object::OffloadFile>&, llvm::opt::InputArgList const&, char**, int)::'lambda'(auto&)::operator()<llvm::SmallVector<llvm::object::OffloadFile, 3u>>(auto&) const ClangLinkerWrapper.cpp:0:0 #24 0x0000559b75e14215 (anonymous namespace)::linkAndWrapDeviceFiles(llvm::SmallVectorImpl<llvm::object::OffloadFile>&, llvm::opt::InputArgList const&, char**, int) ClangLinkerWrapper.cpp:0:0 #25 0x0000559b75d5a746 main (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x288746) #26 0x00007f74d131a083 __libc_start_main /build/glibc-SzIz7B/glibc-2.31/csu/../csu/libc-start.c:342:3 #27 0x0000559b75df66de _start (/home/dobercea/rocm/trunk_1.0/bin/clang-linker-wrapper+0x3246de) ```
shiltian commented 11 months ago

This is a known issue. We don't really support dynamic scheduling on the device because of the current way that the device runtime is implemented.

doru1004 commented 11 months ago

We may not support the full semantics of the dynamic schedule but a program with a dynamic schedule specified was compiling and running correctly before the commit I outlined above.

shiltian commented 11 months ago

It is not "correct" per se, and the reason is exactly the one shown in the crash information. We don't support thread local global variable currently. We did some tricks to hide the issue. However it is fair to say w/o the patch there is no compiler crash. I think the malloc support on AMDGPU makes some code not be able to eliminate.

doru1004 commented 11 months ago

It would be great if we didn't break the compilation of source code we can compile today.

JonChesterfield commented 11 months ago

The code pointed to by the stacktrace is in Workshare.cpp,

static DynamicScheduleTracker *THREAD_LOCAL(ThreadDSTPtr);

where that macro is in types.h

  [[clang::address_space(5)]] NAME [[clang::loader_uninitialized]]

On amdgpu, that's nonsense. You don't get to say that a global static variable is on the stack, as it isn't. I don't know whether that works on nvptx. As in this is only working by luck, because our IR verifier doesn't complain that stack variables are not global variables, and the corresponding variable presumably got optimised out before the backend flatly rejected it.

If we actually had thread_local variables that would be a syntax for using them, but we don't, so it isn't.

jhuber6 commented 11 months ago

The code pointed to by the stacktrace is in Workshare.cpp,

static DynamicScheduleTracker *THREAD_LOCAL(ThreadDSTPtr);

where that macro is in types.h

  [[clang::address_space(5)]] NAME [[clang::loader_uninitialized]]

On amdgpu, that's nonsense. You don't get to say that a global static variable is on the stack, as it isn't. I don't know whether that works on nvptx. As in this is only working by luck, because our IR verifier doesn't complain that stack variables are not global variables, and the corresponding variable presumably got optimised out before the backend flatly rejected it.

If we actually had thread_local variables that would be a syntax for using them, but we don't, so it isn't.

I believe this is nonsense on NVPTX as well. It is most likely passing currently because we use -mlink-builtin-bitcode so this variable is either always optimized out or never used. If you take the following trivial file and try to run it through ptxas it will fail with an appropriate error message.

[[clang::loader_uninitialized]] int [[clang::address_space(5)]] x;
$ clang++ local.cpp --target=nvptx64-nvidia-cuda -mcpu=sm_89
ptxas /tmp/local-f1058b.s, line 9; error   : Module-scoped variables in .local state space are not allowed with ABI
ptxas fatal   : Ptx assembly aborted due to errors
JonChesterfield commented 11 months ago

Ah good stuff, broken on all our targets. Introduced with the devicertl rewrite as far as I can tell from git. Rocm is probably going to work around by reverting the bump allocator introduction, thus sending this back down the dead code elimination path.

jhuber6 commented 11 months ago

If we need to pretend like we have TLS behavior we can do the same trick I did for srand and burn 1/4th or so of our LDS / shared memory budget on it if it ends up used https://github.com/llvm/llvm-project/blob/main/libc/src/stdlib/rand_util.h#L22.