doc/UsingNVCC.dox

namespace Eigen {

/** \page TopicCUDA Using Eigen in CUDA kernels

\section TopicCUDA_Overview 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 \c .cu file compiled by \c nvcc,
most of %Eigen's functions and methods are prefixed by the \c __device__ \c __host__ keywords
making them callable from both host and device code.

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

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

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

\section TopicCUDA_HIP HIP support

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

\section TopicCUDA_KnownIssues Known issues

 - On 64-bit systems %Eigen uses \c 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 \c int.
   However, to keep host and device code compatible, this cannot be done automatically by %Eigen,
   and the user is thus required to define \c EIGEN_DEFAULT_DENSE_INDEX_TYPE to \c int throughout
   their code (or only for device code if there is no interaction between host and device code
   through %Eigen's objects).

*/

}