Overview

Eigen's fixed-size matrices, vectors, and arrays can be used inside CUDA and HIP kernels. This is especially useful when working on numerous but small problems. By default, when Eigen's headers are included within a .cu file compiled by nvcc, most of Eigen's functions and methods are prefixed by the __device__ __host__ keywords making them callable from both host and device code.

This support can be disabled by defining EIGEN_NO_CUDA (or EIGEN_NO_HIP for HIP) before including any Eigen header. This might be useful to disable some warnings when a .cu file makes use of Eigen on the host side only.

Eigen requires CUDA 11.8 or later and compute capability 6.0 (sm_60) or later; CUDA 13 itself requires sm_75. For HIP the floor is ROCm 5.6.

Warning

Host SIMD vectorization must be disabled in .cu files. It is strongly recommended to move all costly host-side computation from .cu files to regular .cpp files.

HIP support

Eigen also supports AMD's HIP compiler (hipcc). The same EIGEN_DEVICE_FUNC mechanism applies: functions are annotated with __device__ __host__ when compiling with HIP. Define EIGEN_NO_HIP to disable this. Internally, both CUDA and HIP are unified under the EIGEN_GPUCC macro.

Known issues

• On 64-bit systems Eigen uses std::ptrdiff_t, a 64-bit type, as the default type for indexes and sizes. On CUDA/HIP devices, it would make sense to default to 32-bit int. However, to keep host and device code compatible, this cannot be done automatically by Eigen, and the user is thus required to define EIGEN_DEFAULT_DENSE_INDEX_TYPE to int throughout their code (or only for device code if there is no interaction between host and device code through Eigen's objects).