并行计算、费林分类与CUDA：AI基础设施核心概念
摘要
异构计算通过CPU与GPU协同工作应对算力膨胀，并行计算将问题拆解为独立任务。弗林分类法
聊这个话题之前，先看看这些年发生了什么。人工智能、高性能数据分析、金融计算……这些计算密集型领域对算力的需求，几乎是在以肉眼可见的速度膨胀。传统的通用CPU虽然“什么活儿都能干”，但架不住数据量越来越大、运算越来越复杂。于是，异构计算从一个实验室里的概念，逐步变成了产业界的香饽饽。
所谓异构计算，简单说就是把不同类型、不同指令集和体系结构的计算单元组合在一起干活。相比传统CPU那种“单打独斗”的模式，异构计算在效率和延迟上的优势是碾压级的。目前主流的异构计算引擎，无非就是GPU、FPGA、ASIC这几类。而在这其中，GPU又是科研和工程领域最常见的硬件翻跟斗。
回头来看，通用CPU其实已经做得相当复杂了——几十个核心，运行频率动辄几GHz，每个核心还有自己的一级、二级、三级缓存。但GPU呢？它的计算核心数通常是CPU的上百倍，虽然频率略低，可架不住核心多。用一句话来概括：CPU更适合“多面手”的角色，兼顾数据读写、计算、逻辑控制；而GPU的基因决定了它天生适合那些计算密集型的任务，比如视频处理、AI训练，甚至传统的科学计算和工程计算，也正在大规模地向GPU迁移。
并行计算
过去几十年，业界对并行计算的热情一直在升温。目标很明确——更快。从最纯粹的计算视角来看，并行计算就是把一个大问题拆成很多可以同时解的小问题。这项工作涉及两个完全不同的技术领域：一个是计算机架构（硬件层面），一个是并行程序设计（软件层面）。
硬件负责提供支持并行性的结构基础，软件则要想办法充分“榨干”这些硬件的能力。现代处理器大多采用的是哈佛体系结构，如下图所示，核心组件包括：内存（指令内存和数据内存）、中央处理单元（控制单元和算数逻辑单元），以及输入/输出接口。
高性能计算的关键，说白了还是CPU这个“大脑”。早期的芯片是单核处理器，一个芯片上只有一个CPU。现在呢？趋势是把多个核心集成到一个处理器上，也就是多核处理器，目的就是在架构层面支持并行。对程序员来说，并行程序设计就是把一个问题的计算拆开，分给这些可用的核心去跑。
写串行程序的时候，你几乎不需要了解计算机架构的细节。但当你面对多核环境，情况就完全不同了——要写出既正确又高效的并行程序，就必须对多核体系结构有基本的认识。
串行编程
用计算机程序解决问题时，我们很自然地会把问题划分成许多运算块，每个块执行一个指定的任务。这就是串行程序。
那么问题来了，如何区分两个计算单元之间的关系？有些需要排好次序，必须串行执行；另外一些则没有次序上的约束，可以并发执行。所有包含并发执行任务的程序，都可以叫并行程序。但要注意，一个并行程序里很可能还藏着一些串行部分。
从程序员的角度看，一个程序由两部分组成：指令和数据。当一个计算问题被划分成许多小任务后，每个任务里都有指令负责处理输入并产生输出。这里的关键在于“数据相关性”——如果一个任务的输入依赖另一个任务的输出，它们就是相关的，否则就是独立的。
理解数据相关性，是并行算法实现的基本功。因为相关性是限制并行性的主要因素。反过来看，那些没有依赖关系的独立任务链，恰恰为并行化提供了绝佳的机会。
并行性
如今并行性的应用已经相当普遍，并行编程正在成为主流。多层次的并行性设计，是架构设计的核心驱动力。在应用程序中，有两种基本的并行类型：任务并行和数据并行。
任务并行，着眼点在于把多个独立的任务或函数分配给多核系统并行执行。数据并行，则是同时处理大量数据，侧重于将数据合理分配给多核系统。很多处理大数据集的应用，都可以借助数据并行模型来提速。
数据并行程序设计的第一步，是把数据依据线程进行划分。常见的方法有两种：块划分和周期划分。块划分是将一组连续的数据分成一个块，每个线程处理一块；周期划分则是把更少的数据分到一个块，相邻的线程处理相邻的数据块，每个线程可以处理多个块。
下图展示了对一维数据进行划分的两个例子：块划分和周期划分。
下图则是对二维数据进行划分的三种情况：沿y轴的块划分、沿x轴和y轴的块划分，以及沿x轴的周期划分。
需要注意，数据本质上是在一维空间中存储的，即便是多维逻辑数据，最终也要映射到一维物理地址空间。如何在线程中分配数据，不仅与数据的物理存储方式有关，还和每个线程的执行次序密切相关。线程的组织方式，对程序性能有至关重要的影响。
计算机架构
计算机架构有多种分类方法。最广为人知的是弗林分类法，它根据指令和数据进入CPU的方式，将架构分为四种类型：SISD、SIMD、MISD和MIMD。
SISD就是传统计算机那种串行架构，一个核心，同一时间只有一个指令流处理一个数据流。SIMD则是并行架构，多个核心在同一时间执行同一个指令流，但处理不同的数据。向量机就是典型的SIMD架构，而如今大多数计算机都包含了SIMD指令集。SIMD最大的优势在于：程序员可以继续按串行逻辑思考，但对并行数据操作实现加速，剩下的细节交给编译器处理。
MISD架构比较少见，在这种架构中，每个核心使用多个指令流处理同一个数据流。MIMD则是另一种并行架构，多个核心用多个指令流异步处理多个数据流，实现空间上的并行性。很多MIMD架构内部还包含了SIMD执行的子组件。
在架构层面，降低延迟、提高带宽、提高吞吐量是永恒的目标。延迟是一个操作从开始到完成所需的时间，常用微秒表示；带宽是单位时间内可处理的数据量，常用MB/s或GB/s表示；吞吐量则是单位时间内成功处理的运算数量，在科学计算领域常用gflops来表示。
计算机架构还可以根据内存组织方式来划分，主要有两种：分布式内存的多节点系统和共享内存的多处理器系统。
在多节点系统中，大量处理器通过网络连接，每个处理器都有自己独立的本地内存，彼此通过网络通信。下图就是一个典型的分布式内存多节点系统，也就是常说的集群。
多处理器架构的规模从双处理器到几百个处理器不等，它们要么共享同一个物理内存，要么共用一条低延迟的链路（如PCI-E）。共享内存意味着共享地址空间，并不等于只有一个物理内存。这类系统既包括单片机上的多核系统，也包括多个芯片组成的系统，每个芯片又可能是多核的。如今，多核架构早已取代了单核架构。
“众核”通常指包含几十个甚至几百个核心的多核架构。近年来，计算机架构正在从多核向众核演化。GPU就是众核架构的代表，它几乎囊括了前文提到的所有并行结构：多线程、MIMD、SIMD，以及指令级并行。NVIDIA把这套架构叫做SIMT（单指令多线程）。
GPU和CPU的出身不同。GPU最初是图形翻跟斗，直到最近才演变成一个强大的、多用途的、完全可编程的任务和数据并行处理器，特别适合解决大规模并行计算问题。
GPU核心和CPU核心
虽然可以用多核和众核来区分CPU和GPU的架构，但这两种核心本质上完全不同。CPU核心比较“重”，专门处理复杂的控制逻辑，优化串行程序的执行；GPU核心则比较“轻”，针对控制逻辑简单的数据并行任务优化，核心目标是大吞吐量。
GPGPU
通用图形处理器这个概念，最早由NVIDIA公司的Mark J. Harris在2002年提出。基于图形渲染管线的流水线特征，GPU本质上就是一个可以同时处理多个计算任务的硬件翻跟斗。由于GPU包含大量的计算资源，Mark J. Harris从2002年就开始尝试在GPU上做通用并行计算。不过那个阶段受限于架构和编程平台，研究人员不得不把目标算法转换成图形运算算法，才能借用GPU来实现通用计算。
转折点出现在NVIDIA推出Tesla统一渲染架构和CUDA编程模型之后。从那时起，NVIDIA的GPU开始全面支持通用并行计算。CUDA发布近两年后，OpenCL 1.0标准也来了——这意味着利用GPU做通用并行计算已经基本成熟。如今，市面上的GPU芯片除了完成高质量的图形渲染，通用并行计算也成了主流应用。
GPGPU通常由成百上千个架构相对简易的基本运算单元组成，多采用极简的流水线设计，不提供复杂的控制优化技术。每个基本运算单元可以同时执行一至多个线程，由GPGPU中的调度器控制。尽管各家厂商的芯片架构各异，但几乎走的都是众核处理器阵列路线，在一个GPU芯片里集成成百上千个核心，以获得更高的计算性能和更大的数据带宽。
GPU中执行的线程对应的程序通常被称为“内核”。需要注意的是，这个“内核”和操作系统的内核是完全不同的概念。GPU中的线程也比CPU中的线程更简单、更轻量。在GPU众核架构中，多个处理核心通常被组织成一个线程组来调度执行，比如NVIDIA的流多处理器、AMD的SIMD执行单元。一个线程组内的线程执行相同的程序指令，同步执行，每个线程处理不同的数据，实现数据级并行。不同架构对线程组的命名也不一样，NVIDIA叫warp，AMD叫wa vefront，每个组包含的线程数从4个到128个不等。执行模式常见的有SIMT和SIMD两种。
在GPU程序中，数据加载存储和条件分支跳转这两类指令，往往会引发不可预测的执行情况。对于前者，如果一级高速缓存命中缺失，指令执行周期将不可预测。为了避免执行单元因数据加载造成资源浪费，GPU的每个执行单元通常设有线程组缓冲区，支持同时执行多个线程组。线程组之间的调度由硬件调度器负责，这个过程一般是零负载的。准备执行的线程组会先进入缓冲区，以队列方式组织，调度器从中选择一个准备好的启动执行。这种调度方式可以有效解决长延时操作导致的停顿问题。而对于条件分支指令，无论是SIMD还是SIMT模式，当一组线程都执行相同代码路径时性能最佳；如果每个线程各自走不同的分支，为了保证正确性，指令发送单元将串行发送所有指令，效率会大打折扣。
异构计算
最初，计算机只靠CPU完成所有编程任务。但近年来，高性能计算领域的主流计算机开始引入其他处理单元，其中最主要的就是GPU。GPU原本是为并行图形计算而生的，但随着时间的推移，它已经进化成一个更强大、更通用的处理器，在大规模并行计算中展现出优越的性能和效率。
CPU和GPU是两个独立的处理器，在单个计算节点中通过PCI-E总线相连。GPU在这种架构中是一个离散的设备。从同构系统到异构系统的转变，可以说是高性能计算史上的一个里程碑。同构计算使用同一架构下的一个或多个处理器执行应用；异构计算则不同，它会根据任务的特点，为它选择最适合的架构，从而提升整体性能。
异构系统比传统的高性能计算系统优势明显，但问题也很现实——它增加了应用程序设计的复杂性。近期备受关注的并行计算，也因为引入了异构资源而变得更复杂了。对于刚接触并行编程的人来说，这些特性改进和软件工具会很有帮助；如果你已经是并行编程的老手，那么适应异构架构的并行编程其实并不难。
GPU 异构计算
CPU-GPU协同工作，是实现高性能计算的必要条件，这就是所谓的CPU-GPU异构计算。它的核心思路是：把应用程序中计算密集的部分卸载到GPU上，其余代码仍然由CPU运行。这样一来，既能利用CPU和GPU各自的最佳特性，又能避免两个处理单元出现空闲，最终实现高计算增益。当然，要充分发挥异构计算的潜力，还需要新的优化技术。
理解CPU和GPU的差异，一个简单的办法是看它们处理任务的方式。CPU的几个核心是针对顺序串行处理优化的；而GPU拥有大规模并行架构，由数千个更小、更高效的核心组成，专门用来同时处理多个任务。
在GPU上解决计算问题，和用多个CPU解决问题在原则上类似：把任务拆分成小任务，每个小任务由单个GPU核心执行。GPU核心之间的通信依靠芯片上的内部寄存器和内存。程序员不再需要用消息传递来处理通信，而是通过CUDA或OpenCL这类编程语言，由平台提供主机CPU和GPU之间的数据交换与同步机制。
一个现代超级计算系统，通常由大量节点组成，每个节点包含2到32颗常规CPU以及1到16个GPU，再加上高速网络和数据存储系统。软件可以用C/C++、Fortran等传统语言，结合用于CPU并行化的消息传递系统以及用于GPU的CUDA或OpenCL来编写。所有组件都需要精心调整和优化，才能实现系统的最佳性能。
异构架构
一个典型的异构计算节点，包含两个多核CPU插槽和两个或更多个众核GPU。GPU并不是一个独立运行的平台，而是CPU的协处理器，必须通过PCIe总线与基于CPU的主机端相连。这就是为什么CPU所在的位置叫主机端，GPU所在的位置叫设备端。
一个异构应用包括两部分：主机代码（在CPU上运行）和设备代码（在GPU上运行）。异构平台上执行的应用通常由CPU初始化，在加载计算密集型任务到设备端之前，CPU代码负责管理设备端的环境、代码和数据。
在计算密集型应用中，往往有大量并行数据的程序段。GPU就是用来加速这些数据并行任务的。当一个硬件组件在物理上独立于CPU，却用来加速计算密集部分时，它就成了所谓的硬件翻跟斗——而GPU可以说是最常见的那个。
衡量GPU的能力，有两个核心指标：CUDA核心数量和内存大小。相应地，评估GPU性能也有两个标准：峰值计算性能和内存带宽。峰值计算性能通常用GFlops或TFlops来表示，内存带宽则用GB/s表示。下表展示了Fermi架构和Kepler架构的一些性能指标。
异构计算范例
GPU计算并不是要取代CPU计算。每种计算方式都有它独特的优势。CPU适合控制密集型任务，GPU适合包含数据并行的计算密集型任务。两者结合，能有效提高大规模计算问题的处理速度和性能。CPU针对动态工作负载做了优化，这些负载的特点是短序列的计算操作和不可预测的控制流程；而GPU的战场在另一头——那些由计算任务主导、控制流相对简单的工作负载。
具体怎么选？如果一个问题的数据规模小、控制逻辑复杂、并行性也不高，用CPU处理更合适；反过来，如果问题涉及大规模数据、数据并行性极高，那么GPU就是最佳选择——它有大量可编程的核心可以支持大规模多线程运算，峰值带宽也远超CPU。
正是CPU和GPU功能的互补性，催生了CPU+GPU的异构并行计算架构。想让应用程序获得最佳效果，可以同时在CPU和GPU上执行：串行部分或任务并行部分交给CPU，数据密集型并行部分交给GPU。
这种代码编写方式，能确保GPU和CPU相辅相成，充分发挥异构系统的计算能力。为了支持这种架构，NVIDIA设计了一个编程模型——CUDA。
CPU线程与GPU线程
CPU上的线程通常是重量级实体，操作系统必须在CPU执行通道上不断交换线程来提供多线程功能，上下文切换慢且开销大。GPU上的线程则非常轻量——在典型系统中，数千个线程排队等待工作，如果GPU必须等待一组线程，它会直接开始执行另一组。CPU核心致力于最小化一两个线程的延迟，而GPU核心的目标是用大量并发的轻量级线程来最大化吞吐量。
举个例子：一个四核CPU（支持超线程）可以同时运行16到32个线程；而现代NVIDIA GPU，每个多处理器可以同时支持最多1536个活跃线程——一台有16个多处理器的GPU，可以并发支持超过24000个同时活跃的线程。
CUDA：一种异构计算平台
CUDA是一个通用的并行计算平台和编程模型，它利用NVIDIA GPU中的并行计算引擎，更高效地解决复杂的计算问题。通过CUDA，你可以像在CPU上一样，在GPU上执行计算任务。
CUDA平台通过多种方式提供服务：CUDA加速库、编译器指令、应用编程接口以及行业标准程序语言的扩展（包括C、C++、Fortran、Python）。
CUDA C是标准ANSI C的扩展，它用少量语言扩展实现了异构编程，同时提供了API来管理设备、内存和其他任务。CUDA还是一个可扩展的编程模型，对不同核心数的GPU都能显著扩展并行性，而对熟悉C语言的人来说也比较容易上手。
CUDA提供了两层API用于管理GPU设备和组织线程：驱动API和运行时API。
驱动API属于底层，编程难度较高，但对于GPU设备的使用提供了更精细的控制。运行时API则是高层API，构建在驱动API之上。每个运行时API函数都会分解成多个传给驱动API的基本运算。需要注意的是，运行时API和驱动API之间并没有明显的性能差异——真正影响性能的，是程序员如何在设备上使用内存以及如何组织线程。这两套API是互斥的，使用其中之一就不能混用另一种。
一个CUDA程序包含了两个部分：在CPU上运行的主机代码和在GPU上运行的设备代码。更具体地说，CUDA的操作可以概括为6个步骤：
- CPU在GPU上分配内存：cudaMalloc
- CPU把数据发送到GPU：cudaMemcpy
- CPU在GPU上启动kernel（自写程序，在每个线程上运行）
- CPU等待GPU端完成任务：cudaDeviceSynchronize
- CPU把数据从GPU取回：cudaMemcpy
- CPU释放GPU上的内存：cudaFree
/**
* Vector addition: C = A + B.
*
* This sample is a very basic sample that implements element by element
* vector addition. It is the same as the sample illustrating Chapter 2
* of the programming guide with some additions like error checking.
*/
#include
// For the CUDA runtime routines (prefixed with "cuda_")
#include
#include
/**
* CUDA Kernel Device code
*
* Computes the vector addition of A and B into C. The 3 vectors ha ve the same
* number of elements numElements.
*/
__global__ void vectorAdd(const float *A, const float *B, float *C,
int numElements) {
int i = blockDim.x * blockIdx.x + threadIdx.x;
if (i < numElements) {
C[i] = A[i] + B[i] + 0.0f;
}
}
/**
* Host main routine
*/
int main(void) {
// Error code to check return values for CUDA calls
cudaError_t err = cudaSuccess;
// Print the vector length to be used, and compute its size
int numElements = 50000;
size_t size = numElements * sizeof(float);
printf("[Vector addition of %d elements]\n", numElements);
// Allocate the host input vector A
float *h_A = (float *)malloc(size);
// Allocate the host input vector B
float *h_B = (float *)malloc(size);
// Allocate the host output vector C
float *h_C = (float *)malloc(size);
// Verify that allocations succeeded
if (h_A == NULL || h_B == NULL || h_C == NULL) {
fprintf(stderr, "Failed to allocate host vectors!\n");
exit(EXIT_FAILURE);
}
// Initialize the host input vectors
for (int i = 0; i < numElements; ++i) {
h_A[i] = rand() / (float)RAND_MAX;
h_B[i] = rand() / (float)RAND_MAX;
}
// Allocate the device input vector A
float *d_A = NULL;
err = cudaMalloc((void **)&d_A, size);
if (err != cudaSuccess) {
fprintf(stderr, "Failed to allocate device vector A (error code %s)!\n",
cudaGetErrorString(err));
exit(EXIT_FAILURE);
}
// Allocate the device input vector B
float *d_B = NULL;
err = cudaMalloc((void **)&d_B, size);
if (err != cudaSuccess) {
fprintf(stderr, "Failed to allocate device vector B (error code %s)!\n",
cudaGetErrorString(err));
exit(EXIT_FAILURE);
}
// Allocate the device output vector C
float *d_C = NULL;
err = cudaMalloc((void **)&d_C, size);
if (err != cudaSuccess) {
fprintf(stderr, "Failed to allocate device vector C (error code %s)!\n",
cudaGetErrorString(err));
exit(EXIT_FAILURE);
}
// Copy the host input vectors A and B in host memory to the device input
// vectors in device memory
printf("Copy input data from the host memory to the CUDA device\n");
err = cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
if (err != cudaSuccess) {
fprintf(stderr,
"Failed to copy vector A from host to device (error code %s)!\n",
cudaGetErrorString(err));
exit(EXIT_FAILURE);
}
err = cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);
if (err != cudaSuccess) {
fprintf(stderr,
"Failed to copy vector B from host to device (error code %s)!\n",
cudaGetErrorString(err));
exit(EXIT_FAILURE);
}
// Launch the Vector Add CUDA Kernel
int threadsPerBlock = 256;
int blocksPerGrid = (numElements + threadsPerBlock - 1) / threadsPerBlock;
printf("CUDA kernel launch with %d blocks of %d threads\n", blocksPerGrid,
threadsPerBlock);
vectorAdd<<>>(d_A, d_B, d_C, numElements);
err = cudaGetLastError();
if (err != cudaSuccess) {
fprintf(stderr, "Failed to launch vectorAdd kernel (error code %s)!\n",
cudaGetErrorString(err));
exit(EXIT_FAILURE);
}
// Copy the device result vector in device memory to the host result vector
// in host memory.
printf("Copy output data from the CUDA device to the host memory\n");
err = cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);
if (err != cudaSuccess) {
fprintf(stderr,
"Failed to copy vector C from device to host (error code %s)!\n",
cudaGetErrorString(err));
exit(EXIT_FAILURE);
}
// Verify that the result vector is correct
for (int i = 0; i < numElements; ++i) {
if (fabs(h_A[i] + h_B[i] - h_C[i]) > 1e-5) {
fprintf(stderr, "Result verification failed at element %d!\n", i);
exit(EXIT_FAILURE);
}
}
printf("Test PASSED\n");
// Free device global memory
err = cudaFree(d_A);
if (err != cudaSuccess) {
fprintf(stderr, "Failed to free device vector A (error code %s)!\n",
cudaGetErrorString(err));
exit(EXIT_FAILURE);
}
err = cudaFree(d_B);
if (err != cudaSuccess) {
fprintf(stderr, "Failed to free device vector B (error code %s)!\n",
cudaGetErrorString(err));
exit(EXIT_FAILURE);
}
err = cudaFree(d_C);
if (err != cudaSuccess) {
fprintf(stderr, "Failed to free device vector C (error code %s)!\n",
cudaGetErrorString(err));
exit(EXIT_FAILURE);
}
// Free host memory
free(h_A);
free(h_B);
free(h_C);
printf("Done\n");
return 0;
}
在后续的文章中，我们会继续介绍CUDA编程示例。
NVIDIA的CUDA nvcc编译器在编译过程中，会将设备代码从主机代码中分离出来。主机代码是标准C代码，使用C编译器编译；设备代码（也就是核函数）用扩展的CUDA C语言编写，由nvcc编译。在链接阶段，CUDA运行时库会加入内核程序调用和GPU设备操作。
来源：互联网
本网站新闻资讯均来自公开渠道，力求准确但不保证绝对无误，内容观点仅代表作者本人，与本站无关。若涉及侵权，请联系我们处理。本站保留对声明的修改权，最终解释权归本站所有。
