Блог

 __kernel void Test(__global uchar* A, __global uchar* B, int W)
{
    int i = get_global_id(0);
    uchar c = 0;
    for (int di=-1; di<2;   di)
        //  W%W is too loop the indeces around the array
        c  = A[(i di W)%W];
    // Make sure all the values are read, before nullifying A
    // The values I pass to work_group_barrier are overkill - having one of them should've been enough
    work_group_barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
    // Without this nullification all the values of B are 3, as expected
    A[i] = 0;
    B[i] = c;
}
 

A и B являются буферами размера W .
Перед выполнением A заполняется всеми байтами 1, в то время B как заполняется всеми нулями.
Я хочу прочитать A локальную переменную во всех потоках, прежде чем ее переопределят, в данном случае, просто нулями. Разве не для этого все это work_group_barrier ?
Но когда W переваливает за ~5000 — я начинаю видеть байты 2 B ввода после некоторых выполнений ядра. Когда W 10000000 — я вижу сотни 2 , каждый раз разное количество.
Также, что интересно, если я помещу B[i] = c; перед барьером — только последнее значение B is 2 , все остальное будет 3 так, как ожидалось. Но размещение его раньше A[i] = 0 не имеет никакого эффекта.

Для этого теста на стороне процессора я создавал новую cl_command_queue для каждой новой команды и ждал события последней команды, прежде чем продолжить.
Я также проверил, что перед запуском буферы ядра находятся в том состоянии, в котором я их ожидаю.
Поэтому код на стороне процессора не отвечает за то, что я вижу.

Это упрощенный пример, я на самом деле создавал Игру Жизни. Там следующее состояние текущего значения определяется по состоянию соседних значений на предыдущем шаге.
Поэтому мне нужно прочитать и сохранить предыдущие значения, а затем записать новые. Использование настройки двойного буфера-это вариант, и именно так я делал это раньше, но потом я вспомнил об work_group_barrier этом и решил очистить код, но обнаружил эту проблему.

Я бы подумал, что проблема в том, что потоки разделены на группы, и каждая группа выполняется только после завершения предыдущей. Но тогда я ожидал бы иметь значения 2 по краям таких групп, однако, когда W=1000000 индексы 2 » s:

 13499 13750 14249 21249 262249 306250 315999 365249 377750 542250 617749 699499 717749 804249 846000 868499 898249 901500 907750 910250 934249 967500 987250 987999 999999
17749 21250 47499 48000 48249 82499 146250 202749 213499 270500 339000 382749 402999 524750 539249 600499 671749 690250 739500 800499 826999 899249 931249 961499 999999
12000 53750 58750 126249 138750 282250 299500 408999 449750 453749 539249 576749 581999 609000 621000 682249 715500 724250 736500 830249 889999 896999 900999 966999 972499 988749 990499 998000 999250 999999
 

(каждая новая строка-это новое выполнение моей программы, так что это 3 отдельных примера)
Я не вижу ничего единообразного в том, где находятся сломанные значения, так что, вероятно, проблема не в этом.

У меня есть NVidia GeForce GTX 750 Ti , если это может помочь.
Может быть, это ошибка NVidia или я просто неправильно понимаю, как работают барьеры?