CUDA

from：http://163n.blog.163.com/blog/static/56035552201112043528767/

Since Nov. 4

 

概述
2006年，NVIDIA推出了CUDA，一种通用的并行计算架构。
CUDA支持多种编程语言或应用编程接口。
CUDA的核有三个关键特性：层次线程组（a hierarchy of thread groups）、共享内存（shared memories）和屏障同步（barrier synchronization）。
CUDA程序可以在任意数量的处理器核上运行。
CPU与GPU的差异：
CPU线程与GPU线程：CPU的一个核心通常在一个时刻只能运行一个线程的指令，CPU切换线程的代价十分高昂，通常需数百个时钟周期。GPU采用的则是由硬件管理的轻量级线程，可实现零开销的线程切换。
多核与众核：当前主流CPU中一般有2~8个核心，每个核心中有3~6条执行流水线。这些核心采用了很多提高指令级并行的技术。当前的NVIDIA GPU中有1~30个包含完整前端的流多处理器，每个流多处理器可看成一个包含8个1D流处理器的SIMD处理器。CUDA利用了多个流处理器间的粗粒度任务级或数据级并行，以及流多处理器内的细粒度数据并行。更多的执行单元数量使GPU能够在浮点处理能力上获得优势，主流GPU的性能可达到同时期主流GPU性能的10倍左右。
外部存储器：GT200 GPU的显存带宽达到了140GB/s，是同期CPU最高内存带宽的5倍。造成这种差异的主要原因有：
显存中使用的GDDR存储器颗粒与内存的DDR存储器颗粒在技术上基本相同，但显存颗粒直接固化在显卡的PCB板上，而内存为了兼顾可扩展性的需要，须通过DIMM插槽与主板相连。因此，显存的信号完整性问题比内存更容易解决，显存的工作频率也比使用相同技术的内存要高一些。
目前的CPU存储器控制器一般基于双通道或三通道技术，每个通道 位宽64bit；而GPU中则存在数个存储器控制单元，如GTX280 CPU中就有8个存储器控制器，每个控制两片位宽32bit的显存芯片，使总的存储器位宽达到512bit。

缓存：CPU中的缓存主要用于减小访存延迟和节约带宽，在多线程环境下会发生失效反应：每次线上下文切换之后，都需重建缓存上下文，一次缓存失效的代价是几十到上百个时钟周期。且为了实现缓存与内存中数据的一致性，还需要复杂的逻辑进行控制。而在GPU中则没有复杂的缓存体系与替换机制。GPU缓存是只读的，因此，也不用考虑缓存一致性问题。GPU缓存的主要功能是用于过滤对存储器控制器的请求，减少对显存的访问，从而节约显存带宽。

设备计算能力：设备计算能力的版本描述了一种GPU对CUDA功能的支持程度。计算能力版本中小数点前的第一位用于表示设备核心架构，小数点后的第二位则表示更加细微的进步，包换对核心架构的改进以及功能的完善等。例如，计算能力1.0的设备能够CUDA，而计算能力1.1设备加入了对全局存储器原子操作的支持。

CUDA编程模型
CUDA编程模型将CPU作为主机，GPU作为协处理器（co-processor）或设备。在这个模型中，CPU负责逻辑性强的事务处理和串行计算，GPU则专注于高度线程化的并行处理任务。CPU、GPU各自拥有相互独立的存储器地址空间。
一旦确定了程序中的并行部分，就可以考虑把这部分计算工作交给GPU。
kernel：运行在GPU上的C函数称为kernel。一个kernel函数并不是一个完整的程序，而是整个CUDA程序中的一个可以被并行执行的步骤。当调用时，通过N个不同的CUDA线程执行N次。
一个完整的CUDA程序是由一系列的设备端kernel函数并行步骤和主机端的串行处理步骤共同组成的。
一个kernel函数中存在两个层次的并行，即Grid中的block间并行和block中的thread间并行。

Kernel函数的定义与调用

 
内核函数必须通过__global__函数类型限定符定义，并且只能在主机端代码中调用。在调用时，必须声明内核函数的执行参数。例如：
// Define kernel
__global__ void VecAdd(float * A, float * B, float * C)
{

int i = threadIdx.x;
C[i] = A[i] + B[i];

}
 
int main
{

// Call kernel
VecAdd<<<1, N>>>(A, B, C);

}

