Pull conversion throws error: "bulk_kernel_by_value: illegal memory access was encountered"
I am new to CUDA / Thrust and have a code snippet problem. To make it easier, I cut it to a minimum. The code is as follows:
struct functor{
functor(float (*g)(const float&)) : _g{g} {}
__host__ __device__ float operator()(const float& x) const {
return _g(x);
}
private:
float (*_g)(const float&);
};
__host__ __device__ float g(const float& x){return 3*x;}
int main(void){
thrust::device_vector<float> X(4,1);
thrust::transform(X.begin(), X.end(), X.begin(), functor(&g));
}
The idea is that I can pass any function to the functor, so I can apply that function to every element in the vector. Unfortunately I don't understand why I got the described error. I Compile with-w -O3 -shared -arch=sm_20 -std=c++11 -DTHRUST_DEBUG
I am grateful for any help you can give me :)
source to share
The function address __device__
(or __host__ __device__
) cannot be taken in the main code for use on the device:
thrust::transform(X.begin(), X.end(), X.begin(), functor(&g));
^
You will not get the
__device__ function
address here
There are many questions on stackoverflow that discuss the use of CUDA device function addresses passed through kernel calls. This answer links to several that might be of interest.
One possible approach to fix this is to get the device function address in the device code and pass it to the host for use, as you describe:
$ cat t1057.cu
#include <thrust/device_vector.h>
#include <thrust/transform.h>
#include <thrust/copy.h>
#include <iostream>
struct functor{
functor(float (*g)(const float&)) : _g{g} {}
__host__ __device__ float operator()(const float& x) const {
return _g(x);
}
private:
float (*_g)(const float&);
};
__host__ __device__ float g(const float& x){return 3*x;}
__device__ float (*d_g)(const float&) = g;
int main(void){
float (*h_g)(const float&) = NULL;
cudaMemcpyFromSymbol(&h_g, d_g, sizeof(void *));
thrust::device_vector<float> X(4,1);
thrust::transform(X.begin(), X.end(), X.begin(), functor(h_g));
thrust::copy_n(X.begin(), X.size(), std::ostream_iterator<float>(std::cout, ","));
std::cout << std::endl;
}
$ nvcc -o t1057 t1057.cu -std=c++11
$ ./t1057
3,3,3,3,
$
Another possible approach using @ms always smart work here uses templates:
$ cat t1057.cu
#include <thrust/device_vector.h>
#include <thrust/transform.h>
#include <thrust/copy.h>
#include <iostream>
typedef float(*fptr_t)(const float&);
template <fptr_t F>
struct functor{
__host__ __device__ float operator()(const float& x) const {
return F(x);
}
};
__host__ __device__ float g(const float& x){return 3*x;}
int main(void){
thrust::device_vector<float> X(4,1);
thrust::transform(X.begin(), X.end(), X.begin(), functor<g>());
thrust::copy_n(X.begin(), X.size(), std::ostream_iterator<float>(std::cout, ","));
std::cout << std::endl;
}
$ nvcc -o t1057 t1057.cu -std=c++11
$ ./t1057
3,3,3,3,
$
source to share