学步园

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编程模型将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间并行。

• 内核函数必须通过__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函数中调用。

• 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 中的资源限制。