cannot create device_vector with enum types

[GoogleCodeExporter]

nathan@rabota:~/Desktop$ cat foo2.cu 
#include <thrust/device_vector.h>
#include <vector>

enum my_enum {red, blue, green};

int main(void)
{
    thrust::device_vector<my_enum> v(10, red);
}

nathan@rabota:~/Desktop$ nvcc foo2.cu
/tmp/tmpxft_0000696c_00000000-12_foo2.o: In function `void
thrust::detail::device::cuda::detail::launch_closure_by_value<thrust::detail::de
vice::cuda::for_each_n_closure<thrust::device_ptr<unsigned
long long>, long,
thrust::detail::device::generic::detail::generate_functor<thrust::detail::device
::cuda::detail::fill_functor<unsigned
long long> > >
>(thrust::detail::device::cuda::for_each_n_closure<thrust::device_ptr<unsigned
long long>, long,
thrust::detail::device::generic::detail::generate_functor<thrust::detail::device
::cuda::detail::fill_functor<unsigned
long long> > >)':
tmpxft_0000696c_00000000-1_foo2.cudafe1.cpp:(.text._ZN6thrust6detail6device4cuda
6detail23launch_closure_by_valueINS2_18for_each_n_closureINS_10device_ptrIyEElNS
1_7generic6detail16generate_functorINS3_12fill_functorIyEEEEEEEEvT_[void
thrust::detail::device::cuda::detail::launch_closure_by_value<thrust::detail::de
vice::cuda::for_each_n_closure<thrust::device_ptr<unsigned
long long>, long,
thrust::detail::device::generic::detail::generate_functor<thrust::detail::device
::cuda::detail::fill_functor<unsigned
long long> > >
>(thrust::detail::device::cuda::for_each_n_closure<thrust::device_ptr<unsigned
long long>, long,
thrust::detail::device::generic::detail::generate_functor<thrust::detail::device
::cuda::detail::fill_functor<unsigned
long long> > >)]+0x11): undefined reference to `void
thrust::detail::device::cuda::detail::__wrapper__device_stub_launch_closure_by_v
alue<thrust::detail::device::cuda::for_each_n_closure<thrust::device_ptr<unsigne
d
long long>, long,
thrust::detail::device::generic::detail::generate_functor<thrust::detail::device
::cuda::detail::fill_functor<unsigned
long long> > >
>(thrust::detail::device::cuda::for_each_n_closure<thrust::device_ptr<unsigned
long long>, long,
thrust::detail::device::generic::detail::generate_functor<thrust::detail::device
::cuda::detail::fill_functor<unsigned
long long> > >&)'
/tmp/tmpxft_0000696c_00000000-12_foo2.o: In function `void
thrust::detail::device::cuda::detail::launch_closure_by_value<thrust::detail::de
vice::cuda::for_each_n_closure<thrust::device_ptr<unsigned
long long>, unsigned int,
thrust::detail::device::generic::detail::generate_functor<thrust::detail::device
::cuda::detail::fill_functor<unsigned
long long> > >
>(thrust::detail::device::cuda::for_each_n_closure<thrust::device_ptr<unsigned
long long>, unsigned int,
thrust::detail::device::generic::detail::generate_functor<thrust::detail::device
::cuda::detail::fill_functor<unsigned
long long> > >)':
tmpxft_0000696c_00000000-1_foo2.cudafe1.cpp:(.text._ZN6thrust6detail6device4cuda
6detail23launch_closure_by_valueINS2_18for_each_n_closureINS_10device_ptrIyEEjNS
1_7generic6detail16generate_functorINS3_12fill_functorIyEEEEEEEEvT_[void
thrust::detail::device::cuda::detail::launch_closure_by_value<thrust::detail::de
vice::cuda::for_each_n_closure<thrust::device_ptr<unsigned
long long>, unsigned int,
thrust::detail::device::generic::detail::generate_functor<thrust::detail::device
::cuda::detail::fill_functor<unsigned
long long> > >
>(thrust::detail::device::cuda::for_each_n_closure<thrust::device_ptr<unsigned
long long>, unsigned int,
thrust::detail::device::generic::detail::generate_functor<thrust::detail::device
::cuda::detail::fill_functor<unsigned
long long> > >)]+0x11): undefined reference to `void
thrust::detail::device::cuda::detail::__wrapper__device_stub_launch_closure_by_v
alue<thrust::detail::device::cuda::for_each_n_closure<thrust::device_ptr<unsigne
d
long long>, unsigned int,
thrust::detail::device::generic::detail::generate_functor<thrust::detail::device
::cuda::detail::fill_functor<unsigned
long long> > >
>(thrust::detail::device::cuda::for_each_n_closure<thrust::device_ptr<unsigned
long long>, unsigned int,
thrust::detail::device::generic::detail::generate_functor<thrust::detail::device
::cuda::detail::fill_functor<unsigned
long long> > >&)'
/tmp/tmpxft_0000696c_00000000-12_foo2.o: In function `void
thrust::detail::device::cuda::detail::launch_closure_by_value<thrust::detail::de
vice::cuda::for_each_n_closure<thrust::device_ptr<my_enum>,
long,
thrust::detail::device::generic::detail::generate_functor<thrust::detail::device
::cuda::detail::fill_functor<my_enum>
> >
>(thrust::detail::device::cuda::for_each_n_closure<thrust::device_ptr<my_enum>,
long,
thrust::detail::device::generic::detail::generate_functor<thrust::detail::device
::cuda::detail::fill_functor<my_enum>
> >)':
tmpxft_0000696c_00000000-1_foo2.cudafe1.cpp:(.text._ZN6thrust6detail6device4cuda
6detail23launch_closure_by_valueINS2_18for_each_n_closureINS_10device_ptrI7my_en
umEElNS1_7generic6detail16generate_functorINS3_12fill_functorIS7_EEEEEEEEvT_[voi
d
thrust::detail::device::cuda::detail::launch_closure_by_value<thrust::detail::de
vice::cuda::for_each_n_closure<thrust::device_ptr<my_enum>,
long,
thrust::detail::device::generic::detail::generate_functor<thrust::detail::device
::cuda::detail::fill_functor<my_enum>
> >
>(thrust::detail::device::cuda::for_each_n_closure<thrust::device_ptr<my_enum>,
long,
thrust::detail::device::generic::detail::generate_functor<thrust::detail::device
::cuda::detail::fill_functor<my_enum>
> >)]+0x11): undefined reference to `void
thrust::detail::device::cuda::detail::__wrapper__device_stub_launch_closure_by_v
alue<thrust::detail::device::cuda::for_each_n_closure<thrust::device_ptr<my_enum
>,
long,
thrust::detail::device::generic::detail::generate_functor<thrust::detail::device
::cuda::detail::fill_functor<my_enum>
> >
>(thrust::detail::device::cuda::for_each_n_closure<thrust::device_ptr<my_enum>,
long,
thrust::detail::device::generic::detail::generate_functor<thrust::detail::device
::cuda::detail::fill_functor<my_enum>
> >&)'
/tmp/tmpxft_0000696c_00000000-12_foo2.o: In function `void
thrust::detail::device::cuda::detail::launch_closure_by_value<thrust::detail::de
vice::cuda::for_each_n_closure<thrust::device_ptr<my_enum>,
unsigned int,
thrust::detail::device::generic::detail::generate_functor<thrust::detail::device
::cuda::detail::fill_functor<my_enum>
> >
>(thrust::detail::device::cuda::for_each_n_closure<thrust::device_ptr<my_enum>,
unsigned int,
thrust::detail::device::generic::detail::generate_functor<thrust::detail::device
::cuda::detail::fill_functor<my_enum>
> >)':
tmpxft_0000696c_00000000-1_foo2.cudafe1.cpp:(.text._ZN6thrust6detail6device4cuda
6detail23launch_closure_by_valueINS2_18for_each_n_closureINS_10device_ptrI7my_en
umEEjNS1_7generic6detail16generate_functorINS3_12fill_functorIS7_EEEEEEEEvT_[voi
d
thrust::detail::device::cuda::detail::launch_closure_by_value<thrust::detail::de
vice::cuda::for_each_n_closure<thrust::device_ptr<my_enum>,
unsigned int,
thrust::detail::device::generic::detail::generate_functor<thrust::detail::device
::cuda::detail::fill_functor<my_enum>
> >
>(thrust::detail::device::cuda::for_each_n_closure<thrust::device_ptr<my_enum>,
unsigned int,
thrust::detail::device::generic::detail::generate_functor<thrust::detail::device
::cuda::detail::fill_functor<my_enum>
> >)]+0x11): undefined reference to `void
thrust::detail::device::cuda::detail::__wrapper__device_stub_launch_closure_by_v
alue<thrust::detail::device::cuda::for_each_n_closure<thrust::device_ptr<my_enum
>,
unsigned int,
thrust::detail::device::generic::detail::generate_functor<thrust::detail::device
::cuda::detail::fill_functor<my_enum>
> >
>(thrust::detail::device::cuda::for_each_n_closure<thrust::device_ptr<my_enum>,
unsigned int,
thrust::detail::device::generic::detail::generate_functor<thrust::detail::device
::cuda::detail::fill_functor<my_enum>
> >&)'
collect2: ld returned 1 exit status

