Пользователь
0,0
рейтинг
22 января 2009 в 20:15

HPC: nVidia, AMD, Sony Cell, x86

Многие слышали о мифических словах — nVidia CUDA, AMD Brook, IBM/Sony Cell… Ходят слухи, что одно лишь упоминание оных заставляет вашу программу работать в сотни раз быстрее. Попробуем разобраться, что они из себя представляют, как выглядит это магическое высокопроизводительное программирование в общих чертах, и какой выигрыш они могут обеспечить в сравнении со старыми добрыми процессорорами x86.


Увлекся я недавно высокопроизводительными вычислениями (язык сломать можно, HPC — High-performance computing), и хочу поделится кратким обзором, как обстоят сейчас дела с программированием и производительностью различных аппаратных платформ, а именно nVidia CUDA, AMD Brook+, Sony/Toshiba/IBM Cell (основная мощь PS3), ну и для сравнения — старичок x86 в лице Core2Quad/i7, AMD Phenom/Phenom II. Сразу хочу отметить, что тут много нюансов влияющих на производительность — размеры и латентности кэшей разных уровней, совмещаемость команд различных типов и изложить все в одной статье никак не выйдет, так что сейчас я покажу ситуацию в общих чертах.

Разговор везде идет про 32-х битные числа. Для оценки: 64-х битных числах производительност GTX285 падает примерно в 10-12 раз, AMD 4870 — в 4 раза, x86 и Cell- в 2 раза.

На затравку сразу показываю результат:

nVidia GTX285 AMD 4870 Sony Cell X86 i7 3.2Ghz
Реальная скорость, GFLOP* 500 600 153 153
Площадь кристалла 470mm2 256mm2 120mm2 263mm2
Техпроцесс 55nm 55nm 65nm 45nm
*GFLOP — миллиардов операций в секунду

P.S.: Cell на 45nm еще как следует
не запустили, и в соньках он не используется
пока.

nVidia CUDA


nVidia пожалуй больше всех усилий прикладывает к продвижению своей технологии. Волшебное слово CUDA я услышал задолго до того, как решил попробовать, что это такое. CUDA — Это C++ компилятор (достаточно прямой и стабильный), который специально написанную программу компилирует для исполнения SIMD процессорами внутри видеокарты. Например в GTX285 — 240 процессоров, которые делают по 1-3(умножение+сложение+ специальная, в графических шейдерах 2-3 реально, но в обычных вычислениях лучше ориентироваться на 1.5) операции за такт на частоте 1476Mhz, маркетинговая производительность чуть выше 1 TFLOP (триллион операций в секунду), реальная достижимая скорость около 500 GFLOP.

Программа выглядит примерно так:

__device__ unsigned int *data_d;
__global__ void gpu_thing(int somevalue)
{
const
int ix = blockDim.x * blockIdx.x + threadIdx.x;

data_d[ix]
*= somevalue;
}

Запускается на GPU так:

dim3
threads(256);
dim3 grid(256);
int
src[256*256], result[256*256]; //копируем данные в память видеокарты
cudaMemcpyAsync(data_d, src, sizeof(int)*256*256, cudaMemcpyHostToDevice,0); //Запускаем на видеокарте, 256*256 потоков
одновременно
gpu_thing <<<grid, threads>>>( 125); //Забираем результат из памяти
cudaMemcpyAsync(result, data_d, sizeof(int)*256*256, cudaMemcpyDeviceToHost,
NULL);

Ограничения:
Потоки выполняются группами по 16, все инструкции должны быть одинаковыми(иначе снижается производительность, за этим следит компилятор/железо).

