Блог

Я использую NUMBA и cupy для выполнения кодирования GPU. Теперь я переключил свой код с карты NVIDIA V100 на A100, но затем я получил следующие предупреждения:

1. Предупреждение о numbaperformance: размер сетки (27) <2 * SM count (216), вероятно, приведет к недоиспользованию GPU из-за низкой занятости.
2. NumbaPerformanceWarning: массив хостов, используемый в ядре CUDA, повлечет за собой накладные расходы на копирование на устройство / с устройства.

Кто-нибудь знает, что на самом деле предлагают два предупреждения? Как мне тогда улучшить свой код?

Предупреждение о numbaperformance: размер сетки (27) <2 * SM count (216), вероятно, приведет к недоиспользованию GPU из-за низкой занятости.

Графический процессор подразделяется на SMS. Каждый SM может содержать набор блоков потоков (что все равно, что сказать, что он может содержать набор потоков). Чтобы «полностью использовать» графический процессор, вы бы хотели, чтобы каждый SM был «заполнен», что примерно означает, что у каждого SM достаточно блоков потоков для заполнения своего набора потоков. Графический процессор A100 имеет 108 SMS. Если в вашем ядре меньше 108 блоков потоков при запуске ядра (т. Е. Сетки), То ваше ядро не сможет полностью использовать графический процессор. Некоторые SMS будут пустыми. Потоковый блок не может быть резидентным на 2 или более SMS одновременно. Даже 108 (по одному на SM) может быть недостаточно. A100 SM может содержать 2048 потоков, что составляет как минимум два блока потоков по 1024 потока в каждом. Все, что меньше 2 * 108 блоков потоков при запуске вашего ядра, может не полностью использовать графический процессор. Когда вы не используете графический процессор в полной мере, ваша производительность может быть не настолько хорошей, насколько это возможно.

Решение состоит в том, чтобы обеспечить достаточный параллелизм (достаточное количество потоков) при запуске вашего ядра, чтобы полностью «занимать» или «использовать» графический процессор. 216 блоков потоков по 1024 потока в каждом достаточно для A100. Чего-то меньшего может и не быть.

Для дополнительного понимания здесь я рекомендую первые 4 раздела этого курса.

NumbaPerformanceWarning: массив хостов, используемый в ядре CUDA, повлечет за собой накладные расходы на копирование на устройство / с устройства.

Одна из замечательных особенностей запуска ядра numba заключается в том, что я могу передать ему массив данных хоста:

 a = numpy.ones(32, dtype=numpy.int64)
my_kernel[blocks, threads](a)
 

и numba «поступит правильно». В приведенном выше примере это будет:

1. создайте массив устройств, предназначенный для хранения a в памяти устройства, назовем это d_a
2. скопируйте данные из a d_a (хост-> Устройство)
3. запустите свое ядро, где ядро фактически использует d_a
4. когда ядро будет завершено, скопируйте содержимое d_a обратно в a (Устройство-> Хост)

Все это очень удобно. Но что, если бы я делал что-то подобное:

 a = numpy.ones(32, dtype=numpy.int64)
my_kernel1[blocks, threads](a)
my_kernel2[blocks, threads](a)
 

Что будет делать numba, так это то, что он выполнит шаги 1-4 выше для запуска my_kernel1 , а затем снова выполнит шаги 1-4 для запуска my_kernel2 . В большинстве случаев это, вероятно, не то, что вам нужно как программисту numba cuda.

Решение в этом случае состоит в том, чтобы «взять под контроль» перемещение данных:

 a = numpy.ones(32, dtype=numpy.int64)
d_a = numba.cuda.to_device(a)
my_kernel1[blocks, threads](d_a)
my_kernel2[blocks, threads](d_a)
a = d_a.to_host()
 

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

Для дополнительного понимания, вероятно, будет полезен любой онлайн-учебник, такой как этот, или просто документы numba cuda.