Compilation with CUDA fails on PacketMathHalf.h for __ldg

Summary

Using the CUDA toolkit version 11.4, when I compile a simple test cuda program that includes <Eigen/Dense>, despite defining EIGEN_NO_CUDA, the compilation still breaks with 2 errors.

Environment

  • Operating System : Linux
  • Architecture : x64
  • Eigen Version : 3.3.9
  • Compiler Version : Gcc9.3.0
  • Compile Flags :

Minimal Example

A runnable example is available here: Example The code is very minimal:

#define EIGEN_NO_CUDA 1

#include <Eigen/Dense>

int main() {
    return 0;
}

Steps to reproduce

  1. Compile the code above
  2. See the error occur.

What is the current bug behavior?

When we compile, even if we define EIGEN_NO_CUDA, which should disable CUDA support, we still get warnings about device functions, and worse, we get an error about an __ldg function not existing at Eigen/src/Core/arch/CUDA/PacketMathHalf.h(81)

What is the expected correct behavior?

I should see no CUDA code compiled because of the EIGEN_NO_CUDA macro, and also even if I do not define that macro, I should not see any errors within the headers.

Warning Messages

/usr/local/cuda-11.4/bin/../targets/x86_64-linux/include/host_defines.h:54:2: warning: #warning "host_defines.h is an internal header file and must not be used directly.  This file will be removed in a future CUDA release.  Please use cuda_runtime_api.h or cuda_runtime.h instead." [-Wcpp]
   54 | #warning "host_defines.h is an internal header file and must not be used directly.  This file will be removed in a future CUDA release.  Please use cuda_runtime_api.h or cuda_runtime.h instead."
      |  ^~~~~~~
/usr/include/eigen3/Eigen/src/Core/util/XprHelper.h(108): warning: __host__ annotation is ignored on a function("no_assignment_operator") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/util/XprHelper.h(108): warning: __device__ annotation is ignored on a function("no_assignment_operator") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/util/XprHelper.h(109): warning: __host__ annotation is ignored on a function("no_assignment_operator") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/util/XprHelper.h(109): warning: __device__ annotation is ignored on a function("no_assignment_operator") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/util/XprHelper.h(109): warning: __host__ annotation is ignored on a function("~no_assignment_operator") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/util/XprHelper.h(109): warning: __device__ annotation is ignored on a function("~no_assignment_operator") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/MathFunctions.h(1081): warning: calling a constexpr __host__ function("real") from a __host__ __device__ function("abs") is not allowed. The experimental flag '--expt-relaxed-constexpr' can be used to allow this.

/usr/include/eigen3/Eigen/src/Core/MathFunctions.h(1081): warning: calling a constexpr __host__ function("imag") from a __host__ __device__ function("abs") is not allowed. The experimental flag '--expt-relaxed-constexpr' can be used to allow this.

/usr/include/eigen3/Eigen/src/Core/MathFunctions.h(1081): warning: calling a constexpr __host__ function from a __host__ __device__ function is not allowed. The experimental flag '--expt-relaxed-constexpr' can be used to allow this.

/usr/include/eigen3/Eigen/src/Core/MathFunctions.h(1081): warning: calling a constexpr __host__ function from a __host__ __device__ function is not allowed. The experimental flag '--expt-relaxed-constexpr' can be used to allow this.

/usr/include/eigen3/Eigen/src/Core/MathFunctions.h(1086): warning: calling a constexpr __host__ function("real") from a __host__ __device__ function("abs") is not allowed. The experimental flag '--expt-relaxed-constexpr' can be used to allow this.

/usr/include/eigen3/Eigen/src/Core/MathFunctions.h(1086): warning: calling a constexpr __host__ function("imag") from a __host__ __device__ function("abs") is not allowed. The experimental flag '--expt-relaxed-constexpr' can be used to allow this.

/usr/include/eigen3/Eigen/src/Core/MathFunctions.h(1086): warning: calling a constexpr __host__ function from a __host__ __device__ function is not allowed. The experimental flag '--expt-relaxed-constexpr' can be used to allow this.

/usr/include/eigen3/Eigen/src/Core/MathFunctions.h(1086): warning: calling a constexpr __host__ function from a __host__ __device__ function is not allowed. The experimental flag '--expt-relaxed-constexpr' can be used to allow this.

/usr/include/eigen3/Eigen/src/Core/arch/CUDA/Half.h(303): warning: calling a __host__ function from a __host__ __device__ function is not allowed

/usr/include/eigen3/Eigen/src/Core/arch/CUDA/Half.h(306): warning: calling a __host__ function from a __host__ __device__ function is not allowed

