并行编程实战——CUDA编程的整体流程

一、CUDA的异步并行编程

我们在前面学习了CUDA编程的基本知识和环境的安装，对CUDA编程已经有了一个较为全面的基础的认知。不过大家一定要明白，即使CUDA程序，其主要的控制应用仍然是在CPU（HOST主机端）中运行，也只有遇到大规模的并行运算等任务时，才会将其转移到GPU端（DEVICE设备端）进行工作。CUDA右异构并行编程中起到了重要的作用，所以CUDA编程被称为异构并行编程 (heterogeneous parallel programming)，即共同使用主机和设备进行编程的一个重要原因就因为如此。

三、CUDA的代码编译处理

在普通的C++编译器中，是无法识别.cu后缀的CUDA程序代码的，所以在CUDA中必须使用nvcc作为编译器来处理cu文件，nvcc本质将cu文件的代码进行区分（即__gloabl__等关键字），主要是主机端和设备端的代码，即GPU的代码交由ptxas编译器编译而CPU的代码交由C++编译器如如gcc，msvc等编译，最后再将二者链接成为可执行文件或库。
和C++编译一样，可以通过编译选项–keep来保留中间编译单元用于分析。
其编译过程如下：
1、对cu文件进行预处理
2、分离主机端和设备端的代码
3、编译设备端代码
4、编译主机端代码
5、链接
详细的编译过程这里不再展开，有兴趣的可以查看一下相关的资料书籍，这个有点略微超纲了。

三、CUDA代码的运行

CUDA代码是从HOST端开始执行的，然后调用Kernel函数来启动向GPU操作的任务。在这个过程中，CUDA会调用相关的API（如cudaMalloc等）对其Runtime进行初始化。初始化的内容数据结构和相关的资源（如内存）等的分配。每个CUDA设备都会有一个上下文（与OS中的上下文类似），它包括了很多方面的描述如内存等。如果开发者拥有多个设备端则可以通过API cudaSetDevice来指定哪个设备为默认设备（家里的机器就算了，一般就一个显卡即0设备）。
当然，也可以通过设置环境变量来对GPU进行屏蔽和选择，这个在后期会进行展开说明。
CUDA的线程层次在前面进行了分析，但在实际的运行中，一定要明白前面几个关键字的相关细节和关系。特别在对kernel要明白其运行的逻辑流程。首先CUDA程序的启动也是一个main函数启动的，要想启动一个由GPU工作的任务，也就是将并行计算的代码从CPU加载到GPU进行工作，就是通过启动一个kernel函数来实现。核函数在设备端执行时会启动很多的线程，它们统称为Grid（网格），同一个Grid中的所有线程共享一个全局内存空间，网络下又为分块（块簇，这个前面分析过），每个块中又有N个线程。
和CPU类似，物理的GPU中其实包含多个硬件处理核心SM即Streaming Multiprocessor，中文叫做流式多处理器。它包括CUDA核心、共享内存以及一系列的寄存器等。SM用来并发的处理几百个线程，也就是说一个GPU中的SM越多则并行处理的能力越大。
当核函数执行时，Grid中的线程块被分配到SM，一个线程块只能在一个SM上调度，而一个SM则可以调度多个线程块。一个Kernel的多个块被分配到多个SM中。可以这样理解，SM是执行的物理层而Grid是执行的逻辑层。SM使用的是SIMD架构，single instruction multiple data，这个大家应该明白，不明白的就得回头复习一下计算体系的知识了。
SM使用的线程束（Wraps）为执行的基本单元，每个束包括32个线程。明白了这些就知道了CUDA的代码是如何进行正确的调度和运行了。
在此之后（注意是异步，所以立即返回），CUDA程序继续返回主机端进行运行并串行执行主机端代码（这里指在没有其它线程的情况下），直到另外一个核函数被启动。由main函数不会等待核函数的返回，所以如果需要处理核函数的结果，就需要明确的在main函数中进行一些同步处理，这个有过多线程编程经验的开发者应该非常清楚。
所以，自己的CUDA代码到最终能够执行够不够优秀，就要看一看硬件的执行能力，特别是GPU中SM的数量，从理论上讲，这个东西越多，执行的效率就越高。可以通过下面的代码进行查看自己的GPU中的SM的数量：

