"ilegal memory access" with thrust::generate on device

Accelerated Computing CUDA CUDA Programming and Performance
1

I feel a bit stupid right now, because I somehow expected, that this would run:

#include <thrust/device_vector.h>
#include <thrust/fill.h>
#include <thrust/generate.h>
#include <thrust/execution_policy.h>

__device__ unsigned zero() { return 0; }

int main()
{
	thrust::device_vector<unsigned> d(1024);
//	thrust::fill(thrust::device,d.begin(), d.end(),0);
	thrust::generate(thrust::device,d.begin(),d.end(),zero);
}

But:

$ nvcc -O -gencode arch=compute_50,code=sm_50 -gencode arch=compute_61,code=sm_61 ../main.cu -o m
$ ./m
terminate called after throwing an instance of 'thrust::system::system_error'
  what():  device free failed: an illegal memory access was encountered
Aborted

What am I doing wrong here? Replacing generate with fill abive works. Cards are GTX1060 and GTX750Ti, CUDA is 9.1 for Ubuntu.

2

Since this works:

#include <thrust/device_vector.h>
#include <thrust/fill.h>
#include <thrust/generate.h>
#include <thrust/execution_policy.h>

__device__ unsigned zero() { return 1; }

template<unsigned T()> struct wrap
{
 public:
  	__device__ unsigned operator () () const { return T(); }	
};

int main()
{
	thrust::device_vector<unsigned> d(1024);
	thrust::generate(thrust::device,d.begin(),d.end(),wrap<zero>());
}

I’m inclined to believe, that thrust does not handle naked functions and/or can’t handle function pointers. Is that true?

3

That’s correct. Thrust expects a functor, not a naked function. This is evident if you read the thrust quick start guide or study any of the sample codes. As an alternative to a functor, you can also use a lambda.

[url]https://devblogs.nvidia.com/new-compiler-features-cuda-8/[/url]

4

Unfortunately, that’s not evident at all. Thrust expects a “function object”, which according to the wording of the C++ standard means (20.10) “1 A function object type is an object type (3.9) that can be the type of the postfix-expression in a function call (5.2.2, 13.3.1.1)[233]”. Footnote 233 clarifies “233) Such a type is a function pointer or a class type which has a member operator() or a class type which has a conversion to a pointer to function.”. Also, they use rand from in their example for thrust::generate:

https://thrust.github.io/doc/group__transformations.html#ga12910ae45ed109cf3eef7df63573e063

5

If you look at the thrust quick start guide or any of the thrust sample codes, you will find numerous examples of functor usage and no examples of bare function usage. I’m not going to parse through or argue about a language definition, but if only a bare function were sufficient, you would find examples like that and there would be no reason to use the “function object” terminology.

If you would like to see a change to the thrust documentation, I suggest filing a thrust issue or pull request.