Problem passing templatized structs by value


I have encountered some problems with passing templatized structs by value… I have code similar to the snippet below, which used to work, but no longer does with CUDA 2.0

template<unsigned int N>

struct valarray {

  float zzz[N];


template<unsigned int N>

__global__ void

moo( valarray<N> foo )



int main()


  valarray<3> bar = {0.0f, 0.0f, 0.0f};


  return 0;


Compiling the snippet above, I get the following output:

[font=“Times”]modesty ~ $ nvcc -V

nvcc: NVIDIA ® Cuda compiler driver

Copyright © 2005-2007 NVIDIA Corporation

Built on Tue_Jun_10_05:42:45_PDT_2008

Cuda compilation tools, release 1.1, V0.2.1221

modesty ~ $ nvcc --keep

reproduce.cudafe1.stub.h:17: error: template argument required for ‘struct valarray’

reproduce.cudafe1.stub.c:17: error: template-id ‘moo<3u>’ for ‘void moo(_Z8valarrayILj3EE)’ does not match any template declaration

reproduce.cudafe1.stub.c:17: error: invalid function declaration

modesty ~ $ [/font]

Lines 16 and 17 of reproduce.cudafe1.stub.h are


void moo(struct valarray foo);[/font]

While lines 16 and 17 of reproduce.cudafe1.stub.cpp are

[font=“Times”]template <>

void moo<3U> (_Z8valarrayILj3EE foo)[/font]

Seems like somehow the “struct valarray” doesn’t get mangled correctly in the stub.h-file?



I hate to bump, but I’m seeing quite a few errors in this vein in my code development. Chris, it looks like you wrote up a reproducible set of code, which implies you submitted a bug report: have you heard anything back?

The effects seem to be widespread. It seems like a large percentage of templated code is not being treated properly by the compiler. I’ve seen namespace clashes from two different classes declaring the same typedef within, which makes no sense since the typedefs should namespace into the classes. I’ve found it nearly impossible to access and use typedefs within a templated class (which are rather useful as they can bring various template arguments together.)

Hi John,

No, unfortunately I haven’t filed a bug report beyond this post for this problem.