并行计算之路<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代码。

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

架构

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 8uusing 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 0Hello from thread 1Hello from thread 2Hello from thread 3Hello from thread 4Hello from thread 5Hello from thread 6Hello 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应用设计与开发》♥♥♥♥