Какой самый быстрый способ обновить одно значение с плавающей точкой в ​​GPU для доступа к нему в ядре CUDA?

Question

У меня есть симуляция частиц opengl, где положение каждой частицы вычисляется в ядре CUDA. Большая часть памяти находится в памяти графического процессора, но есть одно значение с плавающей запятой, я должен обновлять с ЦП каждый кадр.

В настоящее время я использую cudaMemcpyAsync() для копирования значения с плавающей запятой в графический процессор, но (по крайней мере из того, что я могу сказать), это немного замедляет производительность. Я попытался использовать nvproof, чтобы увидеть, какие вызовы занимают больше всего времени, с такими результатами:

Calls    Avg       Min       Max   Name
477  2.9740us  2.8160us  4.5440us  simulation(float3*, float*, float3*, float*)
477  89.033us  18.600us  283.00us  cudaLaunchKernel
477  47.819us  10.200us  120.70us  cudaMemcpyAsync

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

Я также пытался использовать закрепленную память и cudaHostGetDevicePointer() , как описано здесь , однако по какой-то причине это увеличивает время запуска ядра еще больше, увеличивая время, сэкономленное на не нуждается в функции memcopy.

Я думаю, должен быть лучший / более быстрый способ обновления моей единственной переменной с плавающей точкой в ​​GPU?

Answer 1

Самый простой способ заключается в том, что вы можете добавить дополнительный параметр в функцию ядра моделирования в виде значения простого с плавающей точкой, но не как указатель для смещения, чтобы данные передавались напрямую. структурой параметров запуска ядра, которую CUDA отправляет в GPU при запуске ядра. Затем вы полностью уклоняетесь от этой команды копирования данных. (Я предполагаю, что CUDA упаковывает все данные дескриптора параметра функции ядра в одну команду копирования, поскольку пространство дескриптора параметра ядра ограничено несколькими килобайтами или меньше).

simulation(fooPointer, 
           barPointer, 
           fooBarPointer, 
           floatVariable 
);

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

Если это не интерактивное моделирование, должно работать скрытие задержек вычислений / визуализации / данных с помощью двойной или тройной буферизации.

Если вы после минимизации синхронизации по кадрам (более быстрый отклик на пользовательский ввод в симуляцию?), То вам следует встроить переменную с плавающей точкой в ​​конец массива, который вы уже отправляете / используете в симуляции или что-то еще структура, которую вы используете. Если у вас уже есть буфер 1MB + float для отправки в GPU, то добавление 4B (float) к его концу не должно иметь большого значения, тогда вы можете получить к нему доступ оттуда. 1 операция копирования должна выполняться быстрее, чем 2 операции копирования с одинаковым общим размером.

Если вы буквально посылаете только 4B в GPU в каждом кадре (с простой функцией для генерации этих данных), то (как сказал 3Dave в комментариях) вы можете попробовать добавить дополнительную функцию ядра для обновления значения в GPU и просто накладные расходы на команду запуска ядра вместо накладных расходов на копирование и копирование данных. Положительным моментом является то, что дополнительные издержки ядра могут быть скрыты, если есть «граф» ядер, запущенных для каждого кадра автоматически, без постановки в очередь их всех снова и снова.

Здесь

https://devblogs.nvidia.com/cuda-graphs/

Деталь

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

cudaGraphLaunch(instance, stream);

Говорят, что затраты на запуск ядра в этой функции «графа» составляют всего 3-4 микросекунды, когда в алгоритме много (20) ядер.

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

(Может быть) Вам даже не нужно менять механизм данных вообще. Просто попробуйте отправить данные с плавающей точкой в ​​виде двоичного представления в значение указателя и только прочитать значение указателя (но не значение данных) из ядра и преобразовать его обратно во float. Я не знаю, возвращает ли CUDA ошибку для этого, если вы не пытаетесь достичь (неправильного) адреса указателя, который представляют данные с плавающей точкой, в ядре.

simulation(fooPointer, 
           barPointer, 
           fooBarPointer, 
           toPtr(floatData) // <----- float to 64/32 bit pointer value
);

и в ядре

float val = fromPtrToFloat(parameter4); // converts pointer itself, not the data

Но это не может быть предпочтительной практикой, если вы можете просто использовать параметры типа «значение».