C 使用类调用 CUDA 核函数

(给CPP开发者加星标，提升C/C++技能)

来源：小小一步

导读：CUDA是用于GPU编程的框架，在深度学习高速发展的今天，应用十分普遍。然而虽然CUDA对C语言有很好的支持，但是在C++的类函数中使用CUDA 核函数会遇到很多问题。文章介绍了如何通过封装的方式在C++类中调用CUDA的核函数，结合具体的实例。推荐给对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 200class 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();}

部分结果：

- EOF -