Ultimately, you are asking to take a __device__
function argument in host code, and then pass it as a (function) pointer, in what is ultimately (under the hood) a kernel argument, generated by thrust.
It is illegal to take the address of a __device__
function argument in host code, so passing a __device__
function pointer as an argument this way won't work.
It might be possible to work around this by creating additional __device__
variables (pointers) to store function pointers on the device. Then use cudaGetSymbolAddress
to build a table of pointers-to-pointers to functions. This would necessitate running a precursor kernel to set up the function pointers on the device. It seems rather messy.
It might be simpler to parameterize the functor to select a device function based on the parameter. Lke this:
#include <thrust/device_vector.h>
#include <thrust/transform.h>
#include <thrust/sequence.h>
#include <thrust/copy.h>
#include <thrust/fill.h>
#include <thrust/replace.h>
#include <thrust/functional.h>
#include <thrust/for_each.h>
#include <thrust/iterator/zip_iterator.h>
#include <iostream>
#include <math.h>
__host__ __device__ float f1(float x)
{
return sinf(x);
}
__host__ __device__ float f2(float x)
{
return cosf(x);
}
struct euler_functor
{
unsigned h;
euler_functor(unsigned _h) : h(_h) {};
__host__ __device__
void operator()(float &y) const {
if (h == 1) y = f1(y);
else if (h == 2) y = f2(y);
}
};
int main(void)
{
const unsigned N = 8;
// allocate three device_vectors with 10 elements
thrust::device_vector<float> X(N);
// initilaize to random vaues
thrust::sequence(X.begin(), X.end(), 0.0f, (float)(6.283/(float)N));
// apply euler for each element of X
thrust::for_each(X.begin(),X.end(),euler_functor(1));
// print the values
for(int i = 0; i < N; i++) std::cout<< X[i]<< std::endl;
std::cout << "******************" << std::endl;
thrust::sequence(X.begin(), X.end(), 0.0f, (float)(6.283/(float)N));
// apply euler for each element of X
thrust::for_each(X.begin(),X.end(),euler_functor(2));
// print the values
for(int i = 0; i < N; i++) std::cout<< X[i]<< std::endl;
}