条件内的OpenCL屏障使用

Question

按规格:

如果障碍是在一个条件语句,那么所有工作项必须输入条件如果任何工作项进入条件语句并执行障碍。

如果barrier在循环中，所有的工作项必须在循环的每次迭代中执行barrier，然后才允许任何工作项在barrier之外继续执行。

在我的理解中，这意味着在任何内核中:

if(0 == (get_local_id(0)%2)){
//block a
work_group_barrier(CLK_GLOBAL_MEM_FENCE);
//block a part 2
}else{
//block b
work_group_barrier(CLK_GLOBAL_MEM_FENCE);
//block b part 2
}
//common operations

如果一个工人到达//block a，其他工人也需要到达它。

通过这种逻辑，不可能正确地同步每个奇数本地worker和每个偶数本地worker(块a和b)同时运行。

这种理解正确吗?

在这种情况下，什么是好的同步策略?

• 其中每个其他工人需要做不同的逻辑，但通过//block a part 2和//block b part 2，一个工作组内的工人被同步。
• 的要求usecase有更多的阶段,比2,我想让每个阶段同步。

这样的逻辑是一个可接受的解决方案吗?

__local number_finished = 0;
if(0 == (get_local_id(0)%2)){
//block a
atomic_add(&number_finished, 1);
while(number_finished < get_local_size(0));
//block a part 2
}else{
//block b
atomic_add(&number_finished, 1);
while(number_finished < get_local_size(0));
//block b part 2
}
work_group_barrier(CLK_GLOBAL_MEM_FENCE);
//common operations

Answer 1

通过这种逻辑，不可能正确同步每个奇数本地worker和每个偶数本地worker(块a和b)同时运行。

这种理解正确吗?

不可能。当使用work_group_barrier时，您必须确保工作组中的所有工作项都达到它。如果你的代码中有路径没有到达屏障，它可能会导致未定义的行为。

对于这种情况，什么是好的同步策略?

通常，屏障在任何条件节/循环之外使用:

if (cond)
{
// all work items perform work that needs to be synchronized later 
// such as local memory writes
}
barrier(CLK_LOCAL_MEM_FENCE); // note the CLK_LOCAL_MEM_FENCE flag
// now every work item can read the data the other work items wrote before the barrier. 
for (...)
{
}
barrier(CLK_LOCAL_MEM_FENCE);

这样的逻辑是一个可接受的解决方案吗?

这可能会起作用，但是在条件部分之外设置一个屏障会比繁忙的等待更有效。

if(0 == (get_local_id(0)%2)){
//block a
}else{
// block b
}
barrier(CLK_LOCAL_MEM_FENCE); // ensures number_finished == get_local_size(0)