Outline of Programming Massively Parallel Processors Introduction Part

写在最前

本页面内容仅对于毫无基础的初学者想要快速了解可供参考,深度学习者仍建议自行阅读Programming Massively Parallel Processors书籍原文.

在本页面的所有内容中,我不会对CUDA代码进行详细介绍,我只会针对整体内容进行讲解,为读者构建基础的印象感知,同时方便温习使用.

以下内容为基于自身理解的梗概,内容难免有错漏,欢迎各位阅读者批评指正!

再次声明,以下仅供参考

异构并行计算与现代CPU/GPU架构

首先,我们需要知道,现代计算机的两个基础部件:CPU与GPU

CPU,即Central Processing Unit,作为现代计算机的基础组成部分,是大部分计算控制的核心.在片上,它拥有极大的控制单元(Control Unit),以及多级缓存结构(Cache),这也是它延迟导向(Latency-oriented design)的设计思路的表现.片上大的控制单元,使其及其擅长流程控制以及分支判断.相较而言,由于大量的缓存和控制占用,使之片上的计算单元(ALU,Arithmetic Logical Unit)占比较少,因而对于大规模计算,CPU通常难以胜任,尤其是现代图形处理、3D渲染、机器学习等领域中,CPU更多扮演控制的角色而非计算,计算的任务则主要交给GPU.

GPU,即Graphics Processing Unit,作为现代计算机的又一重大组成部分,承担起大规模并行计算的重任.其片上大量面积分配给ALU,而缓存与控制部分占比极少,这充分展现了GPU强大的计算能力.其与显存(DRAM)之间的通道数量较多,可以提供更高的内存带宽.以上两点,则是GPU吞吐导向(Throughput-oriented design)设计思路的展现.

在GPU中,所有的计算都是以线程(Thread)为最基础的单位在软件层面上进行的,每个线程都会负责一部分计算,这一点在下一部分会讲到.

尽管当下,多核CPU已经存在,同样提高了计算机的性能,但是,对于大规模的计算而言,CPU仅仅支持串行计算,其少量的ALU限制了它的及计算运用.对于GPU,大量的ALU为并行计算提供了可能.英伟达CUDA为GPU编程提供了高效快捷的工具,在如今人工智能高速发展、图形渲染要求更多的情境下,GPU异构并行计算成为处理大规模计算的主流方式.

下图为CPU/GPU的结构示意图

数据并行与GPU编程基础概念

请看如下图片灰度处理的例子

对于每一个图片像素,实际上都是由RGB(红绿蓝)三个值来决定的.对于灰度计算,我们有如下公式:

在实际应用中,我们总是会把二维结构一维线性化,以一维数组的形式储存.实际上,在CUDA编程和GPU计算中,这是很常见也是基本的操作.

对于输入数组,我们对每个像素用上述公式计算得到对应的输出,如下

对于串行计算,并行计算可以在几乎同一时间来完成计算,因为一位数组天然有可并行性.这是一种数据并行化的方法,就是转换数据的存储形式.

数据并行 vs 任务并行
数据并行并不是并行编程中唯一使用的并行类型。任务并行也被广泛应用于并行编程。任务并行通常通过应用程序的任务分解来体现。例如，一个简单的应用程序可能需要进行向量加法和矩阵-向量乘法。这两者各自就是一个任务。也就是说,如果这两个任务可以独立完成，那么就存在任务并行。在大型应用中，通常会有更多的独立任务，因此任务并行也更多。一般来说,我们主要通过数据并行来提升并行程序性能.

在GPU编程中,我们在软件层面将GPU分为grid-block-thread三个层次,每一个grid/block皆有xyz三个维度,因而形成了一个三位的结构.thread不存在方向和维度,因为thread已是基本的执行单元.

每个gird方向上支持的最大block数量,以及每个block方向上支持的最大thread数量,可以通过相关CUDA代码进行查询(由于设备不同,其返回数值有所不同).

下图可以看到形象的展示图

需要注意到,对于二维结构,通常用x方向代表水平轴方向,y方向代表垂直轴方向.三维结构的z轴与我们的认知基本一致.

正如上图图片灰度计算的例子相似,我们对于多维结构,一般会采用一维数组的方式存储.在CUDA中,一般采用行为主(Row-major)的展示形式

对应有如下的映射法则:

其中,是矩阵行元素数量(即宽度,Width)

本部分书中讲到的矩阵乘法,图形模糊处理,在后续Parallel Concepts部分会有讲解,故不在这一部分讲述.

GPU计算架构和调度

在现代GPU编程中,除了如上的grid-block-thread结构,实际上,在硬件上,存在SM(Streaming Multiprocessor,流式多处理器)和warp(线程束)的概念.

