推力变换抛出错误：“bulk_kernel_by_value：遇到非法存储器访问” [英] Thrust transform throws error: "bulk_kernel_by_value: an illegal memory access was encountered"

问题描述

我是一个新的CUDA / Thrust，并有一个代码段的问题。
为了使它更容易我已经修剪它到最低限度。
代码如下：

I'm rather new to CUDA/Thrust and have a problem with a code snippet. To make it easier I have trimmed it down to the bare minimum. The code is the following:

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));
}

这个想法是我可以将任何函数传递给函子，将该函数应用于向量中的每个元素。
不幸的是，我不确定为什么我得到描述的错误。
I编译 -w -O3 -shared -arch = sm_20 -std = c ++ 11 -DTHRUST_DEBUG

The idea is that I can pass any function to the functor, so I can apply that function to every element in a Vector. Unfortunately I'm uncertain to why I get the described error. I Compile with -w -O3 -shared -arch=sm_20 -std=c++11 -DTHRUST_DEBUG

感谢您能给我的任何帮助：）

I'm thankful for any help you all can give me :)

推荐答案

不能在主机代码中使用$ c> __ device __ 函数（或 __ host__ __device __ ），以便在设备上使用：

The address of a __device__ function, (or __host__ __device__) cannot be taken in host 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

在stackoverflow有很多问题，讨论通过内核调用传递的CUDA设备函数地址的使用。 此答案
链接到可能感兴趣的几个。

There are many questions on stackoverflow which discuss usage of CUDA device function addresses passed via kernel calls. This answer links to several which may be of interest.

一个可能的解决方法是获取设备代码中的设备功能地址，它到主机，用于像你描述：

One possible approach to fix this would be to acquire the device function address in device code, and pass it to the host, for usage like you are describing:

$ 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,
$

另一种可能的方法是利用@ms永远聪明的工作这里使用模板：

Another possible approach, leveraging the always-clever work by @m.s. here uses templating:

$ 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,
$

这篇关于推力变换抛出错误：“bulk_kernel_by_value：遇到非法存储器访问”的文章就介绍到这了，希望我们推荐的答案对大家有所帮助，也希望大家多多支持IT屋！