2012-01-24 3 views
4

Я пишу функцию, которая выполняет много операций BLAS gemv.Накладные расходы на вызов OpenCL или CUDA?

Я хотел бы иметь возможность сделать это на GPU, и я пробовал cuBlas.

Моя проблема в том, что мои матрицы и векторы довольно малы, матрица 100x100 и 100 вектор. CuBlas занимает много времени по сравнению с процессором, и я понимаю, почему смесь быстрого кеша на процессоре и большие накладные расходы при выполнении вызовов на GPU.

Поэтому я пытаюсь найти разумный способ измерения времени, необходимого для передачи вызова на GPU.

Это время, которое требуется CUDA для настройки вызова и отправки его на графический процессор - не считая времени, которое фактически требуется для умножения матрицы-вектора.

Как бы я это сделал?

+0

Чтобы вы могли динамически выбирать, отправлять ли задание в CUDA или просто ради интереса? – Rup

+0

@Rup: Просто интересно узнать, что на самом деле стоит звонок. и определить, является ли медленный код моей ошибкой или просто продуктом архитектуры ;-) –

+1

Для небольших объемов данных не только накладные расходы на вас, но и на отсутствие возможного параллелизма. Графический процессор зависит от наличия достаточного количества потоков, чтобы скрыть задержки (что намного хуже на gpu, а затем на CPU). Даже без накладных расходов на вызовы графический процессор, вероятно, будет медленнее, чем процессор, если работа не будет разбита на ** lots ** потоков. Лоты могут легко означать тысячи потоков. – Grizzly

ответ

8

Обновление: Следующие результаты для рукописного FFT алгоритма GPU на 2005 аппаратных средств (Nvidia 7800 GTX), но показывает принцип CPU-GPU Tranfer узкие

Накладные не вызов, но компиляция программы GPU и передача данных между GPU и хостом. ЦП очень оптимизирован для функций, которые могут выполняться полностью в кеше, а латентность памяти DDR3 намного ниже, чем шина PCI-Express, обслуживающая GPU. Я сам это испытал при написании подпрограмм GPU FFT (до CUDA). См. this related question.

 
N  FFTw (s) GPUFFT (s) GPUFFT MFLOPS GPUFFT Speedup 
8  0   0.00006  3.352705  0.006881 
16  0.000001 0.000065 7.882117  0.010217 
32  0.000001 0.000075 17.10887  0.014695 
64  0.000002 0.000085 36.080118  0.026744 
128  0.000004 0.000093 76.724324  0.040122 
256  0.000007 0.000107 153.739856  0.066754 
512  0.000015 0.000115 320.200892  0.134614 
1024 0.000034 0.000125 657.735381  0.270512 
2048 0.000076 0.000156 1155.151507  0.484331 
4096 0.000173 0.000215 1834.212989  0.804558 
8192 0.000483 0.00032  2664.042421  1.510011 
16384 0.001363 0.000605 3035.4551  2.255411 
32768 0.003168 0.00114  3450.455808  2.780041 
65536 0.008694 0.002464 3404.628083  3.528726 
131072 0.015363 0.005027 3545.850483  3.05604 
262144 0.033223 0.012513 3016.885246  2.655183 
524288 0.072918 0.025879 3079.443664  2.817667 
1048576 0.173043 0.076537 2192.056517  2.260904 
2097152 0.331553 0.157427 2238.01491  2.106081 
4194304 0.801544 0.430518 1715.573229  1.861814 

В приведенной выше таблице показаны тайминги реализации FFT GPU и реализации ЦП на основе размера ядра. Для меньших размеров передача данных в/из графического процессора доминирует. Меньшие ядра могут выполняться на ЦПУ, некоторые реализации/размеры полностью в кеше. Это делает CPU лучшим выбором для небольших операций.

Если, с другой стороны, вам необходимо выполнить большие партии работы с данными с минимальными перемещениями на/из графического процессора, тогда GPU будет бить CPU вниз.

Что касается измерения эффекта в вашем примере, я бы предложил выполнить такой эксперимент, как выше. Попробуйте выработать FLOPS, рассчитанные для каждого размера матрицы, и выполнить тест на CPU и GPU для разных размеров матрицы. Выведите в CSV-файл размер, время и FLOPS для GPU и CPU. Для любого профилирования убедитесь, что вы запускаете несколько сот итераций своего кода и времени, а затем разделите общее время на итерации, чтобы получить время цикла. Попробуйте различные формы матрицы также, если ваш алгоритм позволяет (например, 10x100, а не 100x10).

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

Надеется, что это помогает,

1

Вы можете получить время в наносекундах от устройства, когда событие было очереди, представленным, начало и закончил с использованием clGetEventProfilingInfo на вашем случае передачи буфера.

подробнее, и как установить его здесь: http://www.khronos.org/registry/cl/sdk/1.0/docs/man/xhtml/clGetEventProfilingInfo.html

