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 拥有了两个层次的并行：

1) 线程级并行实现的细粒度数据并行;

2) 任务级并行实现的粗粒度并行。

2. linux下 CUDA 开发环境构建

         前一节已经简单了说了一下 CUDA，为了能够使用 CUDA 开放，这一节将说明怎样构建 CUDA 开发环境。

         本节讲解在 ubuntu9.04 操作系统和 gcc 前端的基础上安装 CUDA 开发环境。

   

1) 要保证自己机器上的 gcc 能够使用，因为 ubuntu 缺少 gcc 的一些包和 g++，所以这些得自己安装。安装命令: sudo apt-get install g++，待其完成后，弄个 C 代码试试看;  当然你得保证你的显卡支持 CUDA。

   

2) 到 nVidia 官方网站 (http://www.nvidia.cn/object/cuda_get_cn.html) 上下载对应操作系统的驱动 (driver) 和工具包 (toolkit)。

3) 转换到控制台，命令为 Ctrl+Alt+F1/F2/F3/F4，关掉 gdm，命令为：sudo /etc/init.d/gdm stop，要确定已经关闭，否则在安装时会提示你有ｘserver 程序在运行。

   

4) 进入 driver 和 toolkit 目录，执行安装命令，为了方便，请一定按照默认安装。

   

5) 打开个人目录下的 .bashrc 文件或者 /etc/profile 文件，在其中加入命令：

     export  PATH=${PATH}:/usr/local/cuda/bin/   

     export  LD_LIBRARY_PATH=${LD_LIBRARY_PATH}:/usr/local/cuda/lib  

保存后执行 source .bashrc 或者 /etc/profile，依据你添加 PATH 和 LD_LIBRARY_PATH 时修改了哪个文件确定。

6) 执行 nvcc 命令，看看，如果提示你没有输入文件，就安装完成了。

   

         如果你要安装 SDK 的话，还得安装一些包，依据 make 时的提示，google 和 新力得应该可以搞定一切，现在你可以享受CUDA 了！

3. CUDA 架构

         本来一直都想把自己对  CUDA 架构的理解写出来，但是一方面自己是个懒人，不太愿意动笔，另一方面又感觉不太好组织语言，今天终于借着这个第四期有奖讨论的机会，把我对 CUDA 架构的理解写下来吧！

    

         至于题目中提到的有人将 CUDA 架构和 x86 的一样理解，这一定是错误的，如果说把 GPU 的架构和 x86 架构类比，这还有点道理。

        

         要理解 CUDA 架构，个人认为应当先理解 CUDA，CUDA 是 compute unified 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 的扩展主要包括以下四个方面：

1)  函数类型限定符，用来确定函数是在 CPU 还是在 GPU 上执行，以及这个函数是从 CPU 调用还是从 GPU 调用。

(a) __device__，表示从 GPU 上调用，在 GPU上 执行；

(b) __global__，表示在 CPU 上调用，在 GPU 上执行，也就是所谓的内核 (kernel) 函数；内核主要用来执行多线程调用。

(c) __host__，表明在 CPU 上调用，在 CPU 上执行，这是默认时的情况，也就是传统的 C 函数。CUDA 支持 __host__ 和__device__ 的联用，表示同时为主机和设备编译。此时这个函数不能出现多线程语句。

2) 变量类型限定符，用来规定变量存储什么位置上。在传统的 CPU 程序上，这个任务由编译器承担。在 CUDA 中，不仅要使用主机端的内存，还要使用设备端的显存和 GPU 片上的寄存器、共享存储器和缓存。在 CUDA 存储器模型中，一共抽象出来了 8 种不同的存储器。复杂的存储器模型使得必须要使用限定符要说明变量的存储位置。

(a) __device__，表明声明的数据存放在显存中，所有的线程都可以访问，而且主机也可以通过运行时库访问；

(b) __shared__，表示数据存放在共享存储器中，只有在所在的块内的线程可以访问，其它块内的线程不能访问；

(c) __constant__，表明数据存放在常量存储器中，可以被所有的线程访问，也可以被主机通过运行时库访问；

