Низкое использование графического процессора в CUDA

Question

Я реализовал программу, которая использует разные потоки CUDA из разных потоков ЦП. Копирование памяти осуществляется через cudaMemcpyAsync с использованием этих потоков. Ядро запускает также те потоки. Программа выполняет вычисления с двойной точностью (и я подозреваю, что это преступник, однако, cuBlas достигает 75-85% использования ЦП для умножения матриц double с). Существуют также операции сокращения, однако они реализованы через if(threadIdx.x < s) с s, уменьшающимися 2 раза на каждой итерации, поэтому блокированные перекосы должны быть доступны для других блоков. Приложение является графическим процессором и интенсивным процессором, оно начинается с другой части работы, как только предыдущий закончен. Поэтому я ожидаю, что он достигнет 100% от процессора или графического процессора.

Проблема в том, что моя программа генерирует 30-40% нагрузки на графический процессор (и около 50% загрузки процессора), если доверяет GPU-Z 1.9.0. Нагрузка контроллера памяти составляет 9-10%, загрузка интерфейса шины - 6%. Это касается количества потоков ЦП, равных количеству ядер ЦП. Если я удваиваю количество потоков ЦП, нагрузки остаются примерно одинаковыми (включая загрузку ЦП).

Так почему это? Где узкое место?

Я использую GeForce GTX 560 Ti, CUDA 8RC, MSVC++ 2013, Windows 10.

Одна моя догадка, что для Windows 10 применяет некоторые агрессивные экономии энергии, даже если GPU и CPU температура низкая, план питания установлен на «Высокая производительность», а блок питания составляет 700 Вт, а потребляемая мощность с максимальным процессором и графическим процессором составляет около 550 Вт.

Еще одна догадка заключается в том, что скорость с двойной точностью составляет 1/12 от скорости с одной точностью, поскольку на моей карте имеется 1 ядро ​​CUDA с двойной точностью на 12 ядер CUDA с одной точностью, а GPU-Z занимает 100% ситуация, когда используются все ядра с одинарной точностью и двойной точностью. Однако цифры не совсем совпадают.

Answer 1

По-видимому, причиной было низкое заполнение из-за потоков CUDA с использованием слишком большого количества регистров по умолчанию. Чтобы сообщить компилятору, можно использовать ограничение на количество регистров на поток, __launch_bounds__, как описано here. Таким образом, чтобы иметь возможность запускать все 1536 потоки в 560 Ti, для размера блока 256 следующее может быть указано:

_global__ void __launch_bounds__(256, 6) MyKernel(...) { ... } 

После того, как ограничение количества регистров в CUDA нити, использование графического процессора повысилась до 60% для меня ,

Кстати, карты серии 5xx по-прежнему поддерживаются NSight v5.1 для Visual Studio. It can be downloaded от archive.

EDIT: следующие флаги дополнительно увеличенное использование графического процессора до 70% в приложение, которое использует несколько потоков GPU из нескольких потоков CPU:

cudaSetDeviceFlags(cudaDeviceScheduleYield | cudaDeviceMapHost | cudaDeviceLmemResizeToMax); 

• cudaDeviceScheduleYield позволяет другим потокам выполнения, когда центральный процессор нить ожидая работы GPU, вместо того, чтобы вращать GPU для результата .
• cudaDeviceLmemResizeToMax, как я понял, делает ядро ​​ запускает себя асинхронно и избегает избыточной локальной памяти ассигнования & deallocations.