NVidia CUDA: вычисления на видеокарте или смерть CPU? Что такое NVIDIA CUDA? карты с поддержкой cuda.

Согласно Дарвинской теории эволюции, первая человекообразная обезьяна (если
быть точным – homo antecessor, человек-предшественник) превратилась впоследствии
в нас. Многотонные вычислительные центры с тысячью и больше радиоламп,
занимающие целые комнаты, сменились полукилограммовыми ноутами, которые, кстати,
не уступят в производительности первым. Допотопные печатные машинки превратились
в печатающие что угодно и на чем угодно (даже на теле человека)
многофункциональные устройства. Процессорные гиганты вдруг вздумали замуровать
графическое ядро в «камень». А видеокарты стали не только показывать картинку с
приемлемым FPS и качеством графики, но и производить всевозможные вычисления. Да
еще как производить! О технологии многопоточных вычислений средствами GPU, и пойдет речь.

Почему GPU?

Интересно, почему всю вычислительную мощь решили переложить на графический
адаптер? Как видно, процессоры еще в моде, да и вряд ли уступят свое теплое
местечко. Но у GPU есть пара козырей в рукаве вместе с джокером, да и рукавов
хватает. Современный центральный процессор заточен под получение максимальной
производительности при обработке целочисленных данных и данных с плавающей
запятой, особо не заботясь при этом о параллельной обработке информации. В то же
время архитектура видеокарты позволяет быстро и без проблем «распараллелить»
обработку данных. С одной стороны, идет обсчет полигонов (за счет 3D-конвейера),
с другой – пиксельная обработка текстур. Видно, что происходит «слаженная
разбивка» нагрузки в ядре карты. Кроме того, работа памяти и видеопроцессора
оптимальнее, чем связка «ОЗУ-кэш-процессор». В тот момент, когда единица данных
в видеокарте начинает обрабатываться одним потоковым процессором GPU, другая
единица параллельно загружается в другой, и, в принципе, легко можно достичь
загруженности графического процессора, сравнимой с пропускной способностью шины,
однако для этого загрузка конвейеров должна осуществляться единообразно, без
всяких условных переходов и ветвлений. Центральный же процессор в силу своей
универсальности требует для своих процессорных нужд кэш, заполненный
информацией.

Ученые мужи задумались насчет работы GPU в параллельных вычислениях и
математике и вывели теорию, что многие научные расчеты во многом схожи с
обработкой 3D-графики. Многие эксперты считают, что основополагающим фактором в
развитии GPGPU (General Purpose computation on GPU – универсальные
расчеты средствами видеокарты ) стало появление в 2003 году проекта Brook GPU.

Создателям проекта из Стэндфордского университета предстояло решить непростую
проблему: аппаратно и программно заставить графический адаптер производить
разноплановые вычисления. И у них это получилось. Используя универсальный язык C,
американские ученые заставили работать GPU как процессор, с поправкой на
параллельную обработку. После Brook появился целый ряд проектов по VGA-расчетам,
таких как библиотека Accelerator, библиотека Brahma, система
метапрограммирования GPU++ и другие.

CUDA!

Предчувствие перспективности разработки заставило AMD и NVIDIA
вцепиться в Brook GPU, как питбуль. Если опустить маркетинговую политику, то,
реализовав все правильно, можно закрепиться не только в графическом секторе
рынка, но и в вычислительном (посмотри на специальные вычислительные карты и
серверы Tesla с сотнями мультипроцессоров), потеснив привычные всем CPU.

Естественно, «повелители FPS» разошлись у камня преткновения каждый по своей
тропе, но основной принцип остался неизменным – производить вычисления
средствами GPU. И сейчас мы подробнее рассмотрим технологию «зеленых» – CUDA
(Compute Unified Device Architecture ).

Работа нашей «героини» заключается в обеспечении API, причем сразу двух.
Первый – высокоуровневый, CUDA Runtime, представляет собой функции, которые
разбиваются на более простые уровни и передаются нижнему API – CUDA Driver. Так
что фраза «высокоуровневый» применима к процессу с натяжкой. Вся соль находится
именно в драйвере, и добыть ее помогут библиотеки, любезно созданные
разработчиками NVIDIA : CUBLAS (средства для математических расчетов) и
FFT (расчет посредством алгоритма Фурье). Ну что ж, перейдем к практической
части материала.

Терминология CUDA

NVIDIA оперирует весьма своеобразными определениями для CUDA API. Они
отличаются от определений, применяемых для работы с центральным процессором.

Поток (thread) – набор данных, который необходимо обработать (не
требует больших ресурсов при обработке).

Варп (warp) – группа из 32 потоков. Данные обрабатываются только
варпами, следовательно варп – это минимальный объем данных.

Блок (block) – совокупность потоков (от 64 до 512) или совокупность
варпов (от 2 до 16).

Сетка (grid) – это совокупность блоков. Такое разделение данных
применяется исключительно для повышения производительности. Так, если число
мультипроцессоров велико, то блоки будут выполняться параллельно. Если же с
картой не повезло (разработчики рекомендуют для сложных расчетов использовать
адаптер не ниже уровня GeForce 8800 GTS 320 Мб), то блоки данных обработаются
последовательно.

Также NVIDIA вводит такие понятия, как ядро (kernel) , хост (host)
и девайс (device) .

Работаем!

Для полноценной работы с CUDA нужно:

1. Знать строение шейдерных ядер GPU, так как суть программирования
заключается в равномерном распределении нагрузки между ними.
2. Уметь программировать в среде C, с учетом некоторых аспектов.

Разработчики NVIDIA раскрыли «внутренности» видеокарты несколько
иначе, чем мы привыкли видеть. Так что волей-неволей придется изучать все
тонкости архитектуры. Разберем строение «камня» G80 легендарной GeForce 8800
GTX .

