Копирование данных GPU CPU

Лимитирующим фактором в любых вычислениях на видеокарте является передача данных из ОЗУ в видеокарту и обратно. Попробуем оценить это время.

Скорость передачи ОЗУ ↔ видеопамять

Попробуем измерить скорость передачи данных «в» и «из» видеопамяти. Код очень прост — используем только функцию cudaMemcpy.

int count = 65536 * 1024 * 4;
int size = sizeof(int) * count;

int *cpu_a = (int *)malloc(size); int *gpu_a; cudaMalloc((void**)&gpu_a, size);

for(int i = 0; i < count; i++) cpu_a[i]=1;

printf("size: %.3f GBn", size/1024.0/1024.0/1024.0);

long long tm1 = -gettimeus();
cudaMemcpy(gpu_a, cpu_a, size, cudaMemcpyHostToDevice);
tm1 += gettimeus();
printf("DRAM -> GDRAM: %lld ms. speed: %.3f GB/sn", tm1/1000, size/tm1/1000.0);

long long tm2 = -gettimeus();
cudaMemcpy(cpu_a, gpu_a, size, cudaMemcpyDeviceToHost);
tm2 += gettimeus();
printf("DRAM <- GDRAM: %lld ms. speed: %.3f GB/sn", tm2/1000, size/tm2/1000.0);

free(cpu_a); cudaFree(gpu_a);

Скорость копирования между памятью компьютера и видеокарты

Как видим, скорость примерно равна 4 гигабайтам в секунду или 32 гигабитам/с (при измерении скорости интерфейсов всегда используют только десятичные приставки). Много это или мало? Да честно говоря, не очень много. Однако, на пути данных есть несколько интерфейсов, и какой-то из них может оказаться бутылочным горлышком. Перечислим их все:

Видеопамять GDDR5 имеет скорость 80 Гб/с;

Интерфейс PCI-Express 3.0 x16 — 128 Гб/с;

Измеренная скорость передачи DRAM -> DRAM составила 27 Гб/с;

Похоже, производительность упёрлась в скорость ОЗУ.

UPD:

Очень большую роль играет выбор порта PCI-Express. Эксперименты проводились на чипсете Z77 и видеокарте GTX650. Если вставить её в «PCI-Express1″ — скорость обмена будет около 4 ГБ/с и выше. Если же использовать порт «PCI-Express2″ — скорость упадёт до 1.5-1.7 ГБ/с. Для работы с CUDA используйте только PCI-Express1!

Аудио–кроссовер на CUDA (через КИХ–фильтры)

Один из красивых и интересных примеров применения CUDA представлен здесь: реализация очень точного кроссовера (полосового аудио-фильтра) для домашней акустики на маломощной видеокарте. Человек использует фильтр порядка 8192! Такие порядки нечасто встретишь в подобных алгоритмах :)

КИХ-фильтр — классический пример map/reduce алгоритма: проходим по всему массиву предыдущих значений входного параметра, умножая их на коэффициенты фильтра, потом редуцируем массив, складывая все полученные произведения. Как и любой map-алгоритм, он легко поддаётся параллелизации.

Hello world на CUDA

В сети сравнительно мало информации о CUDA, особенный дефицит helloworld`ов. Без теории перейдём сразу к делу (теория потом). Для начала нужно установить Visual studio версии 2008, 2010 или 2012, а так же CUDA SDK. Всё нужное эти парни установят сами, и даже создадут папку с примерами. Поехали, самый первый пример!