Блокировка не работает с OpenCL

Question

Я застрял с проблемой в моем коде OpenCL, где я пытаюсь синхронизироваться внутри ядра:

__kernel void pdiffs (__global const long2 *inData, __global const long2 *inData2, __global long2 *outData) {
  long2 diffSum = 0;
  uint idx0 = get_local_size(0)*get_group_id(0);
  for (uint idx=idx0; idx<idx0+get_local_size(0); idx += 1) {
    diffSum += inData[idx] -  inData2[idx];
  outData[get_group_id(0)] = diffSum;
  printf("%d %d %d %d/%d\n", get_group_id(0), get_num_groups(0), get_local_size(0), diffSum.x, diffSum.y);
  barrier(CLK_GLOBAL_MEM_FENCE|CLK_LOCAL_MEM_FENCE);
  if (get_group_id(0) == 0) {
    for (size_t i = 1; i < get_num_groups(0); i++){
      outData[0] += outData[i];
      printf("v(%d): %d/%d\n", i, outData[i].x, outData[i].y);
    }
  }
}

(Я знаю, что этот кусок кода просто плохо...)

Я просто подумал, что барьер будет синхронизировать отдельные группы, поэтому значения в outData определены. Но моя трассировка показывает, что некоторые различия не рассчитываются и содержат ноль (я настраиваю свои данные, чтобы все возвращали значение 1, но некоторые отображают как 0). Кроме того, имеет значение, есть ли у меня printf заявления или нет. Без printf еще больше различий кажется неправильным.