#include <Eigen/Dense>

struct Test
{
typedef std::complex<float> cpx;
typedef Eigen::Array<cpx, Eigen::Dynamic, 1> Vc;
typedef Eigen::Array<float, Eigen::Dynamic, 1> Vr;
Vc z;
Vr a;

void run()
{
float k=1;
a.resize(1); z.resize(1);
a.setZero(); z.setZero();

// Always OK
a *= k;

// Trouble
z *= k;
// > nvcc eigentest.cc -o eigentest -I eigen
// Works fine
//
// > nvcc eigentest.cu -o eigentest -I eigen
// Fails with message:
//
// AssignmentFunctors.h(92): error: a value of type "int"
// cannot be assigned to an entity of type
// "_ZNSt7complexIfE9_ComplexTE"
//
// No diffference to fail if k cpx
// += does not fail
//
// Other complex array operations also fail with similar
// messages.
}
};

int main(int argc, char **argv) { Test t; t.run(); return 0; }

PS: should have said: this does not happen for Eigen 3.2.4 and earlier, does in 3.3.1 and 3.3.3 beat (ubuntu latest).
Versions 7.5 and 8.0 of the cuda toolkit show the same result.

This is probably a template handling problem in nvcc processing, but it is nasty as I have a great deal of Eigen and cuda mixed code (no Eigen in device code though).

I will have to live with 3.2.4 for now .

I'm using nvcc 8 and your example compiles fine for me (with some unrelated warnings though).

Thanks for replying .

It happens on my home machine running Ubuntu (latest) and on my work machine running centos.

Did you notice it had to be a *.cu suffix to fail? .cc suffix gets passed unexamined to gcc.
If you did, I really want to know what you are doing different.

lots of mpy and divide ops involving complex get fatal compile errors e.g.

Code: Select all

Eigen::Array<cpx, Eigen::Dynamic, 1> x(1),y(1),z(1);
x<<1; z=x; float s=x.sum(); // OK
float m=x.mean(); // compile error
x = x*1; // compile error
y = x*z; // compile error


The previous example should have used cpx (std::complex<float>) scalars (rather than float) - perils of typing from memory, sorry.

PS

I have tried it with the current Ubuntu default Eigen (3.3.3 pre release) and with last night's mercurial version (in addition to 3.3.1 and 3.3.3 at work).

I need the .cu file to compile with nvcc as this file type (normally) mixes CUDA code and PC code, to be split by nvcc.

The nvcc compiler also fails on BOOST expressive for example - I had to ensure that was placed in pure .cc code - but this is not attractive for the much more tightly integrated Eigen stuff.

I also tried the CUDA #define to limit the C++ version to 03 - made no difference.

If this points to an nvcc fault NVIDIA may be persuaded to fix it, eventually.

Gail,

You need to name the file "something.cu" to get nvcc to process it and get the error.

If nvcc gets a .cc or .cpp or etc. suffix, it will pass the file unexamined to gcc, and you will not see the compilation errors.

The problem is very reliable on all Eigen 3.3 variants I have tried and on 7.5 and 8.0 NVIDIA toolkits. It does not happen on 3.2 and earlier versions of Eigen.

Any complex Eigen array involved in multiply or divide generates the odd syntax error - even v.mean().

nvcc -c something.cu -I eigen
will fail

nvcc -c something.cc -I eigen
will not fail.

yes, I did named the file .cu !!!

nvcc -version gives me:

Code: Select all

$ nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2016 NVIDIA Corporation
Built on Tue_Jan_10_13:22:46_CST_2017
Cuda compilation tools, release 8.0, V8.0.61


The content of the .cu file is:

Code: Select all

#include <Eigen/Dense>

struct Test
{
typedef std::complex<float> cpx;
typedef Eigen::Array<cpx, Eigen::Dynamic, 1> Vc;
typedef Eigen::Array<float, Eigen::Dynamic, 1> Vr;
Vc z;
Vr a;

void run()
{
float k=1;
a.resize(1); z.resize(1);
a.setZero(); z.setZero();

// Always OK
a *= k;

// Trouble
z *= k;
}
};

int main(int argc, char **argv) { Test t; t.run(); return 0; }


I updated nvcc on my home machine to the version you used (and had to change my gcc version down to 5.4.1 as a result)

I am using eigen -> eigen-eigen-67e894c6cd8f currently, but all 3.3.x versions have so far failed for me, and all 3.2.4 and earlier worked.

It still fails to compile . Thanks for looking into it - if it compiles for you there must be something that can be used to make it work for me....

Here is a small bash script which generates and uses the test files and shows versions.

Code: Select all

cat > t.cc <<END
#include <Eigen/Dense>

int main(int argc, char **argv)
{
    Eigen::Array<std::complex<float>, 1, 1> z;
    z.setZero();
    z *= 1.0f;
    return 1;
}
END

ln -s t.cc t.cu

echo Testing program:
cat t.cu

echo nvcc version:
nvcc --version

echo
echo gcc version:
gcc --version
echo g++ version:
g++ --version
echo Linux version:
uname -a
echo

( nvcc -c t.cc -I eigen 2>/dev/null ) || echo .cc compile failed
( nvcc -c t.cu -I eigen 2>/dev/null ) || echo .cu compile failed

rm t.c[cu]


On my Ubuntu 16.04 Intel PC machine this generates:

Code: Select all

Testing program:
#include <Eigen/Dense>

int main(int argc, char **argv)
{
    Eigen::Array<std::complex<float>, 1, 1> z;
    z.setZero();
    z *= 1.0f;
    return 1;
}
nvcc version:
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2016 NVIDIA Corporation
Built on Tue_Jan_10_13:22:03_CST_2017
Cuda compilation tools, release 8.0, V8.0.61

gcc version:
gcc-5 (Ubuntu 5.4.1-8ubuntu1) 5.4.1 20170304
Copyright (C) 2015 Free Software Foundation, Inc.
This is free software; see the source for copying conditions.  There is NO
warranty; not even for MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.

g++ version:
g++-5 (Ubuntu 5.4.1-8ubuntu1) 5.4.1 20170304
Copyright (C) 2015 Free Software Foundation, Inc.
This is free software; see the source for copying conditions.  There is NO
warranty; not even for MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.

Linux version:
Linux stuart-Z87-HD3 4.10.0-21-generic #23-Ubuntu SMP Fri Apr 28 16:14:22 UTC 2017 x86_64 x86_64 x86_64 GNU/Linux

.cu compile failed


ho hum

Have you seen my answer there: http://eigen.tuxfamily.org/bz/show_bug.cgi?id=1430

I do can reproduce on Linux with gcc as host compiler, and I even don't need Eigen to make it fails.

I looked at your results after seeing this post. (been away without password, hence delay)

Your pure C++ example needs __device__ to fail. It seems not unreasonable for nvcc to fail device code which would require processing unusable data?

My Eigen example did not mention __device__. Does Eigen 3.3.x add __host__ __device__ to its templates now, even when called from __host__ code? Can I stop it?

I have written an awful lot of GPU/PC code at work relying on Eigen/PC to manipulate small arrays (e.g. eigenvectors, covariancee etc. generated on GPU) and GPU for heavy processing. It would take a lot of re-working to split this up.

Yes, when a file is compile by nvcc, we have to add __DEVICE__ in front all all functions/method that can be called (recursively) from CUDA device code. But there is no way to add __DEVICE__ on demand when the method is truly called from a device code. In a similar vein, we cannot put/remove __DEVICE__ depending on the templated scalar type (e.g., std::complex is not CUDA friendly, so we could remove all the __DEVICE__ when Matrix<> is instantiated with std::complex, but that's obviously not possible).

I'm open to any suggestions to fix this issue but I'm afraid the only solution is to better split C++/CUDA code in separate .cpp/.cu files on the user side.

The insertion of __host__ __ device__ is only needed if Eigen is to be used in GPU code, as far as I can see.

This action could be disabled by a user #define EIGEN_NO_CUDA or similar and remain conditional on nvcc parsing.

If I ever wanted to use Eigen in GPU code, I could enable it for the compilation of the relevant pure GPU modules.

You already have similar #define flags in Eigen for other conditions.

PS. I am curious as to what Eigen can usefully do in GPU code - most of my GPU code does all processing in registers, or shared memory and arrays are spread across threads. As far as I can see Eigen can only operate on memory or stack arrays?

It is useful for computations on small vectors and matrices, like 2x2 to 6x6 for space transformation or local linear regressions.