Skip to Content

Численные методы

AMD/ATI и GPGPU

Я как-то не уследил, потому что AMD/ATI-шными видеокартами начал интересоваться с выходом HD5xxx, а оказывается все очень весело. На gpgpu.ru это уже пообсуждали, ну я сюда наброшу, в более концентрированном виде.

Раньше высокоуровневым средством для разработки считалок на видеокартах у ATI был Brook+. Однако начиная с какой-то беты ATI Stream SDK 2.0 Brook из SDK исчез.

Читаем в ATI-шном форуме (это август-2009):

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 вообще не поддерживает никаких видеокарт, смешно.

Угадал, блин

Довелось оказаться пророком.

В комментариях к записи про анонс NVidia работающего OpenCL я предположил, что

Конечно, сейчас начнется, они вполне могут начать с драйверов для Линукса (или для 32-битной XP, что от меня столь же далеко)
И угадал, блин. Именно XP-32 и Linux-32. XP-шные бинарники на Висте не работают, несмотря на драйвер нужной версии, ругаются что не могут создать OpenCL context

А у меня Vista (32/64) и MacOS. Ну в Маке, ладно, обещали в заснеженном леопарде, а с вистой что? Руки же чешутся....

На закуску: согласно спекам OpenCL, его можно/нужно кормить исходником прямо на этом самом OpenCL (а это практически C). То бишь компилятор этого самого C сидит прямо в драйвере....

Интересно, насколько успешно это пойдет в индустрию, ведь получается что computing kernel засекретить не выйдет, можно же подсунуть приложению драйвер похожий на настоящий и почитать им исходника. Видятся мне OpenCL-обфускаторы.

Opera of the Phantom или сказание о канделябре

452_1.jpg Вся эта история с PhantomOS интересна тем, что будучи еще vapourware оно всколыхнуло в людях пласты и позволило заглянуть в многочисленные бездны. Ссылок на дискуссии не даю, кому интересно - уже все видели.

Вместе с тем, dz продолжает утверждать, что всякий JIT на виртуальной машине рулит примерно не хуже, чем макроассемблер, известный нам под именем C/C++.

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

Да, действительно, судя по этим бенчмаркам Java временами рулит. Но что мы можем сказать про эти бенчмарки:

  1. Они старые, 2003-й год, updated в 2004. Конечно, JIT за это время мог развиться на неимоверные высоты, однако и процессоростроители на месте не стояли, в i7, как я слышал, уже 3 SSE-юнита на ядро, а нормальное программирование SSE - это никакая не автовекторизация компилятором, а натурально фигурное выпиливание прямо на его ассемблере.
  2. Что это за бенчмарки - мне с первого взгляда ясно не стало. Если это обращение матрицы 100x100, которая влезает в кэш L1 - это одно. Если это относительно большие данные, то плохие (не cache-aware) алгоритмы будут одинаково плохи, ибо cache miss будет стоить гораздо дороже, чем эффективность или неэффективность JIT. И, соответственно, выигрыши-провалы на тех графиках - это косвенное указание на то, попали или не попали в кэш.

Вместе с тем, быстрое пользование гуглом навело на реальную имплементацию FFT и реальные бенчмарки в сравнении не с каким-то левым sample/микробенчмаркой, а с кодом, который действительно доводится годами и считается оптимальным или близким к этому.

NVidia GTX280: бенчмарки с плавающей точкой

gtx280.jpg Каждые полгода мы с друзьями я бенчмаркаю вычисления на видеокартах. В этот раз изучалась 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.

Всякие соображения про масштабируемость решения - в самой статье.

OpenMP и C++ class members

Что-то у меня затык и гугление не помогает.

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 2.0 (beta)

Не прислав мне никакого уведомления (хотя раньше присылали), NVidia выпустила CUDA 2.0 beta.

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

Из важного

  • Поддержка Vista (32 и 64 бита);
  • Нет поддержки GeForce 9800GTX (вышедшей на пару недель раньше этой беты), что довольно странно.
  • С двойной точностью какая-то непонятная совсем история:
    1. В CuBLAS она заявлена в документации, символы в библиотеке имеются (собирать еще ничего не пробовал).
    2. В документации (programming guide) слово double встречается 8 раз (на 99 страниц текста), что как-то безобразно мало.
    3. Времена вычислений для double в соответствующей секции не описаны (но я подозреваю, что они другие, чем у float).
    4. Таблица с описанием double-функций (на которую есть указание в тексте) - отсутствует.
    Другими словами, работы ведутся и довольно скоро все может появиться.
Да, dgemm я, естественно, в ближайшие дни попробую опробовать. Интересно же.

А вот кому книжечек по вычметодам?

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

books1.jpg
увеличить фото

Отдаются бесплатно т.е. даром первому обратившемуся. Но только все сразу, а не по одной штучке.

Москва, Коньково - Юго-Западная, пишите lexa@lexa.ru, договоримся.

