This bugzilla service is closed. All entries have been migrated to https://gitlab.com/libeigen/eigen
Bug 1396 - Unable to call operator== and operator!= for matrices inside a CUDA kernel
Summary: Unable to call operator== and operator!= for matrices inside a CUDA kernel
Status: RESOLVED FIXED
Alias: None
Product: Eigen
Classification: Unclassified
Component: Core - general (show other bugs)
Version: 3.4 (development)
Hardware: GPU (CUDA) Linux
: Normal Compilation Problem
Assignee: Nobody
URL:
Whiteboard:
Keywords:
Depends on:
Blocks:
 
Reported: 2017-02-27 15:14 UTC by Patrick Stotko
Modified: 2019-12-04 16:50 UTC (History)
3 users (show)



Attachments
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

Hi,

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

<path>/main.cu(12): 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>/main.cu(28): 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,
Patrick
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.
Comment 3 Nobody 2019-12-04 16:50:46 UTC
-- GitLab Migration Automatic Message --

This bug has been migrated to gitlab.com's GitLab instance and has been closed from further activity.

You can subscribe and participate further through the new bug through this link to our GitLab instance: https://gitlab.com/libeigen/eigen/issues/1396.

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