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();
}

部分結果：