Я думаю, что для 100x100 матриц, вы можете быть лучше придерживаться центрального процессора для хруст. Если вам не удастся много размножаться в одно и то же время, преимущество gpu будет едва заметным из-за (небольших) расходов на передачу и обычно намного более низких тактовых частот. Убедитесь, что вы настроили свое ядро ​​для использования как можно большего количества локальных данных - на моем оборудовании есть 32 КБ на рабочую группу, и это должно быть много, чтобы удерживать две матрицы 100x100. Также должны быть очень удобны функции встроенных функций точек.

Был удивительный разговор об этом на ADFS в прошлом году (см SESSIONID: 2908) http://developer.amd.com/afds/pages/OLD/sessions.aspx Они подробно поговорим об оптимизации ядра и жесткого кодирования оптимальных размеров.

1

Являются ли ваши матрицы уже на графическом процессоре? Если нет, CUBLAS может передать их вам (известный как thunking), что является дополнительным накладным расходами.

Кроме того, графические процессоры на самом деле не сияют для таких небольших вычислений, то есть они, вероятно, будут медленнее, чем процессоры, так как вам нужно вернуть свой результат обратно. Если вы можете, используйте более крупные матрицы. В противном случае вы можете использовать потоки (cudaStream_t) для запуска нескольких параллельных вычислений на графическом процессоре.

Если вы хотите измерить время выполнения ядра в CUDA, вам необходимо заключить, что (или что-нибудь еще, что рассчитывает на GPU) в событиях, как это при использовании CUDA выполнения API:

cudaEvent_t start, stop; 

cudaEventRecord(&start); 

struct timeval cpuStart, cpuEnd; 

gettimeofday(&cpuStart, 0); // get start time on CPU 

// Do something with CUDA on the GPU, e.g. call kernels, transfer memory, ... 

gettimeofday(&cpuEnd, 0); // get end time on CPU 

double seconds = cpuEnd.tv_sec - cpuStart.tv_sec; 
double microseconds = cpuEnd.tv_usec - cpuStart.tv_usec; 
double cpuDuration = (seconds * 1.0e6 + microseconds)/1.0e3; // in milliseconds 

cudaEventRecord(&stop); 

// Wait until the stop event occurred 
cudaError_t eventResult; 

do 
{ 
    eventResult = cudaEventQuery(stop); 
} 
while (eventResult == cudaErrorNotReady); 

// Assert there was no error; check the CUDA Toolkit Reference for further info 
assert(cudaSuccess == eventResult); // requires #include <assert.h> or <cassert> 

// Retrieve the time 
float gpuDuration = 0.0; // in milliseconds 
cudaEventElapsedTime(&gpuDuration, start, stop); 

// Release the event objects 
cudaEventDestroy(stop); 
cudaEventDestroy(start); 

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

(Примечание: в основном я использую API-интерфейс драйвера CUDA, поэтому это может не сработать. Извините за это.)

EDIT: Просто увидел, что вы хотите измерить сам вызов, а не продолжительность ядра. Вы можете сделать это, просто измерив время на CPU для вызова - см. Обновленный код выше. Это работает только в Linux, потому что gettimeofday недоступен для Windows (AFAIK).

+1

В Windows вы можете использовать [QueryPerformanceCounter] (http: // msdn. microsoft.com/en-us/library/ms644904) или [GetSystemTime] (http://msdn.microsoft.com/en-us/library/windows/desktop/ms725473.aspx) и т. д. – Rup

+0

У меня есть все данные на устройстве, и только нужно сделать простой Ax-> y, а затем сохранить y на устройстве. –

+1

В этом случае вы можете измерить время, необходимое CUBLAS для запуска реального ядра, путем добавления gettimeofday() (или аналогичного метода в Windows) вокруг вашего вызова cublasDgemm(). Хотя я не пробовал это самостоятельно, вы могли бы изучить использование Parallel Nsight (в Windows) или Visual Compute Profiler (входит в набор инструментов в Linux). Я не могу найти его прямо сейчас, но я уверен, что видел что-то о профилирующих крючках в CUDA 4 ... EDIT: нашел этот PDF-файл, содержащий некоторую интересную информацию о профилировании CUDA: http: // bit .ly/zn6jbP –

1

Чтобы найти накладные расходы на вызов, вызовите ядро ​​CUDA, которое сделает как можно меньше.

for (int i=0; i<NLoops; i++) { 
    gettimeofday(&cpuStart, 0); // get start time on CPU 

    // Call minimal CUDA kernel 

    gettimeofday(&cpuEnd, 0); // get end time on CPU 

    // save elapsed time 
} 

Следуйте коду Alex P. выше.

Чем меньше обработка вы выполняете в ядре, тем больше разница во времени будет только накладными расходами.

Проведите небольшое экспериментирование, чтобы найти хорошее значение для NLoops (возможно, 1 000 000). Убедитесь, что прошедшее время больше, чем интервал вашего таймера, или вы получите все нули. Если это произойдет, напишите некоторый код ядра, который выполняется в фиксированный временной интервал, который вы можете предсказать: (n циклов из x циклов каждый).

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

Смежные вопросы