1 简介

1.1 从图像处理到通用并行计算

由于市场对实时高清3D图形计算的强烈需求，可编程图像处理器(programmable Graphic Processor Unit)或简称GPU，已发展成高并行、多核、多线程的处理器，拥有强大算力和高内存带宽。

以下图一和图二比较了CPU和GPU之间浮点运算能力和内存带宽间方面的差异：

GPU卡间带宽 gpu带宽有什么用_CUDA

图一：CPU和GPU每秒浮点运算次数比较

GPU卡间带宽 gpu带宽有什么用_机器学习_02

图二：CPU与GPU之间内存带宽比较

GPU强大的浮点数运算能力源自它的结构是专为并发数据计算而设计的，这正是图像渲染所需要的，相比于CPU的设计，GPU将更多的晶体管用于数据处理，而不像CPU用于数据缓存和流程控制。

GPU卡间带宽 gpu带宽有什么用_机器学习_03

图三：GPU中更多的晶体管用于数据处理

这种设计理论上对于并行计算是有效的，因为GPU可以用并行计算弥补内存访问的延迟，而不用通过缓存和流控制来避免内存访问。

数据并行化处理将数据映射到并行运行的线程。很多需要处理大量数据的应用都可以采用数据并行编程模型来加速计算。比如在3D渲染中像素和顶点被映射到并行运行的不同线程。另外，比如图像处理、视频编解码、立体视觉、模式识别等图像和多媒体处理等等，都可以采用类似的并行编程模型。实际中除图形渲染和处理之外很多其他应用可通过数据并发来加速计算。从信号处理、物理模拟到计算金融学和计算生物学等等。

1.2 CUDA: 通用并行计算平台和编程模型

2006年NVIDIA推出CUDA，通用并行计算平台和编程模型，利用NVIDIA GPU的并行计算能力解决复杂的计算问题。

CUDA包含软件环境，开发者可以使用C++作为高层的开发语言进行并行程序开发。如图四所示，除C++之外，CUDA还支持一些其他的语言、应用编程接口和基于指令的操作方式。

GPU卡间带宽 gpu带宽有什么用_神经网络_04

图四：CUDA支持很多的编程语言和应用编程接口

1.3 可扩展编程模型

多核CPU和多核GPU的出现，意味着主流的处理芯片都已是并行系统。随之而来的挑战是，如何设计能根据不断增加的处理器内核数量透明的扩展自身并行度的应用软件。

CUDA的设计初衷就是克服这一挑战，同时对熟悉标准开发语言（如C语言）的开发者保持一个低的学习曲线。

CUDA的核心包括三个抽象概念：线程的分层结构，共享内存和同步机制。这三个概念以C++语言的扩展的形式暴露给开发者。

这些概念提供了精细粒度的数据并行和线程并行，以及粗粒度的数据并行和任务并行。它们指导开发者将问题分解为粗粒度的可由线程组并行独立解决的子问题，每个子问题再分解为细粒度的可由组内线程合作解决的最终问题。

这种分解方式允许了线程间的合作，同时保证了自动扩展性。每一个线程组可以独立的被GPU内的任何内核调度执行，以任何顺序，并行或线性顺序。因此编译后的CUDA程序可以运行在任何内核数量的GPU上，如图五所示，只有运行时系统才需要知道准确的物理内核数量。

GPU卡间带宽 gpu带宽有什么用_GPU卡间带宽_05

图五：自动扩展

注：GPU由一组流式多处理器（Streaming Multiprocessors, SMs)构成。多线程程序会被划分为多个线程组，彼此独立运行在不同的SM中，因此更多的SM数量意味着更短的运行延迟。

2 编程模型

本章我们介绍CUDA编程模型背后的主要概念。更详细的介绍将在后面的编程接口中给出。

本章使用的vectorAdd的源代码在CUDA sample里可以找到。

2.1 内核

CUDA对C++语言进行了扩展，开发者可以定义被称为内核的C++函数，当被调用时内核函数会被CUDA线程并行执行N次，而不像普通函数那样只被执行一次。

内核定义需要使用修饰符GPU卡间带宽 gpu带宽有什么用_机器学习_06，CUDA线程数量在执行配置<<<…>>>中指定。执行内核函数的CUDA线程绑定一个唯一的ID，可以通过内建变量threadIdx获取。

例如下面的代码使用内建变量threadIdx，进行数组A和B相加，并将结果存储到数组C中。

// Kernel definition
__global__ void VecAdd(float* A, float* B, float* C)
{
    int i = threadIdx.x;
    C[i] = A[i] + B[i];
}

int main()
{
    ...
    // Kernel invocation with N threads
    VecAdd<<<1, N>>>(A, B, C);
    ...
}

其中每个cuda线程执行N个加法操作的一个。

2.2 线程层次结构

内建变量threadIdx为一个包含三个元素的向量，因此每个CUDA线程可以通过一个一维，二维或三维的索引唯一标识，整体的CUDA线程则构成一个一维，二维或三维的线程组，我们称之为线程块（thread block）。这种表示方式可以很方便的操作和计算向量、矩阵或分区内的元素。

线程索引与线程ID的关系也比较简单：对于一维线程块，二则相等；对于二维线程块，GPU卡间带宽 gpu带宽有什么用_神经网络_07表示线程块形状，GPU卡间带宽 gpu带宽有什么用_CUDA_08表示线程索引，则线程ID为GPU卡间带宽 gpu带宽有什么用_GPU卡间带宽_09；对于三维线程块，GPU卡间带宽 gpu带宽有什么用_深度学习_10表示线程块形状，GPU卡间带宽 gpu带宽有什么用_机器学习_11表示线程索引，则线程ID表示为GPU卡间带宽 gpu带宽有什么用_神经网络_12.

