CUDA入门笔记（三）GPU编程基础——一个典型GPU程序

参考：

优达学城：https://classroom.udacity.com/courses/cs344/lessons/55120467/concepts/670611900923

 

一、典型GPU程序构成

一个典型GPU程序有如下几个部分：

①CPU在GPU上分配内存

②CPU将CPU中的数据copy到GPU中

③调用内核函数来处理数据

④CPU将GPU中的数据copy到CPU中

 

*可以看出，四个步骤中有两个是数据的copy，因此如果你的程序需要不断地进行copy，那么运行效率会比较低，不适合利用GPU运算。

一般情况下，最好的方式是，让GPU进行大量运算，同时保证计算量与通信的比值很高。

 

二、运行程序时GPU是如何运作的？

前面提到，我们将运算看作是一个由一系列或多个内核组成的结构。GPU有很多并行的计算单元，那么当我们写kernels的时候，应该充分利用这些硬件并行。

 

★Big Idea（one of the very core concept of CUDA）：

当你在写程序的时候，你把它写的像是一个串行的程序，就好像是在一个单独的线程里运行一样。当你从CPU调用核的时候，就给核规定好要启动几个线程，每一个线程都会运行这个程序（的一部分？，存疑）。

 

小结：

GPU擅长：①启动很多的线程   ②并行地运行很多线程

因此，如果你的程序没有利用它进行多线程计算，那么就没有很好的利用GPU。

 

三、一个典型的程序

程序要求：对一个数组当中的所有元素进行平方计算。

 

先看一下C语言版本的程序：

for (int i = 0; i < 64; i++){
    out[i] = in[i] * in[i];
}

C语言的程序有以下两个特点：

①只有一个线程在执行计算，该线程遍历输入数组内所有元素。（将该线程定义为“执行代码的一个独立路径”，这一定义也将适用于GPU代码中。）

②代码中没有显式并行性，是串行代码。这段代码是一个循环64次的线程，每一个循环进行一次平方计算。

 

对于GPU来说，程序要分为两个部分：①在CPU上运行   ②在GPU上运行的

图3.1 GPU程序构成

GPU上运行的程序很简单，可以看作是一个普通的串行运算，在这个例子里就是 OUT = IN * IN，即计算平方。那么这个核函数就可以看作跟“并行”无关。

CPU上的程序则要负责分配内存、从GPU复制数值、将数值传递给GPU、启动Kernel函数。

*CPU启动kernels的过程：用64个线程，启动一个执行平方计算的kernel，即square_kernel。每一个内核的实例（instance）将运算一个平方计算。

*线程索引：分辨每一个独立的内核做的工作。

CPU code：

square_kernel<<<64>>>(input_array, output_array)

完整的代码如下：

#include <stdio.h>

//内核函数，就跟一个串行程序一样
//__global__ : 定义说明符，用来表明这是一个内核，而非普通函数
//kernel的参数都必须是分配在GPU上的，不然会程序崩溃
__global__ void square(float * d_out, float * d_in){
    int idx = threadIdx.x; //用来得到线程的索引值,threadIdx是一个CUDA内置变量，其本质是一个C结构，有三个成员，x、y、z。
    float f = d_in[idx];
    d_out[idx] = f * f;
}

//下面是CPU代码
int main(int argc, char ** argv) {
    const int ARRAY_SIZE = 64;
    const int ARRAY_BYTES = ARRAY_SIZE * sizeof(float);

//generate 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,看起来跟c program一样，是一个普通的数组
    float * d_in;
    float * d_out;

//allocate GPU memory，在这里给它分配GPU内存
    cudaMalloc((void **) &d_in, ARRAY_BYTES);
    cudaMalloc((void **) &d_out, ARRAY_BYTES);

//transfer the array to the GPU，前三个参数跟C Memcpy一样，目标地址、源地址、字节数，最后一个参数表明转移方向cudaMemcpyHostToDevice, cudaMemcpyDeviceToHost,cudaMemcpyDeviceToDevice
    cudaMencpy(d_in, h_in, ARRAY_BYTES,cudaMemcpyHostToDevice);

//launch the kernel, cuda启动运算符，在一个具有64个元素的block上启动一个核，两个参数d_in,d_out
//这句话告诉CPU，在GPU上启动64个线程，运行64个内核副本
//注意，我们只能在调用内核处理GPU数据，而不能处理CPU数据
    square<<<1,ARRAY_SIZE>>>(d_out, d_in);

//copy back the result array to the CPU，完成内核运行以后，结果在d_out中
    cudaMemcpy(h_out, d_out, ATTAY_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;
}

 

需要注意的是，在CUDA程序中，我们以h_开头表示host上的数据，d_开头表示GPU上的数据。

针对代码，有值得注意的几点：

①square<<<1,64>>>(d_out, d_in)

参数说明： 

1：number of blocks

64：Threads Per Block

从硬件角度说，①硬件可以同时运行多个块（blocks），②每个块有线程数限制，较新的GPU一个块可以支持1024个线程，旧的512.

 

②另外，这是一个一维diagram，只在一个维度发展。如果想要处理2/3等维度问题，则需要如下代码：

Kernel<<<GRID OF BLOCKS, BLOCK OF THREADS，shemem>>>(...)

说明： 

GRID OF BLOCKS：线程块网格维数，取值1-3

BLOCK OF THREADS：每个块中线程维数，取值1-3

dim3(x,y,z)：线程/blocks维数 = x*y*z，xyz分别代表三个方向的值

shemem：每一个块共享的内存，字节数

 

③线程索引

threadIdx：线程索引，有threadIdx.x，threadIdx.y，threadIdx.z来找三个方向的索引。

blockIdx:块索引，同上。

*CUDA擅长处理多块、多线程、多维度问题。