NVidia 8800GTX: пропускная способность памяти (при использовании CUDA)

После чтения руководства по NVidia CUDA, остается ощущение сложности модели программирования: треды, блоки тредов, warp-ы, иерархическая память. Непонятно, какие параметры вычислительной задачи оптимальны и какие у них вообще допустимые значения. Само руководство точных рекомендаций не дает, дает лишь приблизительные.

Из общих соображений, понятно что самая медленная часть суперкомпьютера - память. С одной стороны, теоретическая пропускная способность (bandwidth) составляет 900MHz * 384 бита * 2 (DDR) = 86.4 GB/sec. С другой стороны, раздел 6.1.1.3 руководства говорит о 200-300 циклах memory latency (при, по всей видимости,случайном доступе).

К счастью, проблема легко изучается: если взять достаточно много данных (скажем, полгигабайта) и, например, сложить все 4-байтовые значения (как float), то основные затраты времени будут именно на чтение из памяти, а всей прочей арифметикой можно пренебречь (или подсчитать ее отдельно).

Соображения

  1. Чтения из global memory могут быть 64- и 128-битными (раздел 6.1.2.1), руководство по программированию рекомендует выравнивать чтения блока тредов (точнее, warp-а) на 16*sizeof(type) байт - в этом случае обещается одновременная загрузка всех данных для всех тредов.
  2. Количество threads в блоке рекомендовано иметь 64, а лучше 192 или 256 (раздел 6.2). При этом:
    • размер warp-а (набора одновременно выполняемых threads) для Geforce 8800 равен 32 threads
    • блоки threads могут выполняться на одном мультипроцессоре одновременно (как я понимаю, целым числом блоков)
    • задержки при read-after-write регистра не чувствуются (прячутся scheduler-ом) при числе тредов на мультипроцессоре >192.
    • максимальное количество threads на мультипроцессор (одновременно): 768.
    • выбирая количество threads в блоке следует учитывать использование локальной shared memory и локальных регистров. В частности, размер Register File на мультипроцессор составляет 32 килобайта, следовательно для запуска 768 threads (максимальное количество) использование регистров в kernel должно быть не более 10 (регистры 4-байтовые). Посмотреть использование регистров можно в файле с расширением .cubin, который останется в каталоге компиляции, если компилятору дать ключик -keep (там же образуется и крайне интересный файл .ptx с ассемблерным листингом).
  3. Количество thread blocks (CTA) должно быть достаточно большим. Руководство по программированию говорит о "не менее 100" (раздел 6.2), а "1000 будут достаточны на несколько поколений оборудования". Сотрудники NVidia в личной переписке более конкретны:
    • CTA должно быть не меньше чем мультипроцессоров (которых на 8800GTX 16 штук).
    • А лучше больше чем мультипроцессоров или(и?) кратно их числу.
    • Не меньше общего числа CTA, способных одновременно исполняться (см. выше: количество одновременно исполняемых CTA определяется register usage и общим числом threads).
    • Кратно числу CTA, способных одновременно исполняться
    При этом утверждается, что от лишних thread blocks вреда нет

Сама по себе модель исполнения на G80 — "много тредов в блоке, много блоков и все это как-то диспетчеризуется с возможностью синхронизации только внутри блока" — в голове укладывается плохо, нужны эксперименты.

Тестовая задача

Возьмем "квадратный массив", влезающий в память видеокарты. Желательно, чтобы размер был бы кратен степени двойки. Размер 12288x12888 всем хорош: это 600 мегабайт данных (для типа float или int), размер кратен 212. Проинициализируем массив переменным паттерном чисел по порядку величины около 1:
  1. for(int i=0;i < SIZE*SIZE; i+=4)
  2. {
  3.  data[i]= 1.0+(SIZE*SIZE-i-1)/((float)SIZE*SIZE);
  4.  data[i+1]= 1.0-2.0*i/((float)SIZE*SIZE);
  5.  data[i+2]= 1.0+3.0*(SIZE*SIZE-i-1)/((float)SIZE*SIZE);
  6.  data[i+3]= 1.0-2.0*i/((float)SIZE*SIZE);
  7. }