/usr/include/eigen3/Eigen/src/Core/arch/CUDA/PacketMathHalf.h(81): error: no instance of overloaded function "__ldg" matches the argument list
            argument types are: (const Eigen::half *)

/usr/include/eigen3/Eigen/src/Core/arch/CUDA/PacketMathHalf.h(81): error: no instance of overloaded function "__ldg" matches the argument list
            argument types are: (const Eigen::half *)

/usr/include/eigen3/Eigen/src/Core/DenseBase.h(590): warning: __host__ annotation is ignored on a function("DenseBase") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/DenseBase.h(590): warning: __device__ annotation is ignored on a function("DenseBase") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/MatrixBase.h(467): warning: __host__ annotation is ignored on a function("MatrixBase") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/MatrixBase.h(467): warning: __device__ annotation is ignored on a function("MatrixBase") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/MatrixBase.h(468): warning: __host__ annotation is ignored on a function("MatrixBase") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/MatrixBase.h(468): warning: __device__ annotation is ignored on a function("MatrixBase") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/MatrixBase.h(468): warning: __host__ annotation is ignored on a function("~MatrixBase") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/MatrixBase.h(468): warning: __device__ annotation is ignored on a function("~MatrixBase") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/ArrayBase.h(156): warning: __host__ annotation is ignored on a function("ArrayBase") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/ArrayBase.h(156): warning: __device__ annotation is ignored on a function("ArrayBase") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/ArrayBase.h(157): warning: __host__ annotation is ignored on a function("ArrayBase") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/ArrayBase.h(157): warning: __device__ annotation is ignored on a function("ArrayBase") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/ArrayBase.h(157): warning: __host__ annotation is ignored on a function("~ArrayBase") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/ArrayBase.h(157): warning: __device__ annotation is ignored on a function("~ArrayBase") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/CwiseUnaryView.h(70): warning: __host__ annotation is ignored on a function("CwiseUnaryView") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/CwiseUnaryView.h(70): warning: __device__ annotation is ignored on a function("CwiseUnaryView") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/CwiseUnaryView.h(110): warning: __host__ annotation is ignored on a function("CwiseUnaryViewImpl") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/CwiseUnaryView.h(110): warning: __device__ annotation is ignored on a function("CwiseUnaryViewImpl") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/CwiseUnaryView.h(125): warning: __host__ annotation is ignored on a function("CwiseUnaryViewImpl") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/CwiseUnaryView.h(125): warning: __device__ annotation is ignored on a function("CwiseUnaryViewImpl") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/CwiseUnaryView.h(125): warning: __host__ annotation is ignored on a function("~CwiseUnaryViewImpl") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/CwiseUnaryView.h(125): warning: __device__ annotation is ignored on a function("~CwiseUnaryViewImpl") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/MapBase.h(185): warning: __host__ annotation is ignored on a function("MapBase") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/MapBase.h(185): warning: __device__ annotation is ignored on a function("MapBase") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/MapBase.h(186): warning: __host__ annotation is ignored on a function("MapBase") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/MapBase.h(186): warning: __device__ annotation is ignored on a function("MapBase") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/MapBase.h(186): warning: __host__ annotation is ignored on a function("~MapBase") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/MapBase.h(186): warning: __device__ annotation is ignored on a function("~MapBase") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/MapBase.h(300): warning: __host__ annotation is ignored on a function("MapBase") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/MapBase.h(300): warning: __device__ annotation is ignored on a function("MapBase") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/MapBase.h(301): warning: __host__ annotation is ignored on a function("MapBase") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/MapBase.h(301): warning: __device__ annotation is ignored on a function("MapBase") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/MapBase.h(301): warning: __host__ annotation is ignored on a function("~MapBase") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/MapBase.h(301): warning: __device__ annotation is ignored on a function("~MapBase") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/Map.h(162): warning: __host__ annotation is ignored on a function("Map") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/Map.h(162): warning: __device__ annotation is ignored on a function("Map") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/Ref.h(90): warning: __host__ annotation is ignored on a function("RefBase") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/Ref.h(90): warning: __device__ annotation is ignored on a function("RefBase") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/Ref.h(230): warning: __host__ annotation is ignored on a function("Ref") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/Ref.h(230): warning: __device__ annotation is ignored on a function("Ref") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/Block.h(111): warning: __host__ annotation is ignored on a function("Block") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/Block.h(111): warning: __device__ annotation is ignored on a function("Block") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/Block.h(161): warning: __host__ annotation is ignored on a function("BlockImpl") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/Block.h(161): warning: __device__ annotation is ignored on a function("BlockImpl") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/Block.h(181): warning: __host__ annotation is ignored on a function("BlockImpl_dense") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/Block.h(181): warning: __device__ annotation is ignored on a function("BlockImpl_dense") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/Block.h(341): warning: __host__ annotation is ignored on a function("BlockImpl_dense") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/Block.h(341): warning: __device__ annotation is ignored on a function("BlockImpl_dense") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/Transpose.h(66): warning: __host__ annotation is ignored on a function("Transpose") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/Transpose.h(66): warning: __device__ annotation is ignored on a function("Transpose") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/Transpose.h(123): warning: __host__ annotation is ignored on a function("TransposeImpl") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/Transpose.h(123): warning: __device__ annotation is ignored on a function("TransposeImpl") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/Transpose.h(150): warning: __host__ annotation is ignored on a function("TransposeImpl") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/Transpose.h(150): warning: __device__ annotation is ignored on a function("TransposeImpl") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/Transpose.h(150): warning: __host__ annotation is ignored on a function("~TransposeImpl") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/Transpose.h(150): warning: __device__ annotation is ignored on a function("~TransposeImpl") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/Diagonal.h(78): warning: __host__ annotation is ignored on a function("Diagonal") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/Diagonal.h(78): warning: __device__ annotation is ignored on a function("Diagonal") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/TriangularMatrix.h(220): warning: __host__ annotation is ignored on a function("TriangularView") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/TriangularMatrix.h(220): warning: __device__ annotation is ignored on a function("TriangularView") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/TriangularMatrix.h(546): warning: __host__ annotation is ignored on a function("TriangularViewImpl") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/TriangularMatrix.h(546): warning: __device__ annotation is ignored on a function("TriangularViewImpl") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/TriangularMatrix.h(547): warning: __host__ annotation is ignored on a function("TriangularViewImpl") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/TriangularMatrix.h(547): warning: __device__ annotation is ignored on a function("TriangularViewImpl") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/TriangularMatrix.h(547): warning: __host__ annotation is ignored on a function("~TriangularViewImpl") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/TriangularMatrix.h(547): warning: __device__ annotation is ignored on a function("~TriangularViewImpl") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/Reverse.h(90): warning: __host__ annotation is ignored on a function("Reverse") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/Reverse.h(90): warning: __device__ annotation is ignored on a function("Reverse") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/ArrayWrapper.h(47): warning: __host__ annotation is ignored on a function("ArrayWrapper") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/ArrayWrapper.h(47): warning: __device__ annotation is ignored on a function("ArrayWrapper") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/ArrayWrapper.h(145): warning: __host__ annotation is ignored on a function("MatrixWrapper") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Core/ArrayWrapper.h(145): warning: __device__ annotation is ignored on a function("MatrixWrapper") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Geometry/Quaternion.h(202): warning: __host__ annotation is ignored on a function("QuaternionBase") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Geometry/Quaternion.h(202): warning: __device__ annotation is ignored on a function("QuaternionBase") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Geometry/Quaternion.h(203): warning: __host__ annotation is ignored on a function("QuaternionBase") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Geometry/Quaternion.h(203): warning: __device__ annotation is ignored on a function("QuaternionBase") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Geometry/Quaternion.h(203): warning: __host__ annotation is ignored on a function("~QuaternionBase") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Geometry/Quaternion.h(203): warning: __device__ annotation is ignored on a function("~QuaternionBase") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Geometry/Quaternion.h(258): warning: __host__ annotation is ignored on a function("Quaternion") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Geometry/Quaternion.h(258): warning: __device__ annotation is ignored on a function("Quaternion") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Geometry/Quaternion.h(374): warning: __host__ annotation is ignored on a function("Map") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Geometry/Quaternion.h(374): warning: __device__ annotation is ignored on a function("Map") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Geometry/Quaternion.h(411): warning: __host__ annotation is ignored on a function("Map") that is explicitly defaulted on its first declaration

/usr/include/eigen3/Eigen/src/Geometry/Quaternion.h(411): warning: __device__ annotation is ignored on a function("Map") that is explicitly defaulted on its first declaration

2 errors detected in the compilation of "eigen_breaks.cu".

Anything else that might help

The broken code comes from this:

#if defined __CUDACC__
  #define EIGEN_VECTORIZE_CUDA
  #include <vector_types.h>
  #if EIGEN_CUDACC_VER >= 70500
    #define EIGEN_HAS_CUDA_FP16
  #endif
#endif

#if defined EIGEN_HAS_CUDA_FP16
  #include <host_defines.h>
  #include <cuda_fp16.h>
#endif

If I change __CUDACC__ to EIGEN_CUDACC, the code works if EIGEN_NO_CUDA is defined, but still breaks if EIGEN_NO_CUDA is not defined. So making this change to Eigen/Core:49 and Eigen/Core:256, compilation works again.

  • Have a plan to fix this issue.
Edited by rahulaggarwal965