Шейдерное ядро состоит из восьми TPC (Texture Processor Cluster) – кластеров
текстурных процессоров (так, у GeForce GTX 280 – 15 ядер, у 8800 GTS
их шесть, у 8600 – четыре и т.д.). Те, в свою очередь, состоят из двух
потоковых мультипроцессоров (streaming multiprocessor – далее SM). SM (их всего
16) состоит из front end (решает задачи чтения и декодирования инструкций) и
back end (конечный вывод инструкций) конвейеров, а также восьми scalar SP (shader
processor) и двумя SFU (суперфункциональные блоки). За каждый такт (единицу
времени) front end выбирает варп и обрабатывает его. Чтобы все потоки варпа
(напомню, их 32 штуки) обработались, требуется 32/8 = 4 такта в конце конвейера.

Каждый мультипроцессор обладает так называемой общей памятью (shared memory).
Ее размер составляет 16 килобайт и предоставляет программисту полную свободу
действий. Распределяй как хочешь:). Shared memory обеспечивает связь потоков в
одном блоке и не предназначена для работы с пиксельными шейдерами.

Также SM могут обращаться к GDDR. Для этого им «пришили» по 8 килобайт
кэш-памяти, хранящих все самое главное для работы (например, вычислительные
константы).

Мультипроцессор имеет 8192 регистра. Число активных блоков не может быть
больше восьми, а число варпов – не больше 768/32 = 24. Из этого видно, что G80
может обработать максимум 32*16*24 = 12288 потоков за единицу времени. Нельзя не
учитывать эти цифры при оптимизации программы в дальнейшем (на одной чашу весов
– размер блока, на другой – количество потоков). Баланс параметров может сыграть
важную роль в дальнейшем, поэтому NVIDIA рекомендует использовать блоки
со 128 или 256 потоками. Блок из 512 потоков неэффективен, так как обладает
повышенными задержками. Учитывая все тонкости строения GPU видеокарты плюс
неплохие навыки в программировании, можно создать весьма производительное
средство для параллельных вычислений. Кстати, о программировании...

Программирование

Для «творчества» вместе с CUDA требуется видеокарта GeForce не ниже
восьмой серии . С

официального сайта нужно скачать три программных пакета: драйвер с
поддержкой CUDA (для каждой ОС – свой), непосредственно пакет CUDA SDK (вторая
бета-версия) и дополнительные библиотеки (CUDA toolkit). Технология поддерживает
операционные системы Windows (XP и Vista), Linux и Mac OS X. Для изучения я
выбрал Vista Ultimate Edition x64 (забегая вперед, скажу, что система вела себя
просто превосходно). В момент написания этих строк актуальным для работы был
драйвер ForceWare 177.35. В качестве набора инструментов использовался
программный пакет Borland C++ 6 Builder (хотя подойдет любая среда, работающая с
языком C).

Человеку, знающему язык, будет легко освоиться в новой среде. Требуется лишь
запомнить основные параметры. Ключевое слово _global_ (ставится перед функцией)
показывает, что функция относится к kernel (ядру). Ее будет вызывать центральный
процессор, а вся работа произойдет на GPU. Вызов _global_ требует более
конкретных деталей, а именно размер сетки, размер блока и какое ядро будет
применено. Например, строчка _global_ void saxpy_parallel<<>>, где X –
размер сетки, а Y – размер блока, задает эти параметры.

Символ _device_ означает, что функцию вызовет графическое ядро, оно же
выполнит все инструкции. Эта функция располагается в памяти мультипроцессора,
следовательно, получить ее адрес невозможно. Префикс _host_ означает, что вызов
и обработка пройдут только при участии CPU. Надо учитывать, что _global_ и
_device_ не могут вызывать друг друга и не могут вызывать самих себя.

Также язык для CUDA имеет ряд функций для работы с видеопамятью: cudafree
(освобождение памяти между GDDR и RAM), cudamemcpy и cudamemcpy2D (копирование
памяти между GDDR и RAM) и cudamalloc (выделение памяти).

Все программные коды проходят компиляцию со стороны CUDA API. Сначала берется
код, предназначенный исключительно для центрального процессора, и подвергается
стандартной компиляции, а другой код, предназначенный для графического адаптера,
переписывается в промежуточный язык PTX (сильно напоминает ассемблер) для
выявления возможных ошибок. После всех этих «плясок» происходит окончательный
перевод (трансляция) команд в понятный для GPU/CPU язык.

Набор для изучения

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

Специально для новичков разработан CUDA SDK Browser. Любой желающий может
ощутить силу параллельных вычислений на своей шкуре (лучшая проверка на
стабильность – работа примеров без артефактов и вылетов). Приложение имеет
большой ряд показательных мини-программок (61 «тест»). К каждому опыту имеется
подробная документация программного кода плюс PDF-файлы. Сразу видно, что люди,
присутствующие со своими творениями в браузере, занимаются серьезной работой.
Тут же можно сравнить скорости работы процессора и видеокарты при обработке
данных. Например, сканирование многомерных массивов видеокартой GeForce 8800
GT 512 Мб с блоком с 256 потоками производит за 0.17109 миллисекунды.
Технология не распознает SLI-тандемы, так что если у тебя дуэт или трио,
отключай функцию «спаривания» перед работой, иначе CUDA увидит только один
девайс. Двуядерный AMD Athlon 64 X2 (частота ядра 3000 МГц) тот же опыт
проходит за 2.761528 миллисекунды. Получается, что G92 более чем в 16 раз
быстрее «камня» AMD ! Как видишь, далеко не экстремальная система в
тандеме с нелюбимой в массах операционной системой показывает неплохие
результаты.

