Skip to content

Instantly share code, notes, and snippets.

@nwh
Created March 2, 2015 00:12
Show Gist options
  • Save nwh/7cf44c8a9217599b659e to your computer and use it in GitHub Desktop.
Save nwh/7cf44c8a9217599b659e to your computer and use it in GitHub Desktop.
Thrust functor in anonymous namespace

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
#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;
}
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment