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

О вычислениях

lexa - 23/Дек/2011 22:49

clusters

Понастраивал тут кластер, добился для 12 нодов эффективности около 70%, на чем (заказчик) и успокоился.

Но все эти дни не отпускала такая вот простая мысль:

Вот значит кластер, в нем 7U Blade-серверов (набивка неполная, мест на шасси 10) плюс еще 1U "управляющего фронтенда", OpenMPI, подбор топологии и размеров задачи, чтобы коммуникации и вычислительная моща были сбалнсированы (а уже для 12 нод, без опыта - пришлось повозиться, хотя конечно можно было бы просто оставить подбираться на месяц и помешивать).
Ну и стоит - не знаю сколько, но так подозреваю, что заметно за $100k.
Но те же 1.7 терафлопса получаются на половинке юнита (а добив памяти и увеличив задачу - на той машине уже за 2 Tflop/s получили). Ну ладно, на целом юните, если не увлекаться и просто пихнуть две HD6990 в подходящий корпус. И, по ощущениям, гемороя с отладкой какбэ не меньше. Стоить будет - ну скажем $10k за сервер и еще $2k за две видеокарты.
Ну хорошо, пусть даже mainstream-решение: 4-GPU-Tesla (1U) и два 1U-сервера. И даже IB (на две ноды - можно без свитча - будет несколько сотен за 40Gbit порт). Но сбалансировать две ноды сильно проще чем 12, я проверял.
Такой мейнстрим стоить будет тоже не $100k+, а 30-40. Электричества жрать не 6 киловатт, а два. А на 4xTesla те же 1.7 GF

Ну то есть понятно, GPGPU еще не везде мейнстрим, особенно на AMD. Но посмотрев на все это вживую - никакой идеи считать что-то научное на "обычных компьютерах" - у меня больше не возникает.

P.S. С удовольствием приму участие в настройке какого-то кластера с теслами. Чисто за интерес.

I like to move it, move it

lexa - 17/Авг/2011 19:18

Мониторю всяческие новости про OpenCL, CUDA и прочие GPGPU и в последние дни просто засыпан новостью про то, что GEGL is getting GPU-based image rendering and processing.

Довеском к этой новости идет ссылка на OpenCL on GEGL: Results up to now, где сравнивается реализация brightness-contrast фильтра на CPU и на GPU (и не каком-то, а Tesla 2050 ценой 2 килобакса) и получается для 1-мегапиксельного изображения:

526 msec на CPU
483 msec на GPU

Просто гигантский выигрыш, почти 10% !!! Притом, как я понял (может быть и неправильно), во время исполнения на GPU не посчитана пересылка "обратно", с карты в RAM.

При этом, заметим, на CPU оно работает на одном ядре (хоть и на SSE2), а значит на 4 ядрах оно банально будет быстрее разика в три.

Причина тривиальна и прямо в том блог-посте описана, весь пар ушел в свисток, а все время исполнения - на пересылку данных. Собственно исполнение обработки на GPU занимает около 1/10 всего времени.

При этом, само время исполнения - чудовищно. Полсекунды на 1 мегапиксель, даже если пиксели 16-байтные (4 float) - это 32 мегабайта в секунду. Э.... По PCIe обычно ходит 4-5 гигабайт/сек..... Проблема, судя по всему, кроется в тайлововой организации картинки в GEGL, тайл при этом мелкий (128x64) и даже на CPU обрабатывать их эффективно не получается, что уж говорить про GPU, где под каждый тайл аллоцируется текстура.

На эту тему имею рассказать следующую историю:

О цветовой интерполяции

lexa - 14/Май/2011 15:33

профилирование

Навеяно вот этим вот обсуждением, пожалуй запишу.

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

К примеру, Monaco Profiler с удовольствием сделает вам цветовой профиль с таблицами 33x33x33. То бишь почти 108 тысяч коэффициентов (+тоновые кривые). Круто, да.

А на вход ему при этом подадут ну, скажем, 1728 патчей (12 в кубе). Так как жрет он не спектральные данные, а простой рабоче-крестьянский XYZ, то это будет примерно 5200 значений.

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

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

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

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

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

P.S. То что мишени для тех же принтеров генерируются, как правило, путем равномерной расстановки точек по всему пространству координат - отдельная печальная песня.

AMD/ATI и GPGPU

lexa - 12/Янв/2010 21:50

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

