GPU编程 | 并行计算的helloworld！_gpu运行并行程序

在上一篇文章《你了解CUDA，了解GPU吗？》中，我们介绍了什么是CUDA，什么是GPU。那么接下来我们就要来讲解下GPU到底是如何实现并行计算的，以及CUDA是如何通过编程调用GPU的并行计算力的。

1.GPU是如何实现并行计算

从软件层面来说，GPU在进行并行计算时，是以核（kernel）为单位进行的。如下图所示，每个核相当于一个功能函数。每个核由若干线程块（thread block）负责运算，而每个线程块又由若干个线程组成。所有的线程由CPU负责发放给GPU设备。

从硬件层面来说，GPU包含若干个流处理器（SM）（越多GPU性能越好），每个流处理器由若干个处理单元（粉色）和一个存储单元（黄色）组成，每个流处理器独立并行工作。用以计算CPU发送过来的线程。

当我们执行并行计算时，GPU会将核分配给流处理器进行并行计算。该过程中存在无序和有序两个部分，无序部分保证了程序的并行能力，有序部分保证了程序的逻辑能力。其中无序部分为线程与线程之间无序，有序部分为核与核之间有序。

• 无序：
  1 当对一个核进行处理时，用于处理当前核的每个线程块会随机分配给任意一个流处理器中的任意处理单元进行计算。
  2 每个线程的开始结束时间是随机的，不等待原则。需要注意是，由于线程的互相不等待，往往会导致程序出现问题，因此在CUDA编程中会通过设置屏障来解决这个问题，本节不展开讨论。

• 有序：
  1 每个线程块中的线程必定是在同一时间内运行在同一个流处理器中的 。
  2 两个有先后逻辑关系的核，则后一个核要等前一个核完成后才执行。

2.利用CUDA编程操作GPU

在上一篇文章中我们也有介绍，要想用CUDA操纵GPU，首先要以CPU作为宿主，GPU作为一个能够进行并行计算的设备，最终实现并行计算。
因此利用CUDA操作GPU的整体流程如下，对应的示意图如图：

• CPU分配空间给GPU（cudaMalloc）
• CPU复制数据给GPU（cudaMemcpy）
• CPU加载kernels给GPU作计算
• CPU把GPU的运算结果复制回来

上图中不同内存间的访问速度如下图所示，因此要想加快计算速度，我们的核心思想就是将整体计算切分，使得每次的计算规模越小，计算次数越多，那么程序的速度也就越快。这是从内存访问的角度说明了GPU并行计算的优势。

接下来介绍下CUDA编程的HelloWorld！——实现对输入数组求平方的操作。
其实利用CUDA编程操纵GPU和编写C语言程序没有太大的区别，我们只需要严格按照上述流程就可以写出这个HelloWorld.cu了。

• 定义核函数

#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;
}
1
2
3
4
5
6
7
8

• 主函数，流程跟上述类似（h表示host，d表示device）

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

    // 在cpu上产生输入数据
    float h_in[ARRAY_SIZE];
    for(int i=0; i<ARRAY_SIZE; i++)
    {
      h_in[i] = float[i];
    }
    float h_out[ARRAY_SIZE];

    // 返回值定义
    float* d_in;
    float* d_out;

    // 分配GPU内存
    cudaMalloc(void**) &d_in,ARRAY_BYTES);
    cudaMalloc(void**) &d_out,ARRAY_BYTES);

    // 将输入数据传到GPU
    cudaMemcpy(d_in, h_in, ARRAY_BYTES, cudaMemcpyHostToDevice);

    // 在GPU上运行kernel
    square<<<1,ARRAY_SIZE>>>(d_out, d_in);
    //将GPU运算结果传回cpu
    cudaMemcpy(h_out, d_out, ARRAY_BYTES, cudaMemcpyDeviceToHost);

    // 利用cpu将计算结果打印
    for(int i=0; i<ARRAY_SIZE; i++)
    {
        printf("%f",h_out[i]);
        printf(((i%4) != 3) ? "\t":"\n");
    }

    // 释放GPU上的内存空间
    cudaFree(d_in);
    cudaFree(d_out);

    return 0;
}
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42

• 用NVCC编译器编译.cu程序生成可执行文件，并执行
  
  [外链图片转存失败(img-BvQXIjlS-1562393455070)(https://upload-images.jianshu.io/upload_images/5529997-86f12e858bfd54bf.png?imageMogr2/auto-orient/strip|imageView2/2/w/1240)]