风辰的CUDA入门系列教程
1. CUDA简介
GPU是图形处理单元(Graphic Processing Unit)的简称，最初主要用于图形渲染。自九十年代开始，GPU的发展产生了较大的变化，NVIDIA、AMD(ATI)等GPU生产商敏锐的观察到GPU天生的并行性，经过他们对硬件和软件的改进，GPU的可编程能力不断提高，GPU通用计算应运而生。由于GPU具有比CPU强大的计算能力，为科学计算的应用提供了新的选择。
最早的GPU的开发直接使用了图形学的API，将任务映射成纹理的渲染过程，使用汇编或者高级着色器语言Cg，HLSL等编写程序，然后通过图形学API执行（Direct3D和OpenGL），这样的开发不仅难度较大，而且难以优化，对开发人员的要求非常高，因此，传统的GPU计算并没有广泛应用。
2007年6月，NVIDIA公司推出了CUDA(Compute Unified Device Architecture)，CUDA不需要借助图形学API，而是采用了类C语言进行开发。同时，CUDA采用了统一处理架构，降低了编程的难度，使得NVIDIA相比AMD/ATI后来居上。相比AMD的GPU，NVIDIA GPU引入了片内共享存储器，提高了效率。这两项改进使CUDA架构更加适合进行GPU通用计算。由于这些特性，CUDA推出后迅速发展，被应用于石油勘测、天文计算、流体力学模拟、分子动力学仿真、生物计算、图像处理、音视频编解码等领域。
由于采用的是C/C++编译器为前端，以C/C++语法为基础设计，因此对熟悉C系列语言的程序员来说，CUDA的语法比较容易掌握。CUDA只对ANSI C进行了最小的必要扩展，以实现其关键特性--线程按照两个层次进行组织、共享存储器和栅栏同步。
这些关键特性使得CUDA拥有了两个层次的并行：线程级并行实现的细粒度数据并行，和任务级并行实现的粗粒度并行。
2. linux下CUDA开发环境构建
前一节已经简单了说了一下CUDA,为了能够使用CUDA开放，这一节将说明怎样构建CUDA开发环境。
本节讲解在ubuntu9.04操作系统和gcc前端的基础上安装CUDA开发环境。
首先，要保证自己机器上的gcc能够使用，因为ubuntu缺少gcc的一些包和g++，所以这些得自己安装。安装命令:sudo apt-get install g++，待其完成后，弄个C代码试试看:);当然你得保证你的显卡支持CUDA。
其次，到nVidia官方网站(http://www.nvidia.cn/object/cuda_get_cn.html)上下载对应操作系统的驱动(driver)和工具包(toolkit)。
再次，转换到控制台，命令为Ctrl+Alt+F1/F2/F3/F4，关掉gdm，命令为：sudo /etc/init.d/gdm stop，要确定已经关闭，否则在安装时会提示你有ｘ server程序在运行。
再次，进入driver和toolkit目录，执行安装命令,为了方便，请一定按照默认安装。
然后，打开个人目录下的.bashrc文件或者/etc/profile文件，在其中加入命令：PATH=${PATH}:/usr/local/cuda/bin/ export PATH，LD_LIBRARY_PATH=${LD_LIBRARY_PATH}:/usr/local/cuda/lib export LD_LIBRARY_PATH，保存后执行source .bashrc或者/etc/profile，依据你添加PATH和LD_LIBRARY_PATH时修改了那个文件确定。
最后执行nvcc命令，看看，如果提示你没有输入文件，就安装完成了。
如果你要安装SDK的话，还得安装一些包，依据make时的提示，google和新力得应该可以搞定一切，现在你可以享受CUDA了！
3. CUDA架构
本来一直都想把自己对ＣＵＤＡ架构的理解写出来，但是一方面自己是个懒人，不太愿意动笔，另一方面又感觉不太好组织语言，今天终于借着这个第四期有奖讨论的机会，把我对CUDA 架构的理解写下来吧！
至于题目中提到的有人将CUDA架构和x86的一样理解，这一定是错误的，如果说把GPU的架构和x86架构类比，这还有点道理。
要理解CUDA架构，个人认为应当先理解CUDA，CUDA是compute unit device architecture的简称，它统一了NVIDIA用于通用计算的GPU的编程模式，同时引入的共享存储器也是一大特色，大大提高了速度。有时我们说CUDA是指软件，有时又是指硬件，更多时是指CUDA C语言，因此本文将分三个部分简略的说明。
从软件的角度上说，CUDA软件栈包含两个层次，一个是驱动层的API,这类函数以cu开关，一个是运行层的API，以cuda开头，运行层API是建立在驱动层API之上的，是对驱动层API的封装，这种封装方便了使用，但是同时也牺牲了灵活性。一般而言，我们的开发都是优先使用运行时API。要提到的是：NVIDIA自身也提供了一些建立在这两者之上的库，供大家使用。
从另外一个方面说，CUDA的软件架构也应当包含这一系列的语言支持。当然这只是个人观点，欢迎大家讨论。
从硬件的角度看，CUDA架构应当包含TPC-->SM-->SP结构及它们的关系。各种存储器及其组织，以及硬件对线程的创建、执行、调度等，本人博客将会专门细说这个问题，所以也就不详细说了。
从CUDA C语言的角度看，CUDA的架构包括函数修饰符，变量修饰符，以及一些内置函数等，这方面内容本人的博客将会详细说明，因此，此处就不多说了。
4.CUDA C语言
前一节，我已经说了怎样在ubuntu linux上构建CUDA开发环境了，对一个语言来说，最简单的，也是用得最多的当然是它的语法了，下面我简单的介绍一下CUDA的语法。
CUDA C不是C语言，而是对C语言进行扩展。CUDA对C的扩展主要包括以下四个方面：
函数类型限定符，用来确定函数是在CPU还是在GPU上执行，以及这个函数是从CPU调用还是从GPU调用。
- __device__，__device__表示从GPU上调用，在GPU上执行；
- __global__，__global__表示在CPU上调用，在GPU上执行，也就是所谓的内核(kernel)函数；内核主要用来执行多线程调用。
- __host__，__host__表明在CPU上调用，在CPU上执行，这是默认时的情况，也就是传统的C函数。CUDA支持__host__和__device__的联用，表示同时为主机和设备编译。此时这个函数不能出现多线程语句。
变量类型限定符，用来规定变量存储什么位置上。在传统的CPU程序上，这个任务由编译器承担。在CUDA中，不仅要使用主机端的内存，还要使用设备端的显存和GPU片上的寄存器、共享存储器和缓存。在CUDA存储器模型中，一共抽象出来了8种不同的存储器。复杂的存储器模型使得必须要使用限定符要说明变量的存储位置。
- __device__，__device__表明声明的数据存放在显存中，所有的线程都可以访问，而且主机也可以通过运行时库访问；
- __shared__，__shared__表示数据存放在共享存储器在，只有在所在的块内的线程可以访问，其它块内的线程不能访问；
- __constant__，__constant__表明数据存放在常量存储器中，可以被所有的线程访问，也可以被主机通过运行时库访问；
texture，texture表明其绑定的数据可以被纹理缓存加速存取，其实数据本身的存放位置并没有改变，纹理是来源于图形学的一介概念，CUDA使用它的原因一部分在于支持图形处理，另一方面也可以利用它的一些特殊功能。
如果在GPU上执行的函数内部的变量没有限定符，那表示它存放在寄存器或者本地存储器中，在寄存器中的数据只归线程所有，其它线程不可见。
如果SM的寄存器用完，那么编译器就会将本应放到寄存器中的变量放到本地存储器中。
执行配置运算符<<< >>>，用来传递内核函数的执行参数。执行配置有四个参数，第一个参数声明网格的大小，第二个参数声明块的大小，第三个参数声明动态分配的共享存储器大小，默认为0，最后一个参数声明执行的流，默认为0。
五个内建变量，用于在运行时获得网格和块的尺寸及线程索引等信息
- gridDim, gridDim是一个包含三个元素x,y,z的结构体，分别表示网格在x,y,z三个方向上的尺寸，虽然其有三维，但是目前只能使用二维；
- blockDim, blockDim也是一个包含三个元素x,y,z的结构体，分别表示块在x,y,z三个方向上的尺寸，对应于执行配置中的第一个参数，对应于执行配置的第二个参数；
- blockIdx, blockIdx也是一个包含三个元素x,y,z的结构体，分别表示当前线程所在块在网格中x,y,z三个方向上的索引；
- threadIdx, threadIdx也是一个包含三个元素x,y,z的结构体，分别表示当前线程在其所在块中x,y,z三个方向上的索引；
- warpSize，warpSize表明warp的尺寸，在计算能力为1.0的设备中，这个值是24，在1.0以上的设备中，这个值是32。
其它的还有数学函数，原子函数，纹理读取、绑定函数，内建栅栏，内存fence函数等。一般而言，知道这些就应该能够写出CUDA程序了，当然要写好的话，必须知道很多其它的细节。
6. 一个例子
上一节，已经简单的说了一下CUDA C的基本语法；因而在本节，兄弟决定以一个例子为基础说明CUDA程序的基本组成部分，不过说实话兄弟选择的例子并不太好，这个例子就是采用积分法计算圆周率/π的值。其计算原理是：在[0,1]范围内积分1/(1+x*x)
首先，让我们看一下在CPU上的计算流程，其计算流程如下 很明显，我们可以将for循环分解，使用CUDA处理。
有一个问题就是：for内部对sum变量的更新是互斥的，而CUDA中并没有浮点原子函数，对于这个问题的解决方案是：将for循环内部的两个语句拆开，分成两个内核函数来做运算。内核函数英文名为kernel，就是一个能够在GPU上运算的模块。第一个内核计算计算每个小积分块面积，并将每个block内所有线程对应的积分块面积加起来，存入全局存储器；第二个内核将前一个内核存入的数据加起来。下面是kernel代码： 其中__syncthreads()是CUDA的内置命令，其作用是保证block内的所有线程都已经运行到调用__syncthreads()的位置。
由上面的代码可以看出，使用使用CUDA的主要阻碍在于数据相关性。
一般而言，CUDA程序的基本模式是：
- 分配内存空间和显存空间
- 初始化内存空间
- 将要计算的数据从内存上复制到显存上
- 执行kernel计算
- 将计算后显存上的数据复制到内存上
- 处理复制到内存上的数据
这个程序在我的机器(CPU 2.0GHZ,GPU GTX295)上的加速比超过100，不知道在你们的机器上能够加速多少？
7. CUDA编程模式
CUDA支持大量的线程级并行（Thread Level Parallel），并在硬件中动态地创建、调度和执行这些线程，在CPU中，这些操作是重量级的，但是在CUDA中，这些操作是轻量级的。CUDA编程模型将CPU作为主机（Host），而将GPU做为协处理器（Coprocessor），或者设备（Device），以CPU来控制程序整体的串行逻辑和任务调度，而让GPU来运行一些能够被高度线程化的数据并行部分。即让GPU与CPU协同工作，更确切的说是CPU控制GPU工作。GPU只有在计算高度数据并行任务时才发挥作用。
一般而言，ＣＵＤＡ并行程序包括串行计算部分和并行计算部分，并行计算部分称之为内核（Kernel），内核只是一个在ＧＰＵ上执行的数据并行代码段。理想情况下，串行代码的作用应该只是清理上个内核函数，并启动下一个内核函数，但由于目前的GPU的功能仍然十分有限，串行部分的工作量仍然十分可观
8. CUDA线程层次
GPU线程以网格（grid）的方式组织，而每个网格中又包含若干个线程块，在G80/GT200系列中，每一个线程块最多可包含512/768/1024个线程，Fermi架构中每个线程块支持高达1536个线程。同一线程块中的众多线程拥有相同的指令地址，不仅能够并行执行，而且能够通过共享存储器（Shared memory）和栅栏（barrier）实现块内通信。这样，同一网格内的不同块之间存在不需要通信的粗粒度并行，而一个块内的线程之间又形成了允许通信的细粒度并行。这些就是CUDA的关键特性：线程按照粗粒度的线程块和细粒度的线程两个层次进行组织、在细粒度并行的层次通过共享存储器和栅栏同步实现通信，这就是CUDA的双层线程模型。
在执行时，GPU的任务分配单元（global block scheduler）将网格分配到GPU芯片上。启动CUDA 内核时，需要将网格信息从CPU传输到GPU。任务分配单元根据这些信息将块分配到SM上。任务分配单元使用的是轮询策略：轮询查看SM是否还有足够的资源来执行新的块，如果有则给SM分配一个新的块，如果没有则查看下一个SM。决定能否分配的因素有：每个块使用的共享存储器数量，每个块使用的寄存器数量，以及其它的一些限制条件。任务分配单元在SM的任务分配中保持平衡，但是程序员可以通过更改块内线程数，每个线程使用的寄存器数和共享存储器数来隐式的控制，从而保证SM之间的任务均衡。任务以这种方式划分能够使程序获得了可扩展性：由于每个子问题都能在任意一个SM上运行，CUDA程序在核心数量不同的处理器上都能正常运行，这样就隐藏了硬件差异。
对于程序员来说，他们需要将任务划分为互不相干的粗粒度子问题(最好是易并行计算)，再将每个子问题划分为能够使用线程处理的问题。
同一线程块中的线程开始于相同的指令地址，理论上能够以不同的分支执行。但实际上，在块内的分支因为SM构架的原因被大大限制了。内核函数实质上是以块为单位执行的。同一线程块中的线程需要SM中的共享存储器共享数据，因此它们必须在同一个SM中发射。线程块中的每一个线程被发射到一个SP上。
任务分配单元可以为每个SM分配最多8个块。而SM中的线程调度单元又将分配到的块进行细分，将其中的线程组织成更小的结构，称为线程束（warp）。在CUDA中，warp对程序员来说是透明的，它的大小可能会随着硬件的发展发生变化，在当前版本的CUDA中，每个warp是由32个线程组成的。SM中一条指令的延迟最小为4个指令周期。8个SP采用了发射一次指令，执行4次的流水线结构。所以由32个线程组成的Warp是CUDA程序执行的最小单位，并且同一个warp是严格串行的，因此在warp内是无须同步的。在一个SM中可能同时有来自不同块的warp。当一个块中的warp在进行访存或者同步等高延迟操作时，另一个块可以占用SM中的计算资源。这样，在SM内就实现了简单的乱序执行。不同块之间的执行没有顺序，完全并行。无论是在一次只能处理一个线程块的GPU上,还是在一次能处理数十乃至上百个线程块的GPU上，这一模型都能很好的适用。
目前，某一时刻只能有一个内核函数正在执行，但是在Fermi架构中，这一限制已被解除。如果在一个内核访问数据时，另一个内核能够进行计算，则可以有效的提高设备的利用率。
9. CUDA存储器组织
CUDA的存储器由一系列不同的地址空间组成。其中，shared memory和register位于GPU片内，Texture memory和Constant memory可以由GPU片内缓存加速对片外显存的访问，而Local memory和Device memory位于GPU片外的显存中。
最靠近流处理器的是寄存器文件（register file），每个寄存器文件是32bit。对线程来说，寄存器都是私有的，不允许其它线程染指。由于更靠近流处理器，寄存器具有最快的速度，GT200的每个SM拥有64KB的寄存器文件（Register Files），故一个块内最多可分配16K个寄存器，而G80中每个SM只有32KB，故一个块最多可分配8K个寄存器。最新加入的64bit数据类型（双精度浮点和64位整数型）将占用两个相邻的寄存器单元。CUDA的运行环境能够动态的为线程块分配寄存器，而每个线程块中的线程占用的寄存器大小则是静态分配的，在线程块寿命期间都不会更改，因此一个线程占用的寄存器数目是它在运行时占用的最大数目。如果寄存器被消耗完，数据将被存储在本地存储器（local memory）。对每个线程来说，本地存储器也是私有的，但是本地存储器是显存中的一个分区，速度很慢，而且使用本地存储器过多的话，程序也会终止，因此，编程时要尽量保证不能将数据放到本地存储器中，这可以通过修改块大小，使用共享存储器等方法来解决。
共享存储器是可以被同一块中的所有线程访问的可读写存储器，它的生存期就是块的生命期。在没有冲突的情况下，访问共享存储器几乎与访问寄存器一样快，是实现线程间通信的最好方法。共享存储器可以实现许多不同的功能，如用于保存共用的计数器或者块内的公用结果（例如reduction）。在同一个块内，所有的线程都能够读写共享存储器中的数据，相比于AMD的显卡来说，共享存储器是NVIDIA显卡的一项特色。一般而言，在kernel运行时，要先将数据从全局存储器写入共享存储器；计算完成后要将共享存储器中的结果转存入全局存储器。
Tesla的每个SM拥有16KB共享存储器，用于同一个线程块内的线程间通信。为了使一个half-warp内的线程能够在一个内核周期中并行访问，共享存储器被组织成16个bank，每个bank拥有32bit的宽度，故每个bank可保存256个整形或单精度浮点数，或者说目前的bank组织成了256行16列的矩阵。如果一个half-warp中有一部分线程访问属于同一bank的数据，则会产生bank conflict，降低访存效率，在冲突最严重的情况下，速度会比全局显存还慢，但是如果half-warp的线程访问同一地址的时候，会产生一次广播，其速度反而没有下降。在不发生bank conflict时，访问共享存储器的速度与寄存器相同。在不同的块之间，共享存储器是毫不相关的。
在实现中，GPU要把显存中的数据写到共享存储器中，必须先把数据写到寄存器里，再转移到共享存储器中，在编程时，这是隐式实现的。所以如果没有块内的数据共享千万不能用共享存储器，否则会降低速度。但是如果由于寄存器使用过量，那么我们可以使用共享存储器来当寄存器使用，此时比纯使用寄存器慢一点，但是是值得的。
Tesla能够在共享存储器内进行高速的原子操作。这里的原子操作是指保证每个线程能够独占的访问存储器，即只有当一个线程完成对存储器的某个位置的操作以后，其他线程才能访问这一位置。G80只支持对global memory的原子操作。访问global memory需要很长的访存延迟（长达数百个时钟周期），性能很低。在GT200及以后的GPU上，可以支持对shared memory中的原子操作指令（其中包括CAS指令，并且支持64位）。但是，CUDA并不提供对浮点数的原子操作(只有一个赋值的浮点原子指令)，而在科学计算中，浮点数的使用远比整数要多，而且在Fermi的特性列表中，也没有看到加入浮点原子指令的信息，这不能不说是一个遗憾。
除此以外，多处理器上，还有两种只读的存储器：常数存储器（constant memory）和纹理存储器（texture memory），它们是利用GPU用于图形计算的专用单元实现的。常数存储器空间较小（只有64KB），属于片外存储器，其速度比shared要慢，但是它具有缓存，并且无须考虑冲突问题，主要用来加速对常数的访问。
从物理上说，纹理存储器不是存储器，它只是利用了纹理缓存而已。纹理缓存与CPU的缓存有很大的不同。首先，CPU的缓存往往是一维的，因为大多数的架构中的存储器地址是线性的。当访问一个只有4-8Byte的数据字时，会取出一个缓存单元中所有的64B数据。根据局部性原理，CPU处理的数据往往有很强的时空相关性，因此多取出的相邻的数据极有可能会被用到。CPU处理的数据只有一维，因而其缓存也只是在一个维度上是连续的；GPU需要处理的纹理则是连续的二维图像，因此纹理缓存也必须是在两个维度上连续分布的。典型的存储器控制器会将二维的纹理存储器空间映射为一维。其次，纹理缓存是只读的，也不满足数据一致性。当纹理被修改以后，必须更新整个纹理缓存，而不是纹理缓存中被修改的一小部分。第三，纹理缓存的主要功能是为了节省带宽和功耗，而CPU的缓存则是为了实现较低的延迟。第四，纹理可以实现对数据的特殊处理，比如怎样处理越界数据，自动实现插值等。
最后是全局存储器（global memory），使用的是普通的显存。整个网格中的任意线程都能读写全局存储器的任意位置。目前对Global memory的访问没有缓存，因此显存的性能对GPU至关重要。为了能够高效的访问显存，读取和存储必须对齐，宽度为4Byte。如果没有正确的对齐，读写将被编译器拆分为多次操作，极大的影响效率。此外，多个half-warp的读写操作如果能够满足合并访问（coalesced access），那么多次访存操作会被合并成一次完成，从而提高访问效率。
G80的合并访存条件十分严格。首先，访存的开始地址必须对齐：16x32bit的合并必须对齐到64Byte（即访存起始地址必须是64Byte的整数倍）；16x64bit的合并访存起始必须对齐到128Byte；16x128bit合并访存的起始地址必须对齐到128Byte，但是必须横跨连续的两个128Byte区域。其次，只有当第K个线程访问的就是第K个数据字时，才能实现合并访问，否则half warp中的16个访存指令就会被发射成16次单独的访存。
GT200不仅放宽了合并访问条件，而且还能支持对8bit和16bit数据字的合并访问（分别使用32Byte和64Byte传输）。在一次合并传输的数据中，并不要求线程编号和访问的数据字编号相同。其次，当访问128Byte数据时如果地址没有对齐到128Byte，在G80中会产生16次访存指令发射，而在GT200中只会产生两次合并访存。而且，这两次合并访存并不是两次128Byte的。例如，一次128Byte访存中有32Byte在一个区域中，另外一个区域中有96Byte，那么只会产生一次32Byte合并访存（对有32Byte数据的区域）和一次128Byte（对有96Byte数据的区域）。
除了device端存储器外，还有存在于host端的存储器，即内存。在CUDA中，主机端内存分为两种：Pageable host memory和Page-locked host memory，其中Page-locked host memory保证位在于物理内存中，并且能够通过DMA加速与显卡的通信，提高数据传输速度，但是如果主机的内存不够用的话，会减弱系统的性能，但是一般不会出现这种情况。
10.CUDA执行模式
Tesla架构的构建以一个可伸缩的流多处理器（SM）阵列为中心。当主机CPU上的CUDA程序调用内核网格时，网格的块将被枚举并分发到多处理器上。一个线程块的线程在一个多处理器上并发执行。在线程块终止时，调度单元将决定是否启动新块和启动那一个块。
为了管理运行各种不同程序的数百个线程，多处理器利用了一种称为SIMT（单指令、多线程）的新架构。多处理器会将各线程映射到一个标量处理器核心，各标量线程使用自己的指令地址和寄存器状态独立执行。多处理器SIMT单元以warp为单位来创建、管理、调度和执行线程，构成warp的各个线程在同一个程序地址一起启动，严格串行。
为一个多处理器指定了一个或多个要执行的线程块时，它会将其分成warp块，并由SIMT单元进行调度。将块分割为warp的方法总是相同的，每个warp都包含连续的线程，递增线程索引，第一个warp中包含全局线程过索引0-31。每发出一条指令时，SIMT单元都会选择一个已准备好执行的warp块，并将指令发送到该warp块的活动线程。Warp块每次执行一条通用指令，因此在warp块的全部32个线程执行同一条路径时，可达到最高效率。如果一个warp块的线程通过独立于数据的条件分支而分散，warp块将连续执行所使用的各分支路径，而禁用未在此路径上的线程，完成所有路径时，线程重新汇聚到同一执行路径下，其执行时间为各时间总和。分支仅在warp块内出现，不同的warp块总是独立执行的--无论它们执行的是通用的代码路径还是彼此无关的代码路径。
SIMT架构类似于SIMD（单指令流多数据流）向量组织方法，共同之处是使用单指令来控制多个处理元素。一项主要差别在于SIMD向量组织方法会向软件公开SIMD宽度，而SIMT指令指定单一线程的执行和分支行为。与SIMD向量机不同，SIMT允许程序员为独立、标量线程编写线程级的并行代码，还允许为协同线程编写数据并行代码。为了确保正确性，程序员可忽略SIMT行为，但通过维护很少需要使一个warp块内的线程分支的代码，即可实现显著的性能提升。
另外一个重要不同是SIMD中的向量中的元素相互之间可以自由通信，因为它们存在于相同的地址空间（例如，都在CPU的同一寄存器中），而SIMT中的每个线程的寄存器都是私有的，线程之间只能通过shared memory和同步机制进行通信。
在SIMT编程模型中如果需要控制单个线程的行为，必须使用分支，这会大大的降低效率。例如，如果一个warp中需要进行分支（即warp内的线程执行的指令指针指向不同的位置），性能将急剧的下降。如果一个warp内需要执行N个分支，那么SM就需要把每一个分支的指令发射到每一个SP上，再由SP根据线程的逻辑决定需不需要执行。这是一个串行过程，因此SIMT完成分支的时间是多个分支时间之和。
扫一扫
4018
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
