在CUDA设备代码中使用std :: vector

Question

问题是:有没有办法在Cuda内核中使用类"向量"？当我尝试时,我收到以下错误:

error : calling a host function("std::vector<int, std::allocator<int> > ::push_back") from a __device__/__global__ function not allowed

那么有一种方法可以在全局部分使用向量吗？我最近尝试了以下内容:

1. 创建一个新的Cuda项目
2. 转到项目的属性
3. 打开Cuda C/C++
4. 转到设备
5. 将"代码生成"中的值更改为此值:compute_20,sm_20

........之后我能够在我的Cuda内核中使用printf标准库函数.

有没有办法vector在内核代码中支持printf 的方式使用标准库类？这是在内核代码中使用printf的示例:

// this code only to count the 3s in an array using Cuda
//private_count is an array to hold every thread's result separately 

__global__ void countKernel(int *a, int length, int* private_count) 
{
    printf("%d\n",threadIdx.x);  //it's print the thread id and it's working

    // vector<int> y;
    //y.push_back(0); is there a possibility to do this?

    unsigned int offset  = threadIdx.x * length;
    int i = offset;
    for( ; i < offset + length; i++)
    {
        if(a[i] == 3)
        {
            private_count[threadIdx.x]++;
            printf("%d ",a[i]);
        }
    }   
}

Answer 1

您不能在CUDA中使用STL,但您可以使用Thrust库来执行您想要的操作.否则只需将向量的内容复制到设备并正常操作即可.

Answer 2

在cuda库推进中,您可以使用thrust::device_vector<classT>在设备上定义向量,并且主机STL向量和设备向量之间的数据传输非常简单.你可以参考这个有用的链接:http://docs.nvidia.com/cuda/thrust/index.html找到一些有用的例子.

Answer 3

你不能std::vector在设备代码中使用,你应该使用数组.

Answer 4

我认为你可以自己实现一个设备向量，因为CUDA支持设备代码中的动态内存分配。还支持运算符新建/删除。这是 CUDA 中一个非常简单的设备向量原型，但它确实有效。它没有经过充分的测试。

template<typename T>
class LocalVector
{
private:
    T* m_begin;
    T* m_end;

    size_t capacity;
    size_t length;
    __device__ void expand() {
        capacity *= 2;
        size_t tempLength = (m_end - m_begin);
        T* tempBegin = new T[capacity];

        memcpy(tempBegin, m_begin, tempLength * sizeof(T));
        delete[] m_begin;
        m_begin = tempBegin;
        m_end = m_begin + tempLength;
        length = static_cast<size_t>(m_end - m_begin);
    }
public:
    __device__  explicit LocalVector() : length(0), capacity(16) {
        m_begin = new T[capacity];
        m_end = m_begin;
    }
    __device__ T& operator[] (unsigned int index) {
        return *(m_begin + index);//*(begin+index)
    }
    __device__ T* begin() {
        return m_begin;
    }
    __device__ T* end() {
        return m_end;
    }
    __device__ ~LocalVector()
    {
        delete[] m_begin;
        m_begin = nullptr;
    }

    __device__ void add(T t) {

        if ((m_end - m_begin) >= capacity) {
            expand();
        }

        new (m_end) T(t);
        m_end++;
        length++;
    }
    __device__ T pop() {
        T endElement = (*m_end);
        delete m_end;
        m_end--;
        return endElement;
    }

    __device__ size_t getSize() {
        return length;
    }
};