GPGPU.ru

Декабрьские новости

lexa — Wed, 14 Dec 2011 12:20:18 +0000

Накопилось сразу много новостей, мимо которых нельзя пройти мимо.

OpenCL

Месяц назад был объявлен OpenCL 1.2, тут о нем еще не писалось, потому что писать особо нечего: изменения мне кажутся не слишком значительными (ну кроме лучшей интеграции с DX11), а реализаций все еще нет (но вот-вот могут появиться, см. ниже).

Nvidia

Вышел CUDA Toolkit RC2, доступный всем без предварительной регистрации. Судя по Release Notes, есть множество мелких (но существенных) изменений:
- Компилятор теперь основан на LLVM, а не на Open32. Соответственно, возможны заметные изменения в эффективности в обе стороны, надо внимательно смотреть. Обещают ускорение up to 10%, но для некоторого кода жалуются на замедление.
- Существенно расширены библиотеки (NPP, cuSPARSE, cuRAND)
- P2P-communications (между картами, мимо CPU) теперь работают не только для Tesla, а для любых Fermi-карт.
- И много других мелких улучшений.
К CUDA 4.1 вышел соответствующий Parallel NSight 2.1 RC2.
В отличие от CUDA, он доступен пока только зарегистрированным девелоперам, причем не старым зарегистрированным девелоперам (на сайте nvdeveloper.nvidia.com), а на полностью новом developer.nvidia.com
Для доступа к Parallel Nsight требуется одобрение вашей регистрации, в моем случае оно произошло в течение нескольких часов, ночью по калифорнийскому времени, а одобрение CUDA и вовсе прилетело мгновенно.
Помянутый выше LLVM-компилятор обещают заопенсорсить. Пока по этому адресу только форма регистрации, я зарегистрировался, посмотрим что дальше будет. В-принципе, JIT-компилятор CUDA, делающий код, заточенный под конкретное устройство на конкретной машине - полезная же вещь.

AMD

AMD APP SDK 2.6 Preview недоступен с родного сайта AMD, но доступен с GPU-Tech.org (сайт временами лежит). Обещают помянутый выше OpenCL 1.2, пока не смотрел: для работы требуется Catalyst 11.12, который вышел только сегодня, пока не успел.
AMD тоже опенсорсит свой компилятор (из OpenCL в IL). Не в виде анонса, а в виде патчей к LLVM 2.9 которые можно взять тут

Intel

Это не GPGPU, а data parallel, но все равно интересно. Интел зарелизил версию 1.1 своего SPMD Program Compiler. BSD-лицензия, LLVM-based.

Это такой компилятор C-подобного языка, позволяющий легко писать SIMD- и (с нашлепкой для multi-threading) MultiCore/SMP код. В отличие от OpenCL для CPU, этот код может работать с данными прямо по месту, без их копирования в буферы, что для задач с низкой арифметической интенсивностью оказывается очень выгодно.

OpenCL tools & compilers

lexa — Thu, 20 Oct 2011 06:51:57 +0000

Последние дни оказались богаты на всякие штуки для OpenCL:

OpenCL Compiler tools: Небольшая библиотека, позволяющая удобно компилировать OpenCL kernels заранее, сохраняя результат (можно - шифрованый). Цель - поэкономить на времени компиляции и на том, что kernel находится где-то внутри программы в исходных текстах.
Несколько смазывает впечатление то, что библиотека распространяется под GPL3, т.е. цель "спрятать kernel" под такой лицензией, очевидно, не может быть достигнута т.к. использующий библиотеку варез обязан тоже быть под GPL3+
Portable OpenCL: Реализация OpenCL (для CPU) на базе LLVM. OpenSource, BSD-style license.
Не смотрел, но идея мне нравится.
libclc: Реализация OpenCL для CPU в виде библиотеки, тоже поверх LLVM/Clang, хотя поминается и PTC
Анонсирую исключительно до кучи, на настоящий момент проект фактически пустой, но вдруг там что-то разовьется со временем?

Административное: внутренние изменения на сайте

lexa — Thu, 06 Oct 2011 08:19:13 +0000

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

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

Хотя, вроде бы, все эти проблемы были найдены и вычищены, но вдруг что-то не заметили.

Intel OpenCL 1.5: AVX, отладчик, но плохая производительность

lexa — Tue, 27 Sep 2011 10:43:54 +0000

Пощупал вышедший вчера Intel OpenCL SDK 1.5.

Заявленные новые фишки:

Поддержка AVX-векторизации (т.е. векторов в 256 бит)
Отладчик

