OpenCL: ядро ​​цикла?

Question

Я использую ядро ​​OpenCL, которое обрабатывает и повторно обрабатывает один и тот же набор данных снова и снова (это решатель итеративной физики).

В моих тестах нетривиальная стоимость вызова clEnqueueNDRangeKernel.Например, при выполнении 1000 подэтапов моделирования (требующих 1000 идентичных вызовов clEnqueueNDRangeKernel для обработки одних и тех же данных) кажется, что эти вызовы clEnqueueNDRangeKernel фактически становятся узким местом.Мой (псевдо) код выглядит следующим образом:

[create buffers]
[set kernel arguments]

for (int i = 0; i < 1000; i++) //queuing the kernels takes a while
{
    clEnqueueNDRangeKernel(queue, kernel, args...); 
}

clFinish(queue); //waiting for the queue to complete doesn't take much time
[read buffers]

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

У меня также сложилось впечатление, что clEnqueueNDRangeKernel не является блокирующим в том смысле, что он не будет блокироваться, пока ядро ​​не будет завершено, поэтому сложность ядра не должна быть узким местом (в моем случаевыполнение ядра не должно блокироваться до вызова clFinish ()).

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

Мой вопрос : есть ли способ заставить GPU перезапуститьранее помещенное в очередь ядро ​​N раз, вместо того, чтобы вручную ставить ядро ​​N раз?В моей ситуации не нужно менять или обновлять аргументы для ядра каждую итерацию ... ядро ​​просто нужно перезапустить.Можно ли сделать повторные звонки более эффективными?

Answer 1

OpenCL 2.x поддерживает динамический параллелизм, который позволяет 1 рабочему элементу запускать новые ядра.Если при каждом запуске ядра не требуется передача данных gpu-cpu, вы можете иметь 1 рабочий элемент для запуска 1000 ядер и ждать, пока каждое из них не завершится к этому рабочему элементу.Используйте события, чтобы все дочерние ядра работали за другим.

В OpenCL 1.2 вы могли бы использовать атомарность и цикл для выполнения синхронизации ядра "в полете потоков", но это не было бы быстрее, чем запуск нового ядраimo, и это не переносимый способ их синхронизации.

Если каждый запуск ядра стоит больше времени, чем каждый запуск ядра, то на GPU недостаточно работы.Вы можете просто сделать c = a + b на графическом процессоре, и это не будет достаточно быстро только из-за того, что для планирования ядра на конвейерах gpu требуется больше времени, чем для c = a + b.

Но вы все равно можете выполнить следующий подход, используя clEnqueueWaitForEvents и очереди команд по порядку:

thread-0:
    enqueue user event, not triggered
    enqueue 1000 kernels, they don't start yet because of untriggered wait
thread-1:
    nothing

следующий шаг:

thread-0:
    enqueue new user event on a new command queue, not triggered
    enqueue 1000 kernels on new command queue so they don't start yet
thread-1:
    run the old command queue from last timestep by triggering the user event

, чтобы поставить в очередьи бег может быть «перекрыт» как минимум.Если вам нужно больше очереди для запуска коэффициента перекрытия,

thread-0 to thread-N-2:
    enqueue new user event on a new command queue, not triggered
    enqueue 1000 kernels on new command queue so they don't start yet
thread-N-1:
    iterate all command queues
          run currently selected command queue from last timestep  by triggering the user event

теперь, когда у вас в N-1 раз быстрее постановка в очередь, запуск их будет только при накладных расходах планирования на стороне GPU.Если у вас много накладных расходов на планирование на GPU (1M рабочих элементов для 1M c = a + b calcs), вам следует выполнять больше работы на каждый рабочий элемент.

Может быть, лучше сделать запуск ядра в стиле производитель-потребитель, где 7 потоковсоздавать заполненные очереди команд, ожидающие запуска по своим собственным пользовательским событиям, и 8. потреблять их потоком, вызывая их.Это будет работать, даже если им нужно загрузить данные для загрузки в / из GPU.

Даже старые графические процессоры, такие как HD7870, поддерживают 32 + очереди команд (на каждый графический процессор) одновременно, так что вы можете масштабировать производительность очереди высокопроизводительными процессорами.

Если мост pci-e (высокая задержка у стояков?) Вызывает узкое место, то динамический параллелизм OpenCL2.x должен быть лучше, чем шаблон производитель-потребитель на стороне процессора.