cudaMemcpy твърде бавен

Използвам cudaMemcpy() веднъж, за да копирам точно 1 GB данни на устройството. Това отнема 5,9 секунди. Обратно отнема 5.1s. Това нормално ли е?
Самата функция има ли толкова много допълнителни разходи преди копиране? Теоретично трябва да има пропускателна способност от най-малко 4 GB/s за PCIe шината.
Няма припокриващи се прехвърляния на памет, защото Tesla C870 просто не го поддържа. Някакви съвети?

РЕДАКТИРАНЕ 2: моята тестова програма + актуализирани времена; Надявам се, че не е твърде много за четене!
Функциите cutCreateTimer() няма да се компилират за мен: 'грешка: идентификаторът "cutCreateTimer" е недефиниран' - това може да е свързано със старата версия на cuda (2.0), инсталирана на машината

 __host__ void time_int(int print){
static struct timeval t1; /* var for previous time stamp */
static struct timeval t2; /* var of current time stamp */
double time;
if(gettimeofday(&t2, 0) == -1) return;
if(print != 0){
  time = (double) (t2.tv_sec - t1.tv_sec) + ((double) (t2.tv_usec - t1.tv_usec)) / 1000000.0;
  printf(...);
}
t1 = t2;
}

main:
time(0);
void *x;
cudaMallocHost(&x,1073741824);
void *y;
cudaMalloc(&y, 1073741824);
time(1);
cudaMemcpy(y,x,1073741824, cudaMemcpyHostToDevice);
time(1);
cudaMemcpy(x,y,1073741824, cudaMemcpyDeviceToHost);
time(1);

Показаните времена са:
0,86 s разпределение
0,197 s първо копие
5,02 s второ копие
Странното е: въпреки че показва 0,197 s за първо копие, отнема много повече време, ако гледам как програмата работи .


person Callahan    schedule 15.09.2011    source източник
comment
Можете ли да добавите малко информация към въпроса си за това как правите измерванията на времето?   -  person talonmies    schedule 15.09.2011
comment
@talonmies: описаното време в редакцията   -  person Callahan    schedule 15.09.2011


Отговори (3)


Да, това е нормално. cudaMemcpy() прави много проверки и работи (ако паметта на хоста е разпределена от обичайните malloc() или mmap()). Трябва да провери дали всяка страница с данни е в паметта и да премести страниците (една по една) към драйвера.

Можете да използвате cudaHostAlloc функция или cudaMallocHost за разпределяне на памет вместо malloc. Той ще разпредели закрепена памет, която винаги се съхранява в RAM и може да бъде достъпна директно от DMA на GPU (по-бързо cudaMemcpy()). Цитирайки от първия линк:

Разпределя броя байтове от паметта на хоста, която е заключена за страници и достъпна за устройството. Драйверът проследява обхватите на виртуалната памет, разпределени с тази функция, и автоматично ускорява извикванията към функции като cudaMemcpy().

Единственият ограничаващ фактор е, че общото количество фиксирана памет в системата е ограничено (не повече от RAM размер; по-добре е да използвате не повече от RAM - 1Gb):

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

person osgx    schedule 15.09.2011
comment
comment
Добре, качването (на устройство) падна до 1,1 секунди, но изтеглянето (от устройство) все още е на 5 секунди. Това също ли е нормално? - person Callahan; 15.09.2011
comment
Изтеглянето също ли е в Pinned memory? - person osgx; 15.09.2011
comment
да В този тест изтеглям в същата памет, от която качих. - person Callahan; 15.09.2011
comment
Калахан, можеш ли да публикуваш подробности за кода? редовете с: разпределение на паметта; качване; стартиране на ядрото (и източника на ядрото); изтегляне, времена. Правилният метод за определяне на времето за Cuda е да се използват таймерите на Cuda, да се проверяват примерни кодове, в комплект с CUDA. Или тук: cs.virginia.edu/~csadmin/wiki /index.php/CUDA_Support/ - person osgx; 15.09.2011
comment
Не мога да се съглася, че тези показатели за производителност при копиране са нормални, освен ако картата не е в много бавен слот. Всички G80, които използвах, бяха в състояние да достигнат около 1,8 Gb/s при големи трансфери в 16-лентов PCI-e слот, може би 2,5 Gb/s към фиксирана памет. Отчетеното представяне е само около 10% от това, което е изключително ниско. - person talonmies; 15.09.2011
comment
Публикува съответния код. Досега не съм използвал ядро. Както е описано в редакцията, предложените времена на Cuda не работят за мен. - person Callahan; 16.09.2011

