CUDA 4.0

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...

Tags: 

Comments

Не только Apple с OpenCL, но

Не только Apple с OpenCL, но и Intel с альфой OpenCL (пока только CPU, потом интеловские встроенные в процессоры GPU), и AMD со своим OpenCL, который не только GPU, но и x86 CPU поддерживает.

Да, это верное замечание, но

Да, это верное замечание, но под CPU можно оптимизировать и просто локальной ручной векторизацией хот-спотов.
Насчет OpenCL на интеловских CPU/GPU - было бы здорово увидеть это на текущем поколении процессоров, но что-то сомнения у меня.

А так - будем посмотреть. Как мне кажется, если у CUDA4 будет принципиальная легкость в программировании относительно CUDA2-3 и OpenCL (которые практически одинаковы по API), то это серьезное преимущество.

Ну а как векторизировать? Вот

Ну а как векторизировать? Вот вышли процессоры с инструкциями AVX и, получается, что по хорошему надо уже 2 ветки писать - одна для SSE, другая для AVX. И к этому еще писать на CUDA, чтобы ускорять только на Nvidia картах? Время разработчиков и тестеров денег стоит.

Не проще ли написать один раз на OpenCL, и получить сравнительно быстрый код везде? Тот же LLVM (используемый для компиляции кернелов для CPU в AMD APP и вроде как в альфе интела) быстро адаптируется к новым процессорам и про AVX давно знает.

Мне кажется, первый релиз OpenCL драйвера для x86(-64) с redistributable package мгновенно перевесит чашу весов в пользу OpenCL.

Да драйвера такие есть,

Да драйвера такие есть, правда без AVX, называются ATI Catalyst. В релизе уже довольно давно.

Но с OpenCL все не очень радужно. Код переносится, а производительность - нет.

Вот простейший пример, сложение векторов, складываем в несколько потоков, N элементов в потоке, base - номер с которого начинает работать данный поток.
1) Intel

  1. for(j=0;j<N;j++)
  2.        C[j+base] = B[j+base]+C[j+base]

2) NVidia (одномерная сетка, CUDA)

  1.  for(j=0; j<blockDim.x*N; j+=blockDim.x)
  2.        C[j+base] = B[j+base]+C[j+base]

Потому что если на NVidia запустить первый код, то у тредов одного варпа будет не-coalesced доступ.
А если на Intel запустить второй, то оно или вовсе не векторизуется, или векторизуется крайне неэффективно так как загрузка пойдет из разных мест памяти.

OpenCL код 1. const uint j =

OpenCL код

1. const uint j = get_global_id(0);
2. С[j] = B[j] + C[j];

будет работать быстро и на CPU, и на GPU. И coalesced запись на Nvidia тоже будет. Максимум, что придется сделать - это аккуратно задать workgroup size в enqueueNDRangeKernel, но хардкодить ничего не надо, надо высчитывать его из параметров девайса (которые, опять же, можно получить вызывав соответствующий методв в OpenCL). Да и это не обязательно.

Coalesced доступ - это не уникальная "фича" NVidia. AMD карты тоже выигрывают, когда данные пишутся таким образом.

Но это все не важно. Важно вот что: "Код переносится, а производительность - нет". Если писать на CUDA, то не переносится ни код, ни производительность.

У меня тут антиспамовые

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

Я не поленился и запихал эти две строчки в Intel OpenCL offline compiler. До счастья тоже очень далеко: сначала 8 загрузок из разных мест по слову, потом mulps, потом movss и три extractps.
"Жить будете, но хреново...."

Я совершенно не спорю с тем тезисом, что CUDA - это пока вообще не переносимо (пока, потому что PGI обещает CUDA-x86), ну и GPUOcelot есть, хотя я его не пробовал.

Но OpenCL от счастья тоже очень далек, если интересует результат, то придется таки написать по kernel под каждую архитектуру. Ну вот разве что SSE/AVX будут вам автоматом, но опять же во многих (простых) случаях уже SSE-код уперт в память и от AVX выигрыша никакого нет.

Еще интересно (я правда не знаю), что у CPU-OpenCL с clEnqueueWriteBuffer() и подобным. Оно реально копирует куда-то к себе в карман?

Короче, счастья нет.

"сначала 8 загрузок из разных

"сначала 8 загрузок из разных мест по слову, потом mulps, потом movss и три extractps" - это для обработки float или float4?

Весь труд стоит деньги. Полагаю, с выходом в массы, к конечному потребителю параллельных вычислений (а это неизбежно) именно программирование "good enough" с максимальным покрытием железа будет востребовано.

clEnqueueWriteBuffer - работа асинхронного варианта очевидно зависит от реализации. Т.е. никакого кармана может и не быть. Но есть еще набор функций MapBuffer, но я с ними не игрался.

Это float, естественно. С

Это float, естественно.

С float4 код получается какой-то для меня удивительный, какие-то сплошные shuffle, это в векторном варианте. Понять за минуту не смог, заколдобился.
В невекторном тоже интересно, там явно минимизация использования регистров происходит, обходится одним XMM0