Сумма по всем элементам массива равна SIZE*SIZE.

Складывать все числа массива можно разными способами:

  • Порядок обхода:
    • по строкам
    • по столбцам
  • Выборка данных:
    • 4 байта (одно float число)
    • 16 байт (один вектор float4)
Для обхода по строкам и выборки по 4 байта получается такой тривиальный код, выполняемый на GPU:
  1. __global__ void
  2. Sum_h(int run,FTYPE *g_idata, float *g_odata)
  3. {
  4.         const unsigned int blocks = gridDim.x;
  5.         const unsigned int threads = blockDim.x;
  6.         const unsigned int tid = threadIdx.x;
  7.         const unsigned int bx = blockIdx.x;  
  8.         unsigned int rowN,colN;
  9.         float sum = 0.0;
  10.  
  11.         for(rowN=bx; rowN < SIZE; rowN+=blocks){
  12.                 for(colN=tid; colN < SIZE; colN+=threads){
  13.                         sum += g_idata[rowN*SIZE+colN];
  14.                 }
  15.         }
  16.         g_odata[bx*threads+tid]=sum;
  17. }
Каждый thread block обходит несколько строк. Внутри строки, каждый thread суммирует несколько элементов. Результаты суммирования thread-ом записываются в выходной массив g_odata, суммирование элементов этого массива производится уже на CPU. Если заменить тело цикла на sum+=1.0, можно оценить время на исполнение без доступа к глобальной памяти. Вычитать его целиком из времени выполнения тестового кода нельзя (т.к. многие сотни-тысячи threads выполняются одновременно и долгие операции идут "в фоне"), но использовать как оценку снизу времени выполнения - можно.

Перед исполнением тестового кода выполняется вызов пустого kernel (с одним небольшим циклом и записью результатов в выходной массив), время исполнения которого (это 0.03-0.5 миллисекунды, в зависимости от количества блоков и тредов) вычитается из времени исполнения измеряющего вызова.

Первые результаты

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

NVidia 8800GTX: скорость чтения из глобальной памяти (Гбайт/сек)
Тип данныхпорядок обхода
по строкампо столбцам
FLOAT70.520.92
FLOAT438.833.73

Как видим, чтение по 4 байта по строкам дает 82% скорости от теоретического максимума, прекрасный результат как по абсолютному значению (на порядок быстрее того, что удается намерять на Intel Woodcrest и в ~7 раз быстрее memory bandwidth для Opteron), так и в доле от теории (у того же Woodcrest теоретическая полоса около 21Gb/sec, а практическая - около шести).

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

Цифры для FLOAT4 объяснимы лишь частично:

  • доступ по столбцам вчетверо лучше, чем для единичного FLOAT. Все совершенно прозрачно, загрузка 32 бит и 128 бит занимают одинаковое время, но во втором случае читается вчетверо больше.
  • Для чтения по строкам: я грешу на большую потребность thread-а в регистрах (4 регистра на чтение промежуточных значений и до 3 регистров на промежуточные результаты сложения), отчего количество одновременно исполняемых на мультипроцессоре threads должно было упасть.

Предварительные выводы

  1. При последовательном доступе к памяти выгодно не выпендриваться и читать не больше данных, чем нужно.
  2. Неоптимальный (по столбцам) доступ к памяти крайне нежелателен, скорость доступа снижается почти на два порядка.
  3. Если доступ по столбцам все-таки нужен, то выгоднее читать по 128 бит за чтение (т.е. делать это до такой степени "по строкам", до какой это возможно).

Оптимальные параметры мультизадачности

Для рассматриваемой задачи мы можем менять два параметра:
  • Количество threads в блоке (рекомендуют кратно 32).
  • Количество блоков

Количество threads в блоке