Время запуска потоков — не 1мс конечно, но заметно, поэтому мелкие задачи на видеокарту лучше не перекладывать. Во время работы программы
на видеокарте, к которой подключен монитор, windows не может обновить экран, даже курсор мыши не двигается — проблема решается подбором объема работы на 10-20мс, чтобы не слишком плохо было пользователю. Если время работы ядра превышает 5 секунд, Windows решает что драйвер завис, и перезапускает его — результаты плачевны для вашей программы (опять же, это справедливо только если к видеокарте подключен монитор) :-) Linux и MacOS поддерживаются. Поддержка нескольких видеокарт — работает, только если отключен режим SLI. Поддерживает все видеокарты начиная с 8400GS (хотя там обычно медленнее процессора получается). Формально видеокарты разного типа работать не обязаны, но на практике работают любые комбинации (требует ручного разделения задач).

AMD Brook+


AMD постаралась на славу с железом — её 4870 имеет 160 процессоров, каждый из которых может делать по 5 операций (с незначительными ограничениями) за такт на частоте 750 Mhz. Теоретическая и практическая производительность — 600GFlop. Фанаты AMD — следующий абзац не читайте, но знайте — на реальных приложениях (если их удается написать) 4870 делает GTX285 на 10-20%. Ну а теперь о плохом: C Brook+ просто вилы. Нужно быть очень спокойным человеком чтобы с ним работать. Компилятор С(без ++), а потому знакомимся заново как с давно забытыми было ограничениями (вроде определения локальных переменных только в начале тела функции), так и с совершенно незнакомыми: локальные массивы не поддерживаются, структуры поддерживаются, но компилятор практически никогда не может нормально это скомпилировать (крэшится / выдает
неправильный код — это справедливо для Stream SDK 1.3). В общем, приготовьтесь что версия называется 1.3 а работает как 0.3alpha :-D Несмотря на то, что у нас архитектура VLIW и было бы полезно уметь разворачивать циклы (например как в Intel C++ — pragma unroll(5)), сделать этого нельзя.
Да, препроцессора тоже нет, прикручивайте свой. Ну и завершает все невозможность работать с несколькими видеокартами (вернее один процесс(не тред) — одна видеокарта). Стоит отдать должное AMD CAL — программирование на низком уровне (не машинный код, но почти) — там проблем почти нет, но и недостатки очевидны. Программа выглядит примерно так:

kernel void gpu_thing(int input<>,
int somevalue, out int output<>)
{

output
= somevalue*input;

unsigned int md5_dim[] = {800*256};
::brook::Stream<int> input(1,size);
::brook::Stream<int> result(1,size);

int src[800*256], dst[800*256];

input.read(src); //готовим данные для видеокарты
hello_brook_check(input,45,output); //запускаем 800*256 потоков
output.write(dst);

Ограничения: время запуска тоже значительное, опять все замирает кроме курсора мыши, опять 5сек максимальное время выполнения под windows.
Железо поддерживается HD 2xxx (с ограничениями), HD3xxx и 4xxx.

Sony Cell


Многие в восторге от производительности Sony Cell. Она конечно впечатляет, если забыть о том, что x86 тоже продвинулся (не говоря уже о видеокартах). Cell состоит из обычного Power-ядра называемого PPE (с незначительной производительностью по сравнению с вышеупомянутыми монстрами) и 8-и ядер SPE, из которых 1 — потенциально бракованое ("небраковыные" чипы продаются отдельно за гораздо большие
деньги), и 1 зарезервировано операционной системой (если перешить приставку на Linux можно получить доступ к 7, но это для конечных пользователей трудновато ;-) ), итого остается 6. Каждое из ядер делает 8 операций за такт(2*4) на частоте 3.2Ghz, что дает общую реальную
производительность 153.6 GFLOP. С программированием тут не подскажу, т.к. эти 153 GFLOP достаточно непрактично использовать, потому и не
пробовал (с точки зрения выпуска программы не для внутреннего потребления).

x86


Уже сбросили со счетов x86? А ведь он еще может тряхнуть стариной: На i7/Core2Quad каждое из 4-х ядер может выполнять 3 SSE2 операции по 4 числа за такт, итого 48 операций за такт (а ведь еще относительно недавно считалось что преодолеть 1 операцию за такт сложно/невозможно, и начинали идти в сторону VLIW/RISC), что дает нам 153.6 GFLOP (опс, ровно столько же сколько и Cell ) Phenom/PhenomII тоже может, но у него модули SSE2 специализированные (одно умножение, одно сложение, одно универсальное), поэтому производительность оказывается на 20-40% ниже. Если кому-то кажется что использовать 48 операций на такт трудно, вот вам пример:

