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>, так что я его уберу. :(
    Поделиться публикацией
    Похожие публикации
    Реклама помогает поддерживать и развивать наши сервисы

    Подробнее
    Реклама
    Комментарии 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 карточки, и из них одна по шине еще зажата будет…

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