int main()
{
	CPUFunction();
	GPUFunction << <1, 1 >> > ();
  //sm count
	cudadevProp gpuProp;
	cudaGetdevProperties(&gpuProp, 0);
	int smCount = gpuProp.multiProcessorCount;
	std::cout << "sm count:"<< smCount << std::endl;
	cudaDeviceSynchronize();
	system<<"pause");
	return 0;
}

在本文的机器中，GTX960有8个SM。

四、CUDA运行环境的代码

针对上面的分析，大家可以用下面的代码来查看一下自己电脑的GPU的能力：

#include <cuda_runtime_api.h>
#include <iostream>

int main()
{
  // set first
  int dev = 0;
  cudaSetDevice(dev);
  // get GPU info
  cudaDeviceProp devProp;
  cudaGetDeviceProperties(&devProp, dev);
  int dVer = 0, rVer = 0;
  cudaDriverGetVersion(&dVer);
  cudaRuntimeGetVersion(&rVer);

  std::cout << " Device id:" << dev << ",name:"<< devProp.name << std::endl;
  std::cout << " CUDA Driver Version:" << dVer << ", Runtime Version :" << rVer << std::endl;
  std::cout << " global mem:"<<devProp.totalGlobalMem<<std::endl;
  std::cout << " Memory Bus Width: :"<< devProp.memoryBusWidth<<std::endl;

  if (devProp.l2CacheSize) {
    std::cout << " L2 Cache Size:"<<devProp.l2CacheSize<<std::endl;
  }

  std::cout << " Max Texture Dimension Size (x,y,z) 1D="<< devProp.maxTexture1D<<",2D="<<  devProp.maxTexture2D[0]<<","<<
    devProp.maxTexture2D[1]<<std::endl;


  std::cout << " const mem:"<< devProp.totalConstMem<<std::endl;
  std::cout << "shared mem block:"<< devProp.sharedMemPerBlock<<std::endl;
  std::cout << " registor block:"<<devProp.regsPerBlock<<std::endl;
  std::cout << " Warp size: "<<devProp.warpSize<<std::endl;
  std::cout << " processor max thread : "<< devProp.maxThreadsPerMultiProcessor<<std::endl;
  std::cout << " block max thread:"<< devProp.maxThreadsPerBlock<<std::endl;
  std::cout << " thread dim: "<< devProp.maxThreadsDim[0]<<","<< devProp.maxThreadsDim[1]<<","<< devProp.maxThreadsDim[2]<<std::endl;
  std::cout << " grid size:"<< devProp.maxGridSize[0]<<","<< devProp.maxGridSize[1]<<","<< devProp.maxGridSize[2]<<std::endl;
  std::cout << " Maximum memory pitch:"<< devProp.memPitch<<std::endl;

  return 0;
}

运行结果是：

Device id:0,name:GeForce GTX 960
 CUDA Driver Version:9010, Runtime Version :9010
 global mem:4294967296
 Memory Bus Width: :128
 L2 Cache Size:1048576
 Max Texture Dimension Size (x,y,z) 1D=65536,2D=65536,65536
 const mem:65536
shared mem block:49152
 registor block:65536
 Warp size: 32
 processor max thread : 2048
 block max thread:1024
 thread dim: 1024,1024,64
 grid size:2147483647,65535,65535
 Maximum memory pitch:2147483647

六、总结

CUDA开发的一个整体的概念性的理解，到现在应该相对比较完整了。当然，如果想更好的掌握CUDA的开发，更深入的硬件和软件以及二者的关系，往往会起到更好的效果，但这是后面每个人自己发力的结果了。
所以大家对CUDA的理解不能简单认为是一个软件框架，而是一个软硬件结合的一个框架，它与具体的GPU型号和数量有着密切的关系。