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

旧城等待， 2023-07-25 09:57 0阅读 0赞

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

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

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

下面这个图是 CUDA 编程中常见到的软硬件的结构图。  
![在这里插入图片描述][watermark_type_ZmFuZ3poZW5naGVpdGk_shadow_10_text_aHR0cHM6Ly9ibG9nLmNzZG4ubmV0L3dhbmdib3dqMTIz_size_16_color_FFFFFF_t_70]  
**Kernel核**: 可以理解为C/C++中的一个函数function。不过这样的理解其实不够准确，个人认为更准确的理解方式是将其视为一次对在 device 上运行函数的调用，每次调用 kernel 核，都需要指定一些参数，参数的形式也有很多种。下图中绿色的框便是一个 kernel 核（只有一个线程块）。**一个 kernel 核可以具备很多个线程块**。

![在这里插入图片描述][watermark_type_ZmFuZ3poZW5naGVpdGk_shadow_10_text_aHR0cHM6Ly9ibG9nLmNzZG4ubmV0L3dhbmdib3dqMTIz_size_16_color_FFFFFF_t_70_pic_center]  
所以说，Kernel 对应于需要在GPU上执行的程序，并且一个Kernel对应一个**Grid**。  
![在这里插入图片描述][watermark_type_ZmFuZ3poZW5naGVpdGk_shadow_10_text_aHR0cHM6Ly9ibG9nLmNzZG4ubmV0L3dhbmdib3dqMTIz_size_16_color_FFFFFF_t_70 1]

**SM（stream multiprocessor）: 流处理器**  
![在这里插入图片描述][20200414160536418.png]  
**GPU**:每个GPU有若干个SM，最少有1个，每个SM并行而独立运行  
![在这里插入图片描述][watermark_type_ZmFuZ3poZW5naGVpdGk_shadow_10_text_aHR0cHM6Ly9ibG9nLmNzZG4ubmV0L3dhbmdib3dqMTIz_size_16_color_FFFFFF_t_70 2]  
从上面的图可以看出：**CUDA最大的特点：对线程块将在何处、何时运行不作保证。**

**优点：**

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

**后果：**

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

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

1.  所有在同一个线程块上的线程必然会在同一时间运行在同一个SM上
2.  同一个内核的所有线程块必然会全部完成了后，才会运行下一个内核  
    如下图所示：  
    ![在这里插入图片描述][watermark_type_ZmFuZ3poZW5naGVpdGk_shadow_10_text_aHR0cHM6Ly9ibG9nLmNzZG4ubmV0L3dhbmdib3dqMTIz_size_16_color_FFFFFF_t_70 3]

下图是 GPU 工作的**内存模型**。GPU 硬件架构具有三层的存储结构，分别是：**线程私有存储、线程块共享存储、全局存储**。访问速度自然是由快到慢。CUDA 程序的编写的一个重要理念也是尽可能少的访问全局缓存。  
![在这里插入图片描述][watermark_type_ZmFuZ3poZW5naGVpdGk_shadow_10_text_aHR0cHM6Ly9ibG9nLmNzZG4ubmV0L3dhbmdib3dqMTIz_size_16_color_FFFFFF_t_70 4]  
**同步性synchronisation和屏障barrier：不同的线程在共享和全局内存中读写数据需要有先后的控制，所以引入了同步性**的概念。  
\*\*屏障的作用：\*\*用来控制多个线程的停止与等待，当所有线程都到达了屏障点，程序才继续进行。  
![在这里插入图片描述][watermark_type_ZmFuZ3poZW5naGVpdGk_shadow_10_text_aHR0cHM6Ly9ibG9nLmNzZG4ubmV0L3dhbmdib3dqMTIz_size_16_color_FFFFFF_t_70 5]  
**CUDA程序中CPU是主导地位，负责完成以下的事情：**

1.  从CPU同步数据到GPU
2.  从GPU同步数据到CPU（1、2使用cudaMemcpy）
3.  给GPU分配内存（cudaMalloc）
4.  加载Kernel到GPU上，launch kernel on GPU  
    ![在这里插入图片描述][watermark_type_ZmFuZ3poZW5naGVpdGk_shadow_10_text_aHR0cHM6Ly9ibG9nLmNzZG4ubmV0L3dhbmdib3dqMTIz_size_16_color_FFFFFF_t_70 6]

# 第一行 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个进程并行地工作求出答案。

[watermark_type_ZmFuZ3poZW5naGVpdGk_shadow_10_text_aHR0cHM6Ly9ibG9nLmNzZG4ubmV0L3dhbmdib3dqMTIz_size_16_color_FFFFFF_t_70]: /images/20230528/6bb4c37eefe24f81932d85384f60688a.png
[watermark_type_ZmFuZ3poZW5naGVpdGk_shadow_10_text_aHR0cHM6Ly9ibG9nLmNzZG4ubmV0L3dhbmdib3dqMTIz_size_16_color_FFFFFF_t_70_pic_center]: /images/20230528/9b7a05c08747457fb01a1a49dda0223a.png
[watermark_type_ZmFuZ3poZW5naGVpdGk_shadow_10_text_aHR0cHM6Ly9ibG9nLmNzZG4ubmV0L3dhbmdib3dqMTIz_size_16_color_FFFFFF_t_70 1]: /images/20230528/b3754251beb1447f9328cd16ba8098c3.png
[20200414160536418.png]: /images/20230528/c9ddf101bd584aaab55e2c90d5f428b7.png
[watermark_type_ZmFuZ3poZW5naGVpdGk_shadow_10_text_aHR0cHM6Ly9ibG9nLmNzZG4ubmV0L3dhbmdib3dqMTIz_size_16_color_FFFFFF_t_70 2]: /images/20230528/6e5ed5fa2e2a4a53bb0c3bafae519af6.png
[watermark_type_ZmFuZ3poZW5naGVpdGk_shadow_10_text_aHR0cHM6Ly9ibG9nLmNzZG4ubmV0L3dhbmdib3dqMTIz_size_16_color_FFFFFF_t_70 3]: /images/20230528/eeb201484e1443d187917c981b642c28.png
[watermark_type_ZmFuZ3poZW5naGVpdGk_shadow_10_text_aHR0cHM6Ly9ibG9nLmNzZG4ubmV0L3dhbmdib3dqMTIz_size_16_color_FFFFFF_t_70 4]: /images/20230528/dd41117bd43e4505aa91d0973b86201a.png
[watermark_type_ZmFuZ3poZW5naGVpdGk_shadow_10_text_aHR0cHM6Ly9ibG9nLmNzZG4ubmV0L3dhbmdib3dqMTIz_size_16_color_FFFFFF_t_70 5]: /images/20230528/690ddb3f0261402d9487e050db738a18.png
[watermark_type_ZmFuZ3poZW5naGVpdGk_shadow_10_text_aHR0cHM6Ly9ibG9nLmNzZG4ubmV0L3dhbmdib3dqMTIz_size_16_color_FFFFFF_t_70 6]: /images/20230528/9f7dc02b32684c04ac9a1637238b6621.png