指令级并行和线程级并行如何在GPU上运行？

Question

假设我正在尝试对数组大小n进行简单的缩减,比如保持在一个工作单元内...说添加所有元素.一般策略似乎是在每个GPU上生成许多工作项,这会减少树中的项目.天真地看起来似乎采取了log n步骤,但它并不像第一波线程所有这些线程一次性进行,是吗？它们被安排在经线中.

for(int offset = get_local_size(0) / 2;
      offset > 0;
      offset >>= 1) {
     if (local_index < offset) {
       float other = scratch[local_index + offset];
       float mine = scratch[local_index];
       scratch[local_index] = (mine < other) ? mine : other;
     }
     barrier(CLK_LOCAL_MEM_FENCE);
   }

因此,并行添加了32个项目,然后该线程在屏障处等待.另外32人去,我们在障碍物等待.另外32个,我们在屏障处等待,直到所有线程都完成了必要的n/2次添加以进入树的最顶层,然后我们绕过循环.凉.

这看起来不错,但也许很复杂？我理解指令级并行是一个大问题,所以为什么不产生一个线程并做类似的事情

while(i<array size){
    scratch[0] += scratch[i+16]
    scratch[1] += scratch[i+17]
    scratch[2] += scratch[i+17]
    ...
    i+=16
}
...
int accum = 0;
accum += scratch[0]
accum += scratch[1]
accum += scratch[2]
accum += scratch[3]
...

这样所有的添加都发生在经线内.现在你有一个线程可以保持gpu的忙碌程度.

现在假设指令级并行性不是真的.如下所示,工作大小设置为32(warp数).

for(int i = get_local_id(0);i += 32;i++){
    scratch[get_local_id(0)] += scratch[i+get_local_id(0)]
}

然后将前32个项目一起添加.我想这32个线程会一次又一次地开火.

如果你不放弃放弃OpenCL的普遍性,那么当你知道每个循环会有多少加法时,为什么还要在树中减少呢？

Answer 1

一个线程无法使GPU保持忙碌状态.这与说一个线程可以保持8核CPU忙碌大致相同.

为了最大限度地利用计算资源以及可用内存带宽,必须使用整个机器(即可以执行线程的所有可用资源).

对于大多数较新的GPU,通过使线程代码按顺序具有多个独立指令,您当然可以通过指令级并行来提高性能.但是你不能把所有这些都扔进一个线程并期望它能提供良好的性能.

按顺序有2条指令时,如下所示:

scratch[0] += scratch[i+16]
scratch[1] += scratch[i+17]

这对ILP有利,因为这两个操作完全相互独立.但是,由于GPU发出内存事务的方式,第一行代码将参与特定的内存事务,第二行代码必然会参与不同的内存事务.

当我们有一个warp一起工作时,一行代码如下:

float other = scratch[local_index + offset];

将导致warp的所有成员生成请求,但这些请求将全部合并为单个或两个内存事务.这就是如何实现全带宽利用率.

虽然大多数现代GPU具有缓存,并且缓存将倾向于弥合这两种方法之间的差距,但它们绝不会弥补所有warp成员发出组合请求与单个warp之间的事务中的巨大差异.成员按顺序发出一组请求.

您可能想要阅读GPU内存合并.由于您的问题似乎是以OpenCL为中心的,因此您可能对此文档感兴趣.