Имею сказать:

Отладчик: Отладчик - работает. Пошаговая отладка, просмотр значений переменных, все отлично. С каким конкретно потоком работать - задается до запуска программы.
Несколько неудобно то, что путь до файла с исходными текстами передается в флагах компиляции. Казалось бы, компилятор все видит в момент компиляции, однако ж.
AVX и 256 бит: Генерируемый 256-битный код, на мой вкус, безумно странен. Да, 256-битные регистры (YMM) используются, но как-то очень редко.

Для тестирования производительности я взял 6 примеров из AMD APP SDK и запустил их c Intel OpenCL 1.1, а потом с 1.5. Результат меня огорчил:

BitonicSort, EigenValue, Nbody: исполнение под OpenCL 1.5 в 2.5-3 раза медленнее.
RadixSort: новая версия в 1.5 раза медленнее.
SimpleConvolution: время исполнения не изменилось.
Histogram: стало на 15% быстрее.

При этом Historgam новый OpenCL отказывается векторизовать (на старый - не посмотрел), а все остальные - векторизуются и от того стало хуже.

Помимо проблем с производительностью, у меня еще не запускается GUI к 64-битному Offline Compiler (комманд-лайн версия работает без проблем), с диагностикой "не найден файл" (а какой файл - не скажу). Проблема с ioclgui64 решается как написано в README: нужно явно прописать путь к Java Runtime, иначе или 32-битная версия ioclgui не находит 64-битную яву, или, наоборот, 64-битный gui страдает без явы 32-битной.

Мораль: хотите отладчика - ставьте новую версию. Для всех других использований рекомендовать ее трудно. Надеемся на скорый выход фиксов.

Статьи по понедельникам

lexa — Tue, 13 Sep 2011 08:27:33 +0000

Не могу не поделиться очередным источником полезных ссылок по теме GPGPU.

Сайт StreamComputing.eu (за новостями которого лично я слежу через Твиттер) завел очень полезный обычай: по понедельникам они публикуют подборки статей по тематике GPGPU, кластеров и всего подобного.

5-го сентября вроде что-то тоже было, но разрозненно и в твиттере.

Язык, увы, английский.

High Performance Linpack для CUDA/Fermi

lexa — Tue, 23 Aug 2011 08:39:53 +0000

Внезапно обнаружил, что High Performance Linpack для Fermi с начала августа доступен на сайте Nvidia для зарегистрированных девелоперов.

Презентация про эту штуку есть тут: CUDA Accelerated Linpack for Clusters, пишут про 350 GFLOP/s для DGEMM на Tesla M2050 (только GPU) и 435 GFLOP/s для комбинации M2050+8CPU (2xQuad core Xeon 2.66). Для более мощной машины (CPU на 2.8 а не 2.33) показаны 475 GFLOP/s на DGEMM, 401 GFLOP/s для Linpack с одним GPU и 705 GFLOP/s с двумя GPU.
Нашлась еще статья китайцев: Early Linpack Performance Benchmarking on IPE Mole 8.5 Fermi GPU Cluster, где есть цифирка про 200 TFLOP/s на кластере из 320 узлов, по 6 GPU на узел.

С лицензией на этот HPL-Fermi мне непонятно. С одной стороны, при скачивании вылезает грозная EULА с которой надо согласиться. С другой стороны, рядом там лежит BSD-style лицензия, в самом пакете сохранены копирайты университета Теннеси, а в CUDA-специфичных файлах в начале написана BSD-style лицензия с оговоркой об использовании названия NVIDIA только по согласованию.

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

Теоретически, пакет требует Tesla (т.к. требования по памяти GPU от 2Gb), на практике наверное можно уговорить и на что-то поменьше.

Если у кого есть машина с парой новых Fermi, я бы с удовольствием эту штуку бы совместно с вами повоспитывал.

NVidia OpenCL 1.1

lexa — Sat, 06 Aug 2011 11:54:03 +0000

NVidia в рассылке для девелоперов сообщает нам, что в драйверах начиная с 280.13 поддерживается OpenCL 1.1

Я пощупал драйвера 280.18 для Win7/x64, и так оно и оказалось.

Проверил все примеры из AMD APP SDK, существенная часть которых раньше не работала с NVidia т.к. требовала OpenCL 1.1. Ситуация стала лучше, хотя и не идеальной:

Два примера должны бы работать, но не работают по причине разницы синтаксиса:
- FluidSimulation2D ломается c диагностикой "не могу преобразовать вектора разных размеров" на конструкции вида int8 = unsigned + int8.
- GlobalMemoryBandwidth ломается на вызове "overloadable" функции native_divide()
Еще три примера ломаются т.к. NVidia не поддерживает ряд расширений (что вполне ожидаемо):
- DeviceFission - нет соответствующего расширения (хотя на Fermi его хотелось бы иметь, аппаратура то умеет)
- SimpleDX10 - нет интеграции с DX10
- TransferOverlap - работает, если не попросить для памяти флаг CL_MEM_USE_PERSISTENT_MEM_AMD

Но жизнь стала лучше. Например, поддерживаются 2D-изображения 16k-x-16k, вместо совершенно бессмысленных 4k-x-32k в предыдущих драйверах.

код второго кернела

santarus — Wed, 03 Aug 2011 13:49:02 +0000

__global__ void action_2 (  int *out,

                            int *final,

                            size_t pitch,

                            int width,

                            int height)

{

     for (int r = 0; r < height; r++)

     {

          int temp = 0;

          int *row_o = (int*)((char*)out + r * pitch);

          for (int c = 0; c < width; c++)

          {

               temp += row_o[c];

          }

          final [r] = temp;

     }

}

код первого кернела

santarus — Wed, 03 Aug 2011 13:42:51 +0000

__global__ void action (  int *test, 

                          int *target, 

                          int *out, 

                          size_t pitch, 

                          int width, 

                          int height)

{

      for (int r = 0; r < height; r++)

      {

             int *row_t = (int*)((char*)test + r * pitch);

             int *row_o = (int*)((char*)out + r * pitch);

             for (int c = 0; c < width; c++)

             {

                     row_o[c] = row_t[c] + target[c];

             }

      }

}

CUDA by Example на русском

lexa — Fri, 08 Jul 2011 07:56:10 +0000

Нам сообщают, что вышел русский перевод книги CUDA by Example.

ISBN: 978-5-94074-504-4, 978-0-13-138768-3

Доступна в онлайн-магазинах. например в Озоне и Лабиринте.

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

Книга производит очень приятное впечатление как по порядку изложения материала, так и по полноте: помимо базовых вещей не забыты текстуры, не забыто взаимодействие с OpenGL/DX. В ней нет ничего про "продвинутые оптимизации", а вот как базовое изложение - очень понравилось.

О качестве русского перевода пока ничего сказать не могу.

Вышли CUDA 4.0 & Parallel Nsight 2.0 (релизы)

lexa — Sat, 28 May 2011 16:56:36 +0000

Вышли релизные версии:

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

Вот в драйверах разница заметна, копирование хост-видеокарта и обратно в драйвере 270.81 (доступен по первой ссылке) где-то на 10% быстрее, чем в текущем 270.61 WHQL.

Свежие GPGPU tools

lexa — Fri, 08 Apr 2011 06:45:09 +0000

В последние дни оба основных разработчика GPU обновили свои GPGPU toolkits до CUDA 4.0 RC2 и APP SDK 2.4, соответственно.

CUDA 4.0 RC2

CUDA 4.0 доступна для скачивания без пароля и регистрации с Nvidia developer zone.

В RC2 исправлено некоторое количество проблем RC1, подробнее см. Release Notes Errata.

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

AMD Accelerated Parallel Processing (APP) SDK 2.4

Основное новшество новой версии APP SDK (бывший Stream SDK) - поддержка процессоров AMD Fusion, как десктопных (A-series) так и энергоэкономных (C и E-series). Сказать про эту поддержку что-либо практическое мне сложно за неимением таких процессоров. Кроме того, много улучшено по мелочи: скорость FFT, поддержка профайлинга, поддержка OpenCL binary kernels.

Небольшая странность происходящего заключается в том, что в требованиях поминается драйвер (Catalyst) версии 11.4, тогда как найти такой драйвер для 64-битных Windows мне не удалось (в блоге AMD поминается Early Preview для 32-бит), а Catalyst 11.3 вышел буквально чуть больше недели назад.

CUDA 4.0

lexa — Mon, 28 Feb 2011 15:56:31 +0000

NVidia анонсировала CUDA 4.0

С моей колокольни интересны две вещи:

GPU Direct 2.0 т.е. прямая пересылка данных между (нескольки) картами, не задействуя память (и процессор?) хоста.
Унифицированная адресация на GPU и на хосте. Я это понял так, что адресное пространство общее, никаких отдельных сudaMemcpy.