Предложение действительно одну неделю: до вечера 4 апреля 2008, потом вынесу.

Update: все забрали!

Исходники CuBLAS/CuFFT

Программирующим на CUDA может быть интересно: NVidia начала раздавать исходники библиотек CUBLAS/CUFFT.

Я, правда, не очень понимаю статус этого дела:

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

А вот что точно открыто всем желающим, так это визуальный профайлер (beta) для той же CUDA. Пока не смотрел, руки не дошли.

CUDA в массы!

Анонсирована NVidia CUDA 1.1 (beta).

Несмотря на некруглый номер версии, это революция:

  • CUDA-функциональность теперь будет в обычном драйвере;
  • Необходимые DLL-s из SDK можно редистрибутировать вместе с приложением;

По сути, это означает, что из тестовой системы оно стало боевой: в требованиях к программе можно писать просто NVidia 8x00 и минимальную версию (обычных!) драйверов и у любого пользователя оно будет работать.

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

Умножение матриц, серия 5: вычисления на GPU (2)

Почему переделываем тесты?

Предыдущая моя статья на эту тему была написана в феврале 2007 года, сразу после выхода первой публичной бета-версии CUDA Toolkit/CUDA SDK. Представители NVidia предупреждали, что в бета-версии производительность не является оптимальной, к релизу она будет улучшена.

За прошедшие полгода, пока я занимался совсем другими вещами, были выпущены релизы:

  • NVidia CUDA: SDK и библиотеки CUBLAS/CUFFT v1.0;
  • NVidia CUDA Display Driver 162.xx (драйвер, собственно, транслирует псевдокод в реальные программы GPU);
  • RapidMind Platform версий 2.0.0, а затем и 2.0.1.

Интересно посмотреть, стала ли производительность лучше.

О пирамидальном сложении на параллельной архитектуре

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

Для второго шага reduce обычно используется пирамидальная схема: сначала в N/2 потоков сложим N результатов попарно, затем сложим N/2 в N/4 и так далее. Число итераций равно, очевидно, log2N.

Возникает вопрос, «сколько данных складывать на каждой итерации?» Ведь можно складывать в N/4-N/16-N/256 кучек, можно по 1/8-64-512 и так далее. Из общих соображений, складывать по несколько лучше чем по два. Конечно, потоков получается меньше, но меньше и оверхед на создание-завершение потока.

640 килобайт будет достаточно...

Вот смотрите:
  • Radeon X1800 - 512Mb памяти
  • Geforce 8800GTX - 768Mb
  • NVidia Quadro 5600 - 1.5Gb
  • AMD (ATI) Stream processor - 2Gb
А ведь еще 1-2 поколения и будет 64-битная адресация

Программирование NVidia 8800: вести с веба

В University of Illinois at Urbana, куда я чуть было не уехал заниматься геологией 15 лет назад, в настоящее время читается курс ECE 498 AL : Programming Massively Parallel Microprocessors.

Опуская обычные охи "ну почему этому не учат на ВМК" - все в наших руках и я думаю, что через пару лет такие курсы будут и у нас, тема вкусная - хочу обратить внимание на слайды и транскрипты лекций, доступные вот тут. Читают приглашенные лекторы из NVidia, поэтому основное внимание уделено, сюрприз, NVidia 8800. Курс включает в себя лабы, которые сделаны очень интересно: есть готовая рыба, делающая подготовительную работу (I/O, печать результаты) и студент должен только написать несколько десятков-сотен строк изучаемой функциональности. Что, конечно, экономит кучу непроизводительного времени (смотреть тут)

Я успел посмотреть не все, из того что посмотрел - многое нравится, а многое не нравится:
Не нравится то, что танцы идут от предыдущего поколения видеокарт. Если говорить о программировании вообще (а все примеры - про это), а не про видеоигры, то надо забывать про шейдеры и текстурные буферы. Ведь студентов учат, которые, я надеюсь, всего этого ужаса не нюхали.

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

Вот о подробностях - ниже.

NVidia 8800GTX: скорость чтения текстур

В предыдущей части мы рассматривали чтение из глобальной памяти Geforce 8800 напрямую ("как из массива C"). При этом отсутствует кэширование, но при оптимальной схеме доступа получается (согласно указаниям NVidia) наибольшая производительность.

В то же время, скорость доступа при неоптимальном паттерне очень маленькая. Для решения этой проблемы (помимо оптимизации паттерна) NVidia CUDA предлагает доступ к памяти как к текстуре. При этом работает двумерное кэширование (оптимизированное под локальный доступ), пиковые скорости должны получаться меньше, а наихудшие варианты - наоборот лучше.

Все это, как обычно, нуждается в изучении.

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), то основные затраты времени будут именно на чтение из памяти, а всей прочей арифметикой можно пренебречь (или подсчитать ее отдельно).

Syndicate content