第一行 CUDA 代码之GPU架构理解与代码编写

首先是对 CUDA 编程中可能会遇到的各种概念进行简单总结。如对 Kernel、Grid、Device、Host、Thread、Thread
Block、SM 等部件进行梳理，并且牵涉到CUDA编程的基本理念与基本方法。

CUDA 编程之软硬件结构的相关概念

首先是对 CUDA 编程中可能会遇到的各种概念进行简单总结。

下面这个图是 CUDA 编程中常见到的软硬件的结构图。

Kernel核: 可以理解为C/C++中的一个函数function。不过这样的理解其实不够准确，个人认为更准确的理解方式是将其视为一次对在 device 上运行函数的调用，每次调用 kernel 核，都需要指定一些参数，参数的形式也有很多种。下图中绿色的框便是一个 kernel 核（只有一个线程块）。一个 kernel 核可以具备很多个线程块。

所以说，Kernel 对应于需要在GPU上执行的程序，并且一个Kernel对应一个Grid。

SM（stream multiprocessor）: 流处理器

GPU:每个GPU有若干个SM，最少有1个，每个SM并行而独立运行

从上面的图可以看出：CUDA最大的特点：对线程块将在何处、何时运行不作保证。

优点：

1. 硬件真正有效的运行，灵活
2. 无需要线程间互相等待
3. 可扩展性强

后果：

1. 对于那个块在那个SM上运行无法进行任何假设
2. 无法获取块之间的明确通讯（hard to get communications between blocks）
3. dead lock（并行死锁）
4. 线程退出

不过也有一些确定的概念：

1. 所有在同一个线程块上的线程必然会在同一时间运行在同一个SM上
2. 同一个内核的所有线程块必然会全部完成了后，才会运行下一个内核
   如下图所示：
   

下图是 GPU 工作的内存模型。GPU 硬件架构具有三层的存储结构，分别是：线程私有存储、线程块共享存储、全局存储。访问速度自然是由快到慢。CUDA 程序的编写的一个重要理念也是尽可能少的访问全局缓存。

同步性synchronisation和屏障barrier：不同的线程在共享和全局内存中读写数据需要有先后的控制，所以引入了同步性的概念。
**屏障的作用：**用来控制多个线程的停止与等待，当所有线程都到达了屏障点，程序才继续进行。

CUDA程序中CPU是主导地位，负责完成以下的事情：

1. 从CPU同步数据到GPU
2. 从GPU同步数据到CPU（1、2使用cudaMemcpy）
3. 给GPU分配内存（cudaMalloc）
4. 加载Kernel到GPU上，launch kernel on GPU
   

第一行 CUDA 代码

接下来便是对 CUDA 代码的编写。

GPU程序一般步骤

1. CPU分配空间给GPU（cudaMalloc）
2. CPU复制数据给GPU（cudaMemcpy）
3. CPU加载kernels给GPU做计算
4. CPU把GPU计算结果复制回来

过程中，一般要尽量降低数据通讯的消耗，所以如果程序需要复制大量的数据到GPU，显然不是很合适使用GPU运算，最理想的情况是，每次复制的数据很小，然后运算量很大，输出的结果还是很小，复制回CPU。

第一个 CUDA 程序——并行地求平方

#include <stdio.h>

__global__ void square(float* d_out,float* d_in){
  int idx = threadIdx.x;
  float f = d_in[idx];
  d_out[idx] = f * f;
}

int main(int argc,char** argv){
  const int ARRAY_SIZE = 8;
  const int ARRAY_BYTES = ARRAY_SIZE * sizeof(float);

  // generate the input array on the host
  float h_in[ARRAY_SIZE];
  for(int i=0;i<ARRAY_SIZE;i++){
    h_in[i] = float(i);
  }
  float h_out[ARRAY_SIZE];

  // declare GPU memory pointers
  float* d_in;
  float* d_out;

  // allocate GPU memory
  cudaMalloc((void**) &d_in,ARRAY_BYTES);
  cudaMalloc((void**) &d_out,ARRAY_BYTES);

  // transfer the array to GPU
  cudaMemcpy(d_in,h_in,ARRAY_BYTES,cudaMemcpyHostToDevice);

  // launch the kernel
  square<<<1,ARRAY_SIZE>>>(d_out,d_in);

  // copy back the result array to the GPU
  cudaMemcpy(h_out,d_out,ARRAY_BYTES,cudaMemcpyDeviceToHost);

  // print out the resulting array
  for(int i=0;i<ARRAY_SIZE;i++){
    printf("%f",h_out[i]);
    printf(((i%4) != 3) ? "\t" : "\n");
  }

  // free GPU memory allocation
  cudaFree(d_in);
  cudaFree(d_out);
  return 0;
}

这段代码实现了8个数组并行的求平方。并行在哪里体现出来了呢？自然是在调用 kernel 核的时候，我们给 GPU 指定了1个线程块，8个进程并行地工作求出答案。