Помимо браузера существует ряд полезных обществу программ. Adobe
адаптировала свои продукты к новой технологии. Теперь Photoshop CS4 в полной
мере использует ресурсы графических адаптеров (необходимо скачать специальный
плагин). Такими программами, как Badaboom media converter и RapiHD можно
произвести декодирование видео в формат MPEG-2. Для обработки звука неплохо
подойдет бесплатная утилита Accelero. Количество софта, заточенного под CUDA API,
несомненно, будет расти.

А в это время…

А пока ты читаешь сей материал, трудяги из процессорных концернов
разрабатывают свои технологии по внедрению GPU в CPU. Со стороны AMD все
понятно: у них есть большущий опыт, приобретенный вместе с ATI .

Творение «микродевайсеров», Fusion, будет состоять из нескольких ядер под
кодовым названием Bulldozer и видеочипа RV710 (Kong). Их взаимосвязь будет
осуществляться за счет улучшенной шины HyperTransport. В зависимости от
количества ядер и их частотных характеристик AMD планирует создать целую ценовую
иерархию «камней». Также планируется производить процессоры как для ноутбуков (Falcon),
так и для мультимедийных гаджетов (Bobcat). Причем именно применение технологии
в портативных устройствах будет первоначальной задачей для канадцев. С развитием
параллельных вычислений применение таких «камней» должно быть весьма популярно.

Intel немножко отстает по времени со своей Larrabee. Продукты AMD ,
если ничего не случится, появятся на прилавках магазинов в конце 2009 – начале
2010 года. А решение противника выйдет на свет божий только почти через два
года.

Larrabee будет насчитывать большое количество (читай – сотни) ядер. Вначале
же выйдут продукты, рассчитанные на 8 – 64 ядера. Они очень сходны с Pentium, но
довольно сильно переработаны. Каждое ядро имеет 256 килобайт кэша второго уровня
(со временем его размер увеличится). Взаимосвязь будет осуществляться за счет
1024-битной двунаправленной кольцевой шины. Интел говорит, что их «дитя» будет
отлично работать с DirectX и Open GL API (для «яблочников»), поэтому никаких
программных вмешательств не потребуется.

А к чему я все это тебе поведал? Очевидно, что Larrabee и Fusion не вытеснят
обычные, стационарные процессоры с рынка, так же, как не вытеснят с рынка
видеокарты. Для геймеров и экстремалов пределом мечтаний по-прежнему останется
многоядерный CPU и тандем из нескольких топовых VGA. Но то, что даже
процессорные компании переходят на параллельные вычисления по принципам,
аналогичным GPGPU, говорит уже о многом. В частности о том, что такая
технология, как CUDA, имеет право на существование и, по всей видимости, будет
весьма популярна.

Небольшое резюме

Параллельные вычисления средствами видеокарты – всего лишь хороший инструмент
в руках трудолюбивого программиста. Вряд ли процессорам во главе с законом Мура
придет конец. Компании NVIDIA предстоит пройти еще длинный путь по
продвижению в массы своего API (то же можно сказать и о детище ATI/AMD ).
Какой он будет, покажет будущее. Так что CUDA will be back:).

P.S. Начинающим программистам и заинтересовавшимся людям рекомендую посетить
следующие «виртуальные заведения»:

официальный сайт NVIDIA и сайт
GPGPU.com . Вся
предоставленная информация – на английском языке, но, спасибо хотя бы, что не на
китайском. Так что дерзай! Надеюсь, что автор хоть немного помог тебе в
захватывающих начинаниях познания CUDA!

И другие. Однако, поиск комбинации «CUDA scan » выдал всего 2 статьи никак не связанные с, собственно, алгоритмом scan на GPU - а это один из самых базовых алгоритмов. Поэтому, вдохновившись только что просмотренным курсом на Udacity - Intro to Parallel Programming , я и решился написать более полную серию статей о CUDA. Сразу скажу, что серия будет основываться именно на этом курсе, и если у вас есть время - намного полезнее будет пройти его. На данный момент планируются следующие статьи:
Часть 1: Введение.
Часть 2: Аппаратное обеспечение GPU и шаблоны параллельной коммуникации.
Часть 3: Фундаментальные алгоритмы GPU: свертка (reduce), сканирование (scan) и гистограмма (histogram).
Часть 4: Фундаментальные алгоритмы GPU: уплотнение (compact), сегментированное сканирование (segmented scan), сортировка. Практическое применение некоторых алгоритмов.
Часть 5: Оптимизация GPU программ.
Часть 6: Примеры параллелизации последовательных алгоритмов.
Часть 7: Дополнительные темы параллельного программирования, динамический параллелизм.

Задержка vs пропускная способность

Первый вопрос, который должен задать каждый перед применением GPU для решения своих задач - а для каких целей хорош GPU, когда стоит его применять? Для ответа нужно определить 2 понятия:
Задержка (latency) - время, затрачиваемое на выполнение одной инструкции/операции.
Пропускная способность - количество инструкций/операций, выполняемых за единицу времени.
Простой пример: имеем легковой автомобиль со скоростью 90 км/ч и вместимостью 4 человека, и автобус со скоростью 60 км/ч и вместимостью 20 человек. Если за операцию принять перемещение 1 человека на 1 километр, то задержка легкового автомобиля - 3600/90=40с - за столько секунд 1 человек преодолеет расстояние в 1 километр, пропускная способность автомобиля - 4/40=0.1 операций/секунду; задержка автобуса - 3600/60=60с, пропускная способность автобуса - 20/60=0.3(3) операций/секунду.
Так вот, CPU - это автомобиль, GPU - автобус: он имеет большую задержку но также и большую пропускную способность. Если для вашей задачи задержка каждой конкретной операции не настолько важна как количество этих операций в секунду - стоит рассмотреть применение GPU.

