New user self-registration is currently disabled. Please email eigen-core-team @ if you need an account.
Bug 1396 - Unable to call operator== and operator!= for matrices inside a CUDA kernel
Unable to call operator== and operator!= for matrices inside a CUDA kernel
Product: Eigen
Classification: Unclassified
Component: Core - general
3.4 (development)
GPU (CUDA) Linux
: Normal Compilation Problem
Assigned To: Nobody
Depends on:
  Show dependency treegraph
Reported: 2017-02-27 15:14 UTC by Patrick Stotko
Modified: 2017-02-28 21:58 UTC (History)
3 users (show)

Minimal failing example (3.79 KB, application/zip)
2017-02-27 15:14 UTC, Patrick Stotko
no flags Details

Description Patrick Stotko 2017-02-27 15:14:29 UTC
Created attachment 778 [details]
Minimal failing example


I have observed this error when trying to check if two matrices are equal inside a CUDA kernel (see attached reproducer):

<path>/ error: calling a __host__ function("Eigen::MatrixBase< ::Eigen::Matrix<float, (int)5, (int)3, (int)0, (int)5, (int)3> > ::operator ==< ::Eigen::Matrix<float, (int)5, (int)3, (int)0, (int)5, (int)3> > ") from a __global__ function("testEqual<(int)5, (int)3> ") is not allowed

<path>/ error: calling a __host__ function("Eigen::MatrixBase< ::Eigen::Matrix<float, (int)5, (int)3, (int)0, (int)5, (int)3> > ::operator !=< ::Eigen::Matrix<float, (int)5, (int)3, (int)0, (int)5, (int)3> > ") from a __global__ function("testNotEqual<(int)5, (int)3> ") is not allowed

This fails on both version 3.3.3 and the development version from the default branch.

In Eigen/src/Core/MatrixBase.h lines 296-306 where both functions are defined, there is indeed no EIGEN_DEVICE_FUNC. Furthermore, it is also missing in Eigen/src/Core/DenseBase.h lines 487-488 where all() and any() are defined.

Is there a reason why calling them is not allowed inside kernels? If not, I would suggest to add the EIGEN_DEVICE_FUNC qualifier and inline the definitions of all() and any().

Tested on (K)Ubuntu 14.04 and CUDA 8.0 using the attached example program.

Kind regards,
Comment 1 Gael Guennebaud 2017-02-28 08:49:35 UTC
This should be fixed in both devel and 3.3 branches.
Comment 2 Patrick Stotko 2017-02-28 11:41:49 UTC
I can confirm that the newest version of 3.3 and devel resolved the problem. Thank you.

Note You need to log in before you can comment on or make changes to this bug.