А стоит внутри kernel написать цикл от 0 до 127 - так сразу восторг, благолепие, понятный мне и компактный код для float. И он будет весьма эффективным, упрется в память.

Цикл для float4 остается для меня непонятным (тоже сплошные shuffle), может быть это попытка обработки произвольного выравнивания. Не могу осилить, не выспался сегодня жестоко....

О, я, кстати, понял, что

О, я, кстати, понял, что такое "сначала 8 загрузок из разных мест по слову, потом mulps, потом movss и три extractps". Это заявленная интелом векторизация кернелов. Объединяются 4 work-item для того, чтобы заиспользовать SSE.

Ну да, естественно. Только

Ну да, естественно. Только вот, похоже, что никаких гарантий порядка исполнения нет, поэтому оно и грузит черти как.
Если бы можно было сгруппировать kernels так, чтобы get_global_id гарантированно давал бы результаты "подряд", то две загрузки, умножение (я, для наглядности смотрел на c = a*b, чтобы собственно рабочий код сразу отличался от арифметики с адресами), запись содержимого по одному адресу (тем паче, для i7 выравнивание не обязательно).

А вот с float4 там что-то непостижимое.

Гарантия порядка исполнения

Гарантия порядка исполнения как раз есть. В конце концов, это ж интеловская реализация OpenCL, они вольны ее написать как им заблогорассудится. А вот чего нет - так это того, что соседние work-item-ы будут читать и писать в соседние области памяти. Для того, чтобы это понять, надо ж код функции анализировать.

Когда у меня цикл (for

Когда у меня цикл (for i=0...128) с[i+base]=a[i+base]*b[i+base] - оно код функции анализирует.
А когда цикла нет т.к. он подразумевается снаружи (цикл по get_global_id) - не анализирует.

Ну да, не выкопали еще.

> А вот с float4 там что-то

> А вот с float4 там что-то непостижимое.

Сейчас читаю документ интела по оптимизации работы с интеловским OpenCL.

Vectorization module transforms scalar operations on adjacent work-items into an equivalent vector operation. When vector operations already exist in the kernel source code, they are scalarized (broken down into component operations) and re-vectored. This provides additional performance gain by transforming the memory access pattern of the kernel into structure of arrays (SOA), which is often more cache-friendly than the array of structures (AOS).

Попробуйте поставить аттрибут в описании функции:
__kernel __attribute__((vec_type_hint(float4)))

Полагаю, получившийся код будет существенно проще.

Да а по сути ("good enough")

Да а по сути ("good enough") - целевые функции у всех свои. Какой-нибудь видеодекодер будут вылизывать до полного блеска, кода немного, а пользователей много. А что-то узкоспециальное и дорогое - наоборот (ибо несложно заставить пользователя теслу купить, если его интересует результат).

С многоядерностью на CPU тоже какие-то непонятки. Все хотят и мечтают, что не мешает тому же MS-у сломать OpenMP в VS2010.
В колыбели все еще, что приятно т.к. возможностей много.

Закон Мура реализовывается

Закон Мура реализовывается многоядерностью. Уже код для дескотопов должен быть параллельным, чтобы быть конкурентноспособным. Но, в общем то, я не спорю, что останется резон тратить ресурсы на написание очень эффективного кода для тех же тесл.

А что с OpenMP в VS2010? У меня работает.

Про многоядерность,

Про многоядерность, векторизацию и десктопы я согласен. Но OpenCL слишком уж неуклюж. Если у меня маленький хотспот на 20 строчек кода, то вылить в буфер, запустить kernel, вылить из буфера - может не иметь смысла и надо нести сразу большой кусок кода на "устройство" (даже если это устройство - то же CPU).
OpenMP и всякие мелкие библиотеки/улучшения (как у Интела) - тоже не панацея. Ручное выпиливание векторного кода мне тоже не нравится (хотя память насытить получается).
Все оно пока какое-то кривое, короче.

Про OpenMP/VS2010: http://blog.lexa.ru/2010/12/27/o_stepennoi_funktsii_openmp_i_prochikh_gr...
(ну и дальше там по ссылкам).

Лично я в деталях не разбирался, но странности, когда параллельный код медленнее последовательного - уже видел (и это много данных т.е. все нормально и у того же Intel C++ OpenMP работает).

Cпинлок на 100мс - это они

Cпинлок на 100мс - это они переборщили, да, на порядки :) Думаю, кто то там ошибся просто в коде. С нолями. Спасибо за инфу, буду в курсе. Впрочем, я ж сейчас могу попробовать проверить эту новость.

Попробовал. Добиться точно

Попробовал. Добиться точно описываемой ситуации (100% загрузка) у меня не получилось, но странности увидел. 50% получилось. Выглядт так, как будто на моем 2-ядерном Core 2 Duo, OpenMP реализация от MS оставляет вторичный поток OpenMP работать (спинлок) в течение 50мс. Основной поток, понятное дело, такими делами не занимается.

В общем, при таком раскладе 100% процессорных ресурсов сама прога вполне может сожрать, однако замедление работы самой программы наблюдаться не должно.

В моих экспериментах

