转自: http://blog.csdn.net/augusdi/article/details/12529247
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 中放入多个线程块是为了隐藏延迟(latency),更好地利用执行单元的资源。当一个 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条指令的第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 kernelVecAdd<<<1, N>>>(A, B, C);}
- 必须先为Kernel中用到的数组或变量分配好足够的空间,再调用kernel函数。否则,在GPU计算时会发生错误。
- 在设备端运行的线程之间是并行执行的,其中的每个线程按指令的顺序串行执行一次kernel函数。每一个线程有自己的block ID和thread ID用于与其他线程相区分。block ID和thread ID只能在kernel中通过内建变量访问。内建变量是由设备中的专用寄存器提供的,是只读的,且只能在GPU端的kernel函数中调用。