CUDA线程模型二

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

                                               CUDA学习笔记(二)——CUDA线程模型

 

 

在启动内核时指定:

       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] }

                  CUDA学习笔记(二)——CUDA线程模型

 

 

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

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

 

      CUDA学习笔记(二)——CUDA线程模型 这样的话,同样的程序,可以以多种不同的方式执行。每种执行方式可能有不同的开销,功耗和性能。例如,一个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

阅读更多
个人分类: CUDA编程资料
上一篇CUDA编程模型一
下一篇CUDA内存三
想对作者说点什么? 我来说一句

没有更多推荐了,返回首页

关闭
关闭
关闭