计算模型

        如果计算模型中使用了多块同类型架构的芯片进行运算，那么这种计算模型就叫做同构计算模型，比如使用了多块一样的CPU进行矩阵运算；如果计算模型中使用了多块不同类型架构的芯片进行协同运算，那么这种计算模型就叫做异构计算模型，比如用一块CPU加一块GPU进行矩阵运算。
        实际上GPU并行计算模型中，光靠GPU是无法完成所有工作的，数据的初始化和最终结果的处理都需要CPU去完成，所以GPU并行计算模型是一种异构计算模型，需要CPU和GPU协同工作。
        异构计算技术从80年代中期产生，由于它能经济有效地获取高性能计算能力、可扩展性好、计算资源利用率高、发展潜力巨大，已成为当今信息时代并行计算，分布计算领域中的研究热点。

执行过程

        在CUDA中，host和device是两个重要的概念，我们用host指代CPU及其内存，而用device指代GPU及其内存。CUDA程序中既包含host程序，又包含device程序，它们分别在CPU和GPU上运行。同时，由于两者之间存在数据通路，host与device可以进行通信，实现CPU与GPU数据拷贝。典型的CUDA程序的执行流程如下：

  1. 分配host内存，CPU执行指令进行数据的初始化
  2. 分配device内存，并将host数据拷贝到device
  3. 调用CUDA的核函数在device上完成指定的运算
  4. 将device上的最终运算结果拷贝到host上处理
  5. 任务全部结束后释放device和host上分配的内存

函数限定

        由于GP计算模型U实际上是异构模型，所以需要区分host和device上的代码，在CUDA中是通过函数类型限定词开区别host和device上的函数，主要的三个函数类型限定词如下：

  • global：从host中调用，在device上执行，返回类型必须是void。
  • device：从device中调用，在device上执行，不可在host中使用。
  • host：从host上调用，在host上执行，一般程序中可以省略不写。

参考如下代码：

#include<stdio.h>
 //CPU中调用，GPU中执行
__global__ void helloFromGPU(){
  printf("Hello World from GPU！\n");
}
//CPU中调用，CPU中执行
__host__ int main(void){
  printf("Hello World from CPU!\n");
  //调用GPU执行
  helloFromGPU<<<1,2>>>();
  cudaDeviceReset();
  return 0;
}

线程结构

        用__global__符号声明的函数为核函数(kernel)，在调用时需要用<<<grid, block>>>来指定kernel要执行的线程数量。在CUDA中，每一个线程都要执行核函数，并且每个线程会分配一个唯一的线程ID。在上述HelloWorld程序中，__global__ void helloFromGPU()为一个和函数，CPU在执行主函数__host__ int main(void)中语句helloFromGPU<<<1,2>>>();时便启用了GPU执行__global__ void helloFromGPU()函数，此时在终端就可以看到GPU执行打印的输出：Hello World from GPU！。
        要深刻理解kernel(方便交流，简称核)，必须要对kernel的线程层次结构有一个清晰的认识。首先GPU上很多并行化的轻量级线程。kernel在device上执行时实际上是启动很多线程，一个kernel所启动的所有线程称为一个网格，简称格(grid)，grid是线程结构的第一层次；网格又可以分为很多线程块，简称块(block)，一个线程块里面包含很多线程，这是第二个层次。线程英文为thread，简称线，简称只是为了方便交流而用，原始概念还是参考英文原版资料。这样就形成了格–>块–>线三级逻辑结构，后面了解硬件结构后就会明白为什么需要这样的概念了。

层级结构

        线程结构主要分为两个层次结构，第一个层次结构为格到块的映射，第二个层次结构为块到线的映射。如果单独分析一个层次结构，常规思维是直接线性化处理，比如一个格里面有6个块，那么常规思维可能这6个块是线性排列的，但实际上我们还可以进行二维排列，如下图的2行3列矩阵排列：

gpu的成像逻辑 gpu逻辑计算_CUDA


        为了表示这种结构，我们用一种叫dim3的数据结构表示，这种数据结构，可以看成是包含三个无符号整数（x，y，z）成员的结构体变量，在定义时，缺省值初始化为1。如果层次结构采用一维结构，即线性结构，并且长度为6，那么dim3为（6，1，1）。如果采用了上述的二维结构，即矩阵结构，那么dim3为（3，2，1）。由此很容易联想到如果采用三维结构，即立体结构，那么对应的z值大于1即可。

        如果格到块之间映射为二维，块到线之间的映射为二维，那么可能的排列如下图：

gpu的成像逻辑 gpu逻辑计算_CUDA_02


        所以，一个线程需要两个内置的坐标变量（blockIdx，threadIdx）来唯一标识，它们都是dim3类型变量，其中blockIdx指明线程所在grid中的位置，而threaIdx指明线程所在block中的位置，如图中的Thread (1,1)满足：

threadIdx.x = 1
threadIdx.y = 1
blockIdx.x = 1
blockIdx.y = 1