Нубских вопросов про OpenCL тред

Я не спец конечно, но для:

записать 1 в общую переменную, изначально выставленную в 0

нужно иметь какой то механизм синхронизации, сводящий весь выигрыш к нулю.

Это примерно как с мэп редьюс - привести здоровый набор данных к заданному мелкому индексу, характеризующему эти данные за разумное время.

Типа допустим ты выполнил на 6 ядрах по 1кк операций, а в ответ получил 6 булевских значений. Ты свёл 6кк -> к 6, получил выигрыш в 1кк раз, глупо дальше экономить на спичках.

Если я ошибаюсь, поправьте меня.

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

2. Остановить воркеры на GPU не получится, потому что это GPU. Если нужен результат в виде одного значения, то есть следующие варианты: а) на хосте определить наличие единицы в массиве, б) с помощью атомиков записать в общую переменную, в) использовать MapReduce. Метод б требует поддержки атомиков платформой и, наверное, самый медленный. Кроме того, он требует повышенной внимательности из-за необходимости синхронизации - считай, что мьютексов нет. MapReduce, очевидно, стоит доводить не до конца - GPU работает медленнее, если потоков мало (меньше примерно тысячи).

Очень рекомендую почитать вот это.

3. Яростно советую почитать про глобальную и локальную память со страницы документации pyopencl, например, это. Без понимания модели памяти эффективно писать под GPU не получится.

4. Сначала запускаешь N локальных групп по M потоков (если железо позволяет столько) на одном кернеле, потом запускаешь второй кернел с N потоками. Синхронизация данных на твоей совести. Также см. п.3.

P.S. Если используешь OpenCL чисто для CPU, часть из вышеизложенного не актуальна. Впрочем, я не думаю, что это осмыслено - есть же OpenMP и много чего еще, более удобного.

2. Во всех примерах, которые я нашел, кернел возвращает массив, количество элементов которого пропорционально количеству воркеров, т. е. каждый поток дает какой-то независимый частный результат, а хост-программа забирает массив этих результатов. Мне же нужно единое булево значение, т. е. если в одном из воркеров получается true, то он должен остановить остальные воркеры, записать 1 в общую переменную, изначально выставленную в 0, и завершиться сам. Можно ли сделать нечто подобное, чтобы не перебирать массив булевых значений в хост-программе?

Время выполнения одного ядра обычно делают достаточно малым чтобы не требовалось останавливать другие воркеры, когда один из вокреров нашёл ответ. Если нужно просто определить наличие результата - то передавать занулённый буфер в ядро, когда какой-то воркер получил true то он записывает в этот буфер результат с помощью atomic_inc(), например, и завершается сам. Хост-программа вычитывает результат и видит, сколько воркеров дали результат true. Если же обязательно надо останавливаться когда соседний воркер записал 1 в буфер, то можно каждому воркеру в цикле своих вычислений регулярно вызывать atomic_max(buf, 0) - и если атомарный максимум единица, значит какой-то воркер уже всё нашёл.

3. Почему при запуске кернела вместо просто целочисленного числа потоков указываются какие-то local worksize и global worksize с размерностями? В чем их суть или где найти подробное объяснение? Я в своем просто указал global = (количество_элементов), local = None, массив сделал плоским, в кернеле сделал вычисление смещений с номером воркера для доступа к элементам конкретного экземпляра. Работает, хотя не уверен, что эффективно. Сколько потоков может на самом деле выполнять GPU? Что происходит, если пытаются запустить больше?

Локальные размеры нужны когда воркерам нужно обмениваться данными через локальную память и синхронизироваться. Глобальной синхронизации в OpenCL нет, только внутри локальной рабочей группы. Все воркеры из рабочей группы запускаются параллельно на одном compute unit'е. Это, например, позволяет им эффективно использовать кэш если воркеры из одной рабочей группы лезут в память по соседним адресам. Если для алгоритма разделение на рабочие группы не нужно - можно просто передать глобальный размер и OpenCL сам себе придумает локальный. Глобальный размер, т.е. и количество потоков, в OpenCL не ограничен - на сколько size_t и терпения хватит. Но размеры локальной группы по понятным причинам ограничены - им надо иметь возможность лезть в одну локальную память, они запущены на одном compute unit и т.д. Максимальный размер - в свойстве устройства CL_DEVICE_MAX_WORK_GROUP_SIZE. Если попросить размер группы больше, чем может устройство - будет ошибка CL_INVALID_WORK_GROUP_SIZE. Если попросить очень-очень большой глобальный размер - реализация OpenCL просто будет запускать его по кусочкам на исполнение и все рабочие группы одна пачка за другой когда-нибудь выполнятся, механизм такой же как если разбивать большой глобальный размер на мелкие самому с помощью параметра global_work_offset у clEnqueueNDRangeKernel().

4. Как реализовать схему «для N объектов запускаются по M воркеров, до определенной точки алгоритма работают параллельно, далее из каждых M остается по 1-му, делающему что-то с M значениями, итого N»? Надеюсь, понятно сформулировал.

Как уже подсказали - если нужна глобальная синхронизация то это значит что то, что делается до синхронизации - делается в первом ядре, а то, что после синхронизации - во втором. Есть один паттерн как в конкретном случае избежать разделения на несколько ядер, но он работает только для случая (сначала работают все рабочие группы)->(работает только одна рабочая группа). Суть его в том, что ядру передается буфер для синхронизации, с нулевым значением внутри. Когда рабочая группа завершает свою часть расчетов, она добавляет единицу к буферу с помощью int old = atomic_inc(buf); Далее: если old не равен (количеству групп - 1), то выходим. Если же old == (get_num_groups()-1) то это значит что мы - последняя оставшаяся рабочая группа, зануляем buf (для следующего запуска) и делаем своими воркерами оставшуюся часть работы с тем, что посчитали все воркеры всех групп. Если количество не делится нацело на количество воркеров в рабочей группе - ничего страшного, можно запустить цикл вида for(int i = get_local_id(); i < data_size; i+= get_local_size()).