В моих экспериментах замедление было по причине лока, причем на read-only переменной. Прошло после замены default(shared) на default(none) и явное расписывание (единственного) обрабатываемого массива как shared.

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

Да, 50мс - это дофига,

Да, 50мс - это дофига, согласен абсолютно. Интересно, есть такое же поведение в нововведениях в C++, которые пока все никак не могут вылиться в стандарт?

Ага. То ли расширения языка,

Ага. То ли расширения языка, то ли расширение stl, то ли и то, и другое. Я в этом особо не разбирался, жалко было тратить время на то, что стандартом не является. При наличии альтернатив типа OpenMP.

OpenMP + (просто) C++ - это

OpenMP + (просто) C++ - это тот еще уродец, прямо скажем:
* обработка исключений
* невозможность член класса описать как shared или там firstprivate

И вообще, стандартные data-parallel конструкции всяко лучше, чем vendor extensions (cilk plus и так далее).

Невозможность пробрасывать

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

Время разработчиков и

Время разработчиков и тестеров денег стоит.

Я вот посмотрел полчаса на CUDA4 (и вчера еще час послушал вебинар). Совершенно не споря с процитированным тезисом, хочу заметить, что nvidia сделала офигенный шаг вперед именно с точки зрения разработчиков и тестеров
1) втащили Thrust в релиз. Непонятным опенсорсом не все хотят пользоваться, библиотека с поддержкой вендора - совсем другое дело. То есть если нужно что-то быстро отсортировать или там scan/reduce какой несложный, то разработчику про OpenCL или CUDA и вовсе думать не надо, простые понятные вызовы и отлично. Сортировка реально быстро работает
2) NPP (очень похожа на Intel IPP снаружи) - аналогично, думать не надо, хочешь гистограмму, хочешь RGB2Luv хочешь DCT-таблицы, все типичные такие задачи.

Что толку от ускорения

Что толку от ускорения написания кода, если этот код будет работать лишь на очень небольшом проценте компьютеров?

Это очень сильно зависит от

Это очень сильно зависит от программы.

Скажем, на Adobe Premiere CS5 совершенно не стесняются написать "если вам медленно - идите и купите себе NV Quadro за пару штук".
То же самое с играми: не нравится как играет наш Кризис? Вам в магазин!

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

Adobe стоит в разы больше

Adobe стоит в разы больше видеокарты. Можно и посоветовать именно NVidia, что уж тут. Стоил бы Адоб еще подороже немного, можно было бы в коробку ускоритель докладывать.

А с другой стороны, у Адоба же есть конкуренты, верно? Которые вполне могут заиспользовать OpenCL и с большой улыбкой на лице говорить клиентам "Наш продукт использует ускорители всех ведущих производителей".

А игры - нет. В играх все равно какую видеокарту покупать. Впрочем, есть припритарный PhysX, но с ним Nvidia такой стыд учудила, что просто ужас же.

Мой поинт в том, что если

Мой поинт в том, что если какой-то продукт используется для профессиональной деятельности т.е. за ним сидят 8 часов в день, 5 дней в неделю, то написать в пожеланиях "если вы хотите ускорения таких-то операций в 10 раз - купите карточку из этого списка за $200" - не запретительно. И даже за $2000 - не запретительно, "окупится за месяц"
И в больших пакетах: Adobe, Matlab, ANSYS - так уже делают.

При наличии конкуренции пользователю будет легче. Пока ее по факту практически нет, а те кто пишут конкурентов Адобу - вынуждены выбирать между time to market (NVidia NPP) и поддержкой оборудования (OpenCL и руками программировать). Это грубое приближение, конечно.

Повторюсь, в области imaging эппл здорово поднасрал NVidia новыми ноутами.

Я согласен с этой Вашей

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

> в области imaging эппл здорово поднасрал NVidia новыми ноутами

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

Если попытаться перейти от

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

Как мы уже выяснили тут, "производительность не переносится" и в способность одного куска кода, написанного "просто прикладным программистом" работать хотя бы на 50% эффективности на трех видах железа (CPU, 2xGPU, а скоро еще всякие ARM-ы добавятся) - я верю слабо. Тут и coalesced access и много всякой другой веселухи (вроде 24/32 bit int на разных версиях nvidia).

Казалос бы, тот же Thrust или CUFFT - отличная же нашлепка сверху, которая может выбирать оптимальный execution path для разного железа. Пусть даже с OpenCL intermediate. И пусть только для типовых задач: sort, map, reduce, FFT, BLAS, гистограммы (несть их числа).

Потому что сейчас ужасный же бардак. Ну вот тот же BLAS: MKL на одних CPU, ACML на других (и одном виде GPU), CUBLAS - на третьих. Для вычислительных кластеров это терпимо, а для end-user - тоже нет.

Как правильно ходить, от псевдокода (вроде PTX) с трансляцией на лету, или же от языка или вовсе от темплейтов C++ - это я не знаю (но и занимаюсь я другим). Но вот взять какие-то very integrated performance primitives, которые на всех видеокартах и всех CPU посчитают мне гистограмму у HD-потока видео (к примеру) - очень бы хотелось.

Хотя библиотеки плохо продаются, это да.