Ядро OpenCL START задерживается при использовании нескольких графических процессоров?

Question

У меня есть приложение, которое я разработал для работы на графических процессорах AMD с OpenCL. Наконец, вчера было запущено приложение и ошибка (ха-ха), нацеленная на один графический процессор. Теперь, когда приложение работает, пришло время масштабировать его на несколько графических процессоров.

Читайте много о том, как настроить его. Мы используем единый контекст, метод с несколькими очередями.

Я вытаскиваю список устройств и выбираю 2 из графических процессоров и создаю контекст, содержащий их, а затем одну BuildProgram, содержащую оба устройства. Создайте две отдельные очереди.

ПСЕВДОКОД оригинала, работающего приложения, теперь конвертируется обрабатывать 2 GPUs:

context = clCreateContext(0, 2, device_list, NULL, NULL, &ret); 
for(x = 0; x < 2; x++) 
    queue[x] = clCreateCommandQueue(context, device_list[x], ENABLE_PROFILING, &ret); 
clBuildProgram(program, 2, device_list, options, NULL, NULL); 

create kernels.. 

run... 
for(outer_loop = 0; outer_loop < 10; outer_loop++) { 
    clEnqueueNDRangeKernel(queue[0], kernel_init, offset, &event[0]); 
    clEnqueueNDRangeKernel(queue[1], kernel_init, different_offset, &event[1]); 
    clFinish(queue[0]); 
    clFinish(queue[1]); 

    get profiling data and printf results 
} 

Вот в основном, как код выглядит. Аргументы задаются, и записи выполняются ПЕРЕД циклом - ядро ​​init не полагается на ввод, чтобы начать работу. После этого он выполняет async_work_group_copy своих сгенерированных данных в глобальном буфере.

Теперь, прежде чем я изменил код для 2-х графических процессоров, ядро ​​побежал в 27ms (для каждого цикла)

После того как я изменил код, если я закомментировать один или другой из прогонов ядра 2 (EnqueueNDRangeKernel и связанный clFinish), они оба будут работать в течение 27 мс.

Если я запускаю код для запуска на обоих графических процессорах параллельно, я получаю очень странное поведение.

Первый запуск в цикле, оба они выполняются примерно в 37-42 мс отдельно. Я в порядке с небольшим замедлением, так как я получаю в два раза больше работы. Но после первого запуска одно или другое ядро ​​будет случайным образом иметь задержку 4-5 секунд между очередью и запуском.

Вот результат моего профилирования/времени для него. Все числа указаны в мс.

Q0: til sub: 8.8542 til start: 9.8594 til fin: 47.3749 
Q1: til sub: 0.0132 til start: 13.4089 til fin: 39.2364 

Q0: til sub: 0.0072 til start: 0.2310 til fin: 37.1187 
Q1: til sub: 0.0122 til start: 4152.4638 til fin: 4727.1146 

Q0: til sub: 0.0302 til start: 488.6218 til fin: 5049.7233 
Q1: til sub: 0.0179 til start: 5023.9310 til fin: 5049.7762 

Q0: til sub: 0.0190 til start: 2.0987 til fin: 39.4356 
Q1: til sub: 0.0164 til start: 3996.2654 til fin: 4571.5866 

Q0: til sub: 0.0284 til start: 488.5751 til fin: 5046.3555 
Q1: til sub: 0.0176 til start: 5020.5919 til fin: 5046.4382 

Машина, в которой я работаю, имеет в ней 5 графических процессоров. Независимо от того, какой из двух я использую, один из двух графических процессоров (его не всегда один и тот же) получает 4-5 секундную задержку при запуске. Используйте один GPU - без задержки.

Что может быть причиной этого? Есть идеи? Я не блокирую - clFinish только для получения информации профилирования. Даже если бы это было блокирование, это не было бы 5-секундной задержкой.

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

Фактически, я добавил возврат; как первая строка ядра, так что он абсолютно ничего не делает. 40 мс упали до 0,25, но 5-секундная задержка все еще была.

Answer 1

Драйвер OpenCL не заботится о том, что происходит в ядре. Если ядро ​​записывает/читает или является нулевым ядром или записывает только один раздел буфера. Он заботится о флажках параметров буфера и обеспечивает согласованность данных между графическими процессорами, блокируя ядра, если они имеют ЛЮБУЮ зависимость в других ядрах. Передача GPU на GPU происходит прозрачно и может быть очень дорогостоящей.

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

Если ваши ядра могут работать параллельно (поскольку GPU1 работает с разными данными, которые находятся на GPU2 и т. Д.), Тогда вы должны создавать разные буферы для каждого графического процессора. Или, если данные одинаковые, правильно установите типы CL_READ_ONLY/CL_WRITE_ONLY, чтобы обеспечить правильное поведение OpenCL. И минимальные операции копирования/согласования.

---

Например для этих ядер:

kernel Sum(read_only A, read_only B, write_only C); 
kernel Sum_bad(read_write A, read_write B, write_only C); 

Если вы используете один GPU, как будет вести себя точно так же, потому что вся память находится в том же GPU. Но использование нескольких графических процессоров может вызвать ужасные проблемы, например:

Queue 1/GPU 1: Sum_Bad(A,B,C); 
Queue 2/GPU 2: Sum_Bad(A,D,E); 

события будут происходить следующим образом:

1. памяти A, B будут скопированы в память GPU1 (если это не было уже). C, выделенной в GPU1.
2. GPU 1 запустит ядро.
3. Память A будет скопирована с GPU1 на GPU2. Память D будет скопирована на GPU2. Память E выделена.
4. GPU2 будет запускать ядро.

Как вы видите, GPU2 должен дождаться завершения первого и дополнительно дождаться возврата всех параметров. (Может ли это быть 5s может быть, в зависимости от размеров?)

---

Однако при использовании правильного подхода:

Queue 1/GPU 1: Sum(A,B,C); 
Queue 2/GPU 2: Sum(A,D,E); 

события будут происходить следующим образом:

1. памяти A, B будет скопирован в память GPU1 (если его там уже не было). C, выделенной в GPU1.
2. GPU 1 запустит ядро.

Параллельно (потому что нет зависимостях)

1. памяти А, D будут скопированы в GPU2 (если это не было уже). Память E выделена.
2. GPU2 будет запускать ядро.