我正在使用 OpenCL。我感兴趣的是在下面的示例中工作项将如何执行。
我的一维范围为 10000,工作组大小为 512。内核如下:
__kernel void
doStreaming() {
unsigned int id = get_global_id(0);
if (!isExecutable(id))
return;
/* do some work */
}
这里检查是否需要处理具有以下 id 的元素。
假设执行从第一个大小为 512 的工作组开始,其中 20 个被 isExecutable
拒绝。 GPU 是否会继续执行其他 20 个元素而不等待前 492 个元素?
不涉及任何障碍或其他同步技术。
最佳答案
当某些工作项的分支远离通常的/*做一些工作*/时,它们可以通过从下一个wavefront(amd)或下一个warp(nvidia)获取指令来利用管道占用优势,因为当前的warp/wavefront工作项正忙于做其他事情。但这会导致内存访问序列化并清除工作组的访问顺序,从而降低性能。
避免发散的扭曲/波前:如果你在循环中执行 if 语句,那真的很糟糕,所以你最好找到另一种方法。
如果工作组中的每个工作项都有相同的分支,那么就可以了。
如果每个工作项在数百次计算中只执行很少的分支,那就可以了。
尝试为所有工作项(极其并行的数据/算法)生成相同的条件,以利用 GPU 拥有的能力。
我知道摆脱最简单的分支与计算情况的最佳方法是使用全局是/否数组。 0=是,1=否:始终计算,然后将结果与工作项的是/否元素相乘。通常,为每个核心添加 1 字节元素内存访问比为每个核心执行一个分支要好得多。实际上,在添加这个 1 字节后,将对象长度设置为 2 的幂可能会更好。
关于opencl - 工作项执行顺序,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/17517138/