(d) texture，表明其绑定的数据可以被纹理缓存加速存取，其实数据本身的存放位置并没有改变，纹理是来源于图形学的一介概念，CUDA 使用它的原因一部分在于支持图形处理，另一方面也可以利用它的一些特殊功能。

         如果在 GPU 上执行的函数内部的变量没有限定符，那表示它存放在寄存器或者本地存储器中，在寄存器中的数据只归线程所有，其它线程不可见。

         如果 SM 的寄存器用完，那么编译器就会将本应放到寄存器中的变量放到本地存储器中。

3) 执行配置运算符 <<< >>>，用来传递内核函数的执行参数。

         执行配置有四个参数：

第一个参数声明网格的大小；

第二个参数声明块的大小；

第三个参数声明动态分配的共享存储器大小，默认为 0；

最后一个参数声明执行的流，默认为 0。

4) 五个内建变量，用于在运行时获得网格和块的尺寸及线程索引等信息

(a) gridDim, 一个包含三个元素 x, y, z 的结构体，表示网格在 x, y, z 三个方向上的尺寸，虽有三维，但目前只能使用二维；

(b) blockDim, 也是一个包含三个元素  x, y, z 的结构体，分别表示块在  x, y, z 三个方向上的尺寸；

(c) blockIdx, 也是一个包含三个元素  x, y, z 的结构体，分别表示当前线程所在块在网格中  x, y, z 三个方向上的索引；

(d) threadIdx, 也是一个包含三个元素  x, y, z 的结构体，分别表示当前线程在其所在块中  x, y, z 三个方向上的索引；

(e) warpSize，表明 warp 的尺寸，在计算能力为 1.0 的设备中，这个值是 24，在 1.0 以上的设备中，这个值是 32。

        其它的还有数学函数 、原子函数 、纹理读取、绑定函数 、内建栅栏 、内存 fence 函数等。一般而言，知道这些就应该能够写出 CUDA 程序了，当然要写好的话，必须知道很多其它的细节。

5. 一个例子

         上一节，已经简单的说了一下 CUDA C 的基本语法；因而在本节，兄弟决定以一个例子为基础说明 CUDA 程序的基本组成部分，不过说实话兄弟选择的例子并不太好，这个例子就是采用积分法计算圆周率 π 的值。其计算原理是：在 [0,1] 范围内积分 1/(1+x*x)

         首先，让我们看一下在 CPU 上的计算流程，其计算流程如下

 /* 串行计算 PI 的程序，基本思想为：将积分区间均分为 num 小块，将每小块的面积加起来。*/ 

   float cpuPI(int num)

{

      float sum = 0.0f;

      float temp;

      for (int i = 0; i < num; i ++ )

     {

          temp = (i + 0.5f) / num;

          sum += 4 / (1 + temp*temp);

      }

     return sum / num;

 }

复制代码很明显，我们可以将 for 循环分解，使用 CUDA 处理。

         有一个问题就是：for 内部对 sum 变量的更新是互斥的，而 CUDA 中并没有浮点原子函数，对于这个问题的解决方案是：将 for 循环内部的两个语句拆开，分成两个内核函数来做运算。内核函数英文名为 kernel，就是一个能够在 GPU 上运算的模块。第一个内核计算计算每个小积分块面积，并将每个 block 内所有线程对应的积分块面积加起来，存入全局存储器；第二个内核将前一个内核存入的数据加起来。下面是 kernel 代码： 

/*

在 GPU 上计算 PI 的程序，要求块数和块内线程数都是 2 的幂

前一部分为计算 block 内归约，最后大小为块数

后一部分为单个 block 归约，最后存储到 *pi 中。

*/

__global__ void reducePI1( float *d_sum, int num )

{

   int   id  = blockIdx.x*blockDim.x + threadIdx.x; //线程索引

   int   gid = id;

   float temp;

   extern float __shared__ s_pi[]; // 动态分配，长度为 block 线程数

   s_pi[threadIdx.x] = 0.0f;

   while ( gid < num )

   {

      temp = (gid + 0.5f) / num;  // 当前 x 值

      s_pi[threadIdx.x] += 4.0f / (1 + temp*temp);

      gid += blockDim.x*gridDim.x;

   }

   

   for (int i = (blockDim.x >> 1); i > 0; i >>= 1)

   {

      if (threadIdx.x < i)

      {

         s_pi[threadIdx.x] += s_pi[threadIdx.x + i];

      }

      __syncthreads();

   }

   if (threadIdx.x == 0)

   {

      d_sum[blockIdx.x] = s_pi[0];

   }

}

