Thrust functor in anonymous namespace

readme.md

Beginning of compiler error

/usr/local/cuda/bin/..//include/thrust/detail/function.h(60): error: calling a __device__ function("operator()") from a __host__ __device__ function("operator()") is not allowed
          detected during:
            instantiation of "Result thrust::detail::wrapped_function<Function, Result>::operator()(const Argument &) const [with Function=<unnamed>::RngInit, Result=void, Argument=signed int]" 
/usr/local/cuda/bin/..//include/thrust/system/detail/sequential/for_each.h(83): here
            instantiation of "InputIterator thrust::system::detail::sequential::for_each_n(thrust::system::detail::sequential::execution_policy<DerivedPolicy> &, InputIterator, Size, UnaryFunction) [with DerivedPolicy=thrust::detail::seq_t, InputIterator=thrust::counting_iterator<int, thrust::use_default, thrust::use_default, thrust::use_default>, Size=signed long, UnaryFunction=<unnamed>::RngInit]" 

System

• CUDA 7.0 RC

thrust_curand.cu

#include <thrust/iterator/counting_iterator.h>
#include <thrust/functional.h>
#include <thrust/for_each.h>
#include <thrust/device_vector.h>
#include <thrust/execution_policy.h>
#include <curand_kernel.h>

// fix: remove this namespace
namespace {

struct RngInit : public thrust::unary_function<int, void> {
  int seed;
  curandState_t* rng;
  __device__
  void operator()(int id) {
    curandState_t* rng_item = rng + id;
    curand_init(seed,id,0,rng_item);
  }
};

} // end of anonymous namespace
 
int main() {
  const int N = 10;
  thrust::device_vector<curandState_t> rng(N);
  // create counting iterators
  thrust::counting_iterator<int> first(0);
  thrust::counting_iterator<int> last = first + N;
  // set up functor
  RngInit rng_init;
  rng_init.seed = 0;
  rng_init.rng = thrust::raw_pointer_cast(rng.data());
  // call for each
  thrust::for_each(thrust::device,first,last,rng_init);
  return 0; 
}