并行计算之路<3>——CUDA与CPP文件联姻

原内容来源于《GPGPU编程技术——从GLSL、CUDA到OpenCL》的7.1.3节。 有修改。

目的

之前都是单一的*.cu文件。一个较为复杂的工程，往往需要多个文件来实现，合理安排这些文件的结构会使得工程文件尽然有序。《GPGPU编程技术——从GLSL、CUDA到OpenCL》提供了三种方法来实现将CUDA C代码集成到C++中。，

内核代理：将CUDA C程序集成在面向对象的CPP项目中。在CPP语言的帮助下，CUDA C的程序模块也可以呈现面向对象的特性。使用内核代理的基本原理如下：

把CUDA C代码从CPP程序中提取出来，并用文件分割两者，使CPP中类的成员函数“看不到”CUDA C代码的存在。被提取出来的CUDA C代码被一些代理函数封装。于是，代理函数调用CUDA内核，而同时CPP函数调用代理函数。这些代理函数并不包含任何对功能的实现，只是起到了将调用重定向的作用，这样一来，CUDA内核和CPP代码就可以被隔离开。这样做的目的是，尽量封装CUDA C代码，并用nvcc来编译它们，而用CPP编译剩余的CPP代码。

可能画个图会比较好理解。

Created with Raphaël 2.1.0CUDACUDA代理函数代理函数C++C++帮我把礼物转交给女神。代理：那屌丝也配得上。你吖的不知道自己去啊。我精心准备了一份礼物。谢谢代理函数。CUDA:各种梦

架构

1）C++应用程序：application.cpp

#include <iostream>

#include "class.cuh"

using namespace std;

int main(int argc, char **argv)
{
CHelloWorld* hello = new CHelloWorld(argc, argv);
hello->sayHello();

delete hello;

return 0;
}

2）类的定义：class.cuh

#ifndef CLASS_CUH
#define CLASS_CUH

#include <cstdlib>
#include <iostream>

#include "kernel.cuh"

#define BLOCKSIZE 8u

using namespace std;

class CHelloWorld {
public:
CHelloWorld(int argc, char **argv) {
_argc = argc;
_argv = argv;
init();
}

~CHelloWorld();

void sayHello(void);

private:
void init(void);
int _argc;
char** _argv;
unsigned _nBlockSize;
};

#endif

3）类的实现：class.cu

#include "class.cuh"
#include "kernel.cuh"

void CHelloWorld::init()
{
_nBlockSize = (unsigned)BLOCKSIZE;

//TODO:...
}

CHelloWorld::~CHelloWorld()
{
//TODO:...
}

void CHelloWorld::sayHello(void)
{
sayHello_agent(_nBlockSize);
}

4）内核头文件：kernel.cuh

#ifndef KERNEL_CUH
#define KERNEL_CUH

#include <stdio.h>
#include <stdlib.h>
#include <math.h>

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

__global__ void sayHello_kernel(void);

void sayHello_agent(unsigned unBlockSize);

#endif

5）内核实现：kernel.cu

#include "kernel.cuh"

__global__ void sayHello_kernel(void)
{
printf("Hello from thread %d\n", threadIdx.x);
}

void sayHello_agent(unsigned unBlockSize)
{
dim3 dimBlock;
dim3 dimGrid;

dimBlock.x = unBlockSize;
dimBlock.y = 1;
dimBlock.z = 1;
dimGrid.x = 1;
dimGrid.y = 1;
dimGrid.z = 1;

// 代理函数调用内核
sayHello_kernel <<< dimGrid, dimBlock >>>();
cudaThreadExit();
}

输出结果如下：

Hello from thread 0
Hello from thread 1
Hello from thread 2
Hello from thread 3
Hello from thread 4
Hello from thread 5
Hello from thread 6
Hello from thread 7

总结

科技的日新月异促使我们不能再单一的*.cu文件上停留太久。必须将CUDA C代码带到项目中，来加速一个需要依靠GPGPU来获得性能提升的模块，如矩阵运算、机器学习、图像处理等。

在《GPGPU编程技术——从GLSL、CUDA到OpenCL》书中提到过extern “C”，for more info, please visit C++项目中的extern “C”

参考：

《GPGPU编程技术——从GLSL、CUDA到OpenCL》♥♥♥♥♥

《数字图像处理高级应用——基于MATLAB与CUDA的实现》♥♥♥

《基于CUDA的并行程序设计》♥♥♥

《CUDA专家手册》♥♥♥♥♥

《高性能CUDA应用设计与开发》♥♥♥♥