一个类可以简单地传递给CUDA内核进行并行评估吗

Can a class simply be passed to a CUDA kernel for parallel evaluation?

本文关键字：内核 CUDA 评估 并行 一个 简单      更新时间：2023-10-16

我有一个相当复杂的c++类，它具有myObj.fun(x,y,z)类的函数。我想在3D网格上的GPU上调用这个。

我的高级问题是：将myObj和大量点传递给GPU是一件简单的事情吗？由于工作原因，我避免创建该函数的CUDA实现，但我突然想到这可能非常简单。

同样，这是一个非常高层次的问题，所以"是的，有那么容易"或"不是，更复杂"是受欢迎的，尽管如果方向也有帮助的话。

可以在GPU上使用类/对象，包括它们的方法（例如fun()）。这样的类至少必须具有用__host__ __device__修饰的方法，但代码重构可能不会比这更复杂。

然而，这样的方法（就像任何其他没有重构的代码一样）可能不会访问GPU的任何并行功能。最基本的比较是，在单个CPU线程中运行的方法将在单个GPU线程中运行。如果你所做的只是将单个对象传递给GPU，并在GPU上运行等效的单线程代码（在单个GPU线程中），那么这通常不会更快，而且通常会慢得多。

一种可能的策略是，如果你有很多这样的对象，或者等效地，在你的情况下，有很多"点"，代表每个对象上要做的独立工作，那么你可以将每个对象（对象或点）传递给GPU线程，并以这种方式处理它们，从而实现GPU喜欢的大规模多线程操作。理想情况下，你会有10000或更多的点来处理这种方式。

由于各种原因，这通常仍然不是GPU的最有效使用，其中一个原因与高效的数据访问有关，另一个原因则与（可能的）线程分歧有关。尽管如此，有些人还是追求这种"简单"、"令人尴尬的并行"的代码移植方法，偶尔还会有有趣的加速。

根据您的实际代码，如果您以允许相邻线程访问相邻数据的方式将点传递给GPU，对于每个访问点的操作，您可能会看到有吸引力的结果。可以合理地预期，你可能会以这种方式获得有吸引力的速度，也许只需要相对较少的代码重构，但要注意数据组织以实现最佳GPU访问。

下面是一个完整的例子：

$ cat t30.cu
#include <iostream>
#include <cstdlib>
const int dsize = 3;
const int nTPB = 256;
const int rng = 8;
class myclass
{
  int increment;
  public:
    myclass(int _incr): increment(_incr) {};
    // methods callable on the device need the __device__ decoration
    __host__ __device__ void fun(int &x, int &y, int &z){
      x += increment;
      y += increment;
      z += increment;}
};
// this is the actual device routine that is run per thread
__global__ void mykernel(myclass obj, int *dx, int *dy, int *dz, int dsize){
  int idx = threadIdx.x+blockDim.x*blockIdx.x; // figure out which thread we are
  if (idx < dsize)
    obj.fun(dx[idx], dy[idx], dz[idx]);  // apply method
}

int main(){
  // allocate host data
  int *p_x, *p_y, *p_z, *d_x, *d_y, *d_z;
  p_x = new int[dsize];
  p_y = new int[dsize];
  p_z = new int[dsize];
  // allocate device data
  cudaMalloc(&d_x, dsize*sizeof(int));
  cudaMalloc(&d_y, dsize*sizeof(int));
  cudaMalloc(&d_z, dsize*sizeof(int));
  // initialize host data
  std::cout << "Before:" << std::endl;
  for (int i = 0; i < dsize; i++){
    p_x[i] = rand()%rng;
    p_y[i] = rand()%rng;
    p_z[i] = rand()%rng;
    std::cout << p_x[i] << "," << p_y[i] << "," << p_z[i] << std::endl;}
  // copy to device
  cudaMemcpy(d_x, p_x, dsize*sizeof(int), cudaMemcpyHostToDevice);
  cudaMemcpy(d_y, p_y, dsize*sizeof(int), cudaMemcpyHostToDevice);
  cudaMemcpy(d_z, p_z, dsize*sizeof(int), cudaMemcpyHostToDevice);
  // instantiate object on host
  myclass test(1);
  // copy object to device as kernel parameter
  mykernel<<<(dsize+nTPB-1)/nTPB, nTPB>>>(test, d_x, d_y, d_z, dsize);
  // copy data back to host
  cudaMemcpy(p_x, d_x, dsize*sizeof(int), cudaMemcpyDeviceToHost);
  cudaMemcpy(p_y, d_y, dsize*sizeof(int), cudaMemcpyDeviceToHost);
  cudaMemcpy(p_z, d_z, dsize*sizeof(int), cudaMemcpyDeviceToHost);

  std::cout << "After:" << std::endl;
  for (int i = 0; i < dsize; i++){
    std::cout << p_x[i] << "," << p_y[i] << "," << p_z[i] << std::endl;}
  return 0;
}
$ nvcc -o t30 t30.cu
$ ./t30
Before:
7,6,1
3,1,7
2,4,1
After:
8,7,2
4,2,8
3,5,2
$

为了简洁起见，我省略了正确的cuda错误检查，但我始终建议您在开发cuda代码时使用它。

相关文章：