CUDA: Как работает GPU

[narma]

Большое спасибо за статью.

[Klementev]

Пока не всё осилил, но статья интересная.

P.S.
Кстати в фотошопе можно не только nVidia задействовать, новые модели Radeon'a тоже неплохо ускоряют обработку.

[Inco]

Ам… дык было же уже?

Сам только добавить могу, так как по работе уже применял сие =)
Перенёс один метод под видео карту… получил реальное ускорение в 30(!!!) раз
НО спешу добавить, что большой скачок получается только когда я воткнул в принципы группировки потоков в блок. и использование __constant__ и __share__ модификаторов переменных.

Самая идея там если что-то можно объекдинить в блок — объединять, тогда каждый поток может скопировать например сперва в шаред массив свою ячейку, потом синхронизироваться со всеми, и потом работать только с локальной памятью.
Чтение из обычной 400 клоков, чтение из локальной от 4х… так что вот и выйгрыш. Жаль что там только 16к на локальную память на блок, но пока и этого хватает за глаза.
все параметры, что менять не надо писать в константную область памяти (её 64к на блок)!

А так тема реальная!

[Shmakov]

А каким образом с хоста записать данные в константную память? И не подскажите, возможна ли запись двумерных массивов в эту память?

[Inco]

к примеру так:
__constant__ int _threshold;
…
cudaMemcpyToSymbol( _threshold, threshold, sizeof(int));

Для работы с двумерными массивами там все есть, и специальные функции для выделения даже, чтобы выравнивание по адресам было!

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

[MaxFX]

Подробнее о памяти будет в следующей публикации.

[Inco]

Ещё раз отдельное спасибо минусующим! я вас так люблю! =)))))))

[jawbreaker]

Спасибо за статью, как раз начал разбираться с CUDA :)

[Shmakov]

А продолжение будет? )
Очень интересно как реализовывать хотя бы несколько более сложные задачи, чем Вы привели в примере. В частности интересно узнать про деление на блоки (как в этом случае нить узнает над какими данными она работает ?).
Или какой-нибудь реальный пример из жизни перевода программы с CPU на GPU.
За статью спасибо. Полезно.

[Inco]

Поставьте SDK там есть сэмплы, там есть пример умножения матриц, по разберите его и все станет ясно как Божий свет =)

[soloweb]

Спасибо! Очень нужный маттериал.

[Toshas]

MaxFX, так держать.
Однако в Вашу статью вкралась идеологическая ошибка. Размер грида равный 1 никогда не позволит выжать максимум мощи GPU при такой 100% распараллеливаемой задаче как сложение векторов. Да, формально код написан верно, но скорость работы будет на порядок меньше, чем если написать правильно. А правильно будет загрузить все стрим процессоры задачей, задав меньше размер блока и больше размер грида.
Вот например размер блока 32 треда (один варп), размер грида — 16 блоков. В таком случае разные части массива будут обрабатываться параллельно. Пока весь процесс обработки заключен в одном блоке — работает один из десятков процессоров, а остальные простаивают.

[MaxFX]

Знаю. Это сделано для облегчения понимания материала.

[Rational_Yurij]

Картинки не грузятся

[MaxFX]

Все в ваших руках. С удовольствием бы почитал.
У самого под рукой нет видеокарт ATI, сам бы помучил Brook в свободное время.

[Frosty]

По вашему мануалу переписал свою старую программку для рисования фигур лиссажу на OpenGL… все шикарно, но… оно работает слишком быстро, рисует с такой дикой скоростью, что по экрану ползают вверх-вниз горизонтальные полосы :) Как бы это дело замедлить? :)

[superhabra]

по экрану ползают вверх-вниз горизонтальные полосы

двойная буферизация?

[Frosty]

Без двойной буфферизации у меня бы плыли полосы и на старой программе, сейчас fps за 2000 кадров и это выглядит примерно вот так


Тут конечно еще немного артефактов накинул сам скриншот, но суть ясна.

[Frosty]

Ай-ай-ай, извиняюсь, был не прав, когда решил переписать старую прогу, то взял какую то очень давнишнюю версию, где двойной буфер не был включен, включил — все стало отлично.

[superhabra]

Что и требовалось доказать :)

[icoz]

Картинки почините, плиз.