Cuda C++ 设计：编译时大小未知的可重用类答案

【问题标题】：Cuda C++ design: reusable class with unknown compile-time sizeCuda C++ 设计：编译时大小未知的可重用类
【发布时间】：2018-11-03 16:42:48
【问题描述】：

我正在寻找一种方便的设计，以便能够在设备上使用编译时大小未知的类。只需将此类的一个实例发送到设备，就应该对 cudaMalloc 和 cudaMemcpy 进行一次调用（理想情况下）。

该类的宿主版本如下所示：

Class A {
public:
A(int size) : table(size) {
 // some useful initialization of table
}
double get(int i) const {
  // return some processed element from table 
}
private:
std::vector<int> table;
};

内核：

__global__ void kernel(const A *a){
  int idx = threadIdx.x + blockDim.x * blockIdx.x;
  a->get(idx);  // do something useful with it
}

到目前为止，我设计类的设备版本的方式是这样的：

const int sizeMax = 1000;
Class A {
public:
A(int size) {
 // size checking + some useful initialization of table
}
__host__ __device__
double get(int i) const {
  // 
}
private:
int table[sizeMax];
};

以及客户端代码：

A a(128);
A* da;
cudaMalloc((void**)&da, sizeof(A));
cudaMemcpy(da, &a, sizeof(A), cudaMemcpyHostToDevice);
kernel<<<1, 32>>>(da);
cudaDeviceSynchronize();
cudaFree(da);

这很丑，因为：

它不得不使用太大的 sizeMax 来浪费带宽保持安全
类没有关闭修改，sizeMax的值会不可避免地需要在某个时候提出来

是否有任何其他方法可以以更简洁的方式实现相同的目标，而不会对性能产生负面影响？明确一点，我只需要类的设备版本，第一个版本只是等效的非CUDA代码来说明表大小应该是动态的。

【问题讨论】：

1. table 的独立主机和设备存储，包含在类中，两者都是动态分配的。 2. 在构造函数中动态分配表存储大小，而不是在您的客户端代码中。这也可能包括在必要时调整大小。 3. 区分使用数据的主机副本或设备副本（即指针）的类方法，具体取决于方法是在主机代码还是设备代码中执行 4. 将数据从主机复制到的方法设备，反之亦然，因为类上下文从主机移动到设备，反之亦然。
另一种方法：CUDA managed (i.e. unified) memory
@RobertCrovella 在第 2 点，这意味着用户自己实现 std::vector 功能？使用thrust::device_vector<int> 之类的实现怎么样？
@BAdhi：不，这意味着在运行时分配一些内存并将该内存的地址分配给结构中的指针，并且没有推力::device_vector 不能使用，因为它不是设备代码支持。

标签： c++ class memory-management dynamic cuda

【解决方案1】：

在我的评论中，我说：

表的主机和设备存储分开，包含在类中，两者都是动态分配的。 2. 在构造函数中动态分配表存储大小，而不是在您的客户端代码中。这也可能包括在必要时调整大小。 3. 区分使用数据的主机副本或设备副本（即指针）的类方法，具体取决于方法是在主机代码还是设备代码中执行 4. 将数据从主机复制到的方法设备，反之亦然，因为类上下文从主机移动到设备，反之亦然。

这是我想到的一个例子：

#include <stdio.h>
#include <assert.h>
#include <cuda_runtime_api.h>
#include <iostream>


template <typename T>
class gpuvec{
  private:
    T *h_vec = NULL;
    T *d_vec = NULL;
    size_t vsize = 0;
    bool iscopy;
  public:
    __host__ __device__
    T * data(){
      #ifndef __CUDA_ARCH__
        return h_vec;
      #else
        return d_vec;
      #endif
      }
    __host__ __device__
    T& operator[](size_t i) {
      assert(i < vsize);
        return data()[i];}
    void to_device(){
      assert(cudaMemcpy(d_vec, h_vec, vsize*sizeof(T), cudaMemcpyHostToDevice) == cudaSuccess);}
    void to_host(){
      assert(cudaMemcpy(h_vec, d_vec, vsize*sizeof(T), cudaMemcpyDeviceToHost) == cudaSuccess);}
    gpuvec(gpuvec &o){
      h_vec = o.h_vec;
      d_vec = o.d_vec;
      vsize = o.vsize;
      iscopy = true;}
    void copy(gpuvec &o){
      free();
      iscopy = false;
      vsize = o.vsize;
      h_vec = (T *)malloc(vsize*sizeof(T));
      assert(h_vec != NULL);
      assert(cudaMalloc(&d_vec, vsize*sizeof(T)) == cudaSuccess);
      memcpy(h_vec, o.h_vec, vsize*sizeof(T));
      assert(cudaMemcpy(d_vec, o.d_vec, vsize*sizeof(T), cudaMemcpyDeviceToDevice) == cudaSuccess);}
    gpuvec(size_t ds) {
      assert(ds > 0);
      iscopy = false;
      vsize = ds;
      h_vec = (T *)malloc(vsize*sizeof(T));
      assert(h_vec != NULL);
      assert(cudaMalloc(&d_vec, vsize*sizeof(T)) == cudaSuccess);}
    gpuvec(){
      iscopy = false;
    }
    ~gpuvec(){
      if (!iscopy) free();}
    void free(){
      if (d_vec != NULL) cudaFree(d_vec); 
      d_vec = NULL;
      if (h_vec != NULL) ::free(h_vec);
      h_vec = NULL;}
    __host__ __device__
    size_t size() {
      return vsize;}
};

template <typename T>
__global__ void test(gpuvec<T> d){
  for (int i = 0; i < d.size(); i++){
    d[i] += 1;
    }
}


int main(){
  size_t ds = 10;
  gpuvec<int>  A(ds);
  A.to_device();
  test<<<1,1>>>(A);
  A.to_host();
  for (size_t i = 0; i < ds; i++)
    std::cout << A[i];
  std::cout << std::endl;
  gpuvec<int> B;
  B.copy(A);
  A.free();
  B.to_device();
  test<<<1,1>>>(B);
  B.to_host();
  for (size_t i = 0; i < ds; i++)
    std::cout << B[i];
  std::cout << std::endl;
  B.free();
}

我敢肯定会有不少批评。这可能不符合“向量语法”应该是什么的任何特定意见。此外，我确信它没有涵盖一些用例，并且它可能包含彻底的缺陷。要创建一个健壮的主机/设备向量实现可能需要与thrust 主机和设备向量一样多的工作和复杂性。但是，我并不是说推力矢量是问题似乎要问的问题的直接答案。

【讨论】：

【解决方案2】：

根据 Robert Crovella 的回答，这是一个简化的（仅限设备，因此忽略第 3 点和第 4 点）工作解决方案：

Class A {
public:
A(int size) : table(size) {
 // some useful initialization of table
 cudaMalloc((void**)&dTable, sizeof(int) * size);
 cudaMemcpy(dTable, &table[0], sizeof(int) * size, cudaMemcpyHostToDevice);
}
~A() {
cudaFree(dTable);
}
__device__
double get(int i) const {
  // return some processed element of dTable 
}
private:
std::vector<int> table;
int *dTable; 
};

内核和客户端代码保持完全相同。

【讨论】：