Grid does not compile on Arm with CUDA

[edbennett]

Describe the issue:

Attempting to compile Grid for NVIDIA on Arm fails due to a large number of undefined symbols in arm_neon.h. Following up on @agsunderland's comment on #430 and digging deeper, after some discussion with @RChrHill the issue is that including Eigen with __CUDACC__ undefined causes Eigen to emit code using NEON vector instructions, which according to this post on NVIDIA's forums aren't yet supported in NVCC.

A workaround is to define the EIGEN_DONT_VECTORIZE macro, for example by adding -DEIGEN_DONT_VECTORIZE to the CXXFLAGS; this disables Eigen from using SIMD completely. I'm not sure what performance impact this has compared to being able to use NEON for the things that Eigen is used for on CPU. Upgrading to Eigen 3.4.0 did not fix the problem.

To compile the minimal example below, the following was used:

nvcc -x cu -I../../Grid    -O3 -o eigen_hello eigen_hello.cc

(Replace the ../../Grid with the path to wherever Eigen is available.)

Code example:

#include <iostream>

//uncomment the five commented lines below to be able to compile in the case where EIGEN_DONT_VECTORIZE is defined. These are not needed for the error seen with full Grid to be triggered.
//#undef __CUDA_ARCH__
//#undef __NVCC__
#undef __CUDACC__
#include <Eigen/Dense>
//#define __CUDA_ARCH__
//#define __NVCC__
//#define __CUDACC__

int main()
{
  return 0;
}

Target platform:

This is the Arm GPU testbed in Leicester; there is no model name in the cpuinfo.

Configure options:

N/A

[edbennett]

To add: This also occurs on Grace Hopper.

[paboyle]

Use CXXFLAGS - there are other neon compilers that are happy with Eigen. Or use the -ccbin flag to nvcc to specify a host compiler that actually works for Grace.

[paboyle]

e.g. clang++ should work fine as a host compiler, Ive used it fine on Ampere

[paboyle]

Also -- feel free to contribute a config-command and directory under Systems/GraceHopper or similar. Could save others some hassle.

[edbennett]

Use CXXFLAGS - there are other neon compilers that are happy with Eigen. Or use the -ccbin flag to nvcc to specify a host compiler that actually works for Grace.

I must be missing something—I've used -ccbin g++ on Leicester and now -ccbin=clang++-15 on Grace Hopper, but in both cases the issue persists (with slightly different phrasings in the error), even though both can compile for neon when not running via nvcc.

[paboyle]

I'm puzzled by that, because I thought that nvcc does two things: -- preprocess to host and device code -- run ccbin on the host sequences -- run the device compiler on the device sequences. Is it possible to watch under strace to see if it is the intended host compiler complaining about neon code? Perhaps trapping the intermediates? Have you checked that you g++ and/or clang++-15 are neon aware (might need a -m compiler option passed through to enable)

[paboyle]

  NEONv8)
    AC_DEFINE([NEONV8],[1],[ARMv8 NEON])
    SIMD_FLAGS='-march=armv8-a';;

    So if Grid is targeting NEON as it's SIMD, I'm passing -march=armv8-a to the compiler.
    This is likely missing from the host compiler when you configure for GPU, and then Eigen is failing.
    So perhaps if you pass this through to the host compiler with -Xcompiler it might work?

[paboyle]

You could almost certainly remove Grid from this problem as the challenge appears to be to get Eigen to work with nvcc and your host compile with -DEIGEN_DONT_VECTORIZE set. This is a reasonable expectation of Eigen and nvcc / host compiler.

I'd bet a -Xcompiler -march=armv8-a is needed.

See configure.ac for an example:

case ${CXXTEST} in
  nvcc)
#    CXX="nvcc -keep -v -x cu "
#    CXXLD="nvcc -v -link"
    CXX="${CXXBASE} -x cu "
    CXXLD="${CXXBASE} -link"
    CXXFLAGS="$CXXFLAGS -Xcompiler -fno-strict-aliasing --expt-extended-lambda --expt-relaxed-constexpr"
    if test $ac_openmp = yes; then
       CXXFLAGS="$CXXFLAGS -Xcompiler -fopenmp"
       LDFLAGS="$LDFLAGS -Xcompiler -fopenmp"
    fi

[edbennett]

Have you checked that you g++ and/or clang++-15 are neon aware (might need a -m compiler option passed through to enable)

Yes, the example code upthread compiles without issue in g++ (10.x) and clang++-15.

You could almost certainly remove Grid from this problem as the challenge appears to be to get Eigen to work with nvcc and your host compile with -DEIGEN_DONT_VECTORIZE set. This is a reasonable expectation of Eigen and nvcc / host compiler.

Ah, if the expectation is that we need to use -DEIGEN_DONT_VECTORIZE in order to compile for GPU on Arm, then compilation is already successful (including without -march=armv8-a). (Currently the resulting binary doesn't run on Grace Hopper, for reasons I'm trying to understand, but that is likely a separate issue.) Re-reading my first message, I can see I left open multiple interpretations.

And indeed, this is not necessarily Grid specific, it's general to trying to compile Eigen for an Arm host with nvcc.