Многие слышали о мифических словах — 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 раза.
На затравку сразу показываю результат:
*GFLOP — миллиардов операций в секунду
P.S.: Cell на 45nm еще как следует
не запустили, и в соньках он не используется
пока.
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 постаралась на славу с железом — её 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. Она конечно впечатляет, если забыть о том, что x86 тоже продвинулся (не говоря уже о видеокартах). Cell состоит из обычного Power-ядра называемого PPE (с незначительной производительностью по сравнению с вышеупомянутыми монстрами) и 8-и ядер SPE, из которых 1 — потенциально бракованое ("небраковыные" чипы продаются отдельно за гораздо большие
деньги), и 1 зарезервировано операционной системой (если перешить приставку на Linux можно получить доступ к 7, но это для конечных пользователей трудновато ;-) ), итого остается 6. Каждое из ядер делает 8 операций за такт(2*4) на частоте 3.2Ghz, что дает общую реальную
производительность 153.6 GFLOP. С программированием тут не подскажу, т.к. эти 153 GFLOP достаточно непрактично использовать, потому и не
пробовал (с точки зрения выпуска программы не для внутреннего потребления).
Уже сбросили со счетов 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>, так что я его уберу. :(
Увлекся я недавно высокопроизводительными вычислениями (язык сломать можно, 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 |
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>, так что я его уберу. :(