Рассказы про остальные фишки лично меня не впечатлили (Thrust и так уже был, MPI не волнует), но будем посмотреть вживую (начиная с 4 марта).

"Интересные фишки" тоже будем посмотреть: про GPI Direct на онлайн-конференции был задан вопрос "это только для Tesla", а в момент ответа у меня пропал звук. С унифицированной адресацией у меня есть непонимание, как на эту схему ложится асинхронная передача.

Другими словами, 4-5 марта надо ломиться на nvdeveloper и брать бету на поиграться.

Интересно, что будет дальше с OpenCL: текущие версии его достаточно близки по идеологии к CUDA 2.x-3.x и почти нет разницы на чем программировать. Сейчас может оказаться, что CUDA сильно удобнее в разработке... а с другой стороны вот новые макбуки с ATI какбэ намекают мне про OpenCL...

Конкурс HPC-GPGPU проектов

lexa — Thu, 03 Feb 2011 07:44:52 +0000

Не могу не отрекламировать конкурс:

Эффективное использование GPU-ускорителей при решении больших задач

(дальше просто копированием с сайта конкурса, c его страницы и с пресс-релиза):

Конкурс будет проводиться в три этапа.
На первом этапе (до 28 февраля) участники конкурса проходят процедуру регистрации и предоставляют заявки. По итогам первого этапа будет отобрано до 30 проектов, каждому из которых будет выделено поощрительное финансирование на начальный этап проведения работ в размере 100 тысяч рублей.
На втором этапе (до 25 апреля 2011) участники должны будут предоставить отчеты по результатам пробных работ первого этапа. Подключиться к участию в конкурсе можно на любом этапе его проведения. В частности, к участию во втором этапе допускаются новые проекты, заявка по которым на первом этапе не подавалась. По окончании второго этапа экспертной комиссией будет выбрано 15 заявок и выделено финансирование в размере 200 тысяч рублей на каждую задачу.
На третьем этапе (до 30 сентября 2011) участники должны предоставить отчеты о проделанной работе в виде научно-популярной статьи. Как и на втором этапе, допускается подача заявок по новым проектам, которые не участвовали на первом и втором этапах. По окончании третьего этапа будут подведены итоги конкурса и объявлены восемь проектов, которые получат по 500 тысяч рублей каждый на продолжение своих исследований. Награждение победителей пройдет в конце ноября 2011 года на 7-й суперкомпьютерной конференции RSC 2011.

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

В пресс-релизе написано еще, что подлючиться к конкурсу можно на любом этапе.

Как мне кажется, всем у кого уже есть работающее HPC/GPGPU приложение - имеет смысл подключиться, времени это потребует не так много, а эффект может быть приятным.

Еще немного копипасты:

Критерии конкурентного отбора

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

Востребованность приложения, его соответствие поставленным целям;
Обоснование рациональности использования высокопроизводительных вычислений на графических ускорителях для решения выбранной задачи, в том числе экономический эффект;
Оригинальность реализации;
Уровень сложности приложения.

Подача заявок и порядок их описания - по ссылкам со страницы конкурса.

Новости одной строкой

lexa — Thu, 27 Jan 2011 21:06:13 +0000

Вдруг накопилось много мелких, но довольно существенных новостей, которые анонсирую одним текстом:

Parallel Nsight Pro стал бесплатным

Pro-версия Parallel Nsight теперь раздается бесплатно (раньше - только версия Standard). Берут отсюда, требуется регистрация.

Профессиональная версия отличается наличием отладчика.

Поддержка OpenCL включена в ATI/AMD Catalyst 11.x

До версии Catalyst 10.10, для поддержки OpenCL нужны были девелоперские драйвера, которые входили в Stream SDK. Начиная с 10.10 - появилась специальная версия с включенным OpenCL, которая ставилась как обычные драйвера, но пользователю нужно было специально скачивать эту версию.

В 11.1 (11.0 я пропустил и потому не знаю) все наоборот - в стандартный драйвер поддержка OpenCL входит, но если вам оно не надо, то можно скачать обрезанную версию на 20 мегабайт поменьше.

Другими словами, пользователя можно просить "поставь стандартные драйвера 11.x и будет тебе счастье.

Берут из стандартного места для Catalyst

AMD выпустил готовый комплект преподавателя для курса по OpenCL

AMD/ATI предлагает преподавателям готовый комплект из 13 презентаций и трех примеров кода (два примера - микроскопические, один - NBody), предназначенный для чтения лекций в университетах.

Про качество мне судить трудно, но посмотреть всяко имеет смысл.

Берут отсюда