La transformation de poussée génère une erreur: «bulk_kernel_by_value: un access mémoire illégal a été rencontré»

Je suis plutôt nouveau dans CUDA / Thrust et j’ai un problème avec un extrait de code. Pour faciliter les choses, je l’ai réduit au ssortingct minimum. Le code est le suivant:

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 X(4,1); thrust::transform(X.begin(), X.end(), X.begin(), functor(&g)); }

L’idée est que je peux transmettre n’importe quelle fonction au foncteur, afin de pouvoir appliquer cette fonction à chaque élément d’un vecteur. Malheureusement, je ne sais pas pourquoi j’obtiens l’erreur décrite. Je comstack avec -w -O3 -shared -arch=sm_20 -std=c++11 -DTHRUST_DEBUG

Je suis reconnaissant pour toute aide que vous pouvez tous me donner 🙂

L’adresse d’une fonction __device__ (ou __host__ __device__ ) ne peut pas être prise en code hôte pour être utilisée sur le périphérique:

 thrust::transform(X.begin(), X.end(), X.begin(), functor(&g)); ^ You will not get the __device__ function address here

Plusieurs questions sur stackoverflow traitent de l’utilisation des adresses de fonction de périphérique CUDA transmises via les appels du kernel. Cette réponse renvoie à plusieurs questions pouvant présenter un intérêt.

Une solution possible pour résoudre ce problème consiste à acquérir l’adresse de la fonction de périphérique dans le code de périphérique et à la transmettre à l’hôte, pour une utilisation similaire à celle que vous décrivez:

 $ cat t1057.cu #include  #include  #include  #include  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 X(4,1); thrust::transform(X.begin(), X.end(), X.begin(), functor(h_g)); thrust::copy_n(X.begin(), X.size(), std::ostream_iterator(std::cout, ",")); std::cout << std::endl; } $ nvcc -o t1057 t1057.cu -std=c++11 $ ./t1057 3,3,3,3, $

Une autre approche possible, exploitant le travail toujours intelligent de @ms , utilise ici des modèles:

 $ cat t1057.cu #include  #include  #include  #include  typedef float(*fptr_t)(const float&); template  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 X(4,1); thrust::transform(X.begin(), X.end(), X.begin(), functor()); thrust::copy_n(X.begin(), X.size(), std::ostream_iterator(std::cout, ",")); std::cout << std::endl; } $ nvcc -o t1057 t1057.cu -std=c++11 $ ./t1057 3,3,3,3, $