int data[1024*1024*12];
const int value1 = 123;
for(int i=0;i<1024*1024*12;i++)data[i]^=value1;

Если это скомпилировать в Intel C++, и запустить в 4 потока (впрочем, он сейчас иногда даже может сам догадаться потоки запустить) — он сам объединит операции в группы, удобные для SSE2 (стоит признать, что не всегда у всех проходит гладко, можно подсказывать многочисленными #pragma-ми или использовать SSE-intrinsics, но это тема отдельной статьи)

Какой можно сделать вывод? Безусловно, видеокарты нынче показывают феноменальную производительность на достаточно узкой категории задач (особенно двухчиповые), примерно в 8 раз быстрее хорошего кода на обычном x86 процессоре. 8 это далеко не 30-100 как рекламируют, но всё же достаточно, чтобы иногда сделать невозможное возможным. Кроме того, материнская плата под 4 видеокарты стоит намного дешевле платы под 4 процессора. Эйфория же по поводу производительности Cell в PS3 уже должна спадать, раз уж серийные x86 процессоры догнали его по производительности.
­
А еще парсер хабра не хочет валидно отображать тег <сode>, так что я его уберу. :(
Memento Mori @Alaunquirie
карма
23,5
рейтинг 0,0
Реклама помогает поддерживать и развивать наши сервисы

Подробнее
Реклама

Самое читаемое

Комментарии (66)

  • НЛО прилетело и опубликовало эту надпись здесь
    • –1
      Регистры 128-и битные, влазят 4 single или 2 double
      • НЛО прилетело и опубликовало эту надпись здесь
  • –14
    +1
  • +1
    Intel не позволит кануть x86. Вспомните про Larrabee, который поддерживает и базируется на x86 инструкциях.
  • +2
    эти 153 GFLOP достаточно непрактично использовать, потому и не
    пробовал (с точки зрения выпуска программы не для внутреннего потребления).

    Что имеется в виду?
    • +1
      Доступ к ним тяжеловато получить.
      • 0
        O RLY?
        Не понимаю. Если это готовый сервер или акселератор с интерфейсом PCI-E, то вот они, ресурсы. Если это PS3, то ставим Yellow Dog (или какой-нибудь другой дистрибутив GNU/Linux) и делаем свои расчеты на всех доступных SPE.
        Во всех случаях для программирования используются одни и те же библиотеки/фреймворки, так что разработка отличается не очень кардинально.
        • 0
          Для конечного пользователя релизнуть трудно.
          Объяснять простому юзеру как ему надо на приставку линух ставить — достаточно трудновато.

          Ну а о цене акселей с интерфейсом PCI-E лучше помолчим, чтобы не расстраиваться зря :-)
          • 0
            Речь о HPC. Зачем это «простым юзерам»?
            Простому юзеру достаточно купить PS3 и интересующую его игру, но мы не о том.
            • +1
              На случай если пишеться клиент для распределенных вычислений типа Folding@Home.
              Фолдингу то хорошо, с помощью сони они все без Линуха сделали, и запускать конечным пользователям легко.
              К сожалению для простых смертных этот пусть закрыт :-)
      • +1
        Когда речь идёт о пиковой производительности то получить её на х86 не намного легче чем на CELL.
        Умножение больших матриц даёт на полноценном CELL ожидаемые 200GFlops.
  • 0
    Строка «Реальная скорость, GFLOP» не несёт практической смысловой нагрузки, т.к. не указано каких операций. Тем более у разных операций различное время выполенения.
    • 0
      FLOPS — FLoating point Operations Per Second (операций с плавающей точкой(одинарной точности наверно))

      так что все очень даже понятно: например надо перемножить 2 матрицы NxN, считаешь, сколько надо операций, и считаешь приблизительное время
      • 0
        Вы не правы. Каких операций не указано, это может быть вполе суммирование.
        А умножение матриц требует умножение и аккумулирование — эти операции, к примеру могут выполянтся 10 тактов, а суммирование 1 или 2 такта.
        По этому FLOPS довольно абстрактная хар-ка.
        • 0
          я с вами согласен. тут надо уточнить что считается за одну операцию.
          • 0
            Большинство операций на GPU выполняются за один такт(иногда правда со сниженной точностью) — вот именно такие и учитываются.
  • +1
    Побольше таких статей, если можно :)
    • +1
      Можно ;)
      • 0
        А можно немного наглядных графиком с вычислением чего-нибудь одного?
        Например решения судоку?
  • +2
    Это что-то не так, если при работе CUDA не работает ОС :) (Максимум могут быть проблемы, если меняется графический режим, и новый будет требовать больше памяти, чем предыдущий, а вся свободная использована под приложение.)
    Безусловно, видеокарты нынче показывают феноменальную производительность на достаточно узкой категории задач (особенно двухчиповые), примерно в 8 раз быстрее хорошего кода на обычном x86 процессоре.
    Откуда такая информация? Прирост вполне может составлять 10—50 раз. Если конечно не сравнивать «плохой» код на GPU с «хорошим» кодом CPU.

    А вообще неплохо бы код раскрасить да лишние переносы строк убрать…
    • 0
      Хабрапарсер со статьей чет мудрит, продолжение в хабраредакторе буду ваять.
  • НЛО прилетело и опубликовало эту надпись здесь
    • 0
      В синтетических — да, из одночиповых — на практике — 295ая нвидиа
      • 0
        295 — двухчиповая :]
    • 0
      А по себе скажу, что 4870 страдает от нехватки нормальных дров. Очень.
    • 0
      Это спорно, к тому же заточка многих игр под nVidia нивелирует незначительный прирост в производительности. Ну а вычисления на GPU пока массово идут под брендом CUDA.
      • 0
        Нет заточки под nVidia. Mean to be played — всего лишь шильдик, который значит, что если видяха в списке поддерживаемых есть, и компьютер соответствует минимальным требованиям — оно запустится.
        • 0
          Эта программа и есть для того, чтобы разработчики затачивали игры под GeForce, вообще-то.
          en.wikipedia.org/wiki/The_way_it's_meant_to_be_played
          • 0
            Их просто_проверяют_на_совместимость. Есть аналогичный шильдик и у ATI, просто они отказались от использования, причину не помню, если найду статью — покажу. Те игры, которые реально лучше работают на ATI маркируются соответствующим красно-белым клеймом.
  • 0
    Кстати, а куда делать платформа OpenCL, которую помойму AMD и продвигает, как основного соперника CUDA?
    • 0
      AMD has decided to support OpenCL (and DirectX 11) instead of the now deprecated Close to Metal in its Stream framework.
      • 0
        АМД решила продвигать OpenCL, вместо CTM. Я об этом и говорю. Или от её слов до дела ещё не дошло?
        • 0
          На деле пока все крайне странно себя ведет, это раз, два — CTM — это как раз реализация низкоуровневых команд. (Даже название об этом говорит)
          • 0
            Вы, видимо не верно поняли мой вопрос.

            Я говорю про то, что вы использовали Brook+ и даже похвалили, что в нём реализованны CTM.

            Хотелось увидеть в статье (возможно продолжение этой) OpenCL, как альтернативу CUDA и Broom+. Если можно, конечно.
            Собираюсь просто этим заниматься скоро — OpenCL сейчас вижу, как достаточно продуманную кроссплатформенную (Radeon+Geforce) альтернативу.
  • 0
    Хороший обзор. Скажите, а чем предполагается компилировать код для Phenom?
    • 0
      Intel С++ с соответствующими настройками
  • 0
    Программы, работающие по 5 секунд — это ещё очень далеко от HPC. Может они и высокопроизводительные, но много насчитать не очень-то получится… Раз уж у вас есть под рукой необходимое оборудование, стоило бы проверить производительность на более приближенных к жизни операциях, например двумерном Фурье-преобразовании матрицы размером 8192x8192 :)
    • 0
      5 секунд — это всего лишь время через которое ОС посчитает что вычисления захватили ресурсы видеокарты и снимет процесс.
      Если покупается GPU для специальных рассчетов, то монитор к ней подключать не следует, и в таком случае Windows watchdog не срабатывает, можно считать сколько нужно.
      • 0
        Обойти можно всё, это понятно :) Но это лишний камень в огород CUDA, пресс-релизы которой пестрят заявлениями о том, как же с ней удобно работать даже на офисных PC. Особенно, если не покупать новые топовые карты, а пробовать работать с той, которая стоит на рабочем компьютере.
        • 0
          5 секунд — это лимит на исполнение одного kernel'я, а в консьюмерских задачах кернели отрабатывают зачастую за миллисекунды.
          И никаких камней нет, потому что ATI имеет те же проблемы.
          • 0
            Я имел в виду сравнение не с ATI, а с вычислительными системами на основе CPU :) Не спорю, что есть много задач, где вычисления занимают мало времени, просто я их, видимо, не отношу сильно к области HPC, которая была заявлена в заголовке статьи. Либо я не до конца понимаю терминологии, чтож, буду читать…
    • +1
      Классическое умножение матриц 6000x6000: www.gpgpu.ru/articles/sgemm-7.html
      • 0
        Спасибо за ссылку. Намного более адекватный результат, превосходство GTX280 над Core2Quad 9300 всего полтора раза. И это для перемножения матриц, которое для GPU идеально подходит.
        • 0
          Так написано же, что на double производительность в 10 раз падает. Смотрите float:
          Blas — 87, Cuda — 372. В 4.3 раза быстрее чем Quad, в 17 раз быстрее одного ядра.

          Думаю, быстрое преобразование фурье может дать больший разрыв во floate за счёт хорошего синуса в Nvidia.
          • 0
            Мне кажется сравнивать с одним ядром уже давно пора перестать :)

            В том-то и дело, что в реальных считательных задачах не всегда достаточно float. Если кратко, то этого достаточно для видео, графики, игр, что и неудивительно, но не стоит экономить на точности при расчёте ядерных реакторов. Фурье само по себе имеет тенденцию терять точность даже при преобразовании туда-обратно, то делать это во float не очень хорошо.
  • 0
    Что то FLOPS в вашей табличке для чипов сильно расходятся с тем что приведено в en.wikipedia.org/wiki/FLOPS
    Это потому что у вас «реальные»? :-)
    • +2
      Да, а там — теоритический максимум, практически нереально
  • 0
    На самом деле пиковая производительность x86 i7 @ 3GHz(Nehalem):
    a) Single precision: 8*3Ghz = 24GFlops per core, 24*4 = 96GFlops per CPU;
    b) Double precision: 4*3GHz = 12GFlops per core, 12*4 = 48GFlops per CPU.

    Смотрите например тут:
    icl.cs.utk.edu/hpcc/hpcc_results.cgi

    Подтест DGEMM — как раз этот случай. Высокая локальность, чисто Flops-ы :)
    S-DGEMM — замер на одном процессе, EP-DGEMM — замер на всех ядрах, усредненный.

    Вот, а «Реальная скорость в 153 GFLOPS» — это от лукавого…
    • +1
      В этой статьи не рассматриваются вопросы связанные с падением производительности из-за памяти/кеша. К указанной в таблице цифре я приближался на всех платформах кросе Cell. На x86 все помещалось в регистры и немного L1 кеша, потому и скорость 150 а не 96.
      • 0
        Уважаемый — я вам даю теоретический пик. С которым вообщем-то согласны в индустрии. А вы мне черт знает что — я не могу понять в чем вы измеряете свои флопсы, откуда взялась эта цифра 150??
        • 0
          Этот теоретический пик — на матричных операциях, и я согласен что он тоже важен. В данной статье я привожу максимальную реально достижимую скорость, не учитывая ограничения памяти (например когда все в регистрах или в L1 кеше). Производительность работы с памятью (bandwidth & latency на разных уровнях) можно рассмотреть в отдельной статье, там тоже много интересного :-)
          • 0
            Теоретический пик как раз и не учитывает кэши и память.
            Ну вы приведите код в конце концов, и все ведь станет ясно. Тот кусок кода что в оригинальной статье меряет не Flops-ы а Mips-ы, т.е целочисленные инструкции. Возможно отсюда у автора и такие цифры.

            Про bandwidth и latency — пожалуйста. Только сразу указывайте для какой операции.
            • +1
              Производительность и на 32 битных целых, и синглах(по крайней мере пока мы говорим о простых операциях) у всех одинаковая — все занимает по одному такту.
              Код в статье — лишь для примера.
              • 0
                Не может быть!!! Вас кто-то обманул :)
                4 сингловые операции за такт или 3 целочисленные, а с применением адресной арифметики может быть и 5.
                И вообще х86 сейчас может выполнять до 6-ти или 11-ти (точно не помню) инструкций за такт с помощию буфера циклов.
                • +1
                  Я говорю об SSE операциях, так и получается 12 операций за такт на ядро — 3 операции по 4 числа.
                  Ставить вместо SSE операций обычные пожалуй не стоит :-)
                  • 0
                    Еще раз. В процессоре 1 блок SSE шириной 128бит. Он за 1 такт выполняет 1 SSE инструкцию над 128 битами.
                    «3 операции по 4 числа.» — нихт. Никак. Невозможно.
                    • +1
                      В Core2Duo и выше 3 SSE блока — это и дает очень большую скорость.
                      В Phenom — тоже 3, но 2 из них специализированные.

                      На реальных приложениях (например моя программа BarsWF) выжимает ~35 млрд операций в секунду на ядро, как это возможно если можно делать только 4 операции за такт?
                      • 0
                        Небольшой апдейт:
                        1 блок FP умножения и 1 блок FP сложения — отдельно.
                        Итого 2+2=4 double precision FP операции за такт.
                        Делилка и вычиталка — те же самые блоки.
                        Другие операции — сильно медленнее.
                        Конверсия чисел, load и store за операции не считаются.



                        Вобщем чтобы прекратить спор предлагаю вам написать программу, делающую больше 4 FP операций за такт. Хотспот вашей BarsWF подойдет.

                        • +1
                          Вы правы, а я не прав :-)

                          • +1
                            peak FP throughput for vectorized code of 2 64-bit MUL and 2 64-bit ADD per cycle (8 SP or 4 DP FLOPS)
                            • 0
                              Отлично :)
                              Теперь можете прикинуть как это выглядит на больших машинах, на отлично вылизанном DGEMM (или распределенном Linpack) — сколько реально достигается процентов от пика. Ссылку я давал выше.
                              Данные будут примерно такого порядка: 75% на п4, 82% на Core2Duo и примерно 85-89% на Нехалеме. В этом и сила этого процессора.
  • –2
    фанАТИчно пахнущая статья…
  • 0
    Спасибо!!! Ещё бы вводные курсы по nVidia CUDA и AMD Brook почитать.
  • 0
    Кроме того, материнская плата под 4 видеокарты стоит намного дешевле платы под 4 процессора.
    А можно пример такой материнки? Я максимум слышал про 3 карточки, и из них одна по шине еще зажата будет…
    • +1
      www.newegg.com/Product/Product.aspx?Item=N82E16813130140

      279$

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

      Работающию систему 4x9800GX2 видел и на скриншоте и внешний вид )

Только зарегистрированные пользователи могут оставлять комментарии. Войдите, пожалуйста.