Ошибка заказа размеров OpenCL NDRange на nVidia?

Question

Я знаю, что OpenCL в настоящее время довольно проста, особенно реализация NVIDIA CUDA. Тем не менее, я думаю, что я нашел существенную ошибку в Nvidia, и я хотел бы увидеть, замечает ли кто-то еще то же самое. Использование Linux Версия платформы OpenCL 1.2 CUDA 10.1.0 с привязками C ++ У меня возникли всевозможные проблемы с порядком NDRange, и, наконец, у меня есть простое ядро, которое может окончательно воспроизвести проблему:

void kernel test()
{
    printf("G0:%d   G1:%d   G2:%d   L0:%d   L1:%d   L2:%d\n", 
    get_global_id(0),
    get_global_id(1),
    get_global_id(2),
    get_local_id(0),
    get_local_id(1),
    get_local_id(2));
}

Если я ставлю это ядро ​​в 3 измерения: глобальное (4,3,2) и локальное (1,1,1):

queue.enqueueNDRangeKernel(kernel, cl::NullRange, 
                cl::NDRange(4, 3, 2), 
                cl::NDRange(1, 1, 1), 
                NULL, events);

, то оно случайным образом корректно выводит следующее на AMD / Intel(случайный вывод отсортирован для ясности):

G0:0   G1:0   G2:0   L0:0   L1:0   L2:0
G0:0   G1:0   G2:1   L0:0   L1:0   L2:0
G0:0   G1:1   G2:0   L0:0   L1:0   L2:0
G0:0   G1:1   G2:1   L0:0   L1:0   L2:0
G0:0   G1:2   G2:0   L0:0   L1:0   L2:0
G0:0   G1:2   G2:1   L0:0   L1:0   L2:0
G0:1   G1:0   G2:0   L0:0   L1:0   L2:0
G0:1   G1:0   G2:1   L0:0   L1:0   L2:0
G0:1   G1:1   G2:0   L0:0   L1:0   L2:0
G0:1   G1:1   G2:1   L0:0   L1:0   L2:0
G0:1   G1:2   G2:0   L0:0   L1:0   L2:0
G0:1   G1:2   G2:1   L0:0   L1:0   L2:0
G0:2   G1:0   G2:0   L0:0   L1:0   L2:0
G0:2   G1:0   G2:1   L0:0   L1:0   L2:0
G0:2   G1:1   G2:0   L0:0   L1:0   L2:0
G0:2   G1:1   G2:1   L0:0   L1:0   L2:0
G0:2   G1:2   G2:0   L0:0   L1:0   L2:0
G0:2   G1:2   G2:1   L0:0   L1:0   L2:0
G0:3   G1:0   G2:0   L0:0   L1:0   L2:0
G0:3   G1:0   G2:1   L0:0   L1:0   L2:0
G0:3   G1:1   G2:0   L0:0   L1:0   L2:0
G0:3   G1:1   G2:1   L0:0   L1:0   L2:0
G0:3   G1:2   G2:0   L0:0   L1:0   L2:0
G0:3   G1:2   G2:1   L0:0   L1:0   L2:0

Это соответствует спецификации. Но если я планирую точно такое же ядро ​​с теми же измерениями, используя NVidia I, получим следующий вывод:

G0:0   G1:0   G2:0   L0:0   L1:0   L2:0
G0:0   G1:0   G2:0   L0:0   L1:1   L2:0
G0:0   G1:0   G2:1   L0:0   L1:0   L2:0
G0:0   G1:0   G2:1   L0:0   L1:1   L2:0
G0:0   G1:0   G2:2   L0:0   L1:0   L2:0
G0:0   G1:0   G2:2   L0:0   L1:1   L2:0
G0:1   G1:0   G2:0   L0:0   L1:0   L2:0
G0:1   G1:0   G2:0   L0:0   L1:1   L2:0
G0:1   G1:0   G2:1   L0:0   L1:0   L2:0
G0:1   G1:0   G2:1   L0:0   L1:1   L2:0
G0:1   G1:0   G2:2   L0:0   L1:0   L2:0
G0:1   G1:0   G2:2   L0:0   L1:1   L2:0
G0:2   G1:0   G2:0   L0:0   L1:0   L2:0
G0:2   G1:0   G2:0   L0:0   L1:1   L2:0
G0:2   G1:0   G2:1   L0:0   L1:0   L2:0
G0:2   G1:0   G2:1   L0:0   L1:1   L2:0
G0:2   G1:0   G2:2   L0:0   L1:0   L2:0
G0:2   G1:0   G2:2   L0:0   L1:1   L2:0
G0:3   G1:0   G2:0   L0:0   L1:0   L2:0
G0:3   G1:0   G2:0   L0:0   L1:1   L2:0
G0:3   G1:0   G2:1   L0:0   L1:0   L2:0
G0:3   G1:0   G2:1   L0:0   L1:1   L2:0
G0:3   G1:0   G2:2   L0:0   L1:0   L2:0
G0:3   G1:0   G2:2   L0:0   L1:1   L2:0

Похоже, что интерпретация глобальных / локальных измерений NVidia чередуется, что не соответствует спецификации. Похоже, это также не связано с привязкой C ++. Локальный идентификатор никогда не должен быть ничем, кроме нуля, а get_global_id (1) всегда равен нулю.

Я знаю, что NVidia не особо заботится об OpenCL, но это кажется довольно серьезной проблемой. Кто-нибудь еще сталкивался с чем-то подобным? Это не проблема синхронизации с printf. Я заметил это в реальных случаях использования данных и собрал это ядро ​​только для того, чтобы продемонстрировать это.

Answer 1

^{Хотя это трудно проверить подробно, я опубликую это как ответ, потому что из моих наблюдений, кажется , чтобы объяснить проблему:}

tl; dr : Причина почти наверняка связана с отсутствием синхронизации в printf.

---

Прежде всего, я наблюдал то же поведение, что и вы: на AMDвывод кажется правильным. На NVIDIA это кажется раздражающе неправильным. Поэтому мне было любопытно, и я расширил ядро, чтобы также вывести get_local_size:

void kernel test()
{
    printf("G0:%d   G1:%d   G2:%d   L0:%d   L1:%d   L2:%d  S0:%d  S1:%d  S2:%d\n", 
        get_global_id(0),
        get_global_id(1),
        get_global_id(2),
        get_local_id(0),
        get_local_id(1),
        get_local_id(2),
        get_local_size(0),
        get_local_size(1),
        get_local_size(2));
}

Теперь, get_local_id, безусловно, должен быть меньше, чем размер, иначе большинство ядерпросто потерпел бы крах. На AMD вывод был приятным и чистым:

platform AMD Accelerated Parallel Processing
device Spectre
G0:0   G1:0   G2:0   L0:0   L1:0   L2:0  S0:1  S1:1  S2:1
G0:1   G1:0   G2:0   L0:0   L1:0   L2:0  S0:1  S1:1  S2:1
G0:2   G1:0   G2:0   L0:0   L1:0   L2:0  S0:1  S1:1  S2:1
G0:3   G1:0   G2:0   L0:0   L1:0   L2:0  S0:1  S1:1  S2:1
G0:0   G1:1   G2:0   L0:0   L1:0   L2:0  S0:1  S1:1  S2:1
G0:1   G1:1   G2:0   L0:0   L1:0   L2:0  S0:1  S1:1  S2:1
G0:2   G1:1   G2:0   L0:0   L1:0   L2:0  S0:1  S1:1  S2:1
G0:3   G1:1   G2:0   L0:0   L1:0   L2:0  S0:1  S1:1  S2:1
G0:0   G1:2   G2:0   L0:0   L1:0   L2:0  S0:1  S1:1  S2:1
G0:1   G1:2   G2:0   L0:0   L1:0   L2:0  S0:1  S1:1  S2:1
G0:2   G1:2   G2:0   L0:0   L1:0   L2:0  S0:1  S1:1  S2:1
G0:3   G1:2   G2:0   L0:0   L1:0   L2:0  S0:1  S1:1  S2:1
G0:0   G1:0   G2:1   L0:0   L1:0   L2:0  S0:1  S1:1  S2:1
G0:1   G1:0   G2:1   L0:0   L1:0   L2:0  S0:1  S1:1  S2:1
G0:2   G1:0   G2:1   L0:0   L1:0   L2:0  S0:1  S1:1  S2:1
G0:3   G1:0   G2:1   L0:0   L1:0   L2:0  S0:1  S1:1  S2:1
G0:0   G1:1   G2:1   L0:0   L1:0   L2:0  S0:1  S1:1  S2:1
G0:1   G1:1   G2:1   L0:0   L1:0   L2:0  S0:1  S1:1  S2:1
G0:2   G1:1   G2:1   L0:0   L1:0   L2:0  S0:1  S1:1  S2:1
G0:3   G1:1   G2:1   L0:0   L1:0   L2:0  S0:1  S1:1  S2:1
G0:0   G1:2   G2:1   L0:0   L1:0   L2:0  S0:1  S1:1  S2:1
G0:1   G1:2   G2:1   L0:0   L1:0   L2:0  S0:1  S1:1  S2:1
G0:2   G1:2   G2:1   L0:0   L1:0   L2:0  S0:1  S1:1  S2:1
G0:3   G1:2   G2:1   L0:0   L1:0   L2:0  S0:1  S1:1  S2:1

На NVIDIA вывод был

platform NVIDIA CUDA
device GeForce GTX 970
G0:3   G1:0   G2:2   L0:0   L1:0   L2:0  S0:0  S1:0  S2:0
G0:3   G1:0   G2:1   L0:0   L1:0   L2:0  S0:0  S1:0  S2:0
G0:3   G1:0   G2:0   L0:0   L1:0   L2:0  S0:0  S1:0  S2:0
G0:0   G1:0   G2:2   L0:0   L1:1   L2:0  S0:0  S1:0  S2:0
G0:0   G1:0   G2:1   L0:0   L1:1   L2:0  S0:0  S1:0  S2:0
G0:2   G1:0   G2:0   L0:0   L1:0   L2:0  S0:0  S1:0  S2:0
G0:2   G1:0   G2:1   L0:0   L1:0   L2:0  S0:0  S1:0  S2:0
G0:2   G1:0   G2:2   L0:0   L1:0   L2:0  S0:0  S1:0  S2:0
G0:1   G1:0   G2:1   L0:0   L1:0   L2:0  S0:0  S1:0  S2:0
G0:3   G1:0   G2:0   L0:0   L1:1   L2:0  S0:0  S1:0  S2:0
G0:1   G1:0   G2:0   L0:0   L1:0   L2:0  S0:0  S1:0  S2:0
G0:3   G1:0   G2:1   L0:0   L1:1   L2:0  S0:0  S1:0  S2:0
G0:0   G1:0   G2:2   L0:0   L1:0   L2:0  S0:0  S1:0  S2:0
G0:1   G1:0   G2:2   L0:0   L1:0   L2:0  S0:0  S1:0  S2:0
G0:3   G1:0   G2:2   L0:0   L1:1   L2:0  S0:0  S1:0  S2:0
G0:0   G1:0   G2:1   L0:0   L1:0   L2:0  S0:0  S1:0  S2:0
G0:2   G1:0   G2:1   L0:0   L1:1   L2:0  S0:0  S1:0  S2:0
G0:0   G1:0   G2:0   L0:0   L1:1   L2:0  S0:0  S1:0  S2:0
G0:2   G1:0   G2:0   L0:0   L1:1   L2:0  S0:0  S1:0  S2:0
G0:0   G1:0   G2:0   L0:0   L1:0   L2:0  S0:0  S1:0  S2:0
G0:2   G1:0   G2:2   L0:0   L1:1   L2:0  S0:0  S1:0  S2:0
G0:1   G1:0   G2:2   L0:0   L1:1   L2:0  S0:0  S1:0  S2:0
G0:1   G1:0   G2:1   L0:0   L1:1   L2:0  S0:0  S1:0  S2:0
G0:1   G1:0   G2:0   L0:0   L1:1   L2:0  S0:0  S1:0  S2:0

Теперь, когда не может быть правильным: локальная работаsize всегда равен 0!

После некоторых дальнейших тестов (например, с двумерными ядрами и другими числами) вывод , как правило, , кажется, вообще не имеет никакого смысла. Итак, я попробовал это ядро:

void kernel test()
{
    printf("G0:%d\n", get_global_id(0));
    printf("G1:%d\n", get_global_id(1));
    printf("G2:%d\n", get_global_id(2));
    printf("L0:%d\n", get_local_id(0));
    printf("L1:%d\n", get_local_id(1));
    printf("L2:%d\n", get_local_id(2));
    printf("S0:%d\n", get_local_size(0));
    printf("S1:%d\n", get_local_size(1));
    printf("S2:%d\n", get_local_size(2));
}

На NVIDIA вывод будет

platform NVIDIA CUDA
device GeForce GTX 970
G0:1
G0:1
G0:1
G0:2
G0:2
G0:2
G0:2
G0:2
G0:3
G0:2
G0:3
G0:3
G0:0
G0:3
G0:3
G0:0
G0:0
G0:3
G0:0
G0:0
G0:0
G0:1
G0:1
G0:1
G1:2
G1:2
G1:0
G1:0
G1:1
G1:2
G1:2
G1:1
G1:1
G1:1
G1:0
G1:0
G1:2
G1:1
G1:0
G1:0
G1:2
G1:1
G1:1
G1:0
G1:2
G1:2
G1:0
G1:1
G2:0
G2:0
G2:1
G2:1
G2:0
G2:0
G2:1
G2:0
G2:0
G2:0
G2:0
G2:0
G2:1
G2:1
G2:0
G2:1
G2:1
G2:1
G2:1
G2:0
G2:1
G2:0
G2:1
G2:1
L0:0
L0:0
L0:0
L0:0
L0:0
L0:0
L0:0
L0:0
L0:0
L0:0
L0:0
L0:0
L0:0
L0:0
L0:0
L0:0
L0:0
L0:0
L0:0
L0:0
L0:0
L0:0
L0:0
L0:0
L1:0
L1:0
L1:0
L1:0
L1:0
L1:0
L1:0
L1:0
L1:0
L1:0
L1:0
L1:0
L1:0
L1:0
L1:0
L1:0
L1:0
L1:0
L1:0
L1:0
L2:0
L1:0
L1:0
L1:0
L1:0
L2:0
L2:0
L2:0
L2:0
L2:0
L2:0
L2:0
L2:0
L2:0
L2:0
L2:0
L2:0
L2:0
L2:0
L2:0
S0:1
L2:0
L2:0
L2:0
L2:0
L2:0
L2:0
L2:0
L2:0
S0:1
S0:1
S0:1
S0:1
S0:1
S0:1
S0:1
S0:1
S0:1
S0:1
S0:1
S0:1
S0:1
S1:1
S0:1
S0:1
S0:1
S0:1
S0:1
S1:1
S0:1
S0:1
S0:1
S0:1
S0:1
S1:1
S1:1
S1:1
S1:1
S1:1
S1:1
S1:1
S1:1
S1:1
S1:1
S2:1
S1:1
S1:1
S1:1
S2:1
S1:1
S1:1
S1:1
S1:1
S1:1
S2:1
S1:1
S1:1
S2:1
S2:1
S1:1
S1:1
S2:1
S2:1
S2:1
S2:1
S2:1
S2:1
S2:1
S2:1
S2:1
S2:1
S2:1
S2:1
S2:1
S2:1
S2:1
S2:1
S2:1
S2:1
S2:1

Ключевой момент: Каждый отдельный вывод правильный !. Кажется, проблема в том, что помещение всего в один printf портит некоторый внутренний буфер.

Жаль, конечно. В основном это делает невозможным использование printf для единственной цели, которую он мог бы разумно использовать для внутри ядра, а именно для отладки ...

---

В сторону:Спецификации остаются немного сложными для интерпретации в этот момент - по крайней мере, когда дело доходит до решения, является ли наблюдаемое поведение «правильным» или «неправильным». Из документации Khronos printf:

В случае, когда printf выполняется одновременно из нескольких рабочих элементов, нет никакой гарантии упорядочения в отношении записанных данных. Например, допустимо, чтобы выходные данные рабочего элемента с глобальным идентификатором (0,0,1) смешивались с выходными данными рабочего элемента с глобальным идентификатором (0,0,4) и т. Д. ,

Документация NVIDIA для реализации CUDA printf также содержит некоторые отказы от ответственности и рассказывает о некоторых буферах, которые могут быть перезаписаны, но отображаются на это (на техническом уровне спецификации) к поведению OpenCL сложно ...