教你一步步写一个cuda path tracer：cuda与类

cuda里能不能使用类？答案当然是可以的。
下面给出一个实例：

arrayadder.cuh

#ifndef ARRAYADDER_CUH
#define ARRAYADDER_CUH

#include <cuda_runtime.h>
#include <device_launch_parameters.h>

class ArrayAdder
{
public:
    ArrayAdder(int *x, int *y, int *z, int n);
    ~ArrayAdder();

    void add();

private:
    __device__ void device_add();
    __global__ friend void global_add(ArrayAdder *arrayAdder);

    int *dx, *dy, *dz, *z;
    int n;
};

#endif // ARRAYADDER_CUH

arrayadder.cu

#include "arrayadder.cuh"

ArrayAdder::ArrayAdder(int *x, int *y, int *z, int n) : n(n)
{
    cudaMalloc(&dx, n * sizeof(int));
    cudaMalloc(&dy, n * sizeof(int));
    cudaMalloc(&dz, n * sizeof(int));

    cudaMemcpy(dx, x, n * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(dy, y, n * sizeof(int), cudaMemcpyHostToDevice);

    this->n = n;
    this->z = z;
}

ArrayAdder::~ArrayAdder()
{
    cudaFree(dx);
    cudaFree(dy);
    cudaFree(dz);
}

__device__ void ArrayAdder::device_add()
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    dz[i] = dx[i] + dy[i];
}

void ArrayAdder::add()
{
    int threadNums = 256;
    int blockNums = (n + threadNums - 1) / threadNums;

    ArrayAdder *dthis;
    cudaMalloc(&dthis, sizeof(ArrayAdder));
    cudaMemcpy(dthis, this, sizeof(ArrayAdder), cudaMemcpyHostToDevice);

    global_add<<<blockNums, threadNums>>>(dthis);

    cudaMemcpy(this, dthis, sizeof(ArrayAdder), cudaMemcpyDeviceToHost);

    cudaMemcpy(z, dz, n * sizeof(int), cudaMemcpyDeviceToHost);

    cudaFree(dthis);
}

__global__ void global_add(ArrayAdder *arrayAdder)
{
    arrayAdder->device_add();
}

main.cpp

#include <cstdio>
#include "arrayadder.cuh"

int main()
{
    int n = 1 << 20;
    int *x, *y, *z;
    x = new int[n];
    y = new int[n];
    z = new int[n];

    for (int i = 0; i < n; i++)
    {
        x[i] = 1;
        y[i] = 2;
    }

    ArrayAdder arrayAdder(x, y, z, n);
    arrayAdder.add();

    int error = 0;
    for (int i = 0; i < n; i++)
        error += z[i] - 3;

    printf("%d\n", error);

    delete[] x;
    delete[] y;
    delete[] z;

    return 0;
}

在main函数中，实例化一个ArrayAdder对象，传入参数，调用计算函数，便可以完成计算。类的使用可以把cuda的核函数调用细节封装起来，使得用户能够直接调用一个清晰的接口而不需要知道底层的实现细节。

有一下三点需要注意：
1.__host__ __device__ __global__
__host__：只能被在cpu端运行的程序；
__device__ ：只能被在gpu端运行的程序；
__global__：在cpu端可以调用的、在gpu端运行的程序。

其中__host__ __device__能够用来修饰类的成员函数，__global__不能（类的静态函数也不行）。__global__只能修饰普通的全局函数。

2.ArrayAdder的add函数

void ArrayAdder::add()
{
    int threadNums = 256;
    int blockNums = (n + threadNums - 1) / threadNums;

    ArrayAdder *dthis;
    cudaMalloc(&dthis, sizeof(ArrayAdder));
    cudaMemcpy(dthis, this, sizeof(ArrayAdder), cudaMemcpyHostToDevice);

    global_add<<<blockNums, threadNums>>>(dthis);

    cudaMemcpy(this, dthis, sizeof(ArrayAdder), cudaMemcpyDeviceToHost);

    cudaMemcpy(z, dz, n * sizeof(int), cudaMemcpyDeviceToHost);

    cudaFree(dthis);
}

由于核函数的参数必须是gpu的数据，所以需要把当前对象的数据拷贝到gpu上，即把this指向的数据拷贝到dthis指向的数据。运算结束后再把dthis复制回this。最后再把gpu上的dz复制到cpu上的z。

3.核函数定义成友元函数
__global__可以修饰类的友元函数，因为友元函数本质上并不是成员函数。这样一来，外界就访问不到核函数了，而且核函数调用的设备函数也可以被定义成私有的了。至此，所有和cuda相关的函数都被封装起来，实现了真正意义上的面向对象。