openCL;全局内存同步;

我需要同步二维数组的计算。数组的一个元素是4x4 int的数据块。每个块依赖于上面的块像blk [Y] [X]应该在blk [Y-1] [X]之后计算。我尝试使用全局地图进行同步。 Map是挥发性__全局int的二维数组，其大小对应于blk [Y] [X]之后的map [Y] [X]的内核设置值。openCL;全局内存同步;

每个内核校验标志上面知道它可能开始喜欢

__kernel 
    void kernel_blk_4x4(
     __global __read_only uchar * __restrict _src, 
     __global __write_only uchar * __restrict _dst, 
     volatile __global int * __restrict map 
        ){ 
int gidY = get_global_id(1); 
int gidX = get_global_id(0); 

// --- check flags before starting 
volatile int kkk = 0; 
volatile __global int * const map_ptr0 = map + (gidY)*31 + gidX; 
volatile __global int * const map_ptr1 = map_ptr0 + 1; 

volatile int val = *map_ptr0; 

while(val == 0) { 
    kkk++; 
    val = *map_ptr0; 
} 
computation here... 

volatile __global int *map_ptr = map + (gidY+1)*31 + gidX; 
*map_ptr = 1; 
}

的第一行地图已经由场1所以理论上它应该工作...
真实的生活更有趣..
其实我会陷入僵局。但是如果我添加“printf（）;”某处代码，之前或之后，而（）一切工作的很好的例子显示地图阵...

是任何想法我做错了吗？

感谢您的帮助！

编辑： 同步已归档:)但出现了另一个问题。答：我改变了方法。每个线程都会扫描地图，然后只需一个块就可以继续进行。请注意，8x8仅用于测试。

__global int *map_ptr = map; 
int val = 0; 
while (1) { 
    for(int y=0; y<8; y++) { 
     for(int x=0; x<8; x++) { 
      val = atomic_cmpxchg(map_ptr+x, 1, 2); 
      if(val == 1) { 
       map_ptr += x; 
       break; 
      } 
     } 
     if(val == 1) break; 
     map_ptr+=stride; 
    } 
    if(val == 1) break; 
    map_ptr = map; 
} 

// do some work 

__global int *map_next = map_ptr+stride; 
atomic_inc(map_next);

该内核在工作组大小为1x1的情况下工作良好。当我尝试改变尺寸时，我会得到成交。

这可能是什么原因？

来源

2014-03-28 user3124812

它根本无法完成。除了每个WG案例有一个简单的WI，你总是会陷入僵局，我不知道printf（）在做什么，但也许是以某种方式序列化代码。但这不是一个有用的用例。改变你的算法到一个完全平行的算法。我建议每行启动一批内核，并在列维上完全并行运行。 BTW发布完整的代码以获得更好的帮助。（什么是gidY/gidX，什么是地图，参数等） – DarkZeros

@DarkZeros，谢谢你的回答。我自己发布整个内核w/o计算（目前我也在我的评论）。不幸的是，我不能按行进行，有简化版本的代码，事实上上面的行还有一些依赖关系。 – user3124812

尽管OpenCL具有关键字volatile，但并不意味着全局内存将跨工作组同步。该规范说，不同的工作组读取和写入相同的全局内存区域是未定义的。所以你可能会看到阅读，或者你可能不会，或者你可能会得到垃圾。 – sharpneli

GPU设备上的OpenCL使用执行（单指令多数据）的SIMD。这意味着组中的所有工作线程将始终执行相同的代码，因为硬件不允许以任何其他方式。

在这个SIMD案例中，不可能将“break”，“for”，“while” etc ......或任何在所有工作中没有遇到相同次数的循环操作线程在一个组中。

由于您的代码明确需要只有1个线程退出循环，因此它将无法工作。它将进入死锁状态，至少在SIMD设备（所有存在的GPU）上。

要做到这一点，唯一的方法是每个工作站只有012个工作站（非常低效），或者使用CPU。您的问题无法以其他方式回答，OpenCL无法做到您需要的东西。没有任何技巧或代码可以解决您的问题。你必须改变你使用的算法，使它真正平行。

注意：使用1x1 WG或printf（）将执行顺序化，这可以解决您的问题。但在这种情况下，根本没有使用GPU的意义。

来源

2014-04-03 17:15:09 DarkZeros

感谢您的明确回复。 – user3124812

openCL;全局内存同步;

回答

相关问题