CS344 Intro To Parallel Programming Using CUDA lesson1

我们说,通常有三种方法可以让电脑运行的更快:

  1. 更快的时钟。

  2. 每个时钟周期做更多的工作

  3. 更多的处理器

在本次课中,我们更注重第3条,也就是如何使用更多的处理器来加快我们程序的运行速度。现代的GPU有着几千个ALU,几百个处理器和数以万计的并行单元,是我们本次课的重点与核心。

这些年来,处理器的速度越来越快,但是由于功率与热量的限制,时钟频率陷入了瓶颈,处理器速度的加快更多的是由于晶体管体积的缩小。所以目前设计处理器的一个难题就是功率问题。

在日常生活中,我们会接触到两种处理器,CPU与GPU。它们的特点分别是:

  • CPU: 1.编程灵活性好,单个处理器处理效率高(pros); 2.受功率限制(cons)。

  • GPU: 1.有更多针对计算的硬件(pros);2.有着更好的能源利用率(ops/watt)(pros),3.编程模型受限(cons)

衡量一个一个处理器的指标有两种,分别是latencythroughput。举例来说:A点和B点距离为4500公里,一个小汽车一次能带两个人,时速为200公里/小时;一辆公交车一次能载40个人,时速为50公里/小时,那我们可以说,小汽车的latency为4500/200 = 22.5小时,公交车的latency为4500/50=90小时,我们可以看到,小汽车的latency大于公交车。然而小汽车的throughput为2/22.5 = 0.089 人/小时,公交车的throughput为40/90 = 0.45 人/小时。并行算法优化的指标为throughput,也就是说,虽然公交车的速度慢于小汽车,但是由于一次载的人多(并行单元多),因此公交车throughput的指标优于小汽车。同时并行算法优化的指标为throughput

GPU设计的原则主要如下:

  1. 有着大量简单的计算单元,用简单的控制换取更多的计算资源。
  2. 显式并行计算模型。
  3. 优化throughput而不是latency

在GPU编程时,通常有以下几种操作:

  1. 数据从CPU拷贝到GPU。

  2. 数据从GPU拷贝到CPU。

  3. 分配GPU的内存。

  4. 启动GPU上的内核。

一个典型的CUDA编程的模式如下:首先CPU分配GPU上的内存(cudaMalloc),接着CPU将输入数据拷贝到GPU(cudaMemcpy),然后CPU启动GPU上的内核开始处理数据,最后将GPU处理好的数据拷贝回CPU。

CUDA编程操作 首先什么是内核?内核程序看起来像是串行程序,并且知道自己所处线程的索引。GPU会用多个线程来运行内核。

举例来说,对于一个数组[0, 1, 2, … , 63],我们希望通过处理器计算得到输出数组为[0^2, 1^2, 2^2, … , 63^2]。

我们可以很轻易地利用下面的代码得到结果:

// CPU code : square each element of an array
for (i = 0; i < 64; i++){
    out[i] = in[i] * in[i]
}

我们可以看到,上面的代码只使用了一个线程来计算,没有显式的并行过程。假如一个乘法费时2ns,则上述代码的执行过程大约耗时2 * 64 = 128ns。

然而我们可以将该程序分为64个线程来计算每个元素的平方,部分代码如下所示:

// 内核程序
__global__ void square(float * d_out, float * d_in){
    // 读取当前的线程索引
    int idx = threadIdx.x;
    float f = d_in[idx];
    d_out[idx] = f * f;
}

我们可以看到,内核程序知道自己所在线程的索引值,便可以执行相应线程的任务(平方数组对应的元素)。假如一个乘法耗时10ns,则上述代码的耗时为10ns * 1 = 10ns。完整的过程如下所示:

#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 = 64;
    const int ARRAY_BYTES = ARRAY_SIZE * sizeof(float);
    float h_in[ARRAY_SIZE];
    
    // 在CPU上初始化数组
    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_SIZE);
    cudaMalloc((void **) &d_out, ARRAY_SIZE);
    
    // 数据从CPU拷贝到GPU
    cudaMemcpy(d_in, h_in, ARRAY_BYTES, cudaMemcpyHostToDevice);
    // 启动GPU上的内核
    square<<<1, ARRAY_SIZE>>>(d_out, d_in);
    // 数据从GPU拷贝到CPU
    cudaMemcpy(h_out, d_out, ARRAY_BYTES, cudaMemcpyDeviceToHost);
    
    for (int i = 0; i < ARRAY_SIZE, i++){
        printf("%f", h_out[i]);
        printf(((i % 4) != 3) ? "\t" : "\n");
    }
    
    // 释放显存
    cudaFree(d_in);
    cudaFree(d_out);
    
    return 0;
}

其中上述代码涉及:分配GPU的内存;启动GPU上的内核;数据从CPU拷贝到GPU;数据从GPU拷贝到CPU等过程。GPU的内核启动代码:square«<# of blocks, # of threads per blocks»>(function_param), # of blocks 为block 数量,# of threads per blocks 为每个block所包含的线程数,每个block,每个线程都是独立运行的,例如square«<10, 128»>(function_param)有着1280个线程。对于一般的GPU,# of threads per blocks一般最大为512,现代的GPU一般最大为1024

举例来说,对于一张128*128的图像,我们可以用以下两种处理方式:

  1. 使用128个block,每个block里有128个线程。
  2. 使用64个block,每个block里有16*16个线程。

对于内核配置,我们有着更广义的表达方式:kernel«<grid of blocks, block of threads»>(…)

如上图我们可以看到,grid of blocks,和 block of threads都可以通过一个三维坐标来索引。举例来说:

kernel«<dim3(8,4,2), dim3(16,16)»>有着8 * 4 * 2个block,每个block有着16 * 16个线程,总计 8 * 4 * 2 * 16 * 16个线程。

其实对于每个数组的元素进行平方的操作,在并行计算里称为MAP,MAP的主要特点为:

  1. 处理一个集合的元素。
  2. 函数运行在每个函数上。

如下图所示: