c – Thrust转换抛出错误：“bulk_kernel_by_value：遇到非法内存访问”

2024-03-24 08:57:22

我对CUDA / Thrust很新,并且在代码片段方面存在问题.
为了使它更容易,我把它修剪到最低限度.
代码如下：

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

我的想法是我可以将任何函数传递给仿函数,因此我可以将该函数应用于Vector中的每个元素.
不幸的是,我不确定为什么我会得到描述的错误.
我编译-w -O3 -shared -arch = sm_20 -std = c 11 -DTHRUST_DEBUG

我很感谢你能给我的任何帮助:)

解决方法:

__device__函数的地址(或__host__ __device__)不能在主机代码中使用,以便在设备上使用：

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

*有很多问题,讨论通过内核调用传递的CUDA设备函数地址的用法. This answer
链接到可能感兴趣的几个.

解决此问题的一种可能方法是获取设备代码中的设备功能地址,并将其传递给主机,以便像您所描述的那样使用：

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

另一种可能的方法,利用@ m.s的始终聪明的工作. here使用模板：

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

码农公寓