Базовые понятия и термины CUDA

• Устройство (device) - GPU. Выполняет роль «подчиненного» - делает только то, что ему говорит CPU.
• Хост (host) - CPU. Выполняет управляющую роль - запускает задачи на устройстве, выделяет память на устройстве, перемещает память на/с устройства. И да, использование CUDA предполагает, что как устройство так и хост имеют свою отдельную память.
• Ядро (kernel) - задача, запускаемая хостом на устройстве.

Основные этапы CUDA-программы

1. Хост выделяет нужное количество памяти на устройстве.
2. Хост копирует данные из своей памяти в память устройства.
3. Хост стартует выполнение определенных ядер на устройстве.
4. Устройство выполняет ядра.
5. Хост копирует результаты из памяти устройства в свою память.

Ядра

1. Внутри ядер вы имеете возможность узнать «идентификатор» или, проще говоря, позицию потока, который сейчас выполняется - используя эту позицию мы добиваемся того, что одно и то же ядро будет работать с разными данными в зависимости от потока, в котором оно запущено. Кстати, такая организация параллельных вычислений называется SIMD (Single Instruction Multiple Data) - когда несколько процессоров выполняют одновременно одну и ту же операцию но на разных данных.
2. В некоторых случаях в коде ядра необходимо использовать различные способы синхронизации.

• Сначала задаются размеры так называемой сетки (grid), в 3D координатах: grid_x, grid_y, grid_z . В результате, сетка будет состоять из grid_x*grid_y*grid_z блоков.
• Потом задаются размеры блока в 3D координатах: block_x, block_y, block_z . В результате, блок будет состоять из block_x*block_y*block_z потоков. Итого, имеем grid_x*grid_y*grid_z*block_x*block_y*block_z потоков. Важное замечание - максимальное количество потоков в одном блоке ограничено и зависит от модели GPU - типичны значения 512 (более старые модели) и 1024 (более новые модели).
• Внутри ядра доступны переменные threadIdx и blockIdx с полями x, y, z - они содержат 3D координаты потока в блоке и блока в сетке соответственно. Также доступны переменные blockDim и gridDim с теми же полями - размеры блока и сетки соответственно.

Пишем первую программу на CUDA

main.cpp

#include #include #include #include #include #include #include #include #include "openMP.hpp" #include "CUDA_wrappers.hpp" #include "common/image_helpers.hpp" using namespace cv; using namespace std; int main(int argc, char** argv) { using namespace std::chrono; if(argc != 2) { cout <<" Usage: convert_to_grayscale imagefile" << endl; return -1; } Mat image, imageGray; uchar4 *imageArray; unsigned char *imageGrayArray; prepareImagePointers(argv, image, &imageArray, imageGray, &imageGrayArray, CV_8UC1); int numRows = image.rows, numCols = image.cols; auto start = system_clock::now(); RGBtoGrayscaleOpenMP(imageArray, imageGrayArray, numRows, numCols); auto duration = duration_cast(system_clock::now() - start); cout<<"OpenMP time (ms):" << duration.count() << endl; memset(imageGrayArray, 0, sizeof(unsigned char)*numRows*numCols); RGBtoGrayscaleCUDA(imageArray, imageGrayArray, numRows, numCols); return 0; }

prepareImagePointers