在实际计算运行中,通常是每32个threads构成一个warp,这是GPU执行计算的最小调度单位.其中的32个线程是按索引连续的.

对于SM,通常包含了多个block(若干block),由于通常一个block包含了多个warp,所以也可以认为SM是多个warp的集合.SM中有控制单元和内存单元,所以负责基础的控制分流等操作.

通常,一个GPU上的SM是有限的,且支持的线程也是有限的,并不与之前代码中可以查看的数量一致.通常来说,一个SM能够支持2048threads,调度几十到几百block,以及含有一定内存的寄存器,这些所有的数量限制都会对实际SM可以含有的线程数造成影响,比如如果一个thread使用过多的寄存器内存,则寄存器会成为最大限制,导致一个SM中实际支持的线程数量达不到可支持的最大线程数.这个实际线程数和可支持的最大线程数的比率,我们叫做占用率(Occupancy)

在编程中,我们追求高占用率,这样可以充分利用SM,也就是说,尽量达到.当然,在一些应用中,为了追求高计算算力,我们会牺牲占用率,一个例子就是GEMM的实现.

在编程中,由于我们知道,最小的调度单位是一个warp,所以我们会尽量减少warp内的分支控制差异(Control Divergence),避免warp内的部分线程的等待(latency)

GPU内存结构与数据逻辑

对于GPU,不同内存层级之间,访问数据的速度是不一样的.

在GPU中,内存分为显存(DRAM)、共享内存(Shared memory)、常量内存(Constant memory)、寄存器内存(Register memory).在访问速度上,显存的IO是最慢的,最快的则是寄存器的IO,而共享内存和常量内存居中,但是都远远快于访问显存.

对于GPU内存访问,有一个极其重要的指标——计算密度(Computational intensity),单位FLOP/B,即每B内存读取执行了多少次浮点运算操作(Float-point Operation per Byte),这个指标越高,说明每次访问内存的数据得到了更高的利用,通常可以通过优化内存访问(减少内存访问)或是优化计算方法(增加数据利用率)来提高这个值.

其次,除了计算密度,峰值带宽决定了内存最快的读取速度,单位为GB/s.它与计算密度乘积则是计算吞吐量(Computational throughput),单位GFLOPS,代表着若干内存读取对应的每秒实际计算量.

三者关系如图

在实际编程运用中,为了减少内存重复读取,我们会采用tiling的方式来减少内存冲突,增加一次读取数据利用率,这一方法经常运用于如矩阵乘法等.再次声明,我不会在这一部分提供详细代码,因为后续内容会详细介绍tiling以及thread coarsening、memory coalescing、register tiling这些方法的综合运用.

需要注意的是,既然我们提到了内存合并(memory coalescing),我们需要知道,GPU的多级缓存和内存结构决定了对于连续的地址读取,其速度总是优于跳跃的内存读取,因为每次内存读取,GPU会尝试将一段连续的地址内容放入缓存,如果本身读取的数据就是连续的,那么可以直接从缓存拿去数据,提高数据访问速度,这是远远快于跳跃读取数据的.对于写入数据也是如此,因此,催生了Grid-Stride Loop方法的大量运用.

Appendix

尽管我们不提供详细的代码使用,但是在此附录,我会给出CUDA基本函数及其规范,供于参考,准确内容请访问CUDA官网.

对于代码,GPU上的部分叫做device,在CPU运行部分叫host

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
/*
基本函数
*/
cudaMalloc((void**)&pointer,size);   //为指定数组申请GPU上内存,size即大小/B,需传递指针的地址
cudaMemcoy(pointer_des,pointer_src,size,Type);   //第一个指针代表目标地址,第二个指针代表数据拷贝源地址,size代表拷贝数据大小,Type指定数据类型,例如cudaMemcpyHostToDevice,cudaMemcpyDeviceToHost
cudaFree(pointer);   //释放GPU上内存,传递指针
__syncthreads();   //warp内线程同步阻塞,确保所有线程执行完毕上面的代码后同时开始执行后续代码

/*
核函数部分
*/
__global__   //核函数声明关键字
void kernal_function(args){}

//调用方法
kernal_function<<<gridDim,blockDim,size,S>>>(args);   //gridDim,类型int或dim3,网格维度;blockDim,类型int或dim3,块维度;size,可选,共享内存大小;S,类型cudaStream_t,可选,默认0

/*
GPU函数写法
*/
__device__ //GPU函数关键字
int device_function(args){}

/*
多维grid/block声明写法
*/
dim3 grid(x, y, z);   //三个参数代表三个方向的grid内block数量
dim3 block(x, y, z);   //三个参数代表三个方向的block内线程数量