Original issue reported on code.google.com by wnbell on 21 Jan 2010 at 4:57

[GoogleCodeExporter]

Original comment by wnbell on 19 Mar 2010 at 2:17

• Added labels: Milestone-Release1.3
• Removed labels: Milestone-Release1.2

[GoogleCodeExporter]

One workaround is to not dispatch wide_fill when sizeof(T) == 4 (and change 
unsigned
long long to just unsigned int).  This is the right choice for performance 
reasons
anyway.

Still, we should find out why enums are different.

Original comment by wnbell on 19 Apr 2010 at 9:22

[GoogleCodeExporter]

Original comment by wnbell on 3 Sep 2010 at 6:31

[GoogleCodeExporter]

I think 64-bit writes might actually be faster, so the WAR proposed in Comment 
2 would be less desirable.

Original comment by wnbell on 3 Sep 2010 at 6:34

[GoogleCodeExporter]

I believe this is a minimal reproducer

#include <iostream>

struct true_type  {};

template<typename T, bool = __has_trivial_assign(T)> struct has_trivial_assign;

template<typename T> struct has_trivial_assign<T,true> : true_type {};

void foo(true_type)
{
}

enum my_enum {red,green,blue};

int main(void)
{
  std::cout << "__has_trivial_assign(my_enum): " << __has_trivial_assign(my_enum) << std::endl;

  // commenting out this offending line reveals that __has_trivial_assign(my_enum) is true
  foo(has_trivial_assign<my_enum>());

  return 0;
}

Original comment by jaredhoberock on 3 Sep 2010 at 10:54

• Changed state: Forwarded

[GoogleCodeExporter]

This is now fixed in nvcc.

Original comment by jaredhoberock on 1 Feb 2011 at 7:22