Если не вдаваться сильно в детали, то для 8800GTX начиная с 32 блоков и 256 threads в блоке, либо с 64 блоков и 128 threads мы достигаем практически максимальной скорости.
16 threads - далеко от оптимума (разница в 3-4 раза от оптимума), 32 threads - тоже недостаточно (разница в 1.5-2 раза).

Максимальная скорость была достигнута для 192 threads и 1024 blocks. Максимум на 384 threads (почти такой же, как для 192) показывает, что удалось загрузить мультипроцессоры полным количеством потоков.

Если запускать тестовую задачу с числом потоков, меняющихся на 1 (от 16 до 512), то становится понятной важность выравнивания хоть на что нибудь (картинка слева).

Для количества threads, некратного 16-ти, пропускная способность очень плохая (в сравнении с легко достижимым при правильном выравнивании идеалом) - порядка 10Gb/sec вместо 65-70.
Bandwidth кое-как подрастает для числа threads кратного 4-м (что соответствует выравниванию на 16 байт) и заметно подрастает для количества потоков, кратного 8-ми. Между кратностями 16 и 32 тоже есть заметная разница (можно сравнить, например, 96-112-128 или 256-272-288), но на мой взгляд она уже определяется оптимальностью диспетчиризации по 32 thread.

Весьма интересен резкий спад в производительности сразу после 384 thread, он отвечает ситуации, когда на одном мультипроцессоре начинает работать только один CTA (до 384 включительно их было два и более).

Количество блоков threads

Предпочтения по выбору количества блоков (CTA) описаны выше, но несложно оценить реальное влияние количества thread blocks на скорость работы.

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

Из графика видно, что при числе блоков до ~1024 наблюдаются заметные скачки производительности в зависимости от количества CTA.

Если рассмотреть левую часть графика более детально, то видно что для небольшого количества CTA правильный выбор числа блоков весьма важен, производительность скачет на 15-20%. Чем больше блоков, тем меньше скачки, что и ожидалось.

Помимо пилообразного графика, мы видим пики максимальной производительности для 256, 512, 640, 768 блоков.

Заключение

По мере уменьшения степени влияния на производительность чтения из глобальной памяти:
  • Наиболее важным для получения большой производительности является порядок чтения. Обход по столбцам (т.е. фактически псевдослучайная выборка) крайне неэффективен, потери производительности в сравнении с обходом по строкам составляют почти два десятичных порядка. Правильный порядок обхода позволяет получить скорость чтения более 70 Gb/sec, что крайне неплохо в сравнении с теоретически-возможным значением.
  • Неправильное количество threads в CTA ведет к неправильному выравниванию доступов к данным и, как следствие, к провалу производительности в разы. Количество threads в CTA желательно иметь кратным 32-м, а лучше 64-м, при этом равным 192 и более. Оптимальное количество threads, конечно, ограничивается использованием регистров.
  • Использование 128-битных типов данных ведет к 1.5-2-кратному провалу в производительности при чтении по строкам и к 4-кратному росту производительности при чтении по столбцам. Таким образом, при необходимости читатьпо столбцам float4 может немного помочь.
    В принципе, чтение по столбцам должно быть сильно быстрее при чтении текстур: там есть 2D-кэш, оптимизированный под такое чтение. Но это - тема отдельного исследования т.к. быстро совладать с текстурами не получилось

Оптимальные значения

  1. Много CTA - хорошо. Причем их количество хорошо бы иметь пропорциональным какой-нибудь степени двойки. В рассматриваемой задаче максимальная производительность достигнута при 256, 512, 640 и 768 блоках.
  2. Много threads - тоже хорошо. Для threads крайне важно, чтобы доступ к глобальной памяти был выровнен, количество threads в блоке полезно иметь кратным 32-м, для сокрытия латентности регистров полезно иметь более чем 192 thread на мультипроцессор (это достигается или увеличением числа thread в блоке, либо увеличением числа одновременно исполняемых CTA). Наилучшая производительность была получена для 192, 256, 320, 384 threads per CTA.

Tags: