Безопасно ли ядру OpenCL произвольно записывать в буфер __global?

Question

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

Упрощенный пример:

__kernel void test(__global int *a, __global int *hiddenCounter) {
    if (get_global_id(0) == 0) {
        // do stuff and then increment the appropriate counter (random numbers here)
        hiddenCounter[0] += 3;
    }
    else {
        // do stuff...
        hiddenCounter[1] += 5;
    }
}

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

Мой вопрос: есть ли условия гонки, когданесколько экземпляров ядра пытаются записать в один и тот же индекс буфера hiddenCounter (что обязательно произойдет в моем проекте). Нужно ли применять принудительную синхронизацию? Или это невозможно с __global аргументами, и мне нужно изменить его на __private? Смогу ли я потом агрегировать буферы __private из кода хоста?

Answer 1

Мой вопрос: существуют ли условия гонки, когда несколько экземпляров ядра пытаются записать в один и тот же индекс буфера hiddenCounter

Ответ на этот вопрос решительно да, ваш код будет уязвим к условиям гонки, как написано в настоящее время.

Нужно ли принудительно устанавливать какую-либо синхронизацию?

Да, вы можете использовать для этой цели глобальная атомика . Все, кроме самых древних графических процессоров, будут поддерживать это. (все, что поддерживает OpenCL 1.2 или cl_khr_global_int32_base_atomics и аналогичные расширения)

Обратите внимание, что это приведет к нетривиальным потерям производительности. В зависимости от ваших схем доступа и частоты, сбор промежуточных результатов в памяти private или local и запись их в глобальную память в конце ядра может быть быстрее. (В случае local вся рабочая группа будет использовать только один глобальный атомарный вызов для каждой обновленной ячейки - вам нужно будет использовать локальную атомарность или алгоритм сокращения, чтобы накапливать значения из отдельных рабочих элементов по всей группе.)

Другой вариант - использовать гораздо больший буфер глобальной памяти со счетчиками для каждого рабочего элемента или группы. В этом случае вам не потребуется атомика для записи в них, но впоследствии вам нужно будет объединить значения на хосте. Очевидно, что он использует гораздо больше памяти и, вероятно, большую пропускную способность памяти - современные графические процессоры должны кэшировать доступ к вашему небольшому буферу hiddenCounter. Так что вам нужно решить / попробовать, что является меньшим злом в вашем случае.