|
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间并行。
硬件映射 计算单元 - 计算核心：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 中的资源限制。
线程结构（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中所有线程的数据都必须交由同一处理核心进行处理。
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函数中调用。
|
