cudaMemcpy слишком медленно

Question

Я использую cudaMemcpy() один раз, чтобы скопировать ровно 1 ГБ данных на устройство. Это занимает 5,9 с. Обратный путь занимает 5,1 с. Это нормально?
У самой функции столько копий перед копированием? Теоретически пропускная способность шины PCIe должна составлять не менее 4 ГБ / с.
Передачи памяти не перекрываются, потому что 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 с
0,197 с, первая копия
5,02 с, вторая копия
Странная вещь: несмотря на то, что для первой копии он отображает 0,197 с, для наблюдения за запуском программы требуется гораздо больше времени.

Answer 1

Да, это нормально. cudaMemcpy() выполняет много проверок и работает (, если память хоста была выделена обычным malloc() или mmap()). Следует проверить, что каждая страница данных находится в памяти, и переместить страницы (одну за другой) в драйвер.

Вы можете использовать cudaHostAlloc function или cudaMallocHost для выделения памяти вместо malloc. Он выделит закрепленную память, которая всегда хранится в ОЗУ и может быть доступна напрямую через DMA графического процессора (быстрее cudaMemcpy()). Ссылаясь на первую ссылку:

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

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

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

Answer 2

При условии точной синхронизации передач 1,1 секунды для передачи 1 ГБ из закрепленной памяти кажутся медленными. Вы уверены, что слот PCIe настроен на правильную ширину? Для полной производительности вам нужна конфигурация x16. Некоторые платформы предоставляют два слота, один из которых настроен как x16, а другой как x4. Так что, если у вашей машины есть два слота, вы можете попробовать переместить карту в другой слот. Другие системы имеют два слота, где вы получаете x16, если занят только один слот, но вы получаете два слота x8, если оба заняты. Настройка BIOS может помочь выяснить, как сконфигурированы слоты PCIe.

Tesla C870 - довольно старая технология, но, если я правильно помню, скорость передачи данных около 2 ГБ / с из закрепленной памяти должна быть возможной с этими частями, которые использовали интерфейс PCIe 1-го поколения. Современные графические процессоры класса Fermi используют интерфейс PCIe gen 2 и могут достигать 5+ ГБ / с для передач из закрепленной памяти (для измерений пропускной способности, 1 ГБ / с = 10 ^ 9 байт / с).

Обратите внимание, что PCIe использует пакетированный транспорт, и издержки пакета могут быть значительными при размерах пакетов, поддерживаемых общими наборами микросхем, при этом более новые наборы микросхем обычно поддерживают несколько более длинные пакеты. Один из них вряд ли превысит 70% от номинального максимума для каждого направления (4 ГБ / с для PCIe 1.0 x16, 8 ГБ / с для PCIe 2.0 x16) даже для передач из / в закрепленную память хоста. Вот технический документ, который объясняет проблему издержек и имеет удобный график, показывающий использование, достижимое с различными размерами пакетов:

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

Answer 3

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

Большинство Intel с несколькими сокетамиМатеринские платы класса i7 (Nehalem, Westmere) имеют один концентратор ввода / вывода на сокет.Поскольку системная память напрямую подключена к каждому ЦП, доступ к DMA, который является «локальным» (извлечение памяти из ЦП, подключенного к тому же IOH, что и у GPU, осуществляющего доступ DMA), намного быстрее, чем нелокальный (извлечение памяти из подключенного ЦПдля другой IOH - транзакция, которая должна быть выполнена через межсоединение QPI, связывающее два ЦП).

ВАЖНОЕ ПРИМЕЧАНИЕ: к сожалению, для SBIOS характерно конфигурировать системы для чередования, что вызывает непрерывное распределение памяти длячередоваться между розетками.Это смягчает потери производительности из-за локального / нелокального доступа для процессоров (один способ думать об этом: он делает все обращения к памяти одинаково плохими для обоих сокетов), но наносит ущерб доступу данных к графическому процессору, так как вызывает каждую другую страницу на-норозетная система должна быть нелокальной.

Похоже, что системы классов Nehalem и Westmere не страдают от этой проблемы, если система имеет только один IOH.

(Кстати, процессоры класса Sandy Bridgeсделайте еще один шаг по этому пути, интегрировав поддержку PCI Express в ЦП, поэтому в Sandy Bridge машины с несколькими сокетами автоматически имеют несколько IOH.)

Вы можете проверить эту гипотезу, выполнив тест с помощью инструмента.это прикрепляет его к сокету (numactl в Linux, если он доступен) или с помощью зависимого от платформы кода, чтобы управлять распределением и потоками для запуска на конкретном сокете.Вы можете многому научиться, не задумываясь - просто вызовите функцию с глобальными эффектами в начале функции main (), чтобы перегрузить все данные на тот или иной сокет, и посмотрите, сильно ли это повлияет на производительность передачи по PCIe.