Я как-то не уследил, потому что AMD/ATI-шными видеокартами начал интересоваться с выходом HD5xxx, а оказывается все очень весело. На gpgpu.ru это уже пообсуждали, ну я сюда наброшу, в более концентрированном виде.
Раньше высокоуровневым средством для разработки считалок на видеокартах у ATI был Brook+. Однако начиная с какой-то беты ATI Stream SDK 2.0 Brook из SDK исчез.
Yes, this SDK 2.0 beta is for CPU only. It focuses on OpenCL 1.0 for CPU.
Brook+ is now available on SourceForge: http://sourceforge.net/projects/brookplus
Ну ладно, Stream SDK Beta-1 вообще не поддерживает никаких видеокарт, смешно.
Конечно, сейчас начнется, они вполне могут начать с драйверов для Линукса (или для 32-битной XP, что от меня столь же далеко)
И угадал, блин. Именно XP-32 и Linux-32. XP-шные бинарники на Висте не работают, несмотря на драйвер нужной версии, ругаются что не могут создать OpenCL context
А у меня Vista (32/64) и MacOS. Ну в Маке, ладно, обещали в заснеженном леопарде, а с вистой что? Руки же чешутся....
На закуску: согласно спекам OpenCL, его можно/нужно кормить исходником прямо на этом самом OpenCL (а это практически C). То бишь компилятор этого самого C сидит прямо в драйвере....
Интересно, насколько успешно это пойдет в индустрию, ведь получается что computing kernel засекретить не выйдет, можно же подсунуть приложению драйвер похожий на настоящий и почитать им исходника. Видятся мне OpenCL-обфускаторы.
Вся эта история с PhantomOS интересна тем, что будучи еще vapourware оно всколыхнуло в людях пласты и позволило заглянуть в многочисленные бездны. Ссылок на дискуссии не даю, кому интересно - уже все видели.
К сожалению, у него там в комментариях каптча даже для авторизованных пользователей, я через себя переступить не могу, приходится писать у себя и отдельно.
Да, действительно, судя по этим бенчмаркам Java временами рулит. Но что мы можем сказать про эти бенчмарки:
Они старые, 2003-й год, updated в 2004. Конечно, JIT за это время мог развиться на неимоверные высоты, однако и процессоростроители на месте не стояли, в i7, как я слышал, уже 3 SSE-юнита на ядро, а нормальное программирование SSE - это никакая не автовекторизация компилятором, а натурально фигурное выпиливание прямо на его ассемблере.
Что это за бенчмарки - мне с первого взгляда ясно не стало. Если это обращение матрицы 100x100, которая влезает в кэш L1 - это одно. Если это относительно большие данные, то плохие (не cache-aware) алгоритмы будут одинаково плохи, ибо cache miss будет стоить гораздо дороже, чем эффективность или неэффективность JIT. И, соответственно, выигрыши-провалы на тех графиках - это косвенное указание на то, попали или не попали в кэш.
Вместе с тем, быстрое пользование гуглом навело на реальную имплементацию FFT и реальные бенчмарки в сравнении не с каким-то левым sample/микробенчмаркой, а с кодом, который действительно доводится годами и считается оптимальным или близким к этому.
Каждые полгода мы с друзьями я бенчмаркаю вычисления на видеокартах. В этот раз изучалась NVidia GTX280.
SGEMM/DGEMM на видеокарте и CPU, серия 7: NVidia GTX280
В чипе NVidia G200 появились операции с двойной точностью. Их производительность не феноменальна, но даже с учетом ввода-вывода данных в карту GTX280 обгоняет 4-ядерный 3-гигагерцовый Penryn. Если же рассматривать только скорость вычислений (что актуально, если задача позволяет спрятать затраты на ввод-вывод), то на двойной точности видеокарта быстрее CPU в 1.8 раза.
На одинарной точности разрыв видеокарты и CPU вырос: GTX280 обгоняет Core2Quad впятеро.
Понятно, что Core i7 разницу несколько сократит, но по тем бенчмаркам с плавающей точкой, что я видел (а видел я пока только Linpack, причем не факт что в оптимальном для i7 виде), рост в производительности i7 - процентов 20.
Всякие соображения про масштабируемость решения - в самой статье.
class Bar
{
public:
int a;
Bar();
void run(void);
};
void Bar::run(void)
{
int i,k;
#pragma omp parallel private(i) firstprivate(k,a)
{
}
}
Ругается вот таким нехорошим словом:
omp-test.cpp: In member function 'void Bar::run()':
omp-test.cpp:22: error: 'Bar::a' is not a variable in clause 'firstprivate'
Ну и что с этим делать?
Update Нашел (тоже не очень легко), что возможен такой синтаксис:
class Bar
{
public:
int a;
#pragma omp firstprivate(a)
}
А что делать, если Bar::a в одном месте должна быть lastprivate, а в другом - reduction? Временные переменные заводить?
Программирующим на CUDA может быть интересно: NVidia начала раздавать исходники библиотек CUBLAS/CUFFT.
Я, правда, не очень понимаю статус этого дела:
С одной стороны, все выложено на девелоперском сайте, куда нужна регистрация (и говорят, что стоит большая очередь желающих оной регистрации, хотя меня в прошлом году зарегистрировали за один день).
С другой стороны, в девелоперской рассылке пришли ссылки на незапароленый сайт, бери кто хочет.
Посему, ссылки не публикую, если кому-то нужно и нет терпения ждать (со временем все попадает в полностью открытый доступ, всегда так было) - пишите лично.
А вот что точно открыто всем желающим, так это визуальный профайлер (beta) для той же CUDA. Пока не смотрел, руки не дошли.
Несмотря на некруглый номер версии, это революция:
CUDA-функциональность теперь будет в обычном драйвере;
Необходимые DLL-s из SDK можно редистрибутировать вместе с приложением;
По сути, это означает, что из тестовой системы оно стало боевой: в требованиях к программе можно писать просто NVidia 8x00 и минимальную версию (обычных!) драйверов и у любого пользователя оно будет работать.
Вообще, анонс довольно интересный и требует подробного комментирования.
Предыдущая моя статья на эту тему
была написана в феврале 2007 года, сразу после выхода первой публичной бета-версии
CUDA Toolkit/CUDA SDK. Представители NVidia предупреждали,
что в бета-версии производительность не является оптимальной, к релизу она будет улучшена.
На параллельных архитектурах часто приходится делать операцию reduce (складывать и умножать вектора, считать среднее и так далее). В отличие от однопоточной конструкции, где все тривиально, параллельная reduce разбивается на два этапа: сначала мы всеми исполняющими юнитами обрабатываем куски данных, а потом должны сложить (усреднить, поделить) результаты уже меньшим числом процессоров.
Для второго шага reduce обычно используется пирамидальная схема: сначала в N/2 потоков сложим N результатов попарно, затем сложим N/2 в N/4 и так далее. Число итераций равно, очевидно, log2N.
Возникает вопрос, «сколько данных складывать на каждой итерации?» Ведь можно складывать в N/4-N/16-N/256 кучек, можно по 1/8-64-512 и так далее. Из общих соображений, складывать по несколько лучше чем по два. Конечно, потоков получается меньше, но меньше и оверхед на создание-завершение потока.
В , куда я чуть было не уехал заниматься геологией 15 лет назад, в настоящее время читается курс .
Опуская обычные охи "ну почему этому не учат на ВМК" - все в наших руках и я думаю, что через пару лет такие курсы будут и у нас, тема вкусная - хочу обратить внимание на слайды и транскрипты лекций, доступные вот тут. Читают приглашенные лекторы из NVidia, поэтому основное внимание уделено, сюрприз, NVidia 8800. Курс
включает в себя лабы, которые сделаны очень интересно: есть готовая рыба, делающая подготовительную работу (I/O, печать результаты) и студент должен только написать несколько десятков-сотен строк изучаемой функциональности. Что, конечно, экономит кучу
непроизводительного времени (смотреть тут)
Я успел посмотреть не все, из того что посмотрел - многое нравится, а многое
не нравится: Не нравится то, что танцы идут от предыдущего поколения видеокарт. Если говорить о программировании вообще (а все примеры - про это), а не про видеоигры, то надо забывать про шейдеры и текстурные буферы. Ведь студентов учат, которые, я надеюсь, всего этого ужаса не нюхали.
Нравится, как практику, подробный разбор аппаратуры. Вылезает множество подробностей, которые нигде не опубликованы.
В предыдущей части мы рассматривали чтение из глобальной памяти Geforce 8800 напрямую ("как из массива C"). При этом отсутствует кэширование, но при оптимальной схеме доступа получается (согласно указаниям NVidia) наибольшая производительность.
В то же время, скорость доступа при неоптимальном паттерне очень маленькая. Для решения этой проблемы (помимо оптимизации паттерна) NVidia CUDA предлагает доступ к памяти как к текстуре. При этом работает двумерное кэширование (оптимизированное под локальный доступ), пиковые скорости должны получаться меньше, а наихудшие варианты - наоборот лучше.
После чтения руководства по NVidia CUDA,
остается ощущение сложности модели программирования: треды, блоки тредов, warp-ы, иерархическая память.
Непонятно, какие параметры вычислительной задачи оптимальны и какие у них вообще допустимые значения.
Само руководство точных рекомендаций не дает, дает лишь приблизительные.
Из общих соображений, понятно что самая медленная часть суперкомпьютера - память. С одной стороны,
теоретическая пропускная способность (bandwidth) составляет 900MHz * 384 бита * 2 (DDR) = 86.4 GB/sec.
С другой стороны, раздел 6.1.1.3 руководства говорит о 200-300 циклах memory latency (при, по всей видимости,случайном доступе).
К счастью, проблема легко изучается: если взять достаточно много данных (скажем, полгигабайта) и, например,
сложить все 4-байтовые значения (как float), то основные затраты времени будут именно на чтение из памяти,
а всей прочей арифметикой можно пренебречь (или подсчитать ее отдельно).