Угадал, блин

lexa - 13/Май/2009 11:59

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

В комментариях к записи про анонс 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 или сказание о канделябре

lexa - 14/Фев/2009 13:17

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

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

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

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

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

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

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

lexa - 26/Дек/2008 10:05

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

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

Каждые полгода ~~мы с друзьями~~ я бенчмаркаю вычисления на видеокартах. В этот раз изучалась NVidia GTX280.

SGEMM/DGEMM на видеокарте и CPU, серия 7: NVidia GTX280
В чипе NVidia G200 появились операции с двойной точностью. Их производительность не феноменальна, но даже с учетом ввода-вывода данных в карту GTX280 обгоняет 4-ядерный 3-гигагерцовый Penryn. Если же рассматривать только скорость вычислений (что актуально, если задача позволяет спрятать затраты на ввод-вывод), то на двойной точности видеокарта быстрее CPU в 1.8 раза.
На одинарной точности разрыв видеокарты и CPU вырос: GTX280 обгоняет Core2Quad впятеро.

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

OpenMP и C++ class members

lexa - 07/Сен/2008 20:00

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

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)

lexa - 30/Апр/2008 23:00

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

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

Из важного

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

Да, dgemm я, естественно, в ближайшие дни попробую опробовать. Интересно же.

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

lexa - 28/Мар/2008 18:30

книги

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

увеличить фото

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

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

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

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

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

lexa - 02/Фев/2008 21:14

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

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

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

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

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

CUDA в массы!

lexa - 13/Ноя/2007 22:36

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

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

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

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

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

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

lexa - 16/Окт/2007 22:57

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

Предыдущая моя статья на эту тему была написана в феврале 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.

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

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

lexa - 23/Апр/2007 11:26

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

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

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

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

lexa - 18/Апр/2007 22:00

Вот смотрите:

Radeon X1800 - 512Mb памяти
Geforce 8800GTX - 768Mb
NVidia Quadro 5600 - 1.5Gb
AMD (ATI) Stream processor - 2Gb

А ведь еще 1-2 поколения и будет 64-битная адресация

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

lexa - 14/Мар/2007 09:36

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

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

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

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

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

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

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

lexa - 08/Мар/2007 19:21

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

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

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

NVidia 8800GTX: пропускная способность памяти (при использовании CUDA)

lexa - 08/Мар/2007 12:28

После чтения руководства по NVidia CUDA, остается ощущение сложности модели программирования: треды, блоки тредов, warp-ы, иерархическая память. Непонятно, какие параметры вычислительной задачи оптимальны и какие у них вообще допустимые значения. Само руководство точных рекомендаций не дает, дает лишь приблизительные.

Из общих соображений, понятно что самая медленная часть суперкомпьютера - память. С одной стороны, теоретическая пропускная способность (bandwidth) составляет 900MHz * 384 бита * 2 (DDR) = 86.4 GB/sec. С другой стороны, раздел 6.1.1.3 руководства говорит о 200-300 циклах memory latency (при, по всей видимости,случайном доступе).

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

Читая веб и блоги: CUDA и прочее программирование на NVidia 8800

lexa - 03/Мар/2007 23:08

На удивление мало жизни происходит по запросу 'NVidia CUDA' в поиске по блогам и новостям. Что у Яндекса, что у Google. Мне это сильно удивительно - штука многообещающая, первая версия SDK датирована ноябрем (появилась примерно 1-го декабря), публичный SDK появился практически месяц назад, а большинство упоминаний в духе "вот вышло", в крайнем случае "читал доку". Таких текстов - навалом, маркетинг NVidia работает. Но скучно.

Помимо форумов NVidia, где заводится по 5-6 новых топиков в день, интересных публикаций немного.

Для начала: Beyond3D опубликовал большой текст про CUDA. Более подробный, чем все что я видел до сих пор (ну, кроме собственно документации).

NVidia CUDA: синхронизация и shared memory

lexa - 02/Мар/2007 15:19

Экспериментально выяснилось, что содержимое shared memory не сохраняется между запусками кода на G80. Всякий раз оно инициализировано, причем значения разные, то 10 (float), то 125.

Плакала идея синхронизации между мультипроцессорами путем завершения kernel. Нет, синхронизироваться можно, конечно, но если хочется сохранить результат, то надо писать в глобальную память.

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

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

Pages