必须先为Kernel中用到的数组或变量分配好足够的空间，再调用kernel函数。否则，在GPU计算时会发生错误。
在设备端运行的线程之间是并行执行的，其中的每个线程按指令的顺序串行执行一次kernel函数。每一个线程有自己的block ID和thread ID用于与其他线程相区分。block ID和thread ID只能在kernel中通过内建变量访问。内建变时不时是由设备中的专用寄存器提供的，是只读的，且只能在GPU端的kernel函数中调用。

线程结构（Thread Hierarchy）
CUDA中以线程网格（Grid）的形式组织，每个线程网格由若干个线程块（block）组成，而每个线程块又由若干个线程（thread）组成。
threadIdx：CUDA中使用了dim3类型的内建变量threadIdx和blockIdx。threadIdx是一个包含3个组件的向量，这样线程可以用一维、二维或三维线程索引进行识别，从而形成一个一维、二维或三维线程块。一个线程的索引和它的线程ID之间的关系非常直接：
对于一个一维的块，线程的threadIdx就是threadIdx.x；
对于一个二维的大小为（Dx,Dy）的块，线程的threadIdx就是（threadIdx.x + threadIdx.y * Dx）；
对于一个三维的大小为（Dx,Dy,Dz）的块，线程的threadIdx是（threadIdx.x + threadIdx.y * Dx + threadIdx.z * Dx * Dy）。

一个block中的线程数量不能超过512个。
在同一个block中的线程可以进行数据通信。CUDA中实现block内通信的方法是：在同一个block中的线程通过共享存储器（shared memory）交换数据，并通过栅栏同步保证线程间能够正确地共享数据。具体来说，可以在kernel函数中需要同步的位置调用__syncthreads()函数。
一个block中的所有thread在一个时刻执行指令并不一定相同。例如，在一个block中可能存在这样的情况：有些线程已经执行到第20条指令，而这时其他的线程只执行到第8条vkjsfdsvd第21条语句的位置通过共享存储器共享数据，那么只执行到第8条语句的线程中的数据可能还没来得及更新，就被交给其他线程去处理了，这会导致错误的计算结构。而调用__syncthreads()函数进行栅栏同步（barrier）以后，就可以确保只有当block中的每个线程都运行到第21条指令以后，程序才会继续向下进行。
每个线程块中的线程数量、共享存储器大小和寄存器数量都要受到处理核心硬件资源的限制，其原因是：
在GPU中，共享存储器与执行单元的物理距离必须很小，处于同一个处理核心中，以使得共享存储器的延迟尽可能小，从而保证线程块中的各个线程能够有效协作。
为了在硬件上用很小的代价就能实现__syncthreads()函数，一个block中所有线程的数据都必须交由同一处理核心进行处理。

硬件映射

 

计算单元
计算核心：GPU中有多个流多处理器（Stream Multiprocessor, SM），流多处理器即计算核心。每个流多处理器又包含8个标量流处理器（Stream Processor），以及少量的其他计算单元。SP 只是执行单元，并不是完整的处理核心。拥有完整前端的处理核心，必须包含取指、解码、分发逻辑和执行单元。隶属同一 SM 的8个 SP共用一套取指与射单元，也共用一块共享存储器。
CUDA 中的 kernel 函数是以 block 为单元执行的，同一 block 中的线程需要共享数据，因此必须在同一个 SM 中发射，而 block 中的每一个 thread 则被发射到一个 SP 上执行。
一个 block 必须被分配到一个 SM 中，但一个 SM 中同一时刻可以有多个活动线程块（active block）在等待执行，即在一个 SM 中可同时存在多个 block 的上下文。在一个 SM 中发射多个线程块是为了隐藏延迟，更好地利用执行单元的资源。当一个 block 进行同步或访问显存等高延迟操作时，另一个 block 就可以“乘虚而入”，占用 GPU 执行资源。
限制 SM 中活动线程块数量的因素包括：SM中的活动线程块数量不超过 8 个；所有活动线程块中的 warp 数之和在计算能力 1.0/1.1 设备中不超过 24，在计算能力 1.2/1.3 设备中不超过 32；所有活动线程块使用的寄存器和存储器之和不超过 SM 中的资源限制。

warp

 

在实际运行中，block 会被分割为更小的线程束 （warp）。线程束的大小由硬件的计算能力版本决定。
 

参考文献：

 

张舒, 褚艳利. GPU 高性能运算之CUDA. 北京：中国水利水电出版社，2009.10.

NVIDIA. NVIDIA CUDA C Programming Guide. 2010. 7.