CUDA线程模型二

一个grid中的所有线程执行相同的内核函数，通过坐标进行区分。这些线程有两级的坐标，blockId和threadId，由CUDA runtime system指定。grimDim.x标识block在x维度上的数目，gridDim.y标识block在y维度上的数目。例如,

                                               

 

 

在启动内核时指定：

       dim3 dimBlock(4,2,2);

       dim3 dimGrid(2,2,1);

       KernelFunction<<<dimGrid, dimBlock>>>(…);

       其中，grid是二维的，所以，最后一个参数一般设置为1。

注意：一个block最多能有512个线程。

如果kernel function没有用到blockId，那么所有的线程都属于同一个block。

 

例如：

P=M* N，那么：Pd[Row][Col]=Sum{ Md[Row][k]*Nd[k][Col] }

                  

 

 

同一个block中的线程通过syncthreads()进行同步。当内核函数调用syncthreads()时，block中的所有线程都在调用位置停留，直到block中的其他线程达到这个位置。Barrier synchronization是常用的协调线程并行行为的方法。

       为了避免同一个block中的线程在同步的时候，会有长时间的等待。往往把相同的资源分配给一个block中的线程。不同的block可以以任意顺序执行，因为不同的block之间不存在同步问题。如图：

 

       这样的话，同样的程序，可以以多种不同的方式执行。每种执行方式可能有不同的开销，功耗和性能。例如，一个mobile processor可以以很低的功耗慢速执行，一个desktop processor可以以高一些的功耗快速执行。程序相同，但是这种改变是透明的。

线程分配

       一旦启动一个kernel，CUDA run-time系统就产生对应的grid。Grid中的线程分配到block上。GeForce 8800GTX有16个Streaming Multiprocessor，每个SM最多可以分配8个block，所以最多有128个block同时分配给Streaming Multiprocessor。系统维护一个block的list，当前面的block执行完成后，分配新的block给SM。

       另一个资源限制是能够同时调度的线程数目。因为保留thread, block的ID并跟踪他们的执行状态需要硬件资源。在GeForce 8800GTX上，每个SM最多能分配768个线程。所以，整个GeForce 8800GTZX最多同时有12288（16*768）个线程。

线程调度

       在GeForce 8800