下面的例子将两个NxN的A、B矩阵相加，并将结果存如C:

// Kernel definition
__global__ void MatAdd(float A[N][N], float B[N][N],
                       float C[N][N])
{
    int i = threadIdx.x;
    int j = threadIdx.y;
    C[i][j] = A[i][j] + B[i][j];
}

int main()
{
    ...
    // Kernel invocation with one block of N * N * 1 threads
    int numBlocks = 1;
    dim3 threadsPerBlock(N, N);
    MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
    ...
}

线程块中的线程数量是有限的，因为块中线程属于同一个处理器并共享处理器有限的内存资源。目前每个线程块中线程数量的上限为1024.

内核可以被相同形状的多个线程块执行，因此线程的总数量等于线程块中的线程数量乘以线程块的数量。

线程块一起又构成一维、二维或三维的网格结构，如图六所示。

GPU卡间带宽 gpu带宽有什么用_深度学习_13

图六：网格结构的线程块

块中的线程数量，以及块数量都在<<<…>>>中以int或dim3变量类型指定。图六中的线程块和网格都是二维的。网格中的每个线程块都由一个一维、二维或三维的索引唯一标识，索引可以通过内建的变量blockIdx获取。块的维度可以通过内建变量blockDim获取。

vecAdd的例子，加入多个线程块之后变成：

// Kernel definition
__global__ void MatAdd(float A[N][N], float B[N][N],
float C[N][N])
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    int j = blockIdx.y * blockDim.y + threadIdx.y;
    if (i < N && j < N)
        C[i][j] = A[i][j] + B[i][j];
}

int main()
{
    ...
    // Kernel invocation
    dim3 threadsPerBlock(16, 16);
    dim3 numBlocks(N / threadsPerBlock.x, N / threadsPerBlock.y);
    MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
    ...
}

在本例中线程块的shape为16x16，是任意设置的。块的数量设置保证矩阵的没一个元素都对应到一个线程。例子中我们假设矩阵的维度可以被块的数量整除，实际上不能整除也不会有问题。

线程块会被独立运行：执行顺序可能使任意的顺序，并行或则串行。这种独立性保证了线程块可以以人物顺序在任意数量的处理器上被调度和执行，如图五所示，使开发者的程序能随着处理数量自动的进行扩展。

同一个块内的线程可以通过共享内存共享数据，并听过同步机制协调内存的访问。具体来说，线程可以通过调用内建函数GPU卡间带宽 gpu带宽有什么用_神经网络_14进行同步。除了GPU卡间带宽 gpu带宽有什么用_神经网络_14之外，后面将详细介绍其他线程同步的机制。

为了提高线程间协作的效率，要求共享内存是靠近处理器的低延迟内存（类似L1缓存），GPU卡间带宽 gpu带宽有什么用_神经网络_14必须是轻量级的操作。

2.3 内存层级结构

如图七所示，CUDA线程可以访问多个内存空间。每个线程有自己的私有内存；每个块内的线程共享块内部的共享内存，快内共享内存生命期与块相同；所有的线程共享全局内存。

另外还有两个只读内存空间：常量内存空间和材质内存空间。全局内存、常量内存、材质内存针对不同的场景进行了优化（参考接下来的设备内存访问章节）。材质内存针对一些特定的数据格式提供不了不同的寻址模式，如数据过滤（参考接下来的材质和外观内存章节）。

全局、常量和材质内存的生命周期与程序一致，在程序的不同内核调用之间保持不变。

GPU卡间带宽 gpu带宽有什么用_GPU卡间带宽_17

图七：内存层级结构

2.4 异构编程

如图八所示，CUDA编程模型假定宿主运行C++程序， 独立的设备作为宿主的协处理器运行CUDA线程。

CUDA还假定宿主与独立设备分别维护各自的内存空间，分别称之为宿主内存和设备内存。程序通过调用CUDA运行时接口管理CUDA内核的全局、常量和材质内存。包括内存的申请和释放，以及宿主和设备内存之间的数据传输。

统一内存机制提供了跨越宿主内存和设备内存的方式。统一内存可以被系统中所有的CPU和GPU访问。详细见接下来的统一内存编程章节。简单来说，统一内存机制使得我们不需要显式地调用cudaMemecpy系列的函数来进行数据传输，但是底层的数据传输依然是存在的，所以程序的执行时间不会减少。统一内存机制使得代码更简洁和易于维护。

GPU卡间带宽 gpu带宽有什么用_神经网络_18

图八：异构编程

2.5 计算能力

设备的计算能力表示为一个版本号，有时称之为"SM版本"。这个版本号标识GPU硬件所支持的特性，因此应用程序可以在运行时以此来判断当前GPU支持哪些硬件特性和指令。

版本号包含一个主版本号X和一个小版本号Y，记为X.Y.

主版本号相同的设备内核架构一样。7的设备基于Volta架构，6的基于Pasca架构，5的基于Maxwell架构，3的基于Kepler架构，2的基于Fermi架构，1的基于Tesla架构。

不同的小版本号对应不同的改进，比如新特性的加入。

注意：不要混淆GPU的计算能力与CUDA的版本号。

从cuda7.0和cuda9.0开始，分别不在支持基于Tesla架构和Fermi架构的GPU.