Ако приемем, че прехвърлянията са синхронизирани точно, 1,1 секунди за прехвърляне на 1 GB от фиксирана памет изглежда бавно. Сигурни ли сте, че PCIe слотът е конфигуриран на правилната ширина? За пълна производителност бихте искали конфигурация x16. Някои платформи предоставят два слота, единият от които е конфигуриран като x16, а другият като x4. Така че, ако вашата машина има два слота, може да опитате да преместите картата в другия слот. Други системи имат два слота, където получавате x16, ако е зает само един слот, но получавате два слота x8, ако и двата са заети. Настройката на BIOS може да помогне да разберете как са конфигурирани PCIe слотовете.

Tesla C870 е доста стара технология, но ако си спомням правилно скорости на трансфер от около 2 GB/s от фиксирана памет би трябвало да са възможни с тези части, които използват PCIe интерфейс от първо поколение. Настоящите графични процесори от клас Fermi използват интерфейс PCIe gen 2 и могат да постигнат 5+ GB/s за прехвърляне от фиксирана памет (за измервания на пропускателна способност, 1 GB/s = 10^9 байта/s).

Обърнете внимание, че PCIe използва пакетизиран транспорт и натоварването на пакетите може да бъде значително при размерите на пакетите, поддържани от обикновените чипсети, като по-новите чипсети обикновено поддържат малко по-дълги пакети. Малко вероятно е човек да надхвърли 70% от номиналния максимум за посока (4 GB/s за PCIe 1.0 x16, 8 GB/s за PCIe 2.0 x16), дори за трансфери от/към фиксирана хост памет. Ето бяла книга, която обяснява проблема със служебните разходи и има удобна графика, показваща постижимото използване с различни размери на пакети:

http://www.plxtech.com/files/pdf/technical/expresslane/Choosing_PCIe_Packet_Payload_Size.pdf

person njuffa    schedule 15.09.2011
comment
Какво можете да кажете за бавното изтегляне (карта-›хост), дори и за фиксирана памет? - person osgx; 16.09.2011
comment
За съжаление няма да стигна нито до машината, нито до биоса, защото е дистанционна, върху която работя. - person Callahan; 16.09.2011

Освен система, която просто не е конфигурирана правилно, най-доброто обяснение за ужасната честотна лента на PCIe е несъответствие между IOH/сокет и PCIe слота, в който е включен GPU.

Повечето дънни платки Intel i7-class (Nehalem, Westmere) с множество гнезда имат един I/O хъб на гнездо. Тъй като системната памет е директно свързана към всеки CPU, DMA достъпите, които са "локални" (извличане на памет от CPU, свързан към същия IOH като GPU, извършващ DMA достъп) са много по-бързи от нелокалните (извличане на памет от CPU, свързан към другия IOH, транзакция, която трябва да бъде изпълнена чрез QPI връзка, която свързва двата процесора).

ВАЖНА ЗАБЕЛЕЖКА: за съжаление е обичайно за SBIOS да конфигурират системи за подреждане, което води до разпределение на непрекъсната памет между гнездата. Това смекчава сривовете в производителността от локален/нелокален достъп за процесорите (един начин да го мислим: прави всички достъпи до паметта еднакво лоши и за двата сокета), но сее хаос с достъпа на GPU до данните, тъй като причинява всяка друга страница на 2 -сокет системата да не е локална.

Системите от клас Nehalem и Westmere изглежда не страдат от този проблем, ако системата има само един IOH.

(Между другото, процесорите от клас Sandy Bridge правят още една стъпка по този път, като интегрират поддръжката на PCI Express в процесора, така че с Sandy Bridge машините с множество гнезда автоматично имат множество IOH.)

Можете да проучите тази хипотеза, като изпълните теста си с помощта на инструмент, който го прикрепя към сокет (numactl на Linux, ако е наличен), или като използвате зависим от платформата код, за да управлявате разпределенията и нишките да се изпълняват на конкретен сокет. Можете да научите много, без да се измисляте - просто извикайте функция с глобални ефекти в началото на main(), за да принудите всичко към един или друг сокет и вижте дали това има голямо влияние върху производителността на вашия PCIe трансфер.

person ArchaeaSoftware    schedule 16.09.2011
comment
Изпълних cudaSetDevice(0); в началото като функция за глобален ефект. Времената не се промениха. - person Callahan; 18.09.2011