zjin-lcf
commented
2 years ago For the HIP backend, is there some 'assert()' called in the sycl kernel ?
Open krasznaa opened 2 years ago
zjin-lcf
commented
2 years ago For the HIP backend, is there some 'assert()' called in the sycl kernel ?
krasznaa
commented
2 years ago For the HIP backend, is there some 'assert()' called in the sycl kernel ?
Oh, yes. The problem with the HIP backend is completely separate from what I'm describing here. We use assert(...) statements in our code in a few places, and the HIP backend doesn't seem to support that yet.
Note that for the CUDA backend we've had the following workaround in place for a while:
But I didn't find a similar workaround for HIP yet. Also note that the CUDA workaround is not actually needed with the latest version of the compiler. (This is where it comes handy that our build checks at configuration whether it should use the workaround or not.)
zjin-lcf
commented
2 years ago Thanks.
hdelan
commented
2 years ago Hi @krasznaa . HIP AMD backend has recently gained support for assert(). See https://github.com/intel/llvm/commit/ade1870a3d6bd858b55d50211bbbf6eab0ac4796 . Would be great if you could see if debug compilation works for DPC++ for AMD HIP now. I'm just building for CUDA backend now so will report back.
hdelan
commented
2 years ago Having some build issues. Could you possibly share your cmake invocation for DPC++ with CUDA backend? Thanks
hdelan
commented
2 years ago Also @krasznaa if you can trigger this linking failure, once you do could you post the output of
dmesg -T | grep -E -i -B100 'killed process'Thanks
hdelan
commented
1 year ago Please let us know if you are still having problems with this. If not we will close the ticket. If you have any more issues in the future with this just comment on this ticket and we can reopen it.
krasznaa
commented
1 year ago Sorry for the silence. 😦 I plan to come back to the issue in this coming week. (Unfortunately it still affects us.)
I've built the 2022-06 tag of the project (which still fails on this build) with the following:
source /software/cuda/11.6.2/x86_64/setup.sh
source /software/cmake/3.23.2/x86_64-ubuntu2004-gcc9-opt/setup.sh
cmake -DCMAKE_BUILD_TYPE=Release \
-DCMAKE_INSTALL_PREFIX=/software/intel/clang-2022-06/x86_64-ubuntu2004-gcc9-opt \
-DLLVM_ENABLE_ASSERTIONS=ON -DCMAKE_INSTALL_LIBDIR=lib -DLLVM_TARGETS_TO_BUILD="X86;NVPTX" \
-DLLVM_EXTERNAL_PROJECTS="sycl;llvm-spirv;opencl;libdevice;xpti;xptifw" \
-DLLVM_EXTERNAL_SYCL_SOURCE_DIR=/atlas/krasznaa/intel/llvm/sycl \
-DLLVM_EXTERNAL_LLVM_SPIRV_SOURCE_DIR=/atlas/krasznaa/intel/llvm/llvm-spirv \
-DLLVM_EXTERNAL_XPTI_SOURCE_DIR=/atlas/krasznaa/intel/llvm/xpti \
-DXPTI_SOURCE_DIR=/atlas/krasznaa/intel/llvm/xpti \
-DLLVM_EXTERNAL_XPTIFW_SOURCE_DIR=/atlas/krasznaa/intel/llvm/xptifw \
-DLLVM_EXTERNAL_LIBDEVICE_SOURCE_DIR=/atlas/krasznaa/intel/llvm/libdevice \
-DLLVM_ENABLE_PROJECTS="clang;compiler-rt;sycl;llvm-spirv;opencl;libdevice;xpti;xptifw;libclc;lld;clang-tools-extra;openmp" \
-DLIBCLC_TARGETS_TO_BUILD="nvptx64--;nvptx64--nvidiacl" -DLIBCLC_GENERATE_REMANGLED_VARIANTS=ON \
-DLLVM_BUILD_TOOLS=ON -DSYCL_ENABLE_WERROR=OFF -DSYCL_INCLUDE_TESTS=OFF -DLLVM_ENABLE_DOXYGEN=OFF \
-DLLVM_ENABLE_SPHINX=OFF -DBUILD_SHARED_LIBS=OFF -DSYCL_ENABLE_XPTI_TRACING=ON -DLLVM_ENABLE_LLD=OFF \
-DXPTI_ENABLE_WERROR=OFF -DOpenCL_INSTALL_KHRONOS_ICD_LOADER=TRUE -DLLVM_ENABLE_PIC=ON \
-DLLVM_ENABLE_RTTI=ON -DSYCL_ENABLE_PLUGINS="opencl;level_zero;cuda" ../llvm/llvm/
make -j
make -j sycl-toolchain
cmake --install .I'll try the what happens with a newer versions of the code this coming week.
krasznaa
commented
1 year ago Unfortunately my build still fails. :frowning:
...
[ 95%] Building SYCL object device/sycl/CMakeFiles/traccc_sycl.dir/src/seeding/seed_finding.sycl.o
[ 95%] Linking SYCL shared library ../../lib/libtraccc_sycl.so
ptxas /tmp/clusterization_algorithm-sm_50-a5dcc9.s, line 17806; fatal : Parsing error near '.': syntax error
ptxas fatal : Ptx assembly aborted due to errors
llvm-foreach:
clang-16: error: ptxas command failed with exit code 255 (use -v to see invocation)
clang version 16.0.0 (https://github.com/intel/llvm.git 0f579bae55c48d810e1ed76db29229c854e61d5e)
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /software/intel/clang-2022-09/x86_64-ubuntu2004-gcc9-opt/bin
clang-16: note: diagnostic msg: Error generating preprocessed source(s).
make[3]: *** [device/sycl/CMakeFiles/traccc_sycl.dir/build.make:169: lib/libtraccc_sycl.so.0.2.0] Error 1
make[2]: *** [CMakeFiles/Makefile2:3723: device/sycl/CMakeFiles/traccc_sycl.dir/all] Error 2
make[1]: *** [CMakeFiles/Makefile2:3730: device/sycl/CMakeFiles/traccc_sycl.dir/rule] Error 2
make: *** [Makefile:992: traccc_sycl] Error 2
[bash][atspot01]:build-x86_64_ubuntu2004_llvm_nvidia > dmesg -T | grep -E -i -B100 'killed process'
[bash][atspot01]:build-x86_64_ubuntu2004_llvm_nvidia >As you can see, I tried using the 2022-09 release of the compiler. Which did require me to make a good number of changes in our code to avoid compilation warnings. (In debug builds we treat warnings as errors, so this was fun... :stuck_out_tongue:)
The output now suggests that the failure happens in this code: https://github.com/acts-project/traccc/blob/main/device/sycl/src/clusterization/clusterization_algorithm.sycl
I could put the clusterization_algorithm-sm_50-a5dcc9.s file somewhere public if that may help... The referenced problematic line in question (if I didn't make a mistake navigating to it in this pretty large file) is this:
.b64 __PRETTY_FUNCTION__._ZN6traccc14container_baseIKNS_11cell_moduleEKNS_4cellEN6vecmem13device_vectorENS5_20jagged_device_vectorEN6thrust4pairEEC2INS5_4data11vector_viewIS2_EENSC_18jagged_vector_viewIS4_EESt9enable_ifILb0EvESI_EERKT_RKT0_Yepp... we do some non-trivial templating in our code... :stuck_out_tongue: (https://github.com/acts-project/traccc/blob/main/core/include/traccc/edm/cell.hpp)
hdelan
commented
1 year ago Hi sorry I meant could you share your cmake invocation that you use to build your project (traccc)? Thanks
krasznaa
commented
1 year ago Ahh, sorry. I didn't even consider that you may want to build our entire project. :stuck_out_tongue:
The procedure, especially with the very latest versions of the compiler, is indeed not trivial. You'll need the following as a setup:
export CC="/path/to/intel/compiler/bin/clang"
export CXX="/path/to/intel/compiler/bin/clang++ -Wno-deprecated-builtins"
export SYCLCXX="${CXX} -fsycl -fsycl-targets=nvptx64-nvidia-cuda -Wno-linker-warnings -Wno-sycl-target"The many flags for SYCLCXX (which is used by "our cmake configuration" to figure out which compiler to use for compiling SYCL code) are needed to suppress all the warnings that come up during the build. Because our build by default uses -Werror in Debug builds. So all those warnings are pretty bad for us. :frowning:
nvcc set up in your environment as well, you should be explicit about how CMake should build CUDA code in your environment. Like:export CUDACXX=/path/to/cuda/bin/nvcc
export CUDAHOSTCXX="${CXX}"
export CUDAFLAGS="-allow-unsupported-compiler -Xcompiler -Wno-deprecated-builtins"The flags are meant to disable warnings here as well. :wink:
2022-09 tag of the compiler, you can find these in this branch of mine: https://github.com/krasznaa/traccc/tree/SYCL2020Update-main-20221102[bash][atspot01]:intel-test > git clone -b SYCL2020Update-main-20221102 https://github.com/krasznaa/traccc.git
Cloning into 'traccc'...
remote: Enumerating objects: 7537, done.
remote: Counting objects: 100% (847/847), done.
remote: Compressing objects: 100% (486/486), done.
remote: Total 7537 (delta 364), reused 644 (delta 278), pack-reused 6690
Receiving objects: 100% (7537/7537), 1.96 MiB | 15.53 MiB/s, done.
Resolving deltas: 100% (3817/3817), done.
[bash][atspot01]:intel-test > source /software/intel/clang-2022-09/x86_64-ubuntu2004-gcc9-opt/setup.sh nvidia
Configured CUDA from: /software/cuda/11.7.1/x86_64
Configured Clang from: /software/intel/clang-2022-09/x86_64-ubuntu2004-gcc9-opt
[bash][atspot01]:intel-test > source /software/boost/1.80.0/x86_64-ubuntu2004-gcc9-opt/setup.sh
Configured Boost from: /software/boost/1.80.0/x86_64-ubuntu2004-gcc9-opt
[bash][atspot01]:intel-test > source /software/cmake/3.24.1/x86_64-ubuntu2004-gcc9-opt/setup.sh
Configured CMake from: /software/cmake/3.24.1/x86_64-ubuntu2004-gcc9-opt
[bash][atspot01]:intel-test > cmake -DCMAKE_BUILD_TYPE=Debug -DTRACCC_BUILD_SYCL=TRUE -DTRACCC_BUILD_EXAMPLES=FALSE -S traccc -B build
-- The CXX compiler identification is Clang 16.0.0
-- Detecting CXX compiler ABI info
-- Detecting CXX compiler ABI info - done
-- Check for working CXX compiler: /software/intel/clang-2022-09/x86_64-ubuntu2004-gcc9-opt/bin/clang++ - skipped
-- Detecting CXX compile features
-- Detecting CXX compile features - done
-- Looking for a CUDA compiler
-- Looking for a CUDA compiler - /software/cuda/11.7.1/x86_64/bin/nvcc
-- Looking for a CUDA host compiler - /software/intel/clang-2022-09/x86_64-ubuntu2004-gcc9-opt/bin/clang++
-- Building VecMem as part of the TRACCC project
-- Looking for a HIP compiler
-- Looking for a HIP compiler - NOTFOUND
-- Looking for a SYCL compiler
-- Looking for a SYCL compiler - /software/intel/clang-2022-09/x86_64-ubuntu2004-gcc9-opt/bin/clang++
...
-- Building GoogleTest as part of the TRACCC project
-- Found Python: /usr/bin/python3.8 (found version "3.8.10") found components: Interpreter
-- Building with plugin type: ARRAY
-- The CUDA compiler identification is NVIDIA 11.7.99
-- Detecting CUDA compiler ABI info
-- Detecting CUDA compiler ABI info - done
-- Check for working CUDA compiler: /software/cuda/11.7.1/x86_64/bin/nvcc - skipped
-- Detecting CUDA compile features
-- Detecting CUDA compile features - done
-- Configuring done
-- Generating done
-- Build files have been written to: /home/krasznaa/ATLAS/projects/traccc/intel-test/build
[bash][atspot01]:intel-test > export MAKEFLAGS=-j16
[bash][atspot01]:intel-test > cmake --build build --target traccc_sycl
[ 0%] Copy to /home/krasznaa/ATLAS/projects/traccc/intel-test/build/_deps/vc-build/trigonometric_SSSE3.cpp
[ 0%] Copy to /home/krasznaa/ATLAS/projects/traccc/intel-test/build/_deps/vc-build/sse_sorthelper_AVX.cpp
[ 0%] Copy to /home/krasznaa/ATLAS/projects/traccc/intel-test/build/_deps/vc-build/avx_sorthelper_AVX2+FMA+BMI2.cpp
...
[ 93%] Building CXX object device/common/CMakeFiles/traccc_device_common.dir/src/seeding/make_triplet_buffer.cpp.o
[ 93%] Linking CXX shared library ../../lib/libtraccc_device_common.so
[ 93%] Built target traccc_device_common
Scanning dependencies of target traccc_sycl
[ 95%] Building SYCL object device/sycl/CMakeFiles/traccc_sycl.dir/src/clusterization/clusterization_algorithm.sycl.o
[ 95%] Building SYCL object device/sycl/CMakeFiles/traccc_sycl.dir/src/seeding/spacepoint_binning.sycl.o
[ 95%] Building SYCL object device/sycl/CMakeFiles/traccc_sycl.dir/src/seeding/seed_finding.sycl.o
[ 95%] Building CXX object device/sycl/CMakeFiles/traccc_sycl.dir/src/seeding/seeding_algorithm.cpp.o
[ 97%] Building SYCL object device/sycl/CMakeFiles/traccc_sycl.dir/src/seeding/track_params_estimation.sycl.o
[ 97%] Building SYCL object device/sycl/CMakeFiles/traccc_sycl.dir/src/utils/get_queue.sycl.o
[ 97%] Building CXX object device/sycl/CMakeFiles/traccc_sycl.dir/src/utils/queue_wrapper.cpp.o
[ 97%] Building SYCL object device/sycl/CMakeFiles/traccc_sycl.dir/src/utils/calculate1DimNdRange.sycl.o
[100%] Building SYCL object device/sycl/CMakeFiles/traccc_sycl.dir/src/utils/make_prefix_sum_buff.sycl.o
[100%] Linking SYCL shared library ../../lib/libtraccc_sycl.so
ptxas /tmp/clusterization_algorithm-sm_50-8e4474.s, line 17790; fatal : Parsing error near '.': syntax error
ptxas fatal : Ptx assembly aborted due to errors
llvm-foreach:
clang-16: error: ptxas command failed with exit code 255 (use -v to see invocation)
clang version 16.0.0 (https://github.com/intel/llvm.git 0f579bae55c48d810e1ed76db29229c854e61d5e)
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /software/intel/clang-2022-09/x86_64-ubuntu2004-gcc9-opt/bin
clang-16: note: diagnostic msg: Error generating preprocessed source(s).
make[3]: *** [device/sycl/CMakeFiles/traccc_sycl.dir/build.make:169: lib/libtraccc_sycl.so.0.2.0] Error 1
make[2]: *** [CMakeFiles/Makefile2:3575: device/sycl/CMakeFiles/traccc_sycl.dir/all] Error 2
make[1]: *** [CMakeFiles/Makefile2:3582: device/sycl/CMakeFiles/traccc_sycl.dir/rule] Error 2
make: *** [Makefile:992: traccc_sycl] Error 2
[bash][atspot01]:intel-test >The "examples" of the project explicitly need to be turned off (-DTRACCC_BUILD_EXAMPLES=FALSE), as they would require a dependency (ROOT) that would be too much of a hassle for you to set up just for this build.
krasznaa
commented
1 year ago I only realised today that I have a much easier-to-build example for you... You see, we've been seeing occasional build errors with acts-project/vecmem for some time now. Like:
As silly as I am, I didn't put 1+1 together that that would be the same error that we've been seeing in our other project as well. :clown_face:
So... VecMem is a lot easier to build. For that you just need the Intel compiler, and CMake 3.17+. Then you can trigger the error like:
export CC="/path/to/intel/compiler/bin/clang"
export CXX="/path/to/intel/compiler/bin/clang++ -Wno-deprecated-builtins"
export SYCLCXX="${CXX} -fsycl -fsycl-targets=nvptx64-nvidia-cuda"
git clone https://github.com/acts-project/vecmem.git
cmake -DCMAKE_BUILD_TYPE=Debug -S vecmem -B build
cmake --build buildWith the latest compiler this always crashes. With 2021-09, which we used earlier in our CI, it only crashed sometimes.
hdelan
commented
1 year ago I can't replicate the failure on my machine. Is the failure only happening in docker? Will try in docker now
hdelan
commented
1 year ago I also can't replicate the failures in docker using the CI setup for vecmem. I will try separately for traccc
krasznaa
commented
1 year ago You're building the project with -DCMAKE_BUILD_TYPE=Debug, right? Since the build errors in Docker are quite reproducible for me... 😕
hdelan
commented
1 year ago Actually another thing was building.. please hold!
krasznaa
commented
1 year ago Also, check that the VecMem build would accept your SYCL compiler. By default, when it doesn't find a compiler that it can use, the build of all SYCL code is turned off in the project.
You can force it to build the SYCL code (and fail if it can't do that) with: -DVECMEM_BUILD_SYCL_LIBRARY=TRUE
hdelan
commented
1 year ago Unfortunately I can't reproduce these failures locally when I replicate the CI pipeline for either traccc or vecmem. Can you tell me a bit about the hardware that you are running on? How much RAM do you have? On my PC I have 32GB RAM, which sometimes was over half full during the CI replicate runs. Since this failure is flaky and hard to reproduce it may be due to memory overallocation. vecmem for instance has what looks like 8 or 9 build containers that do not necessarily build sequentially. This could be a cause of the crashes. It would also make sense that the crashes occur at linking step, which is the most memory intensive operation of the build.
We unfortunately don't have the option to split DWARF for PTX backend yet, which would reduce the size of the debug binary. However it might be worthwhile making sure that the debug build is dependent on all other previous tasks, meaning that it does not have to share resources with other builds. You can do this by using the needs: tag in Github actions. See here https://docs.github.com/en/actions/using-jobs/using-jobs-in-a-workflow
krasznaa
commented
1 year ago I am extremely surprised. :confused: I can very reliably produce this crash with vecmem like the following:
[bash][pcadp04]:vecmem > source /software/intel/clang-2022-09/x86_64-centos9-gcc11-opt/setup.sh nvidia
Configured CUDA from: /software/cuda/11.7.1/x86_64
Configured Clang from: /software/intel/clang-2022-09/x86_64-centos9-gcc11-opt
[bash][pcadp04]:vecmem > echo $CXX
/software/intel/clang-2022-09/x86_64-centos9-gcc11-opt/bin/clang++ -Wno-deprecated-builtins
[bash][pcadp04]:vecmem > echo $SYCLCXX
/software/intel/clang-2022-09/x86_64-centos9-gcc11-opt/bin/clang++ -Wno-deprecated-builtins -fsycl -Wno-linker-warnings -Wno-sycl-target -fsycl-targets=nvptx64-nvidia-cuda -Xsycl-target-backend --offload-arch=sm_86
[bash][pcadp04]:vecmem > free
total used free shared buff/cache available
Mem: 131228636 5500996 17837812 86000 109222604 125727640
Swap: 19533820 145256 19388564
[bash][pcadp04]:vecmem > git clone https://github.com/acts-project/vecmem.git
Cloning into 'vecmem'...
remote: Enumerating objects: 4542, done.
remote: Counting objects: 100% (1519/1519), done.
remote: Compressing objects: 100% (645/645), done.
remote: Total 4542 (delta 928), reused 1149 (delta 826), pack-reused 3023
Receiving objects: 100% (4542/4542), 1021.19 KiB | 1.47 MiB/s, done.
Resolving deltas: 100% (2463/2463), done.
[bash][pcadp04]:vecmem > cmake -DCMAKE_BUILD_TYPE=Debug -S vecmem -B build
-- The CXX compiler identification is Clang 16.0.0
-- Detecting CXX compiler ABI info
-- Detecting CXX compiler ABI info - done
-- Check for working CXX compiler: /software/intel/clang-2022-09/x86_64-centos9-gcc11-opt/bin/clang++ - skipped
-- Detecting CXX compile features
-- Detecting CXX compile features - done
-- Looking for a CUDA compiler
-- Looking for a CUDA compiler - /software/cuda/11.7.1/x86_64/bin/nvcc
-- Looking for a CUDA host compiler - /software/intel/clang-2022-09/x86_64-centos9-gcc11-opt/bin/clang++
-- Looking for a HIP compiler
-- Looking for a HIP compiler - /usr/bin/hipcc
-- Looking for a SYCL compiler
-- Looking for a SYCL compiler - /software/intel/clang-2022-09/x86_64-centos9-gcc11-opt/bin/clang++
-- Performing Test COMPILER_HAS_HIDDEN_VISIBILITY
-- Performing Test COMPILER_HAS_HIDDEN_VISIBILITY - Success
-- Performing Test COMPILER_HAS_HIDDEN_INLINE_VISIBILITY
-- Performing Test COMPILER_HAS_HIDDEN_INLINE_VISIBILITY - Success
-- Performing Test COMPILER_HAS_DEPRECATED_ATTR
-- Performing Test COMPILER_HAS_DEPRECATED_ATTR - Success
-- Performing Test VECMEM_HAVE_PMR_MEMORY_RESOURCE
-- Performing Test VECMEM_HAVE_PMR_MEMORY_RESOURCE - Success
-- Using memory resource types from the std::pmr namespace
-- Performing Test VECMEM_HAVE_DEFAULT_RESOURCE
-- Performing Test VECMEM_HAVE_DEFAULT_RESOURCE - Success
-- The SYCL compiler identification is IntelLLVM 16.0.0
-- Check for working SYCL compiler: /software/intel/clang-2022-09/x86_64-centos9-gcc11-opt/bin/clang++
-- Check for working SYCL compiler: /software/intel/clang-2022-09/x86_64-centos9-gcc11-opt/bin/clang++ - works
-- Performing Test VECMEM_HAVE_SYCL_EXT_ONEAPI_PRINTF
-- Performing Test VECMEM_HAVE_SYCL_EXT_ONEAPI_PRINTF - Success
-- Performing Test VECMEM_HAVE_SYCL_ONEAPI_PRINTF
-- Performing Test VECMEM_HAVE_SYCL_ONEAPI_PRINTF - Failed
-- Performing Test VECMEM_HAVE_SYCL_ATOMIC_REF
-- Performing Test VECMEM_HAVE_SYCL_ATOMIC_REF - Success
-- Performing Test VECMEM_HAVE_LZCNT_U64
-- Performing Test VECMEM_HAVE_LZCNT_U64 - Failed
-- Performing Test VECMEM_HAVE_BUILTIN_CLZL
-- Performing Test VECMEM_HAVE_BUILTIN_CLZL - Success
-- Looking for std::aligned_alloc
-- Looking for std::aligned_alloc - found
-- Found CUDAToolkit: /software/cuda/11.7.1/x86_64/include (found version "11.7.99")
-- Looking for C++ include pthread.h
-- Looking for C++ include pthread.h - found
-- Performing Test CMAKE_HAVE_LIBC_PTHREAD
-- Performing Test CMAKE_HAVE_LIBC_PTHREAD - Success
-- Found Threads: TRUE
-- Check for working HIP compiler: /usr/bin/hipcc
-- Check for working HIP compiler: /usr/bin/hipcc - works
-- Found HIPToolkit: /opt/rocm/include (found version "5.3.22061")
-- Performing Test VECMEM_HAVE_SYCL_ASSERT
-- Performing Test VECMEM_HAVE_SYCL_ASSERT - Success
-- Performing Test VECMEM_HAVE_SYCL_MEMSET
-- Performing Test VECMEM_HAVE_SYCL_MEMSET - Success
-- Performing Test VECMEM_SYCL_WERROR_USABLE
-- Performing Test VECMEM_SYCL_WERROR_USABLE - Failed
-- Building GoogleTest as part of the VecMem project
-- The C compiler identification is Clang 16.0.0
-- Detecting C compiler ABI info
-- Detecting C compiler ABI info - done
-- Check for working C compiler: /software/intel/clang-2022-09/x86_64-centos9-gcc11-opt/bin/clang - skipped
-- Detecting C compile features
-- Detecting C compile features - done
-- Found Python: /usr/bin/python3.9 (found version "3.9.14") found components: Interpreter
-- The CUDA compiler identification is NVIDIA 11.7.99
-- Detecting CUDA compiler ABI info
-- Detecting CUDA compiler ABI info - done
-- Check for working CUDA compiler: /software/cuda/11.7.1/x86_64/bin/nvcc - skipped
-- Detecting CUDA compile features
-- Detecting CUDA compile features - done
-- Configuring done
-- Generating done
-- Build files have been written to: /mnt/ssd1/krasznaa/projects/vecmem/build
[bash][pcadp04]:vecmem > cmake --build build --target vecmem_test_sycl
[ 0%] Building CXX object _deps/googletest-build/googletest/CMakeFiles/gtest.dir/src/gtest-all.cc.o
[ 2%] Linking CXX shared library ../../../lib/libgtestd.so
[ 2%] Built target gtest
[ 5%] Building CXX object core/CMakeFiles/vecmem_core.dir/src/memory/allocator.cpp.o
[ 8%] Building CXX object core/CMakeFiles/vecmem_core.dir/src/memory/details/memory_resource_base.cpp.o
[ 11%] Building CXX object core/CMakeFiles/vecmem_core.dir/src/memory/arena.cpp.o
[ 13%] Building CXX object core/CMakeFiles/vecmem_core.dir/src/memory/arena_memory_resource.cpp.o
[ 16%] Building CXX object core/CMakeFiles/vecmem_core.dir/src/memory/identity_memory_resource.cpp.o
[ 19%] Building CXX object core/CMakeFiles/vecmem_core.dir/src/memory/terminal_memory_resource.cpp.o
[ 22%] Building CXX object core/CMakeFiles/vecmem_core.dir/src/memory/host_memory_resource.cpp.o
[ 25%] Building CXX object core/CMakeFiles/vecmem_core.dir/src/memory/binary_page_memory_resource.cpp.o
[ 27%] Building CXX object core/CMakeFiles/vecmem_core.dir/src/memory/binary_page_memory_resource_impl.cpp.o
[ 30%] Building CXX object core/CMakeFiles/vecmem_core.dir/src/memory/contiguous_memory_resource.cpp.o
[ 33%] Building CXX object core/CMakeFiles/vecmem_core.dir/src/memory/instrumenting_memory_resource.cpp.o
[ 36%] Building CXX object core/CMakeFiles/vecmem_core.dir/src/memory/choice_memory_resource.cpp.o
[ 38%] Building CXX object core/CMakeFiles/vecmem_core.dir/src/memory/coalescing_memory_resource.cpp.o
[ 41%] Building CXX object core/CMakeFiles/vecmem_core.dir/src/memory/conditional_memory_resource.cpp.o
[ 44%] Building CXX object core/CMakeFiles/vecmem_core.dir/src/memory/debug_memory_resource.cpp.o
[ 47%] Building CXX object core/CMakeFiles/vecmem_core.dir/src/memory/details/is_aligned.cpp.o
[ 50%] Building CXX object core/CMakeFiles/vecmem_core.dir/src/utils/copy.cpp.o
[ 52%] Building CXX object core/CMakeFiles/vecmem_core.dir/src/utils/memory_monitor.cpp.o
[ 55%] Linking CXX shared library ../lib64/libvecmem_core.so
[ 55%] Built target vecmem_core
Scanning dependencies of target vecmem_sycl
[ 58%] Building SYCL object sycl/CMakeFiles/vecmem_sycl.dir/src/memory/sycl/details/memory_resource_base.sycl.o
[ 61%] Building SYCL object sycl/CMakeFiles/vecmem_sycl.dir/src/memory/sycl/device_memory_resource.sycl.o
[ 63%] Building SYCL object sycl/CMakeFiles/vecmem_sycl.dir/src/memory/sycl/host_memory_resource.sycl.o
[ 66%] Building SYCL object sycl/CMakeFiles/vecmem_sycl.dir/src/memory/sycl/shared_memory_resource.sycl.o
[ 69%] Building SYCL object sycl/CMakeFiles/vecmem_sycl.dir/src/utils/sycl/copy.sycl.o
[ 72%] Building SYCL object sycl/CMakeFiles/vecmem_sycl.dir/src/utils/sycl/queue_wrapper.sycl.o
[ 75%] Building SYCL object sycl/CMakeFiles/vecmem_sycl.dir/src/utils/sycl/get_queue.sycl.o
[ 77%] Linking SYCL shared library ../lib64/libvecmem_sycl.so
[ 77%] Built target vecmem_sycl
[ 80%] Building CXX object tests/CMakeFiles/vecmem_testing_common.dir/common/memory_resource_name_gen.cpp.o
[ 83%] Linking CXX static library ../lib64/libvecmem_testing_common.a
[ 83%] Built target vecmem_testing_common
[ 86%] Building CXX object _deps/googletest-build/googletest/CMakeFiles/gtest_main.dir/src/gtest_main.cc.o
[ 88%] Linking CXX shared library ../../../lib/libgtest_maind.so
[ 88%] Built target gtest_main
Scanning dependencies of target vecmem_test_sycl
[ 91%] Building CXX object tests/sycl/CMakeFiles/vecmem_test_sycl.dir/test_sycl_memory_resources.cpp.o
[ 94%] Building SYCL object tests/sycl/CMakeFiles/vecmem_test_sycl.dir/test_sycl_containers.sycl.o
[ 97%] Building SYCL object tests/sycl/CMakeFiles/vecmem_test_sycl.dir/test_sycl_jagged_containers.sycl.o
[100%] Linking SYCL executable ../../bin/vecmem_test_sycl
ptxas warning : .debug_abbrev section not found
ptxas warning : .debug_info section not found
ptxas warning : .debug_abbrev section not found
ptxas warning : .debug_info section not found
ptxas /tmp/krasznaa/test_sycl_memory_resources-sm_86-05c6bf.s, line 6151; fatal : Parsing error near '.': syntax error
ptxas fatal : Ptx assembly aborted due to errors
llvm-foreach:
ptxas /tmp/krasznaa/test_sycl_memory_resources-sm_86-a9331f.s, line 6109; fatal : Parsing error near '.': syntax error
ptxas fatal : Ptx assembly aborted due to errors
llvm-foreach:
clang-16: error: ptxas command failed with exit code 255 (use -v to see invocation)
clang version 16.0.0 (https://github.com/intel/llvm.git 0f579bae55c48d810e1ed76db29229c854e61d5e)
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /software/intel/clang-2022-09/x86_64-centos9-gcc11-opt/bin
clang-16: note: diagnostic msg: Error generating preprocessed source(s).
gmake[3]: *** [tests/sycl/CMakeFiles/vecmem_test_sycl.dir/build.make:116: bin/vecmem_test_sycl] Error 1
gmake[2]: *** [CMakeFiles/Makefile2:1384: tests/sycl/CMakeFiles/vecmem_test_sycl.dir/all] Error 2
gmake[1]: *** [CMakeFiles/Makefile2:1391: tests/sycl/CMakeFiles/vecmem_test_sycl.dir/rule] Error 2
gmake: *** [Makefile:719: vecmem_test_sycl] Error 2
[bash][pcadp04]:vecmem >Note that here I ran the build using a single process/core, on a machine with 128 GB of memory. So that should really not be the issue...
What version of CUDA did you use in your test? :thinking:
krasznaa
commented
1 year ago But in the end the CUDA version shouldn't matter much. Since even with CUDA 10.2, which we use in our Docker image, I can produce the crash very reliably.
[bash][pcadp04]:vecmem > docker run -it --rm -v $PWD:$PWD -w $PWD ghcr.io/acts-project/ubuntu1804_cuda_oneapi:v32
root@874effd59f61:/mnt/ssd1/krasznaa/projects/vecmem# cmake -DCMAKE_BUILD_TYPE=Debug -S vecmem -B build
-- The CXX compiler identification is Clang 15.0.0
-- Check for working CXX compiler: /usr/local/bin/clang++
-- Check for working CXX compiler: /usr/local/bin/clang++ - works
-- Detecting CXX compiler ABI info
-- Detecting CXX compiler ABI info - done
-- Detecting CXX compile features
-- Detecting CXX compile features - done
-- Looking for a CUDA compiler
-- Looking for a CUDA compiler - /usr/local/cuda/bin/nvcc
-- Looking for a CUDA host compiler - /usr/local/bin/clang++
-- Looking for a HIP compiler
-- Looking for a HIP compiler - NOTFOUND
-- Looking for a SYCL compiler
-- Looking for a SYCL compiler - /usr/local/bin/clang++
-- Performing Test COMPILER_HAS_HIDDEN_VISIBILITY
-- Performing Test COMPILER_HAS_HIDDEN_VISIBILITY - Success
-- Performing Test COMPILER_HAS_HIDDEN_INLINE_VISIBILITY
-- Performing Test COMPILER_HAS_HIDDEN_INLINE_VISIBILITY - Success
-- Performing Test COMPILER_HAS_DEPRECATED_ATTR
-- Performing Test COMPILER_HAS_DEPRECATED_ATTR - Success
-- Performing Test VECMEM_HAVE_PMR_MEMORY_RESOURCE
-- Performing Test VECMEM_HAVE_PMR_MEMORY_RESOURCE - Failed
-- Performing Test VECMEM_HAVE_EXPERIMENTAL_PMR_MEMORY_RESOURCE
-- Performing Test VECMEM_HAVE_EXPERIMENTAL_PMR_MEMORY_RESOURCE - Success
-- Using memory resource types from the std::experimental::pmr namespace
-- Performing Test VECMEM_HAVE_DEFAULT_RESOURCE
-- Performing Test VECMEM_HAVE_DEFAULT_RESOURCE - Success
-- The SYCL compiler identification is IntelLLVM 15.0.0
-- Check for working SYCL compiler: /usr/local/bin/clang++
-- Check for working SYCL compiler: /usr/local/bin/clang++ - works
-- Performing Test VECMEM_HAVE_SYCL_EXT_ONEAPI_PRINTF
-- Performing Test VECMEM_HAVE_SYCL_EXT_ONEAPI_PRINTF - Success
-- Performing Test VECMEM_HAVE_SYCL_ONEAPI_PRINTF
-- Performing Test VECMEM_HAVE_SYCL_ONEAPI_PRINTF - Failed
-- Performing Test VECMEM_HAVE_SYCL_ATOMIC_REF
-- Performing Test VECMEM_HAVE_SYCL_ATOMIC_REF - Success
-- Performing Test VECMEM_HAVE_LZCNT_U64
-- Performing Test VECMEM_HAVE_LZCNT_U64 - Failed
-- Performing Test VECMEM_HAVE_BUILTIN_CLZL
-- Performing Test VECMEM_HAVE_BUILTIN_CLZL - Success
-- Looking for std::aligned_alloc
-- Looking for std::aligned_alloc - found
-- Found CUDAToolkit: /usr/local/cuda/include
-- Looking for C++ include pthread.h
-- Looking for C++ include pthread.h - found
-- Performing Test CMAKE_HAVE_LIBC_PTHREAD
-- Performing Test CMAKE_HAVE_LIBC_PTHREAD - Failed
-- Looking for pthread_create in pthreads
-- Looking for pthread_create in pthreads - not found
-- Looking for pthread_create in pthread
-- Looking for pthread_create in pthread - found
-- Found Threads: TRUE
-- Performing Test VECMEM_HAVE_SYCL_ASSERT
-- Performing Test VECMEM_HAVE_SYCL_ASSERT - Success
-- Performing Test VECMEM_HAVE_SYCL_MEMSET
-- Performing Test VECMEM_HAVE_SYCL_MEMSET - Success
-- Performing Test VECMEM_SYCL_WERROR_USABLE
-- Performing Test VECMEM_SYCL_WERROR_USABLE - Success
-- Building GoogleTest as part of the VecMem project
-- The C compiler identification is Clang 15.0.0
-- Check for working C compiler: /usr/local/bin/clang
-- Check for working C compiler: /usr/local/bin/clang - works
-- Detecting C compiler ABI info
-- Detecting C compiler ABI info - done
-- Detecting C compile features
-- Detecting C compile features - done
-- Found Python: /usr/bin/python3.6 (found version "3.6.9") found components: Interpreter
-- The CUDA compiler identification is NVIDIA 10.2.89
-- Check for working CUDA compiler: /usr/local/cuda/bin/nvcc
-- Check for working CUDA compiler: /usr/local/cuda/bin/nvcc - works
-- Detecting CUDA compiler ABI info
-- Detecting CUDA compiler ABI info - done
-- Detecting CUDA compile features
-- Detecting CUDA compile features - done
-- Found CUDAToolkit: /usr/local/cuda/include (found version "10.2.89")
-- Configuring done
-- Generating done
-- Build files have been written to: /mnt/ssd1/krasznaa/projects/vecmem/build
root@874effd59f61:/mnt/ssd1/krasznaa/projects/vecmem# cmake --build build --target vecmem_test_sycl
Scanning dependencies of target gtest
[ 2%] Building CXX object _deps/googletest-build/googletest/CMakeFiles/gtest.dir/src/gtest-all.cc.o
[ 5%] Linking CXX shared library ../../../lib/libgtestd.so
[ 5%] Built target gtest
Scanning dependencies of target vecmem_core
[ 8%] Building CXX object core/CMakeFiles/vecmem_core.dir/src/memory/allocator.cpp.o
[ 10%] Building CXX object core/CMakeFiles/vecmem_core.dir/src/memory/details/memory_resource_base.cpp.o
[ 13%] Building CXX object core/CMakeFiles/vecmem_core.dir/src/memory/arena.cpp.o
[ 16%] Building CXX object core/CMakeFiles/vecmem_core.dir/src/memory/arena_memory_resource.cpp.o
[ 18%] Building CXX object core/CMakeFiles/vecmem_core.dir/src/memory/identity_memory_resource.cpp.o
[ 21%] Building CXX object core/CMakeFiles/vecmem_core.dir/src/memory/terminal_memory_resource.cpp.o
[ 24%] Building CXX object core/CMakeFiles/vecmem_core.dir/src/memory/host_memory_resource.cpp.o
[ 27%] Building CXX object core/CMakeFiles/vecmem_core.dir/src/memory/binary_page_memory_resource.cpp.o
[ 29%] Building CXX object core/CMakeFiles/vecmem_core.dir/src/memory/binary_page_memory_resource_impl.cpp.o
[ 32%] Building CXX object core/CMakeFiles/vecmem_core.dir/src/memory/contiguous_memory_resource.cpp.o
[ 35%] Building CXX object core/CMakeFiles/vecmem_core.dir/src/memory/instrumenting_memory_resource.cpp.o
[ 37%] Building CXX object core/CMakeFiles/vecmem_core.dir/src/memory/choice_memory_resource.cpp.o
[ 40%] Building CXX object core/CMakeFiles/vecmem_core.dir/src/memory/coalescing_memory_resource.cpp.o
[ 43%] Building CXX object core/CMakeFiles/vecmem_core.dir/src/memory/conditional_memory_resource.cpp.o
[ 45%] Building CXX object core/CMakeFiles/vecmem_core.dir/src/memory/debug_memory_resource.cpp.o
[ 48%] Building CXX object core/CMakeFiles/vecmem_core.dir/src/memory/details/is_aligned.cpp.o
[ 51%] Building CXX object core/CMakeFiles/vecmem_core.dir/src/utils/copy.cpp.o
[ 54%] Building CXX object core/CMakeFiles/vecmem_core.dir/src/utils/memory_monitor.cpp.o
[ 56%] Linking CXX shared library ../lib/libvecmem_core.so
[ 56%] Built target vecmem_core
Scanning dependencies of target vecmem_sycl
[ 59%] Building SYCL object sycl/CMakeFiles/vecmem_sycl.dir/src/memory/sycl/details/memory_resource_base.sycl.o
[ 62%] Building SYCL object sycl/CMakeFiles/vecmem_sycl.dir/src/memory/sycl/device_memory_resource.sycl.o
[ 64%] Building SYCL object sycl/CMakeFiles/vecmem_sycl.dir/src/memory/sycl/host_memory_resource.sycl.o
[ 67%] Building SYCL object sycl/CMakeFiles/vecmem_sycl.dir/src/memory/sycl/shared_memory_resource.sycl.o
[ 70%] Building SYCL object sycl/CMakeFiles/vecmem_sycl.dir/src/utils/sycl/copy.sycl.o
[ 72%] Building SYCL object sycl/CMakeFiles/vecmem_sycl.dir/src/utils/sycl/queue_wrapper.sycl.o
[ 75%] Building SYCL object sycl/CMakeFiles/vecmem_sycl.dir/src/utils/sycl/get_queue.sycl.o
[ 78%] Linking SYCL shared library ../lib/libvecmem_sycl.so
[ 78%] Built target vecmem_sycl
Scanning dependencies of target vecmem_testing_common
[ 81%] Building CXX object tests/CMakeFiles/vecmem_testing_common.dir/common/memory_resource_name_gen.cpp.o
[ 83%] Linking CXX static library ../lib/libvecmem_testing_common.a
[ 83%] Built target vecmem_testing_common
Scanning dependencies of target gtest_main
[ 86%] Building CXX object _deps/googletest-build/googletest/CMakeFiles/gtest_main.dir/src/gtest_main.cc.o
[ 89%] Linking CXX shared library ../../../lib/libgtest_maind.so
[ 89%] Built target gtest_main
Scanning dependencies of target vecmem_test_sycl
[ 91%] Building CXX object tests/sycl/CMakeFiles/vecmem_test_sycl.dir/test_sycl_memory_resources.cpp.o
[ 94%] Building SYCL object tests/sycl/CMakeFiles/vecmem_test_sycl.dir/test_sycl_containers.sycl.o
[ 97%] Building SYCL object tests/sycl/CMakeFiles/vecmem_test_sycl.dir/test_sycl_jagged_containers.sycl.o
[100%] Linking SYCL executable ../../bin/vecmem_test_sycl
ptxas /tmp/test_sycl_memory_resources-sm_50-1e6b62.s, line 5875; fatal : Parsing error near '.': syntax error
ptxas fatal : Ptx assembly aborted due to errors
llvm-foreach:
ptxas /tmp/test_sycl_memory_resources-sm_50-003e59.s, line 5947; fatal : Parsing error near '.': syntax error
ptxas fatal : Ptx assembly aborted due to errors
llvm-foreach:
clang-15: error: ptxas command failed with exit code 255 (use -v to see invocation)
clang version 15.0.0 (https://github.com/intel/llvm.git 4043dda356af59d3d88607037955a0728dc0f466)
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /usr/local/bin
clang-15: note: diagnostic msg: Error generating preprocessed source(s).
tests/sycl/CMakeFiles/vecmem_test_sycl.dir/build.make:122: recipe for target 'bin/vecmem_test_sycl' failed
make[3]: *** [bin/vecmem_test_sycl] Error 1
CMakeFiles/Makefile2:1339: recipe for target 'tests/sycl/CMakeFiles/vecmem_test_sycl.dir/all' failed
make[2]: *** [tests/sycl/CMakeFiles/vecmem_test_sycl.dir/all] Error 2
CMakeFiles/Makefile2:1346: recipe for target 'tests/sycl/CMakeFiles/vecmem_test_sycl.dir/rule' failed
make[1]: *** [tests/sycl/CMakeFiles/vecmem_test_sycl.dir/rule] Error 2
Makefile:698: recipe for target 'vecmem_test_sycl' failed
make: *** [vecmem_test_sycl] Error 2
root@874effd59f61:/mnt/ssd1/krasznaa/projects/vecmem#This is on the same (128 GB memory) machine.
krasznaa
commented
1 year ago Finally, since I wasn't absolutely 100% sure, I tried what happens if I attempt the build as the superuser, with all limits removed. (ulimit -u unlimited) And the build crashes exactly the same way. So I really don't think this is because of some hardware resource limit. Though even at this point I wouldn't bet my life on it. Since even after the ulimit -u unlimited call I see:
[bash][pcadp04]:vecmem > ulimit -a
real-time non-blocking time (microseconds, -R) unlimited
core file size (blocks, -c) 0
data seg size (kbytes, -d) unlimited
scheduling priority (-e) 0
file size (blocks, -f) unlimited
pending signals (-i) 512314
max locked memory (kbytes, -l) 64
max memory size (kbytes, -m) unlimited
open files (-n) 1024
pipe size (512 bytes, -p) 8
POSIX message queues (bytes, -q) 819200
real-time priority (-r) 0
stack size (kbytes, -s) 8192
cpu time (seconds, -t) unlimited
max user processes (-u) unlimited
virtual memory (kbytes, -v) unlimited
file locks (-x) unlimited
[bash][pcadp04]:vecmem >Aha I was pulling the wrong container. I have replicated the error! Sorry for confusion. This is potentially a ptxas problem which might be specific to CUDA 10.2. I will try to see if I can replicate outside of docker with ptxas 10.2 and then compare with another version. If not then I will have to dig some more
krasznaa
commented
1 year ago As I said a few times, I've been seeing this behaviour with various CUDA 11.X versions as well. If you just look at the log I posted earlier today, I demonstrated the issue there with CUDA 11.7.1.
We are in the process of making a new set of Docker images for our CIs that will be based on CUDA 11.6.
https://github.com/acts-project/machines/pull/75
If that helps convince you, I'll up the urgency on that one, so that you could test the failure with such a pre-built image as well.
hdelan
commented
1 year ago Hi yes bug diagnosed. A problem with the PTX debug info generated by DPC++ when assert is used in a kernel. Investigating a fix. Problem not with ptxas or any specific CUDA version.
hdelan
commented
1 year ago Tuns out that it's a problem with the __func__ arg in __assertfail. Working on a fix. A temporary workaround would be something like:
SYCL_EXTERNAL extern "C" void __assertfail(const char * msg, const char * file, unsigned line,
const char * func, size_t char_sz);
#ifdef __NVPTX__
#define MY_ASSERT(COND) \
if (!COND) \
__assertfail("Kernel assert", __FILE__, __LINE__, "", 1);
#else
#define MY_ASSERT(COND) \
assert(COND);
#endifThanks for the info! The problem being with how assertions are implemented makes a lot of sense. (As to why issues would only show up in Debug builds for me.)
It's a pity that right now one can only work around this by introducing a custom assertion macro. :frowning: I was hoping that I could inject some code into our project similar to what's described in https://github.com/intel/llvm/issues/3385, and what we already do in: https://github.com/acts-project/vecmem/blob/main/sycl/src/utils/sycl/cuda_assert_polyfill.sycl
But as I understand it, it's not possible to fix this issue behind the scenes like that. :frowning: (I tried for a bit just to be sure, but failed.)
We'll just be waiting for a fix in the compiler then. Because I'm not keen on replacing all our assert(...) calls with something else.
hdelan
commented
1 year ago OK thanks. I am going on holiday at the end of the week but if I don't have a fix by then I will hand this off to one of my team mates.
The problem is not actually strictly with assert, but with the __func__/__FUNCTION__/__PRETTY_FUNCTION__ macros - you should see the same failure with:
// Called from a SYCL kernel
printf("%s\n", __func__);When compiling a debug build for cuda backend.
ldrumm
commented
1 year ago Hi. I took this off @hdelan's hands while on holiday. A quick update before I go on holiday: I've found the root cause of this bug and I have a fix, but it's hacky and I want to clean it up a bit.
In short: The NVPTXAssignValidGlobalNames pass runs in the same passmanager as the NVPTXAsmPrinter. This is all as part of the TargetMachine::addPassesToEmitFile step executed by the clang driver.
NVPTXAsmPrinter::doInitialization is called by the parent passmanager before any pass contained is run on the module proper. This is normal, and runs once per passmanager run.
NVPTXAsmPrinter::doInitialization does some checks on the module, then runs the parent doInitialization. The parent doInitialization then constructs all the MCSymbolRefExprs referred to by the debug info, and the symbol name strings are interned by the MCContext.
After the pass manager has initialized all passes, it runs them. NVPTXAssignValidGlobalNames::runOnModule runs, and changes the global symbol names as expected. The MCSymbolRefExprs still implicitly referenced do not have their symbol names changed.
Later, the AsmPrinter prints two different symbol names - one converted (the global), and one invalid (the old global), and ptxas terminates with an error.
Here is the quick hack:
diff --git a/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXTargetStreamer.cpp b/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXTargetStreamer.cpp
index b72cea5d03f1..b67e2258868c 100644
--- a/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXTargetStreamer.cpp
+++ b/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXTargetStreamer.cpp
@@ -13,6 +13,7 @@
#include "NVPTXTargetStreamer.h"
#include "llvm/MC/MCAsmInfo.h"
#include "llvm/MC/MCContext.h"
+#include "llvm/MC/MCExpr.h"
#include "llvm/MC/MCObjectFileInfo.h"
using namespace llvm;
@@ -102,6 +103,26 @@ void NVPTXTargetStreamer::changeSection(const MCSection *CurSection,
}
}
+void NVPTXTargetStreamer::emitValue(const MCExpr *Value) {
+ if (Value->getKind() != MCExpr::ExprKind::SymbolRef)
+ return MCTargetStreamer::emitValue(Value);
+
+ SmallString<128> Str;
+ raw_svector_ostream OS(Str);
+ Value->print(OS, Streamer.getContext().getAsmInfo());
+
+ std::string ValidName;
+ raw_string_ostream ValidNameStream(ValidName);
+ for (char C : Str) {
+ if (C == '.' || C == '@') {
+ ValidNameStream << "_$_";
+ } else {
+ ValidNameStream << C;
+ }
+ }
+ Streamer.emitRawText((ValidNameStream.str()));
+}
+
void NVPTXTargetStreamer::emitRawBytes(StringRef Data) {
MCTargetStreamer::emitRawBytes(Data);
// TODO: enable this once the bug in the ptxas with the packed bytes is
diff --git a/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXTargetStreamer.h b/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXTargetStreamer.h
index 8185efadefdb..3fffe0b048b6 100644
--- a/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXTargetStreamer.h
+++ b/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXTargetStreamer.h
@@ -48,6 +48,7 @@ public:
///
/// This is used to emit bytes in \p Data as sequence of .byte directives.
void emitRawBytes(StringRef Data) override;
+ void emitValue(const MCExpr *Value) override;
};
} // end namespace llvm
Bear with me and I'll post a cleaner fix and proper test-case that has a chance of being merged upstream ASAP
ldrumm
commented
1 year ago As far as I can tell, a variation on the above hack is the best way to go due to strangeness with the way debug info is collected before assembly emission.
The upstream review is here
wangzy0327
commented
6 months ago @ldrumm I met the same problem when I compiled the oneDNN(v3.2) with cuda(11.2). My sycl version is 2022-06-release.
I modified the two files llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXTargetStreamer.cpp and llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXTargetStreamer.h as you listed above with the modifications. I modified it and recompiled SYCL, then compiled oneDNN.
The output of make -j as follow.
[100%] Linking CXX shared library libdnnl.so
llvm-foreach: Segmentation fault (core dumped)
clang-15: error: ptxas command failed with exit code 254 (use -v to see invocation)
clang version 15.0.0 (ssh://git@gitlab.gxnzx12729.ict:2222/wangziyang/intel-llvm-new.git 7ecb566e497fa926844521e8df2a2405c7e92e63)
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /home/wzy/sycl_workspace/build-cuda-2022-06/bin
clang-15: note: diagnostic msg: Error generating preprocessed source(s).
src/CMakeFiles/dnnl.dir/build.make:776: recipe for target 'src/libdnnl.so.3.2' failed
make[2]: *** [src/libdnnl.so.3.2] Error 1
CMakeFiles/Makefile2:355: recipe for target 'src/CMakeFiles/dnnl.dir/all' failed
make[1]: *** [src/CMakeFiles/dnnl.dir/all] Error 2
Makefile:159: recipe for target 'all' failed
make: *** [all] Error 2It's still the same error as before.
wangzy0327
commented
5 months ago Describe the bug
This issue has been hurting us for a while already, but since the compiler's behaviour changed a little bit recently, I thought it would finally be time to try to address it.
When trying to build our code for the CUDA backend in debug mode, we've been getting the following failures since a while:
... [100%] Linking SYCL shared library ../../lib/x86_64-linux-gnu/libtraccc_sycl.so llvm-foreach: Segmentation fault (core dumped) llvm-foreach: Segmentation fault (core dumped) ptxas fatal : Memory allocation failure llvm-foreach: clang-14: error: ptxas command failed with exit code 255 (use -v to see invocation) clang version 14.0.0 (https://github.com/intel/llvm.git bd68232bb96386bf7649345c0557ba520e73c02d) Target: x86_64-unknown-linux-gnu Thread model: posix InstalledDir: /atlas/software/intel/clang/2021-09/x86_64-ubuntu1804-gcc8-opt/bin clang-14: note: diagnostic msg: Error generating preprocessed source(s). device/sycl/CMakeFiles/traccc_sycl.dir/build.make:203: recipe for target 'lib/x86_64-linux-gnu/libtraccc_sycl.so.0.1.0' failed make[3]: *** [lib/x86_64-linux-gnu/libtraccc_sycl.so.0.1.0] Error 255 CMakeFiles/Makefile2:14544: recipe for target 'device/sycl/CMakeFiles/traccc_sycl.dir/all' failed make[2]: *** [device/sycl/CMakeFiles/traccc_sycl.dir/all] Error 2 CMakeFiles/Makefile2:14551: recipe for target 'device/sycl/CMakeFiles/traccc_sycl.dir/rule' failed make[1]: *** [device/sycl/CMakeFiles/traccc_sycl.dir/rule] Error 2 Makefile:5945: recipe for target 'traccc_sycl' failed make: *** [traccc_sycl] Error 2And:
... [100%] Linking SYCL shared library ../../lib/x86_64-linux-gnu/libtraccc_sycl.so llvm-foreach: Segmentation fault (core dumped) llvm-foreach: Segmentation fault (core dumped) llvm-foreach: Segmentation fault (core dumped) llvm-foreach: Segmentation fault (core dumped) llvm-foreach: Segmentation fault (core dumped) llvm-foreach: Segmentation fault (core dumped) ptxas fatal : Memory allocation failure llvm-foreach: clang-14: error: ptxas command failed with exit code 255 (use -v to see invocation) clang version 14.0.0 (https://github.com/intel/llvm.git 27f59d8906fcc8aece7ff6aa570ccdee52168c2d) Target: x86_64-unknown-linux-gnu Thread model: posix InstalledDir: /atlas/software/intel/clang/2021-12/x86_64-ubuntu1804-gcc8-opt/bin clang-14: note: diagnostic msg: Error generating preprocessed source(s). device/sycl/CMakeFiles/traccc_sycl.dir/build.make:202: recipe for target 'lib/x86_64-linux-gnu/libtraccc_sycl.so.0.1.0' failed make[3]: *** [lib/x86_64-linux-gnu/libtraccc_sycl.so.0.1.0] Error 255 CMakeFiles/Makefile2:14515: recipe for target 'device/sycl/CMakeFiles/traccc_sycl.dir/all' failed make[2]: *** [device/sycl/CMakeFiles/traccc_sycl.dir/all] Error 2 CMakeFiles/Makefile2:14522: recipe for target 'device/sycl/CMakeFiles/traccc_sycl.dir/rule' failed make[1]: *** [device/sycl/CMakeFiles/traccc_sycl.dir/rule] Error 2 Makefile:5932: recipe for target 'traccc_sycl' failed make: *** [traccc_sycl] Error 2(With 2021-09 and 2021-12 respectively.)
Now, with the latest nightlies, the linker became a bit more chatty, and fails in the following (very long) way:
... [100%] Linking SYCL shared library ../../lib/x86_64-linux-gnu/libtraccc_sycl.so clang-15: /atlas/krasznaa/intel/llvm/llvm/include/llvm/Support/Casting.h:104: static bool llvm::isa_impl_cl<To, const From*>::doit(const From*) [with To = llvm::ConstantAsMetadata; From = llvm::Metadata]: Assertion `Val && "isa<> used on a null pointer"' 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: /atlas/software/intel/clang/nightly-20220402/x86_64-ubuntu1804-gcc8-opt/bin/clang-15 -cc1 -triple nvptx64-nvidia-cuda -aux-triple x86_64-unknown-linux-gnu -fsycl-is-device -fdeclare-spirv-builtins -fenable-sycl-dae -Wno-sycl-strict -sycl-std=2020 -S -disable-free -clear-ast-before-backend -main-file-name counting_grid_capacities.sycl.o -mrelocation-model pic -pic-level 2 -fhalf-no-semantic-interposition -mframe-pointer=all -ffp-contract=on -fno-rounding-math -fno-verbose-asm -no-integrated-as -aux-target-cpu x86-64 -internal-isystem /atlas/software/intel/clang/nightly-20220402/x86_64-ubuntu1804-gcc8-opt/bin/../include/sycl -internal-isystem /atlas/software/intel/clang/nightly-20220402/x86_64-ubuntu1804-gcc8-opt/bin/../include -mlink-builtin-bitcode /atlas/software/intel/clang/nightly-20220402/x86_64-ubuntu1804-gcc8-opt/lib/clang/15.0.0/../../clc/remangled-l64-signed_char.libspirv-nvptx64--nvidiacl.bc -mlink-builtin-bitcode /atlas/software/cuda/11.5.2/x86_64-ubuntu1804/nvvm/libdevice/libdevice.10.bc -target-feature +ptx75 -target-sdk-version=11.5 -target-cpu sm_50 -target-feature +ptx75 -mllvm -treat-scalable-fixed-error-as-warning -mllvm -sycl-enable-local-accessor -debug-info-kind=constructor -dwarf-version=2 -debugger-tuning=gdb -fno-dwarf-directory-asm -resource-dir /atlas/software/intel/clang/nightly-20220402/x86_64-ubuntu1804-gcc8-opt/lib/clang/15.0.0 -Wno-linker-warnings -Wall -Wextra -Wno-unknown-cuda-version -Wshadow -Wunused-local-typedefs -Werror -pedantic -std=c++17 -fdebug-compilation-dir=/home/krasznaa/ATLAS/projects/traccc/build/device/sycl -ferror-limit 19 -fgnuc-version=4.2.1 -fcolor-diagnostics -D__GCC_HAVE_DWARF2_CFI_ASM=1 -o /tmp/counting_grid_capacities-sm_50-eb0ce2.s -x ir /tmp/counting_grid_capacities-46cbbd/counting_grid_capacities-sm_50_0.bc 1. Code generation 2. Running pass 'Function Pass Manager' on module '/tmp/counting_grid_capacities-46cbbd/counting_grid_capacities-sm_50_0.bc'. 3. Running pass 'NVPTX DAG->DAG Pattern Instruction Selection' on function '@_ZTSN6traccc4sycl19CountGridCapacitiesE' #0 0x000055dc444c36ff PrintStackTraceSignalHandler(void*) Signals.cpp:0:0 #1 0x000055dc444c0db4 SignalHandler(int) Signals.cpp:0:0 #2 0x00007f52c969a980 __restore_rt (/lib/x86_64-linux-gnu/libpthread.so.0+0x12980) #3 0x00007f52c834be87 raise /build/glibc-uZu3wS/glibc-2.27/signal/../sysdeps/unix/sysv/linux/raise.c:51:0 #4 0x00007f52c834d7f1 abort /build/glibc-uZu3wS/glibc-2.27/stdlib/abort.c:81:0 #5 0x00007f52c833d3fa __assert_fail_base /build/glibc-uZu3wS/glibc-2.27/assert/assert.c:89:0 #6 0x00007f52c833d472 (/lib/x86_64-linux-gnu/libc.so.6+0x30472) #7 0x000055dc43212780 llvm::NVPTXTargetLowering::getFunctionParamOptimizedAlign(llvm::Function const*, llvm::Type*, llvm::DataLayout const&) const (.constprop.515) NVPTXISelLowering.cpp:0:0 #8 0x000055dc4321299e llvm::NVPTXTargetLowering::getArgumentAlignment(llvm::SDValue, llvm::CallBase const*, llvm::Type*, unsigned int, llvm::DataLayout const&) const (.constprop.514) NVPTXISelLowering.cpp:0:0 #9 0x000055dc4321ac00 llvm::NVPTXTargetLowering::LowerCall(llvm::TargetLowering::CallLoweringInfo&, llvm::SmallVectorImpl<llvm::SDValue>&) const (/atlas/software/intel/clang/nightly-20220402/x86_64-ubuntu1804-gcc8-opt/bin/clang-15+0x1281c00) #10 0x000055dc453acd75 llvm::TargetLowering::LowerCallTo(llvm::TargetLowering::CallLoweringInfo&) const (/atlas/software/intel/clang/nightly-20220402/x86_64-ubuntu1804-gcc8-opt/bin/clang-15+0x3413d75) #11 0x000055dc453b7ad1 llvm::SelectionDAGBuilder::lowerInvokable(llvm::TargetLowering::CallLoweringInfo&, llvm::BasicBlock const*) (/atlas/software/intel/clang/nightly-20220402/x86_64-ubuntu1804-gcc8-opt/bin/clang-15+0x341ead1) #12 0x000055dc453d4cde llvm::SelectionDAGBuilder::LowerCallTo(llvm::CallBase const&, llvm::SDValue, bool, bool, llvm::BasicBlock const*) (/atlas/software/intel/clang/nightly-20220402/x86_64-ubuntu1804-gcc8-opt/bin/clang-15+0x343bcde) #13 0x000055dc453c40ef llvm::SelectionDAGBuilder::visitCall(llvm::CallInst const&) (/atlas/software/intel/clang/nightly-20220402/x86_64-ubuntu1804-gcc8-opt/bin/clang-15+0x342b0ef) #14 0x000055dc453eacd1 llvm::SelectionDAGBuilder::visit(llvm::Instruction const&) (/atlas/software/intel/clang/nightly-20220402/x86_64-ubuntu1804-gcc8-opt/bin/clang-15+0x3451cd1) #15 0x000055dc45457f81 llvm::SelectionDAGISel::SelectBasicBlock(llvm::ilist_iterator<llvm::ilist_detail::node_options<llvm::Instruction, true, false, void>, false, true>, llvm::ilist_iterator<llvm::ilist_detail::node_options<llvm::Instruction, true, false, void>, false, true>, bool&) (/atlas/software/intel/clang/nightly-20220402/x86_64-ubuntu1804-gcc8-opt/bin/clang-15+0x34bef81) #16 0x000055dc4545959b llvm::SelectionDAGISel::SelectAllBasicBlocks(llvm::Function const&) (/atlas/software/intel/clang/nightly-20220402/x86_64-ubuntu1804-gcc8-opt/bin/clang-15+0x34c059b) #17 0x000055dc4545ae7a llvm::SelectionDAGISel::runOnMachineFunction(llvm::MachineFunction&) (.part.1050) SelectionDAGISel.cpp:0:0 #18 0x000055dc43698e45 llvm::MachineFunctionPass::runOnFunction(llvm::Function&) (/atlas/software/intel/clang/nightly-20220402/x86_64-ubuntu1804-gcc8-opt/bin/clang-15+0x16ffe45) #19 0x000055dc43bb8d80 llvm::FPPassManager::runOnFunction(llvm::Function&) (/atlas/software/intel/clang/nightly-20220402/x86_64-ubuntu1804-gcc8-opt/bin/clang-15+0x1c1fd80) #20 0x000055dc43bb8f79 llvm::FPPassManager::runOnModule(llvm::Module&) (/atlas/software/intel/clang/nightly-20220402/x86_64-ubuntu1804-gcc8-opt/bin/clang-15+0x1c1ff79) #21 0x000055dc43bb9b67 llvm::legacy::PassManagerImpl::run(llvm::Module&) (/atlas/software/intel/clang/nightly-20220402/x86_64-ubuntu1804-gcc8-opt/bin/clang-15+0x1c20b67) #22 0x000055dc44812726 (anonymous namespace)::EmitAssemblyHelper::EmitAssemblyWithLegacyPassManager(clang::BackendAction, std::unique_ptr<llvm::raw_pwrite_stream, std::default_delete<llvm::raw_pwrite_stream> >) BackendUtil.cpp:0:0 #23 0x000055dc44816e72 clang::EmitBackendOutput(clang::DiagnosticsEngine&, clang::HeaderSearchOptions const&, clang::CodeGenOptions const&, clang::TargetOptions const&, clang::LangOptions const&, llvm::StringRef, llvm::Module*, clang::BackendAction, std::unique_ptr<llvm::raw_pwrite_stream, std::default_delete<llvm::raw_pwrite_stream> >) (/atlas/software/intel/clang/nightly-20220402/x86_64-ubuntu1804-gcc8-opt/bin/clang-15+0x287de72) #24 0x000055dc4557e285 clang::CodeGenAction::ExecuteAction() (/atlas/software/intel/clang/nightly-20220402/x86_64-ubuntu1804-gcc8-opt/bin/clang-15+0x35e5285) #25 0x000055dc44f2e479 clang::FrontendAction::Execute() (/atlas/software/intel/clang/nightly-20220402/x86_64-ubuntu1804-gcc8-opt/bin/clang-15+0x2f95479) #26 0x000055dc44ec2f61 clang::CompilerInstance::ExecuteAction(clang::FrontendAction&) (/atlas/software/intel/clang/nightly-20220402/x86_64-ubuntu1804-gcc8-opt/bin/clang-15+0x2f29f61) #27 0x000055dc44ffd163 clang::ExecuteCompilerInvocation(clang::CompilerInstance*) (/atlas/software/intel/clang/nightly-20220402/x86_64-ubuntu1804-gcc8-opt/bin/clang-15+0x3064163) #28 0x000055dc42e8346f cc1_main(llvm::ArrayRef<char const*>, char const*, void*) (/atlas/software/intel/clang/nightly-20220402/x86_64-ubuntu1804-gcc8-opt/bin/clang-15+0xeea46f) #29 0x000055dc42e7fb07 ExecuteCC1Tool(llvm::SmallVectorImpl<char const*>&) driver.cpp:0:0 #30 0x000055dc42e14b84 main (/atlas/software/intel/clang/nightly-20220402/x86_64-ubuntu1804-gcc8-opt/bin/clang-15+0xe7bb84) #31 0x00007f52c832ec87 __libc_start_main /build/glibc-uZu3wS/glibc-2.27/csu/../csu/libc-start.c:344:0 #32 0x000055dc42e7f59a _start (/atlas/software/intel/clang/nightly-20220402/x86_64-ubuntu1804-gcc8-opt/bin/clang-15+0xee659a) llvm-foreach: Aborted (core dumped) ... clang-15: /atlas/krasznaa/intel/llvm/llvm/include/llvm/Support/Casting.h:104: static bool llvm::isa_impl_cl<To, const From*>::doit(const From*) [with To = llvm::ConstantAsMetadata; From = llvm::Metadata]: Assertion `Val && "isa<> used on a null pointer"' 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: /atlas/software/intel/clang/nightly-20220402/x86_64-ubuntu1804-gcc8-opt/bin/clang-15 -cc1 -triple nvptx64-nvidia-cuda -aux-triple x86_64-unknown-linux-gnu -fsycl-is-device -fdeclare-spirv-builtins -fenable-sycl-dae -Wno-sycl-strict -sycl-std=2020 -S -disable-free -clear-ast-before-backend -main-file-name counting_grid_capacities.sycl.o -mrelocation-model pic -pic-level 2 -fhalf-no-semantic-interposition -mframe-pointer=all -ffp-contract=on -fno-rounding-math -fno-verbose-asm -no-integrated-as -aux-target-cpu x86-64 -internal-isystem /atlas/software/intel/clang/nightly-20220402/x86_64-ubuntu1804-gcc8-opt/bin/../include/sycl -internal-isystem /atlas/software/intel/clang/nightly-20220402/x86_64-ubuntu1804-gcc8-opt/bin/../include -mlink-builtin-bitcode /atlas/software/intel/clang/nightly-20220402/x86_64-ubuntu1804-gcc8-opt/lib/clang/15.0.0/../../clc/remangled-l64-signed_char.libspirv-nvptx64--nvidiacl.bc -mlink-builtin-bitcode /atlas/software/cuda/11.5.2/x86_64-ubuntu1804/nvvm/libdevice/libdevice.10.bc -target-feature +ptx75 -target-sdk-version=11.5 -target-cpu sm_50 -target-feature +ptx75 -mllvm -treat-scalable-fixed-error-as-warning -mllvm -sycl-enable-local-accessor -debug-info-kind=constructor -dwarf-version=2 -debugger-tuning=gdb -fno-dwarf-directory-asm -resource-dir /atlas/software/intel/clang/nightly-20220402/x86_64-ubuntu1804-gcc8-opt/lib/clang/15.0.0 -Wno-linker-warnings -Wall -Wextra -Wno-unknown-cuda-version -Wshadow -Wunused-local-typedefs -Werror -pedantic -std=c++17 -fdebug-compilation-dir=/home/krasznaa/ATLAS/projects/traccc/build/device/sycl -ferror-limit 19 -fgnuc-version=4.2.1 -fcolor-diagnostics -D__GCC_HAVE_DWARF2_CFI_ASM=1 -o /tmp/counting_grid_capacities-sm_50-b46803.s -x ir /tmp/counting_grid_capacities-46cbbd/counting_grid_capacities-sm_50_8.bc 1. Code generation 2. Running pass 'Function Pass Manager' on module '/tmp/counting_grid_capacities-46cbbd/counting_grid_capacities-sm_50_8.bc'. 3. Running pass 'NVPTX DAG->DAG Pattern Instruction Selection' on function '@_ZNK6traccc4sycl12WeightUpdateclEN2cl4sycl7nd_itemILi1EEE' #0 0x000056034ea226ff PrintStackTraceSignalHandler(void*) Signals.cpp:0:0 #1 0x000056034ea1fdb4 SignalHandler(int) Signals.cpp:0:0 #2 0x00007fcba6c01980 __restore_rt (/lib/x86_64-linux-gnu/libpthread.so.0+0x12980) #3 0x00007fcba58b2e87 raise /build/glibc-uZu3wS/glibc-2.27/signal/../sysdeps/unix/sysv/linux/raise.c:51:0 #4 0x00007fcba58b47f1 abort /build/glibc-uZu3wS/glibc-2.27/stdlib/abort.c:81:0 #5 0x00007fcba58a43fa __assert_fail_base /build/glibc-uZu3wS/glibc-2.27/assert/assert.c:89:0 #6 0x00007fcba58a4472 (/lib/x86_64-linux-gnu/libc.so.6+0x30472) #7 0x000056034d771780 llvm::NVPTXTargetLowering::getFunctionParamOptimizedAlign(llvm::Function const*, llvm::Type*, llvm::DataLayout const&) const (.constprop.515) NVPTXISelLowering.cpp:0:0 #8 0x000056034d77199e llvm::NVPTXTargetLowering::getArgumentAlignment(llvm::SDValue, llvm::CallBase const*, llvm::Type*, unsigned int, llvm::DataLayout const&) const (.constprop.514) NVPTXISelLowering.cpp:0:0 #9 0x000056034d779c00 llvm::NVPTXTargetLowering::LowerCall(llvm::TargetLowering::CallLoweringInfo&, llvm::SmallVectorImpl<llvm::SDValue>&) const (/atlas/software/intel/clang/nightly-20220402/x86_64-ubuntu1804-gcc8-opt/bin/clang-15+0x1281c00) #10 0x000056034f90bd75 llvm::TargetLowering::LowerCallTo(llvm::TargetLowering::CallLoweringInfo&) const (/atlas/software/intel/clang/nightly-20220402/x86_64-ubuntu1804-gcc8-opt/bin/clang-15+0x3413d75) #11 0x000056034f916ad1 llvm::SelectionDAGBuilder::lowerInvokable(llvm::TargetLowering::CallLoweringInfo&, llvm::BasicBlock const*) (/atlas/software/intel/clang/nightly-20220402/x86_64-ubuntu1804-gcc8-opt/bin/clang-15+0x341ead1) #12 0x000056034f933cde llvm::SelectionDAGBuilder::LowerCallTo(llvm::CallBase const&, llvm::SDValue, bool, bool, llvm::BasicBlock const*) (/atlas/software/intel/clang/nightly-20220402/x86_64-ubuntu1804-gcc8-opt/bin/clang-15+0x343bcde) #13 0x000056034f9230ef llvm::SelectionDAGBuilder::visitCall(llvm::CallInst const&) (/atlas/software/intel/clang/nightly-20220402/x86_64-ubuntu1804-gcc8-opt/bin/clang-15+0x342b0ef) #14 0x000056034f949cd1 llvm::SelectionDAGBuilder::visit(llvm::Instruction const&) (/atlas/software/intel/clang/nightly-20220402/x86_64-ubuntu1804-gcc8-opt/bin/clang-15+0x3451cd1) #15 0x000056034f9b6f81 llvm::SelectionDAGISel::SelectBasicBlock(llvm::ilist_iterator<llvm::ilist_detail::node_options<llvm::Instruction, true, false, void>, false, true>, llvm::ilist_iterator<llvm::ilist_detail::node_options<llvm::Instruction, true, false, void>, false, true>, bool&) (/atlas/software/intel/clang/nightly-20220402/x86_64-ubuntu1804-gcc8-opt/bin/clang-15+0x34bef81) #16 0x000056034f9b859b llvm::SelectionDAGISel::SelectAllBasicBlocks(llvm::Function const&) (/atlas/software/intel/clang/nightly-20220402/x86_64-ubuntu1804-gcc8-opt/bin/clang-15+0x34c059b) #17 0x000056034f9b9e7a llvm::SelectionDAGISel::runOnMachineFunction(llvm::MachineFunction&) (.part.1050) SelectionDAGISel.cpp:0:0 #18 0x000056034dbf7e45 llvm::MachineFunctionPass::runOnFunction(llvm::Function&) (/atlas/software/intel/clang/nightly-20220402/x86_64-ubuntu1804-gcc8-opt/bin/clang-15+0x16ffe45) #19 0x000056034e117d80 llvm::FPPassManager::runOnFunction(llvm::Function&) (/atlas/software/intel/clang/nightly-20220402/x86_64-ubuntu1804-gcc8-opt/bin/clang-15+0x1c1fd80) #20 0x000056034e117f79 llvm::FPPassManager::runOnModule(llvm::Module&) (/atlas/software/intel/clang/nightly-20220402/x86_64-ubuntu1804-gcc8-opt/bin/clang-15+0x1c1ff79) #21 0x000056034e118b67 llvm::legacy::PassManagerImpl::run(llvm::Module&) (/atlas/software/intel/clang/nightly-20220402/x86_64-ubuntu1804-gcc8-opt/bin/clang-15+0x1c20b67) #22 0x000056034ed71726 (anonymous namespace)::EmitAssemblyHelper::EmitAssemblyWithLegacyPassManager(clang::BackendAction, std::unique_ptr<llvm::raw_pwrite_stream, std::default_delete<llvm::raw_pwrite_stream> >) BackendUtil.cpp:0:0 #23 0x000056034ed75e72 clang::EmitBackendOutput(clang::DiagnosticsEngine&, clang::HeaderSearchOptions const&, clang::CodeGenOptions const&, clang::TargetOptions const&, clang::LangOptions const&, llvm::StringRef, llvm::Module*, clang::BackendAction, std::unique_ptr<llvm::raw_pwrite_stream, std::default_delete<llvm::raw_pwrite_stream> >) (/atlas/software/intel/clang/nightly-20220402/x86_64-ubuntu1804-gcc8-opt/bin/clang-15+0x287de72) #24 0x000056034fadd285 clang::CodeGenAction::ExecuteAction() (/atlas/software/intel/clang/nightly-20220402/x86_64-ubuntu1804-gcc8-opt/bin/clang-15+0x35e5285) #25 0x000056034f48d479 clang::FrontendAction::Execute() (/atlas/software/intel/clang/nightly-20220402/x86_64-ubuntu1804-gcc8-opt/bin/clang-15+0x2f95479) #26 0x000056034f421f61 clang::CompilerInstance::ExecuteAction(clang::FrontendAction&) (/atlas/software/intel/clang/nightly-20220402/x86_64-ubuntu1804-gcc8-opt/bin/clang-15+0x2f29f61) #27 0x000056034f55c163 clang::ExecuteCompilerInvocation(clang::CompilerInstance*) (/atlas/software/intel/clang/nightly-20220402/x86_64-ubuntu1804-gcc8-opt/bin/clang-15+0x3064163) #28 0x000056034d3e246f cc1_main(llvm::ArrayRef<char const*>, char const*, void*) (/atlas/software/intel/clang/nightly-20220402/x86_64-ubuntu1804-gcc8-opt/bin/clang-15+0xeea46f) #29 0x000056034d3deb07 ExecuteCC1Tool(llvm::SmallVectorImpl<char const*>&) driver.cpp:0:0 #30 0x000056034d373b84 main (/atlas/software/intel/clang/nightly-20220402/x86_64-ubuntu1804-gcc8-opt/bin/clang-15+0xe7bb84) #31 0x00007fcba5895c87 __libc_start_main /build/glibc-uZu3wS/glibc-2.27/csu/../csu/libc-start.c:344:0 #32 0x000056034d3de59a _start (/atlas/software/intel/clang/nightly-20220402/x86_64-ubuntu1804-gcc8-opt/bin/clang-15+0xee659a) llvm-foreach: Aborted (core dumped) clang-15: error: clang frontend command failed with exit code 254 (use -v to see invocation) clang version 15.0.0 (https://github.com/intel/llvm.git 7fe81de64304f5aba4ce4582d4837d24b4d23fe9) Target: x86_64-unknown-linux-gnu Thread model: posix InstalledDir: /atlas/software/intel/clang/nightly-20220402/x86_64-ubuntu1804-gcc8-opt/bin clang-15: note: diagnostic msg: Error generating preprocessed source(s). device/sycl/CMakeFiles/traccc_sycl.dir/build.make:202: recipe for target 'lib/x86_64-linux-gnu/libtraccc_sycl.so.0.1.0' failed make[3]: *** [lib/x86_64-linux-gnu/libtraccc_sycl.so.0.1.0] Error 254 CMakeFiles/Makefile2:14515: recipe for target 'device/sycl/CMakeFiles/traccc_sycl.dir/all' failed make[2]: *** [device/sycl/CMakeFiles/traccc_sycl.dir/all] Error 2 CMakeFiles/Makefile2:14522: recipe for target 'device/sycl/CMakeFiles/traccc_sycl.dir/rule' failed make[1]: *** [device/sycl/CMakeFiles/traccc_sycl.dir/rule] Error 2 Makefile:5932: recipe for target 'traccc_sycl' failed make: *** [traccc_sycl] Error 2(I had to shorten the output, as the same messages seem to repeat 8-10 times...)
To Reproduce
Unfortunately reproducing it is not completely trivial. 😦 This linking error happens while trying to build https://github.com/acts-project/traccc with:
export CC=`which clang` export CXX=`which clang++` export SYCLCXX="${CXX} -fsycl -fsycl-targets=nvptx64-nvidia-cuda -Wno-linker-warnings" export CUDAFLAGS="${CUDAFLAGS} -allow-unsupported-compiler" export CUDAHOSTCXX=${CXX}I can give a longer description if there will be somebody willing to give this a try...
Environment (please complete the following information)
- OS: Ubuntu 18.04 and 20.04
- Target device and vendor: NVIDIA GPU
- DPC++ version: Any version from the last 6 months or so.
- Dependencies version: Tried with both CUDA 11.4 and 11.5.
Additional context
Building the code in Debug mode for Intel (OpenCL and Level-0) backends works fine. Building the code in Debug mode for the HIP backend fails for a very different reason:
... [100%] Linking SYCL shared library ../../lib/libtraccc_sycl.so lld: error: undefined hidden symbol: __assert_fail >>> referenced by container_base.hpp:128 (/data/ssd-1tb/projects/traccc/traccc/core/include/traccc/edm/details/container_base.hpp:128) >>> lto.tmp:(typeinfo name for traccc::sycl::CountGridCapacities) >>> referenced by container_base.hpp:128 (/data/ssd-1tb/projects/traccc/traccc/core/include/traccc/edm/details/container_base.hpp:128) >>> lto.tmp:(typeinfo name for traccc::sycl::CountGridCapacities) >>> referenced by jagged_device_vector.ipp:61 (/data/ssd-1tb/projects/traccc/build/_deps/vecmem-src/core/include/vecmem/containers/impl/jagged_device_vector.ipp:61) >>> lto.tmp:(typeinfo name for traccc::sycl::CountGridCapacities) >>> referenced 5 more times llvm-foreach: ...But that one is for a different issue... 😛
Finally, the compiler flags used by CMake for this Debug build are the following:
[bash][atspot01]:build > more device/sycl/CMakeFiles/traccc_sycl.dir/flags.make # CMAKE generated file: DO NOT EDIT! # Generated by "Unix Makefiles" Generator, CMake Version 3.23 # compile CXX with /atlas/software/intel/clang/nightly-20220402/x86_64-ubuntu1804-gcc8-opt/bin/clang++ # compile SYCL with /atlas/software/intel/clang/nightly-20220402/x86_64-ubuntu1804-gcc8-opt/bin/clang++ CXX_DEFINES = -DALGEBRA_PLUGINS_INCLUDE_ARRAY -DEIGEN_NO_CUDA -DTHRUST_DEVICE_SYSTEM=THRUST_DEVICE_SYSTEM_CUDA -DTHRUST_HOST_SYSTEM=THRUST_HOST_SYSTEM_CPP -DTRACCC_CUSTOM_SCALARTYPE=float -DVECMEM_DEBUG_MSG_LVL=0 -DVECMEM_HAVE_EXPERIMENTAL_PMR_MEMORY_RESOURCE -DVECMEM_HAVE_SYCL_ATOMIC_REF -DVECMEM_SOURCE_DIR_LENGTH=44 -DVECMEM_SYCL_PRINTF_FUNCTION=cl::sycl::ext::oneapi::experimental::printf -Dtraccc_sycl_EXPORTS CXX_INCLUDES = -I/home/krasznaa/ATLAS/projects/traccc/traccc/device/sycl/include -I/home/krasznaa/ATLAS/projects/traccc/build/_deps/vecmem-build/core/CMakeFiles -I/home/krasznaa/ATLAS/projects/traccc/build/_deps/vecmem-src/core/include -I/home/krasznaa/ATLAS/projects/traccc/traccc/core/include -I/home/krasznaa/ATLAS/projects/traccc/build/_deps/detray-src/core/include -I/home/krasznaa/ATLAS/projects/traccc/build/_deps/thrust-src -I/home/krasznaa/ATLAS/projects/traccc/build/_deps/thrust-src/dependencies/cub -I/home/krasznaa/ATLAS/projects/traccc/build/_deps/acts-src/Core/include -I/home/krasznaa/ATLAS/projects/traccc/build/_deps/acts-build/Core -I/home/krasznaa/ATLAS/projects/traccc/traccc/plugins/algebra/include -I/home/krasznaa/ATLAS/projects/traccc/traccc/plugins/algebra/array/include -I/home/krasznaa/ATLAS/projects/traccc/build/_deps/algebraplugins-src/frontend/array_cmath/include -I/home/krasznaa/ATLAS/projects/traccc/build/_deps/algebraplugins-src/common/include -I/home/krasznaa/ATLAS/projects/traccc/build/_deps/algebraplugins-src/storage/array/include -I/home/krasznaa/ATLAS/projects/traccc/build/_deps/algebraplugins-src/math/cmath/include -I/home/krasznaa/ATLAS/projects/traccc/build/_deps/algebraplugins-src/math/common/include -I/home/krasznaa/ATLAS/projects/traccc/traccc/plugins/algebra/vecmem/include -I/home/krasznaa/ATLAS/projects/traccc/build/_deps/algebraplugins-src/frontend/vecmem_cmath/include -I/home/krasznaa/ATLAS/projects/traccc/build/_deps/algebraplugins-src/storage/vecmem/include -I/home/krasznaa/ATLAS/projects/traccc/traccc/plugins/algebra/eigen/include -I/home/krasznaa/ATLAS/projects/traccc/build/_deps/algebraplugins-src/frontend/eigen_eigen/include -I/home/krasznaa/ATLAS/projects/traccc/build/_deps/algebraplugins-src/storage/eigen/include -I/home/krasznaa/ATLAS/projects/traccc/build/_deps/algebraplugins-src/math/eigen/include -I/home/krasznaa/ATLAS/projects/traccc/traccc/plugins/algebra/vc/include -I/home/krasznaa/ATLAS/projects/traccc/build/_deps/algebraplugins-src/frontend/vc_vc/include -I/home/krasznaa/ATLAS/projects/traccc/build/_deps/algebraplugins-src/storage/vc/include -I/home/krasznaa/ATLAS/projects/traccc/build/_deps/algebraplugins-src/math/vc/include -I/home/krasznaa/ATLAS/projects/traccc/build/_deps/vecmem-build/sycl/CMakeFiles -I/home/krasznaa/ATLAS/projects/traccc/build/_deps/vecmem-src/sycl/include -isystem /home/krasznaa/ATLAS/projects/traccc/build/_deps/eigen3-src -isystem /atlas/software/boost/1.78.0/x86_64-ubuntu1804-gcc8-opt/include -isystem /home/krasznaa/ATLAS/projects/traccc/build/_deps/vc-src CXX_FLAGS = -Wall -Wextra -Wshadow -Wunused-local-typedefs -g -Werror -pedantic -fPIC -std=c++17 SYCL_DEFINES = -DALGEBRA_PLUGINS_INCLUDE_ARRAY -DEIGEN_NO_CUDA -DTHRUST_DEVICE_SYSTEM=THRUST_DEVICE_SYSTEM_CUDA -DTHRUST_HOST_SYSTEM=THRUST_HOST_SYSTEM_CPP -DTRACCC_CUSTOM_SCALARTYPE=float -DVECMEM_DEBUG_MSG_LVL=0 -DVECMEM_HAVE_EXPERIMENTAL_PMR_MEMORY_RESOURCE -DVECMEM_HAVE_SYCL_ATOMIC_REF -DVECMEM_SOURCE_DIR_LENGTH=44 -DVECMEM_SYCL_PRINTF_FUNCTION=cl::sycl::ext::oneapi::experimental::printf -Dtraccc_sycl_EXPORTS SYCL_INCLUDES = -I/home/krasznaa/ATLAS/projects/traccc/traccc/device/sycl/include -I/home/krasznaa/ATLAS/projects/traccc/build/_deps/vecmem-build/core/CMakeFiles -I/home/krasznaa/ATLAS/projects/traccc/build/_deps/vecmem-src/core/include -I/home/krasznaa/ATLAS/projects/traccc/traccc/core/include -I/home/krasznaa/ATLAS/projects/traccc/build/_deps/detray-src/core/include -I/home/krasznaa/ATLAS/projects/traccc/build/_deps/thrust-src -I/home/krasznaa/ATLAS/projects/traccc/build/_deps/thrust-src/dependencies/cub -I/home/krasznaa/ATLAS/projects/traccc/build/_deps/acts-src/Core/include -I/home/krasznaa/ATLAS/projects/traccc/build/_deps/acts-build/Core -I/home/krasznaa/ATLAS/projects/traccc/traccc/plugins/algebra/include -I/home/krasznaa/ATLAS/projects/traccc/traccc/plugins/algebra/array/include -I/home/krasznaa/ATLAS/projects/traccc/build/_deps/algebraplugins-src/frontend/array_cmath/include -I/home/krasznaa/ATLAS/projects/traccc/build/_deps/algebraplugins-src/common/include -I/home/krasznaa/ATLAS/projects/traccc/build/_deps/algebraplugins-src/storage/array/include -I/home/krasznaa/ATLAS/projects/traccc/build/_deps/algebraplugins-src/math/cmath/include -I/home/krasznaa/ATLAS/projects/traccc/build/_deps/algebraplugins-src/math/common/include -I/home/krasznaa/ATLAS/projects/traccc/traccc/plugins/algebra/vecmem/include -I/home/krasznaa/ATLAS/projects/traccc/build/_deps/algebraplugins-src/frontend/vecmem_cmath/include -I/home/krasznaa/ATLAS/projects/traccc/build/_deps/algebraplugins-src/storage/vecmem/include -I/home/krasznaa/ATLAS/projects/traccc/traccc/plugins/algebra/eigen/include -I/home/krasznaa/ATLAS/projects/traccc/build/_deps/algebraplugins-src/frontend/eigen_eigen/include -I/home/krasznaa/ATLAS/projects/traccc/build/_deps/algebraplugins-src/storage/eigen/include -I/home/krasznaa/ATLAS/projects/traccc/build/_deps/algebraplugins-src/math/eigen/include -I/home/krasznaa/ATLAS/projects/traccc/traccc/plugins/algebra/vc/include -I/home/krasznaa/ATLAS/projects/traccc/build/_deps/algebraplugins-src/frontend/vc_vc/include -I/home/krasznaa/ATLAS/projects/traccc/build/_deps/algebraplugins-src/storage/vc/include -I/home/krasznaa/ATLAS/projects/traccc/build/_deps/algebraplugins-src/math/vc/include -I/home/krasznaa/ATLAS/projects/traccc/build/_deps/vecmem-build/sycl/CMakeFiles -I/home/krasznaa/ATLAS/projects/traccc/build/_deps/vecmem-src/sycl/include -isystem/home/krasznaa/ATLAS/projects/traccc/build/_deps/eigen3-src -isystem/atlas/software/boost/1.78.0/x86_64-ubuntu1804-gcc8-opt/include -isystem/home/krasznaa/ATLAS/projects/traccc/build/_deps/vc-src SYCL_FLAGS = -fsycl -fsycl-targets=nvptx64-nvidia-cuda -Wno-linker-warnings -std=c++17 -g -Wall -Wextra -Wno-unknown-cuda-version -Wshadow -Wunused-local-typedefs -Werror -pedantic -fPIC [bash][atspot01]:build >Yes... our build setup is not the simplest...
Pinging @ivorobts and @konradkusiak97.
Do you solve the Segementation fault in debug cuda oneDNN？ If so, how did you solve this problem? Thank you!
@krasznaa
Describe the bug
This issue has been hurting us for a while already, but since the compiler's behaviour changed a little bit recently, I thought it would finally be time to try to address it.
When trying to build our code for the CUDA backend in debug mode, we've been getting the following failures since a while:
And:
(With 2021-09 and 2021-12 respectively.)
Now, with the latest nightlies, the linker became a bit more chatty, and fails in the following (very long) way:
(I had to shorten the output, as the same messages seem to repeat 8-10 times...)
To Reproduce
Unfortunately reproducing it is not completely trivial. :frowning: This linking error happens while trying to build https://github.com/acts-project/traccc with:
I can give a longer description if there will be somebody willing to give this a try...
Environment (please complete the following information)
Additional context
Building the code in Debug mode for Intel (OpenCL and Level-0) backends works fine. Building the code in Debug mode for the HIP backend fails for a very different reason:
But that one is for a different issue... :stuck_out_tongue:
Finally, the compiler flags used by CMake for this Debug build are the following:
Yes... our build setup is not the simplest...
Pinging @ivorobts and @konradkusiak97.