template void prepareImagePointers(const char * const inputImageFileName, cv::Mat& inputImage, T1** inputImageArray, cv::Mat& outputImage, T2** outputImageArray, const int outputImageType) { using namespace std; using namespace cv; inputImage = imread(inputImageFileName, IMREAD_COLOR); if (inputImage.empty()) { cerr << "Couldn"t open input file." << endl; exit(1); } //allocate memory for the output outputImage.create(inputImage.rows, inputImage.cols, outputImageType); cvtColor(inputImage, inputImage, cv::COLOR_BGR2BGRA); *inputImageArray = (T1*)inputImage.ptr(0); *outputImageArray = (T2*)outputImage.ptr(0); }

openMP.hpp

#include #include #include void RGBtoGrayscaleOpenMP(uchar4 *imageArray, unsigned char *imageGrayArray, int numRows, int numCols) { #pragma omp parallel for collapse(2) for (int i = 0; i < numRows; ++i) { for (int j = 0; j < numCols; ++j) { const uchar4 pixel = imageArray; imageGrayArray = 0.299f*pixel.x + 0.587f*pixel.y+0.114f*pixel.z; } } }

Скрытый текст

void RGBtoGrayscaleCUDA(const uchar4 * const h_imageRGBA, unsigned char* const h_imageGray, size_t numRows, size_t numCols) { uchar4 *d_imageRGBA; unsigned char *d_imageGray; const size_t numPixels = numRows * numCols; cudaSetDevice(0); checkCudaErrors(cudaGetLastError()); //allocate memory on the device for both input and output checkCudaErrors(cudaMalloc(&d_imageRGBA, sizeof(uchar4) * numPixels)); checkCudaErrors(cudaMalloc(&d_imageGray, sizeof(unsigned char) * numPixels)); //copy input array to the GPU checkCudaErrors(cudaMemcpy(d_imageRGBA, h_imageRGBA, sizeof(uchar4) * numPixels, cudaMemcpyHostToDevice));

Скрытый текст

#include #define checkCudaErrors(val) check((val), #val, __FILE__, __LINE__) template void check(T err, const char* const func, const char* const file, const int line) { if (err != cudaSuccess) { std::cerr << "CUDA error at: " << file << ":" << line << std::endl; std::cerr << cudaGetErrorString(err) << " " << func << std::endl; exit(1); } }

Скрытый текст

dim3 blockSize; dim3 gridSize; int threadNum; cudaEvent_t start, stop; cudaEventCreate(&start); cudaEventCreate(&stop); threadNum = 1024; blockSize = dim3(threadNum, 1, 1); gridSize = dim3(numCols/threadNum+1, numRows, 1); cudaEventRecord(start); rgba_to_grayscale_simple<<>>(d_imageRGBA, d_imageGray, numRows, numCols); cudaEventRecord(stop); cudaEventSynchronize(stop); cudaDeviceSynchronize(); checkCudaErrors(cudaGetLastError()); float milliseconds = 0; cudaEventElapsedTime(&milliseconds, start, stop); std::cout << "CUDA time simple (ms): " << milliseconds << std::endl;

rgba_to_grayscale_simple

Global__ void rgba_to_grayscale_simple(const uchar4* const d_imageRGBA, unsigned char* const d_imageGray, int numRows, int numCols) { int y = blockDim.y*blockIdx.y + threadIdx.y; int x = blockDim.x*blockIdx.x + threadIdx.x; if (x>=numCols || y>=numRows) return; const int offset = y*numCols+x; const uchar4 pixel = d_imageRGBA; d_imageGray = 0.299f*pixel.x + 0.587f*pixel.y+0.114f*pixel.z; }

Скрытый текст

checkCudaErrors(cudaMemcpy(h_imageGray, d_imageGray, sizeof(unsigned char) * numPixels, cudaMemcpyDeviceToHost)); cudaFree(d_imageGray); cudaFree(d_imageRGBA);

Скрытый текст

Процессор: Intel® Core(TM) i7-3615QM CPU @ 2.30GHz.
GPU: NVIDIA GeForce GT 650M, 1024 MB, 900 MHz.
RAM: DD3, 2x4GB, 1600 MHz.
OS: OS X 10.9.5.
Компилятор: g++ (GCC) 4.9.2 20141029.
CUDA компилятор: Cuda compilation tools, release 6.0, V6.0.1.
Поддерживаемая версия OpenMP: OpenMP 4.0.

rgba_to_grayscale_optimized

#define WARP_SIZE 32 __global__ void rgba_to_grayscale_optimized(const uchar4* const d_imageRGBA, unsigned char* const d_imageGray, int numRows, int numCols, int elemsPerThread) { int y = blockDim.y*blockIdx.y + threadIdx.y; int x = blockDim.x*blockIdx.x + threadIdx.x; const int loop_start = (x/WARP_SIZE * WARP_SIZE)*(elemsPerThread-1)+x; for (int i=loop_start, j=0; j

Скрытый текст

threadNum=128; const int elemsPerThread = 16; blockSize = dim3(threadNum, 1, 1); gridSize = dim3(numCols / (threadNum*elemsPerThread) + 1, numRows, 1); cudaEventRecord(start); rgba_to_grayscale_optimized<<>>(d_imageRGBA, d_imageGray, numRows, numCols, elemsPerThread); cudaEventRecord(stop); cudaEventSynchronize(stop); cudaDeviceSynchronize(); checkCudaErrors(cudaGetLastError()); milliseconds = 0; cudaEventElapsedTime(&milliseconds, start, stop); std::cout << "CUDA time optimized (ms): " << milliseconds << std::endl;

Теги: Добавить метки

Ядра CUDA – условное обозначение скалярных вычислительных блоков в видео-чипах NVidia , начиная с G 80 (GeForce 8 xxx, Tesla C-D-S870 , FX4 /5600 , 360M ). Сами чипы являются производными архитектуры. К слову, потому компания NVidia так охотно взялась за разработку собственных процессоров Tegra Series , основанных тоже на RISC архитектуре. Опыт работы с данными архитектурами очень большой.

CUDA ядро содержит в себе один один векторный и один скалярный юнит, которые за один такт выполняют по одной векторной и по одной скалярной операции, передавая вычисления другому мультипроцессору, либо в для дальнейшей обработки. Массив из сотен и тысяч таких ядер, представляет из себя значительную вычислительную мощность и может выполнять различные задачи в зависимости от требований, при наличии определённого софта поддерживающего . Применение может быть разнообразным: декодирование видеопотока, ускорение 2D/3D графики, облачные вычисления, специализированные математические анализы и т.д.

Довольно часто, объединённые профессиональные карты NVidia Tesla и NVidia Quadro , являются костяком современных суперкомпьютеров.

CUDA — ядра не претерпели каких либо значимых изменений со времён G 80 , но увеличивается их количество (совместно с другими блоками — ROP , Texture Units & etc) и эффективность параллельных взаимодействий друг с другом (улучшаются модули Giga Thread ).

К примеру:

GeForce

GTX 460 — 336 CUDA ядер

GTX 580 — 512 CUDA ядер

8800GTX — 128 CUDA ядер

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

Начиная с чипа GK110 (NVidia GeForce GTX 680) — CUDA ядра теперь не имеют удвоенную частоту, а общую со всеми остальными блоками чипа. Вместо этого было увеличено их количество примерно в три раза в сравнении с предыдущим поколением G110 .

Я расскажу о ключевых моментах компилятора CUDA, интерфейсе CUDA runtime API, ну, и в заключение, приведу пример использования CUDA для несложных математических вычислений.

Приступим.

Вычислительная модель GPU:

При использовании GPU вы можете задействовать грид необходимого размера и сконфигурировать блоки под нужды вашей задачи.

CUDA и язык C:

1. Спецификаторы функций, которые показывают, как и откуда буду выполняться функции.
2. Спецификаторы переменных, которые служат для указания типа используемой памяти GPU.
3. Спецификаторы запуска ядра GPU.
4. Встроенные переменные для идентификации нитей, блоков и др. параметров при исполнении кода в ядре GPU .
5. Дополнительные типы переменных.

• __host__ - выполнятся на CPU, вызывается с CPU (в принципе его можно и не указывать).
• __global__ - выполняется на GPU, вызывается с CPU.
• __device__ - выполняется на GPU, вызывается с GPU.

MyKernelFunc<<>>(float* param1,float* param2), где

• gridSize – размерность сетки блоков (dim3), выделенную для расчетов,
• blockSize – размер блока (dim3), выделенного для расчетов,
• sharedMemSize – размер дополнительной памяти, выделяемой при запуске ядра,
• cudaStream – переменная cudaStream_t, задающая поток, в котором будет произведен вызов.

Так же стоит упомянуть о встроенных переменных:

• gridDim – размерность грида, имеет тип dim3. Позволяет узнать размер гридa, выделенного при текущем вызове ядра.
• blockDim – размерность блока, так же имеет тип dim3. Позволяет узнать размер блока, выделенного при текущем вызове ядра.
• blockIdx – индекс текущего блока в вычислении на GPU, имеет тип uint3.
• threadIdx – индекс текущей нити в вычислении на GPU, имеет тип uint3.
• warpSize – размер warp’а, имеет тип int (сам еще не пробовал использовать).

Дополнительные типы переменных и их спецификаторы будут рассмотрены непосредственно в примерах работы с памятью.

CUDA host API:

В CUDA runtime API входят следующие группы функций:

• Device Management – включает функции для общего управления GPU (получение инфор-мации о возможностях GPU, переключение между GPU при работе SLI-режиме и т.д.).
• Thread Management – управление нитями.
• Stream Management – управление потоками.
• Event Management – функция создания и управления event’ами.
• Execution Control – функции запуска и исполнения ядра CUDA.
• Memory Management – функции управлению памятью GPU.
• Texture Reference Manager – работа с объектами текстур через CUDA.
• OpenGL Interoperability – функции по взаимодействию с OpenGL API.
• Direct3D 9 Interoperability – функции по взаимодействию с Direct3D 9 API.
• Direct3D 10 Interoperability – функции по взаимодействию с Direct3D 10 API.
• Error Handling – функции обработки ошибок.

Понимаем работу GPU:

Задача. Требуется вычислить сумму двух векторов размерностью N элементов.

Нам известна максимальные размеры нашего блока: 512*512*64 нитей. Так как вектор у нас одномерный, то пока ограничимся использованием x-измерения нашего блока, то есть задействуем только одну полосу нитей из блока (рис. 3).

Заметим, что x-размерность блока 512, то есть, мы можем сложить за один раз векторы, длина которых N <= 512 элементов. В прочем, при более массивных вычислениях, можно использовать большее число блоков и многомерные массивы. Так же я заметил одну интересную особенность, возможно, некоторые из вас подумали, что в одном блоке можно задействовать 512*512*64 = 16777216 нитей, естественно это не так, в целом, это произведение не может превышать 512 (по крайней мере, на моей видеокарте).

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

1. Получить данные для расчетов.
2. Скопировать эти данные в GPU память.
3. Произвести вычисление в GPU через функцию ядра.
4. Скопировать вычисленные данные из GPU памяти в ОЗУ.
5. Посмотреть результаты.
6. Высвободить используемые ресурсы.

Первым делом напишем функцию ядра, которая и будет осуществлять сложение векторов:

// Функция сложения двух векторов
__global__ void addVector(float * left, float * right, float * result)
{
//Получаем id текущей нити.
int idx = threadIdx.x;

//Расчитываем результат.
result = left + right;
}

Таким образом, распараллеливание будет выполнено автоматически при запуске ядра. В этой функции так же используется встроенная переменная threadIdx и её поле x, которая позволяет задать соответствие между расчетом элемента вектора и нитью в блоке. Делаем расчет каждого элемента вектора в отдельной нити.

Пишем код, которые отвечает за 1 и 2 пункт в программе:

#define SIZE 512
__host__ int main()
{
//Выделяем память под вектора
float * vec1 = new float ;
float * vec2 = new float ;
float * vec3 = new float ;

//Инициализируем значения векторов
for (int i = 0; i < SIZE; i++)
{
vec1[i] = i;
vec2[i] = i;
}

//Указатели на память видеокарте
float * devVec1;
float * devVec2;
float * devVec3;

//Выделяем память для векторов на видеокарте
cudaMalloc((void **)&devVec1, sizeof (float ) * SIZE);
cudaMalloc((void **)&devVec2, sizeof (float ) * SIZE);
cudaMalloc((void **)&devVec3, sizeof (float ) * SIZE);

//Копируем данные в память видеокарты
cudaMemcpy(devVec1, vec1, sizeof (float ) * SIZE, cudaMemcpyHostToDevice);
cudaMemcpy(devVec2, vec2, sizeof (float ) * SIZE, cudaMemcpyHostToDevice);
…
}

* This source code was highlighted with Source Code Highlighter .

Для выделения памяти на видеокарте используется функция cudaMalloc , которая имеет следующий прототип:
cudaError_t cudaMalloc(void** devPtr, size_t count), где

1. devPtr – указатель, в который записывается адрес выделенной памяти,
2. count – размер выделяемой памяти в байтах.

1. cudaSuccess – при удачном выделении памяти
2. cudaErrorMemoryAllocation – при ошибке выделения памяти

1. dst – указатель, содержащий адрес места-назначения копирования,
2. src – указатель, содержащий адрес источника копирования,
3. count – размер копируемого ресурса в байтах,
4. cudaMemcpyKind – перечисление, указывающее направление копирования (может быть cudaMemcpyHostToDevice, cudaMemcpyDeviceToHost, cudaMemcpyHostToHost, cudaMemcpyDeviceToDevice).

1. cudaSuccess – при удачном копировании
2. cudaErrorInvalidValue – неверные параметры аргумента (например, размер копирования отрицателен)
3. cudaErrorInvalidDevicePointer – неверный указатель памяти в видеокарте
4. cudaErrorInvalidMemcpyDirection – неверное направление (например, перепутан источник и место-назначение копирования)

…
dim3 gridSize = dim3(1, 1, 1); //Размер используемого грида
dim3 blockSize = dim3(SIZE, 1, 1); //Размер используемого блока

addVector<<>>(devVec1, devVec2, devVec3);
…

* This source code was highlighted with Source Code Highlighter .

addVector<<<1, SIZE>>>(devVec1, devVec2, devVec3);

* This source code was highlighted with Source Code Highlighter .

Код после вызова ядра:

//Выполняем вызов функции ядра
addVector<<>>(devVec1, devVec2, devVec3);

//Хендл event"а
cudaEvent_t syncEvent;

CudaEventCreate(&syncEvent); //Создаем event
cudaEventRecord(syncEvent, 0); //Записываем event
cudaEventSynchronize(syncEvent); //Синхронизируем event

//Только теперь получаем результат расчета
cudaMemcpy(vec3, devVec3, sizeof (float ) * SIZE, cudaMemcpyDeviceToHost);

* This source code was highlighted with Source Code Highlighter .

Event создается с помощью функции cudaEventCreate , прототип которой имеет вид:
cudaError_t cudaEventCreate(cudaEvent_t* event), где

1. *event – указатель для записи хендла event’а.

1. cudaSuccess – в случае успеха
2. cudaErrorMemoryAllocation – ошибка выделения памяти

1. event – хендл хаписываемого event’а,
2. stream – номер потока, в котором записываем (в нашем случае это основной нулевой по-ток).

1. cudaSuccess – в случае успеха
2. cudaErrorInvalidValue – неверное значение
3. cudaErrorInitializationError – ошибка инициализации
4. cudaErrorPriorLaunchFailure – ошибка при предыдущем асинхронном запуске функции

1. event – хендл event’а, прохождение которого ожидается.

1. cudaSuccess – в случае успеха
2. cudaErrorInitializationError – ошибка инициализации
3. cudaErrorPriorLaunchFailure – ошибка при предыдущем асинхронном запуске функции
4. cudaErrorInvalidValue – неверное значение
5. cudaErrorInvalidResourceHandle – неверный хендл event’а

На рисунке 4 блок «Ожидание прохождения Event’а» и есть вызов функции cudaEventSynchronize.

Ну и в заключении выводим результат на экран и чистим выделенные ресурсы.

//Результаты расчета
for (int i = 0; i < SIZE; i++)
{
printf("Element #%i: %.1f\n" , i , vec3[i]);
}

//
// Высвобождаем ресурсы
//

CudaEventDestroy(syncEvent);

CudaFree(devVec1);
cudaFree(devVec2);
cudaFree(devVec3);

Delete vec1; vec1 = 0;
delete vec2; vec2 = 0;
delete vec3; vec3 = 0;

* This source code was highlighted with Source Code Highlighter .

Заключение

P.S.: Получилось не очень кратко. Надеюсь, что не утомил. Если нужен весь исходный код, то могу выслать на почту.
P.S.S: Задавайте вопросы.

Теги: Добавить метки

– набор низкоуровневых программных интерфейсов (API ) для создания игр и других высокопроизводительных мультимедиа-приложений. Включает поддержку высокопроизводительной 2D - и 3D -графики, звука и устройств ввода.

Direct3D (D3D ) – интерфейс вывода трёхмерных примитивов (геометрических тел). Входит в .

OpenGL (от англ. Open Graphics Library , дословно – открытая графическая библиотека) – спецификация, определяющая независимый от языка программирования кросс-платформенный программный интерфейс для написания приложений, использующих двухмерную и трёхмерную компьютерную графику. Включает более 250 функций для рисования сложных трёхмерных сцен из простых примитивов. Используется при создании видеоигр, виртуальной реальности, визуализации в научных исследованиях. На платформе Windows конкурирует с .

OpenCL (от англ. Open Computing Language , дословно – открытый язык вычислений) – фреймворк (каркас программной системы) для написания компьютерных программ, связанных с параллельными вычислениями на различных графических (GPU ) и ( ). В фреймворк OpenCL входят язык программирования и интерфейс программирования приложений (API ). OpenCL обеспечивает параллелизм на уровне инструкций и на уровне данных и является реализацией техники GPGPU .

GPGPU (сокр. от англ. General-P urpose G raphics P rocessing U nits , дословно – GPU общего назначения) – техника использования графического процессоравидеокарты для общих вычислений, которые обычно проводит .

Шейдер (англ. shader ) – программа построения теней на синтезируемых изображениях, используется в трёхмерной графике для определенияокончательных параметров объекта или изображения. Как правило, включает произвольной сложности описание поглощения и рассеяния света, наложениятекстуры, отражения и преломления, затенения, смещения поверхности и эффекты пост-обработки. Сложные поверхности могут быть визуализированы припомощи простых геометрических форм.

Рендеринг (англ. rendering ) – визуализация, в компьютерной графике процесс получения изображения по модели с помощью программного .

SDK (сокр. от англ. Software Development Kit ) – набор инструментальных средств разработки программного .

CPU (сокр. от англ. Central Processing Unit , дословно – центральное/основное/главное вычислительное устройство) – центральный (микро) ;устройство, исполняющее машинные инструкции; часть аппаратного обеспечения , отвечающая за выполнение вычислительных операций (заданныхоперационной системой и прикладным программным ) и координирующая работу всех устройств .

GPU (сокр. от англ. Graphic Processing Unit , дословно – графическое вычислительное устройство) – графический процессор; отдельное устройство илиигровой приставки, выполняющее графический рендеринг (визуализацию). Современные графические процессоры очень эффективно обрабатывают иреалистично отображают компьютерную графику. Графический процессор в современных видеоадаптерах применяется в качестве ускорителя трёхмернойграфики, однако его можно использовать в некоторых случаях и для вычислений (GPGPU ).

Проблемы CPU

Долгое время повышение производительности традиционных в основном происходило за счёт последовательного увеличения тактовой частоты (около 80% производительности определяла именно тактовая частота) с одновременным увеличением количества транзисторов на одном кристалле. Однако дальнейшее повышение тактовой частоты (при тактовой частоте более 3,8 ГГц чипы попросту перегреваются!) упирается в ряд фундаментальных физических барьеров (поскольку технологический процесс почти вплотную приблизился к размерам атома: , а размеры атома кремния – приблизительно 0,543 нм):

Во-первых, с уменьшением размеров кристалла и с повышением тактовой частоты возрастает ток утечки транзисторов. Это ведёт к повышению потребляемой мощности и увеличению выброса тепла;

Во-вторых, преимущества более высокой тактовой частоты частично сводятся на нет из-за задержек при обращении к памяти, так как время доступа к памяти не соответствует возрастающим тактовым частотам;

В-третьих, для некоторых приложений традиционные последовательные архитектуры становятся неэффективными с возрастанием тактовой частоты из-за так называемого «фон-неймановского узкого места» – ограничения производительности в результате последовательного потока вычислений. При этом возрастают резистивно-ёмкостные задержки передачи сигналов, что является дополнительным узким местом, связанным с повышением тактовой частоты.

Развитие GPU

Параллельно с шло (и идет!) развитие GPU :

…

Ноябрь 2008 г. – Intel представила линейку 4-ядерных Intel Core i7 , в основу которых положена микроархитектура нового поколения Nehalem . Процессоры работают на тактовой частоте 2,6-3,2 ГГц. Выполнены по 45-нм техпроцессу.

Декабрь 2008 г. – начались поставки 4-ядерного AMD Phenom II 940 (кодовое название – Deneb ). Работает на частоте 3 ГГц, выпускается по техпроцессу 45-нм.

…

Май 2009 г. – компания AMD представила версию графического процессора ATI Radeon HD 4890 с тактовой частотой ядра, увеличенной с 850 МГц до 1 ГГц. Это первый графический процессор, работающий на частоте 1 ГГц. Вычислительная мощность чипа, благодаря увеличению частоты, выросла с 1,36 до 1,6 терафлоп. Процессор содержит 800 (!) вычислительных ядер, поддерживает видеопамять GDDR5 , DirectX 10.1 , ATI CrossFireX и все другие технологии, присущие современным моделям видеокарт. Чип изготовлен на базе 55-нм технологии.

Основные отличия GPU

Отличительными особенностями GPU (по сравнению с ) являются:

– архитектура, максимально нацеленная на увеличение скорости расчёта текстур и сложных графических объектов;

– пиковая мощность типичного GPU намного выше, чем у ;

– благодаря специализированной конвейерной архитектуре, GPU намного эффективнее в обработке графической информации, чем .

«Кризис жанра»

«Кризис жанра» для назрел к 2005 г., – именно тогда появились . Но, несмотря на развитие технологии , рост производительности обычных заметно снизился. В то же время производительность GPU продолжает расти. Так, к 2003 г. и кристаллизовалась эта революционная идея – использовать для нужд вычислительную мощь графического . Графические процессоры стали активно использоваться для «неграфических» вычислений (симуляция физики, обработка сигналов, вычислительная математика/геометрия, операции с базами данных, вычислительная биология, вычислительная экономика, компьютерное зрение и т.д.).

Главная проблема заключалась в том, что не было никакого стандартного интерфейса для программирования GPU . Разработчики использовали OpenGL или Direct3D , но это было очень удобно. Корпорация NVIDIA (один из крупнейших производителей графических, медиа- и коммуникационных процессоров, а также беспроводных медиа-процессоров; основана в 1993 г.) занялась разработкой некоего единого и удобного стандарта, – и представила технологию CUDA .

Как это начиналось

2006 г. – NVIDIA демонстрирует CUDA™ ; начало революции в вычислениях на GPU .

2007 г. – NVIDIA выпускает архитектуру CUDA (первоначальная версия CUDA SDK была представлена 15 февраля 2007 г.); номинация «Лучшая новинка» от журнала Popular Science и «Выбор читателей» от издания HPCWire .

2008 г. – технология NVIDIA CUDA победила в номинации «Техническое превосходство» от PC Magazine .

Что такое CUDA

CUDA (сокр. от англ. Compute Unified Device Architecture , дословно – унифицированная вычислительная архитектура устройств) – архитектура (совокупность программных и аппаратных средств), позволяющая производить на GPU вычисления общего назначения, при этом GPU фактически выступает в роли мощного сопроцессора.

Технология NVIDIA CUDA™ – это единственная среда разработки на языке программирования C , которая позволяет разработчикам создавать программное для решения сложных вычислительных задач за меньшее время, благодаря вычислительной мощности графических процессоров. В мире уже работают миллионы GPU с поддержкой CUDA , и тысячи программистов уже пользуются (бесплатно!) инструментами CUDA для ускорения приложений и для решения самых сложных ресурсоёмких задач – от кодирования видео- и аудио- до поисков нефти и газа, моделирования продуктов, вывода медицинских изображений и научных исследований.

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

Уральский, ведущий специалист по технологиям NVIDIA , сравнивая GPU и , говорит так : « – это внедорожник. Он ездит всегда и везде, но не очень быстро. А GPU – это спорткар. На плохой дороге он просто никуда не поедет, но дайте хорошее покрытие, – и он покажет всю свою скорость, которая внедорожнику и не снилась!..».

Возможности технологии CUDA