__global__ void reducePI2( float *d_sum, int num, float *d_pi )

{

   int id=threadIdx.x;

   extern float __shared__ s_sum[];

   s_sum[id]=d_sum[id];

   __syncthreads();

   

   for (int i = (blockDim.x >> 1); i > 0; i >> = 1)

   {

      if ( id < i )

      {

         s_sum[id] += s_sum[id + i];

      }

      _syncthreads();

   }

   

   //printf("%d,%f\n",id,s_sum[id]);

   if (id == 0)

   {

      *d_pi=s_sum[0]/num;

      //printf("%d,%f\n",id,*pi);

   }

}  

复制代码其中__syncthreads() 是 CUDA 的内置命令，其作用是保证 block 内的所有线程都已经运行到调用 __syncthreads() 的位置。

   

        由上面的代码可以看出，使用使用 CUDA 的主要阻碍在于数据相关性。

   

         一般而言，CUDA 程序的基本模式是： 

* 分配内存空间和显存空间 

* 初始化内存空间 

* 将要计算的数据从内存上复制到显存上 

* 执行 kernel 计算 

* 将计算后显存上的数据复制到内存上 

* 处理复制到内存上的数据

         这个程序在我的机器 (CPU 2.0GHZ,GPU GTX295) 上的加速比超过 100，不知道在你们的机器上能够加速多少？

6. CUDA 编程模式

        CUDA 支持大量的线程级并行（Thread Level Parallel），并在硬件中动态地创建、调度和执行这些线程，在 CPU 中，这些操作是重量级的，但是在 CUDA 中，这些操作是轻量级的。CUDA 编程模型将 CPU 作为 主机（Host），而将 GPU 做为 协处理器（Coprocessor）或者 设备（Device）， 以CPU来控制程序整体的串行逻辑和任务调度，而让GPU来运行一些能够被高度线程化的数据并行部分。即让 GPU 与 CPU 协同工作，更确切的说是 CPU 控制 GPU 工作。GPU 只有在计算高度数据并行任务时才发挥作用。

         一般而言， CUDA 并行程序包括 串行计算部分和 并行计算部分，并行计算部分称之为内核（Kernel），内核只是一个在GPU 上执行的数据并行代码段。理想情况下，串行代码的作用应该只是清理上个内核函数，并启动下一个内核函数，但由于目前的 GPU 的功能仍然十分有限，串行部分的工作量仍然十分可观。

7. CUDA 线程层次

         GPU 线程以 网格（grid）的方式组织，而每个网格中又包含若干个 线程块 （thread block），在 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 架构中，这一限制已被解除。如果在一个内核访问数据时，另一个内核能够进行计算，则可以有效的提高设备的利用率。

8. CUDA 存储器组织

         CUDA 的存储器由一系列不同的地址空间组成。其中， shared memory 和  register 位于 GPU 片内， Texture memory 和  Constant memory 可以由 GPU 片内缓存加速对片外显存的访问，而 Local memory 和  Device memory 位于 GPU 片外的显存中。

         最靠近流处理器 （SP）的是寄存器文件（register file），每个寄存器文件是 32 bit。对线程来说，寄存器都是私有的，不允许其它线程染指。由于更靠近流处理器，寄存器具有最快的速度，GT200 的每个 SM 拥有 64KB 的寄存器文件，故一个块内最多可分配 16K 个寄存器，而 G80 中每个 SM 只有 32KB，故一个块最多可分配 8K 个寄存器。最新加入的 64 bit 数据类型（双精度浮点和 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 加速与显卡的通信，提高数据传输速度，但是如果主机的内存不够用的话，会减弱系统的性能，但是一般不会出现这种情况。

9. 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完成分支的时间是多个分支时间之和。