Стратегии для выбора времени ядра CUDA: плюсы и минусы?

При синхронизации ядер CUDA следующее не работает, потому что ядро ​​не блокирует выполнение программы CPU во время выполнения:

start timer
kernel<<<g,b>>>();
end timer

Я видел три основных способа (успешно) синхронизировать ядра CUDA:

(1) Две записи CUDA.

float responseTime; //result will be in milliseconds
cudaEvent_t start; cudaEventCreate(&start); cudaEventRecord(start); cudaEventSynchronize(start);
cudaEvent_t stop;  cudaEventCreate(&stop);
kernel<<<g,b>>>();
cudaEventRecord(stop); cudaEventSynchronize(stop);
cudaEventElapsedTime(&responseTime, start, stop); //responseTime = elapsed time

(2) Одна запись CUDA.

float start = read_timer(); //helper function on CPU, in milliseconds
cudaEvent_t stop;  cudaEventCreate(&stop);
kernel<<<g,b>>>();
cudaEventRecord(stop); cudaEventSynchronize(stop);
float responseTime = read_timer() - start;

(3) deviceSynchronize вместо eventRecord. (Вероятно, полезно только при использовании программирования в одном потоке.)

float start = read_timer(); //helper function on CPU, in milliseconds
kernel<<<g,b>>>();
cudaDeviceSynchronize();
float responseTime = read_timer() - start;

Я экспериментально подтвердил, что эти три стратегии дают одинаковый результат синхронизации.

Вопросы:

  • Каковы компромиссы этих стратегий? Здесь есть какие-то скрытые детали?
  • Помимо синхронизации многих ядер в нескольких потоках, есть ли преимущества использования двух записей событий и cudaEventElapsedTime() функционировать?

Вы можете, вероятно, использовать свое воображение, чтобы выяснить, что read_timer() делает. Тем не менее, это не помешает привести пример реализации:

double read_timer(){
    struct timeval start;
    gettimeofday( &start, NULL ); //you need to include <sys/time.h>
    return (double)((start.tv_sec) + 1.0e-6 * (start.tv_usec))*1000; //milliseconds
}
Источник

2 ответа

Решение

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

Одним из возможных различий будет переносимость между Windows и Linux. Я считаю, что ваш пример функции read_timer ориентирован на Linux. Возможно, вы могли бы создать функцию read_timer, которая является "переносимой", но система событий cuda (метод 1) переносима как есть.

Источник

Вариант (1) использует cudaEventRecord для измерения времени процессора. Это крайне неэффективно, и я бы не рекомендовал использовать cudaEventRecord для этой цели. cudaEventRecord может использоваться для определения времени буфера push-кода графического процессора для выполнения ядра следующим образом:

float responseTime; //result will be in milliseconds
cudaEvent_t start;
cudaEvent_t stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);

cudaEventRecord(start);
kernel<<<g,b>>>();
cudaEventRecord(stop);
cudaEventSynchronize(stop);
cudaEventElapsedTime(&responseTime, start, stop); //responseTime = elapsed time

Код нужно немного изменить, если вы отправляете несколько элементов работы в несколько потоков. Я бы порекомендовал прочитать ответ Разница во времени, сообщаемая НВВП и счетчиками.

Варианты (2) и (3) аналогичны для данного примера. Вариант (2) может быть более гибким.

Источник
Другие вопросы по тегам