CUDA知识点

线程块坐标：blockIdx
线程坐标：threadIdx
在CUDA中，kernel函数是以线程块为单位执行的
对于一维的block，线程的threadID就是threadldx.x;对于大小为（Dx,Dy）二维的block,线程的threadID就是threadIdx.x+threadIdx.y*Dx;每个block之间没有同步与通信机制,其执行顺序也互不影响,因此block之间为互相独立的.虽然block之间不能通信,但属于同一线程块的线程不仅能够并行执行,而且能够通过共享存储器和栅栏同步来实现block内的线程间通信.。这样就可以实现在同一网格内不同block间的不需要通信的粗粒度并行和在同一block内线程允许通信的细粒度并行。这就是CUDA的关键特性:线程按两个层次进行织、在较低层次通过共享存储器和栅栏同步实现通信.
GPU的计算单元为若干个流多处理器SM(Stream Multiprocessor),相当于GPU的计算核心,每个SM中包含8个标量流处理器SP(Stream Processor),同一SM中的8个sP共用同一套取指和发射单元,也共用一块共享存储器.,GPU根据网格信息将block分配到每一个SM上,而block中的每一个thread则被发射至一个SP上执行,因此一个block必须在一个SM中,但一个SM中可以有多个线程块等待执行。每个SM中最多等待执行的活动线程块数为8个,而在实际运行中,block会被分为更小的线程束(warp)来执行,线程束大小因硬件的计算能力版本不同而不同。在CUDA中,warp是对程序员透明的,在Tesla架构的GPU中,一个线程束由连续的32个线程组成,在硬件实际运行中,warp才是真正的执行单位。
一个线程块block由数个warp组成,执行宽度可以在1到512个线程间变化。
每个线程拥有一个私有存储器寄存器(register)和局部存储器(10cal memory),每个线程块拥有一个共享存储器(shared memory),该共享存储器对于线程块内的所有线程都是可见的,Grid中所有的线程都可以访问相同的全局存储器(global memory)。此外,还有两个只读的存储器空间:常数存储器(constant memory)和纹理存储器(texture memory),可以被所有的线程所访问,这两个存储器分别具有不同的用途。全局存储器、常数存储器和纹理存储器空间是持久的,即它们的值在一个Kernel函数执行完成后将继续存在,可以被同一个应用程序其他的Kernel函数调用。
在CUDA中一个最简单的来实现卷积的方法是将图像中的一块数据载入到共享存储器阵列中,对这一块数据中的像素按照滤波核的尺寸与模板逐点计算相乘,然后将计算和写入到设备内存中的输出图像中。每一个线程可计算得到单个的输出像素,可通过下图来清晰的了解这一过程:
（放入到共享存储器中的图像数据，应该与卷积核大小一样）
我们可以通过降低每个块中线程数目的总数并且在每个进程中载入多个像素至共享内存来减少空闲线程。如果使用与待处理图像块宽度相同的垂直列线程,那么对于共享内存中的48×48个像素区域,只需要16X48个线程,可分三列读取数据。
在卷积滤波计算中,坊线程的活跃度自然大于埔线程的活跃度,但因载入时带宽的限制,并不能完全做到上述过程,因此可以将线程块进一步降低至16×16个线程并将图像块分为9个正方形像素块,这样就可以保证计算时所有线程都是活跃的。需要注意的是,每块中的线程数目必须是warp尺寸的整数倍(G80体系下GPU为32个线程)才可得到最佳效率。如果边缘像素并不与线程块一样宽,那么在载入时必须终止一些线程。

每个网格（Grid）可以最多创建65535个线程块，每个线程块（Block）一般最多可以创建512个并行线程，

blockIdx代表线程块在网格中的索引值，blockDim代表线程块的尺寸大小，另外还有gridDim代表网格的尺寸大小。

blockDim.x 表示block在x轴方向线程的数量， blockIdx.x表示block在gridx方向的block号