C++使用类调用CUDA核函数

正如CUDA C所称，CUDA对C语言进行了很好的扩展，直接使用C语言可以非常简单方便的调用CUDA核函数。但是当想使用C++的类成员函数直接调用核函数是不可行的，第一，核函数不能作为类的成员函数，第二，C++的cpp文件和CUDA的cu文件分别经由g++和nvcc编译，当两种代码混合就会编译出错。

因而C++的类和CUDA结合使用需要进行一层封装，借用两个数组相加的例子说明，主要过程如下：

本项目包括4文件，如上图所示，分别为：add.h、add.cpp、kernel.cuh、kernel.cu。在add.h中封装一个函数AddNum()，调用kernel.cuh中的函数

AddKernel(int *a, int *b, int *c, int DX)，然后在kernel.cu文件中使用AddKernel(...)调用相加核函数Add(int *a, int *b, int *c, int DX)

1、add.h文件定义了一个CTest的类，包括3个指针（数组）、4个函数。

#pragma once

#include "kernel.cuh"

#include <iostream>
using namespace std;

#define DX 200


class CTest
{
public:
    int *a;
    int *b;
    int *c;

    void SetParameter();
    void AddNum();
    void Show();
    void Evolution();
};

void CTest::SetParameter()
{
    cudaMallocManaged(&a, sizeof(int) * DX);
    cudaMallocManaged(&b, sizeof(int) * DX);
    cudaMallocManaged(&c, sizeof(int) * DX);

    for (int f = 0; f<DX; f++)
    {
        a[f] = f;
        b[f] = f + 1;
    }

}

void CTest::AddNum()
{
    AddKernel(a, b, c, DX);
}

void CTest::Show()
{
    cout << " a     b    c"  << endl;

    for (int f = 0; f<DX; f++)
    {
        cout << a[f] << " + " << b[f] << "  = " << c[f] << endl;
    }
}

void CTest::Evolution()
{
    SetParameter();
    AddNum();
    Show();
}

2、add.cpp文件执行主函数，创建一个CTest的对象cTest，然后调用Evolution执行相加操作。

#include "add.h"

void main()
{
    CTest cTest;

    cTest.Evolution();

    system("pause");
}

3、kernel.cuh文件定义一个接口函数AddKernel(int *a, int *b, int *c, int DX);

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

void AddKernel(int *a, int *b, int *c, int DX);

4、kernel.cu包括相加Add核函数以及调用核函数的封装函数AddKernel.

#include "kernel.cuh"


__global__ void Add(int *a, int *b, int *c, int DX)
{
    int f = blockIdx.x*blockDim.x + threadIdx.x;

    if (f >= DX) return;

    c[f] = a[f] + b[f];

}

void AddKernel(int *a, int *b, int *c, int DX)
{
    dim3 dimBlock = (128);
    dim3 dimGrid = ((DX + 128 - 1) / 128);
    Add << <dimGrid, dimBlock >> > (a, b, c, DX);
    cudaDeviceSynchronize();
}

部分结果：