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

Однако проходит немного времени, и оказывается, что новичок бегает быстрее, прыгает выше и вообще сильнее. И мух он лопает больше его соседей-ретроградов. И вот тогда эти самые соседи начинают понимать, что ссориться с этим бывшим неуклюжим не стоит. Лучше с ним дружить, а еще лучше организовать симбиоз. Глядишь, и мух перепадет побольше.

Технология GPGPU (General-Purpose Graphics Processing Units — графический процессор общего назначения) долгое время существовала только в теоретических выкладках мозговитых академиков. А как иначе? Предложить кардинально изменить сложившийся за десятилетия вычислительный процесс, доверив расчет его параллельных веток видеокарте, — на это только теоретики и способны.

Логотип технологии CUDA напоминает о том, что выросла она в недрах
3D-графики.

Но долго пылиться на страницах университетских журналов технология GPGPU не собиралась. Распушив перья своих лучших качеств, она привлекла к себе внимание производителей. Так на свет появилась CUDA — реализация GPGPU на графических процессорах GeForce производства компании nVidia.

Благодаря CUDA технологии GPGPU стали мейнстримом. И ныне только самый недальновидный и покрытый толстым слоем лени разработчик систем программирования не заявляет о поддержке своим продуктом CUDA. IT-издания почли за честь изложить подробности технологии в многочисленных пухлых научно-популярных статьях, а конкуренты срочно уселись за лекала и кросскомпиляторы, чтобы разработать нечто подобное.

Публичное признание — это мечта не только начинающих старлеток, но и вновь зародившихся технологий. И CUDA повезло. Она на слуху, о ней говорят и пишут.

Вот только пишут так, словно продолжают обсуждать GPGPU в толстых научных журналах. Забрасывают читателя грудой терминов типа «grid», «SIMD», «warp», «хост», «текстурная и константная память». Погружают его по самую маковку в схемы организации графических процессоров nVidia, ведут извилистыми тропами параллельных алгоритмов и (самый сильный ход) показывают длинные листинги кода на языке Си. В результате получается, что на входе статьи мы имеем свежего и горящего желанием понять CUDA читателя, а на выходе — того же читателя, но с распухшей головой, заполненной кашей из фактов, схем, кода, алгоритмов и терминов.

А между тем цель любой технологии — сделать нашу жизнь проще. И CUDA прекрасно с этим справляется. Результаты ее работы — именно это убедит любого скептика лучше сотни схем и алгоритмов.

Далеко не везде

CUDA поддерживается высокопроизводительными суперкомпьютерами
nVidia Tesla.

И все же прежде, чем взглянуть на результаты трудов CUDA на поприще облегчения жизни рядового пользователя, стоит уяснить все ее ограничения. Точно как с джинном: любое желание, но одно. У CUDA тоже есть свои ахиллесовы пятки. Одна из них — ограничения платформ, на которых она может трудиться.

Перечень видеокарт производства nVidia, поддерживающих CUDA, представлен в специальном списке, именуемом CUDA Enabled Products. Список весьма внушительный, но легко классифицируемый. В поддержке CUDA не отказывают:

Модели nVidia GeForce 8-й, 9-й, 100-й, 200-й и 400-й серий с минимумом 256 мегабайт видеопамяти на борту. Поддержка распространяется как на карты для настольных систем, так и на мобильные решения.

Подавляющее большинство настольных и мобильных видеокарт nVidia Quadro.

Все решения нетбучного ряда nvidia ION.

Высокопроизводительные HPC (High Performance Computing) и суперкомпьютерные решения nVidia Tesla, используемые как для персональных вычислений, так и для организации масштабируемых кластерных систем.

Поэтому, прежде чем применять программные продукты на базе CUDA, стоит свериться с этим списком избранных.

Кроме самой видеокарты, для поддержки CUDA требуется соответствующий драйвер. Именно он является связующим звеном между центральным и графическим процессором, выполняя роль своеобразного программного интерфейса для доступа кода и данных программы к многоядерной сокровищнице GPU. Чтобы наверняка не ошибиться, nVidia рекомендует посетить страничку драйверов и получить наиболее свежую версию.

...но сам процесс

Как работает CUDA? Как объяснить сложный процесс параллельных вычислений на особой аппаратной архитектуре GPU так, чтобы не погрузить читателя в пучину специфических терминов?

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

Архитектурно центральный процессор (CPU) и его графический собрат (GPU) устроены по-разному. Если проводить аналогию с миром автопрома, то CPU — универсал, из тех, которые называют «сарай». Выглядит легковым авто, но при этом (с точки зрения разработчиков) «и швец, и жнец, и на дуде игрец». Выполняет роль маленького грузовика, автобуса и гипертрофированного хечбэка одновременно. Универсал, короче. Цилиндров-ядер у него немного, но они «тянут» практически любые задачи, а внушительная кэш-память способна разместить кучу данных.

А вот GPU — это спорткар. Функция одна: доставить пилота на финиш как можно быстрее. Поэтому никакой большой памяти-багажника, никаких лишних посадочных мест. Зато цилиндров-ядер в сотни раз больше, чем у CPU.

Благодаря CUDA разработчикам программ GPGPU не требуется вникать в сложности программи-
рования под такие графические движки, как DirectX и OpenGL

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

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

Технологически для программиста ничего не меняется. Код CUDA-программ пишется на языке Си. Точнее, на особом его диалекте «С with streams» (Си с потоками). Разработанное в Стэнфорде, это расширение языка Си получило название Brook. В качестве интерфейса, передающего Brook-код на GPU, выступает драйвер видеокарты, поддерживающей CUDA. Он организует весь процесс обработки этого участка программы так, что для программиста GPU выглядит как сопроцессор CPU. Очень похоже на использование математического сопроцессора на заре персональных компьютеров. С появлением Brook, видеокарт с поддержкой CUDA и драйверов для них любой программист стал способен в своих программах обращаться к GPU. А ведь раньше этим шаманством владел узкий круг избранных, годами оттачивающих технику программирования под графические движки DirectX или OpenGL.

В бочку этого пафосного меда — дифирамбов CUDA — стоит положить ложку дегтя, то бишь ограничений. Далеко не любая задача, которую нужно запрограммировать, подходит для решения с помощью CUDA. Добиться ускорения решения рутинных офисных задач не получится, а вот доверить CUDA обсчет поведения тысячи однотипных бойцов в World of Warcraft — пожалуйста. Но это задача, высосанная из пальца. Рассмотрим же примеры того, что CUDA уже очень эффективно решает.

Труды праведные

CUDA — весьма прагматичная технология. Реализовав ее поддержку в своих видеокартах, компания nVidia весьма справедливо рассчитывала на то, что знамя CUDA будет подхвачено множеством энтузиастов как в университетской среде, так и в коммерции. Так и случилось. Проекты на базе CUDA живут и приносят пользу.

NVIDIA PhysX

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

Одна из возможностей физического движка NVIDIA PhysX — реалистичная работа с тканями.

Реализовать алгоритмы компьютерной симуляции базовых физических законов — дело очень трудоемкое. Наиболее известными компаниями на этом поприще являются ирландская компания Havok с ее межплатформенным физическим Havok Physics и калифорнийская Ageia — прародитель первого в мире физического процессора (PPU — Physics Processing Unit) и соответствующего физического движка PhysX. Первая из них, хотя и приобретена компанией Intel, активно трудится сейчас на поприще оптимизации движка Havok для видеокарт ATI и процессоров AMD. А вот Ageia с ее движком PhysX стала частью nVidia. При этом nVidia решила достаточно сложную задачу адаптации PhysX под технологию CUDA.

Возможным это стало благодаря статистике. Статистически было доказано, что, какой бы сложный рендеринг ни выполнял GPU, часть его ядер все равно простаивает. Именно на этих ядрах и работает движок PhysX.

Благодаря CUDA львиная доля вычислений, связанных с физикой игрового мира, стала выполняться на видеокарте. Освободившаяся мощь центрального процессора была брошена на решение других задач геймплея. Результат не заставил себя ждать. По оценкам экспертов, прирост производительности игрового процесса с PhysX, работающем, на CUDA возрос минимум на порядок. Выросло и правдоподобие реализации физических законов. CUDA берет на себя рутинный расчет реализации трения, тяготения и прочих привычных нам вещей для многомерных объектов. Теперь не только герои и их техника идеально вписываются в законы привычного нам физического мира, но и пыль, туман, взрывная волна, пламя и вода.

CUDA-версия пакета сжатия текстур NVIDIA Texture Tools 2

Нравятся реалистичные объекты в современных играх? Стоит сказать спасибо разработчикам текстур. Но чем больше реальности в текстуре, тем больше ее объем. Тем больше она занимает драгоценной памяти. Чтобы этого избежать, текстуры предварительно сжимают и динамически распаковывают по мере надобности. А сжатие и распаковка — это сплошные вычисления. Для работы с текстурами nVidia выпустила пакет NVIDIA Texture Tools. Он поддерживает эффективное сжатие и распаковку текстур стандарта DirectX (так называемый ВЧЕ-формат). Вторая версия этого пакета может похвастаться поддержкой алгоритмов сжатия BC4 и BC5, реализованных в технологии DirectX 11. Но главное то, что в NVIDIA Texture Tools 2 реализована поддержка CUDA. По оценке nVidia, это дает 12-кратный прирост производительности в задачах сжатия и распаковки текстур. А это значит, что фреймы игрового процесса будут грузиться быстрее и радовать игрока своей реалистичностью.

Пакет NVIDIA Texture Tools 2 заточен под работу с CUDA. Прирост производительности при сжатии и распаковке текстур налицо.

Использование CUDA позволяет существенно повысить эффективность видеослежки.

Обработка видеопотока в реальном времени

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

Полноводные реки видеоинформации стекаются в центры ее обработки и... наталкиваются на узкое звено — человека. Именно он в большинстве случаев — последняя инстанция, следящая за видеомиром. Причем инстанция не самая эффективная. Моргает, отвлекается и норовит уснуть.

Благодаря CUDA появилась возможность реализации алгоритмов одновременного слежения за множеством объектов в видеопотоке. При этом процесс происходит в реальном масштабе времени, а видео является полноценным 30 fps. По сравнению с реализацией такого алгоритма на современных многоядерных CPU CUDA дает двух-, трехкратный прирост производительности, а это, согласитесь, немало.

Конвертирование видео, фильтрация аудио

Видеоконвертер Badaboom — первая ласточка, использующая CUDA для ускорения конвертирования.

Приятно посмотреть новинку видеопроката в FullHD-качестве и на большом экране. Но большой экран не возьмешь с собой в дорогу, а видеокодек FullHD будет икать на маломощном процессоре мобильного гаджета. На помощь приходит конвертирование. Но большинство тех, кто с ним сталкивался на практике, сетуют на длительное время конвертации. Оно и понятно, процесс рутинный, пригодный к распараллеливанию, и его выполнение на CPU не очень оптимально.

А вот CUDA с ним справляется на ура. Первая ласточка — конвертер Badaboom от компании Elevental. Разработчики Badaboom, выбрав CUDA, не просчитались. Тесты показывают, что стандартный полуторачасовый фильм на нем конвертируется в формат iPhone/iPod Touch менее чем за двадцать минут. И это при том, что при использовании только CPU этот процесс занимает больше часа.

Помогает CUDA и профессиональным меломанам. Любой из них полцарства отдаст за эффективный FIR-кроссовер — набор фильтров, разделяющих звуковой спектр на несколько полос. Процесс этот весьма трудоемкий и при большом объеме аудиоматериала заставляет звукорежиссера сходить на несколько часов «покурить». Реализация FIR-кроссовера на базе CUDA ускоряет его работу в сотни раз.

CUDA Future

Сделав технологию GPGPU реальностью, CUDA не собирается почивать на лаврах. Как это происходит повсеместно, в CUDA работает принцип рефлексии: теперь не только архитектура видеопроцессоров nVidia влияет на развитие версий CUDA SDK, а и сама технология CUDA заставляет nVidia пересматривать архитектуру своих чипов. Пример такой рефлексии — платформа nVidia ION. Ее вторая версия специально оптимизирована для решения CUDA-задач. А это означает, что даже в относительно недорогих аппаратных решениях потребители получат всю мощь и блестящие возможности CUDA.

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

Приступим.

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

Рассмотрим вычислительную модель GPU более подробно.

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

CUDA и язык C:

Сама технология CUDA (компилятор nvcc.exe) вводит ряд дополнительных расширений для языка C, которые необходимы для написания кода для GPU:

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

Как было сказано, спецификаторы функций определяют, как и откуда буду вызываться функции. Всего в CUDA 3 таких спецификатора:

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

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

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

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

Ну и конечно сама myKernelFunc – функция ядра (спецификатор __global__). Некоторые переменные при вызове ядра можно опускать, например sharedMemSize и cudaStream.

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

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

Кстати, gridDim и blockDim и есть те самые переменные, которые мы передаем при запуске ядра GPU, правда, в ядре они могут быть read only.

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

CUDA host API:

Перед тем, как приступить к непосредственному использованию CUDA для вычислений, необходимо ознакомиться с так называемым CUDA host API, который является связующим звеном между CPU и GPU. CUDA host API в свою очередь можно разделить на низкоуровневое API под названием CUDA driver API, который предоставляет доступ к драйверу пользовательского режима CUDA, и высокоуровневое API – CUDA runtime API. В своих примерах я буду использовать CUDA runtime 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 – при ошибке выделения памяти

Для копирования данных в память видеокарты используется cudaMemcpy, которая имеет следующий прототип:
cudaError_t cudaMemcpy(void* dst, const void* src ,size_t count, enum cudaMemcpyKind kind), где

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

Возвращает:

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

Теперь переходим к непосредственному вызову ядра для вычисления на GPU.

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

Теперь нам остаеться скопировать результат расчета из видеопамяти в память хоста. Но у функций ядра при этом есть особенность – асинхронное исполнение, то есть, если после вызова ядра начал работать следующий участок кода, то это ещё не значит, что GPU выполнил расчеты. Для завершения работы заданной функции ядра необходимо использовать средства синхронизации, например event’ы. Поэтому, перед копированием результатов на хост выполняем синхронизацию нитей GPU через event.

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

//Выполняем вызов функции ядра
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 Managment API.

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

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

Возвращает:

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

Запись event’а осуществляется с помощью функции cudaEventRecord , прототип которой имеет вид:
cudaError_t cudaEventRecord(cudaEvent_t event, CUstream stream), где

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

Возвращает:

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

Синхронизация event’а выполняется функцией cudaEventSynchronize. Данная функция ожидает окончание работы всех нитей GPU и прохождение заданного event’а и только потом отдает управление вызывающей программе. Прототип функции имеет вид:
cudaError_t cudaEventSynchronize(cudaEvent_t event), где

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

Возвращает:

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

Понять, как работает cudaEventSynchronize, можно из следующей схемы:

На рисунке 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 .

Думаю, что описывать функции высвобождения ресурсов нет необходимости. Разве что, можно напомнить, что они так же возвращают значения cudaError_t, если есть необходимость проверки их работы.

Заключение

Надеюсь, что этот материал поможет вам понять, как функционирует GPU. Я описал самые главные моменты, которые необходимо знать для работы с CUDA. Попробуйте сами написать сложение двух матриц, но не забывайте об аппаратных ограничениях видеокарты.

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

Теги:

• CUDA
• gpgpu
• nvidia

Добавить метки

Ядра 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 .

Позвольте обратиться к истории - вернуться в 2003 год, когда Intel и AMD участвовали в совместной гонке за самый мощный процессор. Всего за несколько лет в результате этой гонки тактовые частоты существенно выросли, особенно после выхода Intel Pentium 4.

Но гонка быстро приближалась к пределу. После волны огромного прироста тактовых частот (между 2001 и 2003 годами тактовая частота Pentium 4 удвоилась с 1,5 до 3 ГГц), пользователям пришлось довольствоваться десятыми долями гигагерц, которые смогли выжать производители (с 2003 до 2005 тактовые частоты увеличились всего с 3 до 3,8 ГГц).

Даже архитектуры, оптимизированные под высокие тактовые частоты, та же Prescott, стали испытывать трудности, причём на этот раз не только производственные. Производители чипов просто упёрлись в законы физики. Некоторые аналитики даже предрекали, что закон Мура перестанет действовать. Но этого не произошло. Оригинальный смысл закона часто искажают, однако он касается числа транзисторов на поверхности кремниевого ядра. Долгое время повышение числа транзисторов в CPU сопровождалось соответствующим ростом производительности - что и привело к искажению смысла. Но затем ситуация усложнилась. Разработчики архитектуры CPU подошли к закону сокращения прироста: число транзисторов, которое требовалось добавить для нужного увеличения производительности, становилось всё большим, заводя в тупик.

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

Почему же они не зашли в тот же тупик, как разработчики архитектуры CPU? Причина очень простая: центральные процессоры разрабатываются для получения максимальной производительности на потоке инструкций, которые обрабатывают разные данные (как целые числа, так и числа с плавающей запятой), производят случайный доступ к памяти и т.д. До сих пор разработчики пытаются обеспечить больший параллелизм инструкций - то есть выполнять как можно большее число инструкций параллельно. Так, например, с Pentium появилось суперскалярное выполнение, когда при некоторых условиях можно было выполнять две инструкции за такт. Pentium Pro получил внеочередное выполнение инструкций, позволившее оптимизировать работу вычислительных блоков. Проблема заключается в том, что у параллельного выполнения последовательного потока инструкций есть очевидные ограничения, поэтому слепое повышение числа вычислительных блоков не даёт выигрыша, поскольку большую часть времени они всё равно будут простаивать.

Напротив, работа GPU относительно простая. Она заключается в принятии группы полигонов с одной стороны и генерации группы пикселей с другой. Полигоны и пиксели независимы друг от друга, поэтому их можно обрабатывать параллельно. Таким образом, в GPU можно выделить крупную часть кристалла на вычислительные блоки, которые, в отличие от CPU, будут реально использоваться.

Нажмите на картинку для увеличения.

GPU отличается от CPU не только этим. Доступ к памяти в GPU очень связанный - если считывается тексель, то через несколько тактов будет считываться соседний тексель; когда записывается пиксель, то через несколько тактов будет записываться соседний. Разумно организуя память, можно получить производительность, близкую к теоретической пропускной способности. Это означает, что GPU, в отличие от CPU, не требуется огромного кэша, поскольку его роль заключается в ускорении операций текстурирования. Всё, что нужно, это несколько килобайт, содержащих несколько текселей, используемых в билинейных и трилинейных фильтрах.

Нажмите на картинку для увеличения.

Да здравствует GeForce FX!

Два мира долгое время оставались разделёнными. Мы использовали CPU (или даже несколько CPU) для офисных задач и интернет-приложений, а GPU хорошо подходили лишь для ускорения визуализации. Но одна особенность изменила всё: а именно, появление программируемых GPU. Поначалу центральным процессорам было нечего бояться. Первые так называемые программируемые GPU (NV20 и R200) вряд ли представляли угрозу. Число инструкций в программе оставалось ограниченным около 10, они работали над весьма экзотическими типами данных, такими как 9- или 12-битными числами с фиксированной запятой.

Нажмите на картинку для увеличения.

Но закон Мура вновь показал себя с лучшей стороны. Увеличение числа транзисторов не только позволило повысить количество вычислительных блоков, но и улучшило их гибкость. Появление NV30 можно считать существенным шагом вперёд по нескольким причинам. Конечно, геймерам карты NV30 не очень понравились, однако новые графические процессоры стали опираться на две особенности, которые были призваны изменить восприятие GPU уже не только как графических акселераторов.

• Поддержка вычислений с плавающей запятой одинарной точности (пусть даже это и не соответствовало стандарту IEEE754);
• поддержка числа инструкций больше тысячи.

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

Идея использования графических акселераторов для математических расчётов не нова. Первые попытки были сделаны ещё в 90-х годах прошлого века. Конечно, они были очень примитивными - ограничиваясь, по большей части, использованием некоторых аппаратно заложенных функций, например, растеризации и Z-буферов для ускорения таких задач, как поиск маршрута или вывод диаграмм Вороного .

Нажмите на картинку для увеличения.

В 2003 году, с появлением эволюционировавших шейдеров, была достигнута новая планка - на этот раз выполнение матричных вычислений. Это был год, когда целая секция SIGGRAPH ("Computations on GPUs/Вычисления на GPU") была выделена под новую область ИТ. Эта ранняя инициатива получила название GPGPU (General-Purpose computation on GPU, универсальные вычисления на GPU). И ранним поворотным моментом стало появление .

Чтобы понять роль BrookGPU, нужно разобраться, как всё происходило до его появления. Единственным способом получить ресурсы GPU в 2003 году было использование одного из двух графических API - Direct3D или OpenGL. Следовательно, разработчикам, которые хотели получить возможности GPU для своих вычислений, приходилось опираться на два упомянутых API. Проблема в том, что они не всегда являлись экспертами в программировании видеокарт, а это серьёзно осложняло доступ к технологиям. Если 3D-программисты оперируют шейдерами, текстурами и фрагментами, то специалисты в области параллельного программирования опираются на потоки, ядра, разбросы и т.д. Поэтому сначала нужно было привести аналогии между двумя мирами.

• Поток (stream) представляет собой поток элементов одного типа, в GPU он может быть представлен текстурой. В принципе, в классическом программировании есть такой аналог, как массив.
• Ядро (kernel) - функция, которая будет применяться независимо к каждому элементу потока; является эквивалентом пиксельного шейдера. В классическом программировании можно привести аналогию цикла - он применяется к большому числу элементов.
• Чтобы считывать результаты применения ядра к потоку, должна быть создана текстура. На CPU эквивалента нет, поскольку там есть полный доступ к памяти.
• Управление местоположением в памяти, куда будет производиться запись (в операциях разброса/scatter), осуществляется через вершинный шейдер, поскольку пиксельный шейдер не может изменять координаты обрабатываемого пикселя.

Как можно видеть, даже с учётом приведённых аналогий, задача не выглядит простой. И на помощь пришёл Brook. Под этим названием подразумеваются расширения к языку C ("C with streams", "C с потоками"), как назвали их разработчики в Стэнфорде. По своей сути, задача Brook сводилась к сокрытию от программиста всех составляющих 3D API, что позволяло представить GPU как сопроцессор для параллельных вычислений. Для этого компилятор Brook обрабатывал файл.br с кодом C++ и расширениями, после чего генерировал код C++, который привязывался к библиотеке с поддержкой разных выходов (DirectX, OpenGL ARB, OpenGL NV3x, x86).

Нажмите на картинку для увеличения.

У Brook есть несколько заслуг, первая из которых заключается в выводе GPGPU из тени, чтобы с этой технологией могли знакомиться и широкие массы. Хотя после объявления о проекте ряд ИТ-сайтов слишком оптимистично сообщил о том, что выход Brook ставит под сомнение существование CPU, которые вскоре будут заменены более мощными GPU. Но, как видим, и через пять лет этого не произошло. Честно говоря, мы не думаем, что это вообще когда-либо случится. С другой стороны, глядя на успешную эволюцию CPU, которые всё более ориентируются в сторону параллелизма (больше ядер, технология многопоточности SMT, расширение блоков SIMD), а также и на GPU, которые, напротив, становятся всё более универсальными (поддержка расчётов с плавающей запятой одинарной точности, целочисленные вычисления, поддержка расчётов с двойной точностью), похоже, что GPU и CPU вскоре попросту сольются. Что же тогда произойдёт? Будут ли GPU поглощены CPU, как в своё время произошло с математическими сопроцессорами? Вполне возможно. Intel и AMD сегодня работают над подобными проектами. Но ещё очень многое может измениться.

Но вернёмся к нашей теме. Преимущество Brook заключалось в популяризации концепции GPGPU, он существенно упростил доступ к ресурсам GPU, что позволило всё большим пользователям осваивать новую модель программирования. С другой стороны, несмотря на все качества Brook, предстоял ещё долгий путь, прежде чем ресурсы GPU можно будет использовать для вычислений.

Одна из проблем связана с разными уровнями абстракции, а также, в частности, с чрезмерной дополнительной нагрузкой, создаваемой 3D API, которая может быть весьма ощутима. Но более серьёзной можно считать проблему совместимости, с которой разработчики Brook ничего не могли сделать. Между производителями GPU существует жёсткая конкуренция, поэтому они нередко оптимизируют свои драйверы. Если подобные оптимизации, по большей части, хороши для геймеров, они могут в один момент покончить с совместимостью Brook. Поэтому сложно представить использование этого API в промышленном коде, который будет где-то работать. И долгое время Brook оставался уделом исследователей-любителей и программистов.

Однако успеха Brook оказалось достаточно, чтобы привлечь внимание ATI и nVidia, у них зародился интерес к подобной инициативе, поскольку она могла бы расширить рынок, открыв для компаний новый немаловажный сектор.

Исследователи, изначально вовлечённые в проект Brook, быстро присоединились к командам разработчиков в Санта-Кларе, чтобы представить глобальную стратегию для развития нового рынка. Идея заключалась в создании комбинации аппаратного и программного обеспечения, подходящего для задач GPGPU. Поскольку разработчики nVidia знают все секреты своих GPU, то на графическое API можно было и не опираться, а связываться с графическим процессором через драйвер. Хотя, конечно, при этом возникают свои проблемы. Итак, команда разработчиков CUDA (Compute Unified Device Architecture) создала набор программных уровней для работы с GPU.

Нажмите на картинку для увеличения.

Как можно видеть на диаграмме, CUDA обеспечивает два API.

• Высокоуровневый API: CUDA Runtime API;
• низкоуровневый API: CUDA Driver API.

Поскольку высокоуровневый API реализован над низкоуровневым, каждый вызов функции уровня Runtime разбивается на более простые инструкции, которые обрабатывает Driver API. Обратите внимание, что два API взаимно исключают друг друга: программист может использовать один или другой API, но смешивать вызовы функций двух API не получится. Вообще, термин "высокоуровневый API" относителен. Даже Runtime API таков, что многие сочтут его низкоуровневым; впрочем, он всё же предоставляет функции, весьма удобные для инициализации или управления контекстом. Но не ожидайте особо высокого уровня абстракции - вам всё равно нужно обладать хорошим набором знаний о nVidia GPU и о том, как они работают.

С Driver API работать ещё сложнее; для запуска обработки на GPU вам потребуется больше усилий. С другой стороны, низкоуровневый API более гибок, предоставляя программисту дополнительный контроль, если нужно. Два API способны работать с ресурсами OpenGL или Direct3D (только девятая версия на сегодня). Польза от такой возможности очевидна - CUDA может использоваться для создания ресурсов (геометрия, процедурные текстуры и т.д.), которые можно передать на графическое API или, наоборот, можно сделать так, что 3D API будет отсылать результаты рендеринга программе CUDA, которая, в свою очередь, будет выполнять пост-обработку. Есть много примеров таких взаимодействий, и преимущество заключается в том, что ресурсы продолжают храниться в памяти GPU, их не требуется передавать через шину PCI Express, которая по-прежнему остаётся "узким местом".

Впрочем, следует отметить, что совместное использование ресурсов в видеопамяти не всегда проходит идеально и может привести к некоторым "головным болям". Например, при смене разрешения или глубины цвета, графические данные приоритетны. Поэтому если требуется увеличить ресурсы в кадровом буфере, то драйвер без проблем сделает это за счёт ресурсов приложений CUDA, которые попросту "вылетят" с ошибкой. Конечно, не очень элегантно, но такая ситуация не должна случаться очень уж часто. И раз уж мы начали говорить о недостатках: если вы хотите использовать несколько GPU для приложений CUDA, то вам нужно сначала отключить режим SLI, иначе приложения CUDA смогут "видеть" только один GPU.

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

• CUBLAS, где есть необходимые блоки для вычислений линейной алгебры на GPU;
• CUFFT, которая поддерживает расчёт преобразований Фурье - алгоритм, широко используемый в области обработки сигналов.

Перед тем, как мы погрузимся в CUDA, позвольте определить ряд терминов, разбросанных по документации nVidia. Компания выбрала весьма специфическую терминологию, к которой трудно привыкнуть. Прежде всего, отметим, что поток (thread) в CUDA имеет далеко не такое же значение, как поток CPU, а также и не является эквивалентом потока в наших статьях о GPU. Поток GPU в данном случае является базовый набор данных, которые требуется обработать. В отличие от потоков CPU, потоки CUDA очень "лёгкие", то есть переключение контекста между двумя потоками - отнюдь не ресурсоёмкая операция.

Второй термин, часто встречающийся в документации CUDA - варп (warp) . Здесь путаницы нет, поскольку в русском языке аналога не существует (разве что вы не являетесь фанатом Start Trek или игры Warhammer). На самом деле термин взят из текстильной промышленности, где через основную пряжу (warp yarn), которая растянута на станке, протягивается уточная пряжа (weft yarn). Варп в CUDA представляет собой группу из 32 потоков и является минимальным объёмом данных, обрабатываемых SIMD-способом в мультипроцессорах CUDA.

Но подобная "зернистость" не всегда удобна для программиста. Поэтому в CUDA, вместо работы с варпами напрямую, можно работать с блоками/block , содержащими от 64 до 512 потоков.

Наконец, эти блоки собираются вместе в сетки/grid . Преимущество подобной группировки заключается в том, что число блоков, одновременно обрабатываемых GPU, тесно связано с аппаратными ресурсами, как мы увидим ниже. Группировка блоков в сетки позволяет полностью абстрагироваться от этого ограничения и применить ядро/kernel к большему числу потоков за один вызов, не думая о фиксированных ресурсах. За всё это отвечают библиотеки CUDA. Кроме того, подобная модель хорошо масштабируется. Если GPU имеет мало ресурсов, то он будет выполнять блоки последовательно. Если число вычислительных процессоров велико, то блоки могут выполняться параллельно. То есть, один и тот же код может работать на GPU как начального уровня, так и на топовых и даже будущих моделях.

Есть ещё пара терминов в CUDA API, которые обозначают CPU (хост/host ) и GPU (устройство/device ). Если это небольшое введение вас не испугало, то настало время поближе познакомиться с CUDA.

Если вы регулярно читаете Tom"s Hardware Guide, то архитектура последних GPU от nVidia вам знакома. Если нет, мы рекомендуем ознакомиться со статьёй "nVidia GeForce GTX 260 и 280: новое поколение видеокарт ". Что касается CUDA, то nVidia представляет архитектуру несколько по-другому, демонстрируя некоторые детали, раньше остававшиеся скрытыми.

Как можно видеть по иллюстрации выше, ядро шейдеров nVidia состоит из нескольких кластеров текстурных процессоров (Texture Processor Cluster, TPC) . Видеокарта 8800 GTX, например, использовала восемь кластеров, 8800 GTS - шесть и т.д. Каждый кластер, по сути, состоит из текстурного блока и двух потоковых мультипроцессоров (streaming multiprocessor) . Последние включают начало конвейера (front end), выполняющее чтение и декодирование инструкций, а также отсылку их на выполнение, и конец конвейера (back end), состоящий из восьми вычислительных устройств и двух суперфункциональных устройств SFU (Super Function Unit) , где инструкции выполняются по принципу SIMD, то есть одна инструкция применяется ко всем потокам в варпе. nVidia называет такой способ выполнения SIMT (single instruction multiple threads, одна инструкция, много потоков). Важно отметить, что конец конвейера работает на частоте в два раза превосходящей его начало. На практике это означает, что данная часть выглядит в два раза "шире", чем она есть на самом деле (то есть как 16-канальный блок SIMD вместо восьмиканального). Потоковые мультипроцессоры работают следующим образом: каждый такт начало конвейера выбирает варп, готовый к выполнению, и запускает выполнение инструкции. Чтобы инструкция применилась ко всем 32 потокам в варпе, концу конвейера потребуется четыре такта, но поскольку он работает на удвоенной частоте по сравнению с началом, потребуется только два такта (с точки зрения начала конвейера). Поэтому, чтобы начало конвейера не простаивало такт, а аппаратное обеспечение было максимально загружено, в идеальном случае можно чередовать инструкции каждый такт - классическая инструкция в один такт и инструкция для SFU - в другой.

Каждый мультипроцессор обладает определённым набором ресурсов, в которых стоит разобраться. Есть небольшая область памяти под названием "Общая память/Shared Memory" , по 16 кбайт на мультипроцессор. Это отнюдь не кэш-память: программист может использовать её по своему усмотрению. То есть, перед нами что-то близкое к Local Store у SPU на процессорах Cell. Данная деталь весьма любопытная, поскольку она подчёркивает, что CUDA - это комбинация программных и аппаратных технологий. Данная область памяти не используется для пиксельных шейдеров, что nVidia остроумно подчёркивает "нам не нравится, когда пиксели разговаривают друг с другом".

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

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

Мультипроцессор имеет 8 192 регистра, которые общие для всех потоков всех блоков, активных на мультипроцессоре. Число активных блоков на мультипроцессор не может превышать восьми, а число активных варпов ограничено 24 (768 потоков). Поэтому 8800 GTX может обрабатывать до 12 288 потоков в один момент времени. Все эти ограничения стоило упомянуть, поскольку они позволяют оптимизировать алгоритм в зависимости от доступных ресурсов.

Оптимизация программы CUDA, таким образом, состоит в получении оптимального баланса между количеством блоков и их размером. Больше потоков на блок будут полезны для снижения задержек работы с памятью, но и число регистров, доступных на поток, уменьшается. Более того, блок из 512 потоков будет неэффективен, поскольку на мультипроцессоре может быть активным только один блок, что приведёт к потере 256 потоков. Поэтому nVidia рекомендует использовать блоки по 128 или 256 потоков, что даёт оптимальный компромисс между снижением задержек и числом регистров для большинства ядер/kernel.

С программной точки зрения CUDA состоит из набора расширений к языку C, что напоминает BrookGPU, а также нескольких специфических вызовов API. Среди расширений присутствуют спецификаторы типа, относящиеся к функциям и переменным. Важно запомнить ключевое слово __global__ , которое, будучи приведённым перед функцией, показывает, что последняя относится к ядру/kernel - эту функцию будет вызывать CPU, а выполняться она будет на GPU. Префикс __device__ указывает, что функция будет выполняться на GPU (который, кстати, CUDA и называет "устройство/device") но она может быть вызвана только с GPU (иными словами, с другой функции __device__ или с функции __global__). Наконец, префикс __host__ опционален, он обозначает функцию, которая вызывается CPU и выполняется CPU - другими словами, обычную функцию.

Есть ряд ограничений, связанных с функциями __device__ и __global__: они не могут быть рекурсивными (то есть вызывать самих себя), и не могут иметь переменное число аргументов. Наконец, поскольку функции __device__ располагаются в пространстве памяти GPU, вполне логично, что получить их адрес не удастся. Переменные тоже имеют ряд квалификаторов, которые указывают на область памяти, где они будут храниться. Переменная с префиксом __shared__ означает, что она будет храниться в общей памяти потокового мультипроцессора. Вызов функции __global__ немного отличается. Дело в том, при вызове нужно задать конфигурацию выполнения - более конкретно, размер сетки/grid, к которой будет применено ядро/kernel, а также размер каждого блока. Возьмём, например, ядро со следующей подписью.

__global__ void Func(float* parameter);

Оно будет вызываться в виде

Func<<< Dg, Db >>> (parameter);

где Dg является размером сетки, а Db - размером блока. Две этих переменных относятся к новому типу вектора, появившегося с CUDA.

API CUDA содержит функции для работы с памятью в VRAM: cudaMalloc для выделения памяти, cudaFree для освобождения и cudaMemcpy для копирования памяти между RAM и VRAM и наоборот.

Мы закончим данный обзор весьма интересным способом, которым компилируется программа CUDA: компиляция выполняется в несколько этапов. Сначала извлекается код, относящийся к CPU, который передаётся стандартному компилятору. Код, предназначенный для GPU, сначала преобразовывается в промежуточный язык PTX. Он подобен ассемблеру и позволяет изучать код в поисках потенциальных неэффективных участков. Наконец, последняя фаза заключается в трансляции промежуточного языка в специфические команды GPU и создании двоичного файла.

Просмотрев документацию nVidia, так и хочется попробовать CUDA на неделе. Действительно, что может быть лучше оценки API путём создания собственной программы? Именно тогда большинство проблем должны выплыть на поверхность, пусть даже на бумаге всё выглядит идеально. Кроме того, практика лучше всего покажет, насколько хорошо вы поняли все принципы, изложенные в документации CUDA.

В подобный проект погрузиться довольно легко. Сегодня для скачивания доступно большое количество бесплатных, но качественных инструментов. Для нашего теста мы использовали Visual C++ Express 2005, где есть всё необходимое. Самое сложное заключалось в том, чтобы найти программу, портирование которой на GPU не заняло бы несколько недель, и вместе с тем она была бы достаточно интересная, чтобы наши усилия не пропали даром. В конце концов, мы выбрали отрезок кода, который берёт карту высот и рассчитывает соответствующую карту нормалей. Мы не будем детально углубляться в эту функцию, поскольку в данной статье это вряд ли интересно. Если быть кратким, то программа занимается искривлением участков: для каждого пикселя начального изображения мы накладываем матрицу, определяющую цвет результирующего пикселя в генерируемом изображении по прилегающим пикселям, используя более или менее сложную формулу. Преимущество этой функции в том, что её очень легко распараллелить, поэтому данный тест прекрасно показывает возможности CUDA.

Ещё одно преимущество заключается в том, что у нас уже есть реализация на CPU, поэтому мы можем сравнивать её результат с версией CUDA - и не изобретать колесо заново.

Ещё раз повторим, что целью теста являлось знакомство с утилитами CUDA SDK, а не сравнительное тестирование версий под CPU и GPU. Поскольку это была первая наша попытка создания программы CUDA, мы не особо надеялись получить высокую производительность. Так как данная часть кода не является критической, то версия под CPU была не оптимизирована, поэтому прямое сравнение результатов вряд ли интересно.

Производительность

Однако мы замерили время выполнения, чтобы посмотреть, есть ли преимущество в использовании CUDA даже с самой грубой реализацией, или нам потребуется длительная и утомительная практика, чтобы получить какой-то выигрыш при использовании GPU. Тестовая машина была взята из нашей лаборатории разработки - ноутбук с процессором Core 2 Duo T5450 и видеокартой GeForce 8600M GT, работающей под Vista. Это далеко не суперкомпьютер, но результаты весьма интересны, поскольку тест не "заточен" под GPU. Всегда приятно видеть, когда nVidia демонстрирует огромный прирост на системах с монстрообразными GPU и немалой пропускной способностью, но на практике многие из 70 миллионов GPU с поддержкой CUDA на современном рынке ПК далеко не такие мощные, поэтому и наш тест имеет право на жизнь.

Для изображения 2 048 x 2 048 пикселей мы получили следующие результаты.

• CPU 1 поток: 1 419 мс;
• CPU 2 потока: 749 мс;
• CPU 4 потока: 593 мс
• GPU (8600M GT) блоки по 256 потоков: 109 мс;
• GPU (8600M GT) блоки по 128 потоков: 94 мс;
• GPU (8800 GTX) блоки по 128 потоков/ 256 потоков: 31 мс.

По результатам можно сделать несколько выводов. Начнём с того, что, несмотря на разговоры об очевидной лени программистов, мы модифицировали начальную версию CPU под несколько потоков. Как мы уже упоминали, код идеален для этой ситуации - всё, что требуется, это разбить начальное изображение на столько зон, сколько существует потоков. Обратите внимание, что от перехода от одного потока на два на нашем двуядерном CPU ускорение получилось почти линейное, что тоже указывает на параллельную природу тестовой программы. Весьма неожиданно, но версия с четырьмя потоками тоже оказалась быстрее, хотя на нашем процессоре это весьма странно - можно было, напротив, ожидать падения эффективности из-за накладных расходов на управление дополнительными потоками. Как можно объяснить такой результат? Сложно сказать, но, возможно, виновен планировщик потоков под Windows; в любом случае, результат повторяем. С текстурами меньшего размера (512x512) прирост от разделения на потоки был не такой выраженный (примерно 35% против 100%), и поведение версии с четырьмя потоками было логичнее, без прироста по сравнению с версией на два потока. GPU работал всё ещё быстрее, но уже не так выражено (8600M GT была в три раза быстрее, чем версия с двумя потоками).

Нажмите на картинку для увеличения.

Второе значимое наблюдение - даже самая медленная реализация GPU оказалась почти в шесть раз быстрее, чем самая производительная версия CPU. Для первой программы и неоптимизированной версии алгоритма результат очень даже ободряющий. Обратите внимание, что мы получили ощутимо лучший результат на небольших блоках, хотя интуиция может подсказывать об обратном. Объяснение простое - наша программа использует 14 регистров на поток, и с 256-поточными блоками требуется 3 584 регистра на блок, а для полной нагрузки процессора требуется 768 потоков, как мы показывали. В нашем случае это составляет три блока или 10 572 регистра. Но мультипроцессор имеет всего 8 192 регистра, поэтому он может поддерживать активными только два блока. Напротив, с блоками по 128 потоков нам требуется 1 792 регистра на блок; если 8 192 поделить на 1 792 и округлить до ближайшего целого, то мы получим четыре блока. На практике число потоков будет таким же (512 на мультипроцессор, хотя для полной нагрузки теоретически нужно 768), но увеличение числа блоков даёт GPU преимущество гибкости по доступу к памяти - когда идёт операция с большими задержками, то можно запустить выполнение инструкций другого блока, ожидая поступления результатов. Четыре блока явно снижают задержки, особенно с учётом того, что наша программа использует несколько доступов в память.

Анализ

Наконец, несмотря на то, что мы сказали выше, мы не смогли устоять перед искушением и запустили программу на 8800 GTX, которая оказалась в три раза быстрее 8600, независимо от размера блоков. Можно подумать, что на практике на соответствующих архитектурах результат будет в четыре или более раз выше: 128 АЛУ/шейдерных процессоров против 32 и более высокая тактовая частота (1,35 ГГц против 950 МГц), но так не получилось. Скорее всего, ограничивающим фактором оказался доступ к памяти. Если быть более точным, доступ к начальному изображению осуществляется как к многомерному массиву CUDA - весьма сложный термин для того, что является не более, чем текстурой. Но ест несколько преимуществ.

• доступы выигрывают от кэша текстур;
• мы используем wrapping mode, в котором не нужно обрабатывать границы изображения, в отличие от версии CPU.

Кроме того, мы можем получить преимущество от "бесплатной" фильтрации с нормализованной адресацией между вместо и , но в нашем случае это вряд ли полезно. Как вы знаете, 8600 оснащён 16 текстурными блоками по сравнению с 32 у 8800 GTX. Поэтому между двумя архитектурами соотношение всего два к одному. Добавьте к этому разницу в частотах, и мы получим соотношение (32 x 0,575) / (16 x 0,475) = 2,4 - близко к "трём к одному", что мы получили на самом деле. Данная теория также объясняет, почему размер блоков многое на G80 не меняет, поскольку АЛУ всё равно упирается в текстурные блоки.

Нажмите на картинку для увеличения.

Кроме многообещающих результатов, наше первое знакомство с CUDA прошло очень хорошо, учитывая не самые благоприятные выбранные условия. Разработка на ноутбуке под Vista подразумевает, что придётся использовать CUDA SDK 2.0, всё ещё находящееся в состоянии бета-версии, с драйвером 174.55, который тоже бета-версия. Несмотря на это мы не можем сообщить о каких-либо неприятных сюрпризах - только начальные ошибки во время первой отладки, когда наша программа, всё ещё весьма "глючная" попыталась адресовать память за пределами выделенного пространства.

Монитор начал дико мерцать, затем экран почернел... пока Vista не запустила службу восстановления драйвера, и всё стало в порядке. Но всё же несколько удивительно это наблюдать, если вы привыкли видеть типичную ошибку Segmentation Fault на стандартных программах, подобно нашей. Наконец, небольшая критика в сторону nVidia: во всей документации, доступной для CUDA, нет небольшого руководства, которое бы шаг за шагом рассказывало о том, как настроить окружение разработки под Visual Studio. Собственно, проблема невелика, поскольку в SDK есть полный набор примеров, которые можно изучить для понимания каркаса для приложений CUDA, но руководство для новичков не помешало бы.

Нажмите на картинку для увеличения.

nVidia представила CUDA с выпуском GeForce 8800. И в то время обещания казались весьма соблазнительными, но мы придержали свой энтузиазм до реальной проверки. Действительно, в то время это казалось больше разметкой территории, чтобы оставаться на волне GPGPU. Без доступного SDK сложно сказать, что перед нами не очередная маркетинговая пустышка, из которой ничего не получится. Уже не в первый раз хорошая инициатива была объявлена слишком рано и в то время не вышла на свет из-за недостатка поддержки - особенно в столь конкурентном секторе. Теперь, через полтора года после объявления, мы с уверенностью можем сказать, что nVidia сдержала слово.

SDK довольно быстро появился в бета-версии в начале 2007 года, с тех пор он быстро обновлялся, что доказывает значимость этого проекта для nVidia. Сегодня CUDA весьма приятно развивается: SDK доступен уже в бета-версии 2.0 для основных операционных систем (Windows XP и Vista, Linux, а также 1.1 для Mac OS X), а для разработчиков nVidia выделила целый раздел сайта.

На более профессиональном уровне впечатление от первых шагов с CUDA оказалось очень даже позитивным. Если даже вы знакомы с архитектурой GPU, вы легко разберётесь. Когда API выглядит понятным с первого взгляда, то сразу же начинаешь полагать, что получишь убедительные результаты. Но не будет ли теряться вычислительное время от многочисленных передач с CPU на GPU? И как использовать эти тысячи потоков практически без примитива синхронизации? Мы начинали наши эксперименты со всеми этими опасениями в уме. Но они быстро рассеялись, когда первая версия нашего алгоритма, пусть и весьма тривиального, оказалась существенно быстрее, чем на CPU.

Так что CUDA - это не "палочка-выручалочка" для исследователей, которые хотят убедить руководство университета купить им GeForce. CUDA - уже полностью доступная технология, которую может использовать любой программист со знанием C, если он готов потратить время и усилия на привыкание к новой парадигме программирования. Эти усилия не будут потеряны даром, если ваши алгоритмы хорошо распараллеливаются. Также мы хотели бы поблагодарить nVidia за предоставление полной и качественной документации, где найдут ответы начинающие программисты CUDA.

Что же требуется CUDA, чтобы стать узнаваемым API? Если говорить одним словом: переносимость. Мы знаем, что будущее ИТ кроется в параллельных вычислениях - сегодня уже каждый готовится к подобным изменениям, и все инициативы, как программные, так и аппаратные, направлены в этом направлении. Однако на данный момент, если смотреть на развитие парадигм, мы находится ещё на начальном этапе: мы создаём потоки вручную и стараемся спланировать доступ к общим ресурсам; со всем этим ещё как-то можно справиться, если количество ядер можно пересчитать по пальцам одной руки. Но через несколько лет, когда число процессоров будет исчисляться сотнями, такой возможности уже не будет. С выпуском CUDA nVidia сделала первый шаг в решении этой проблемы - но, конечно, данное решение подходит только для GPU от этой компании, да и то не для всех. Только GF8 и 9 (и их производные Quadro/Tesla) сегодня могут работать с программами CUDA. И новая линейка 260/280, конечно.

Нажмите на картинку для увеличения.

nVidia может хвастаться тем, что продала 70 миллионов CUDA-совместимых GPU по всему миру, но этого всё равно мало, чтобы стать стандартом де-факто. С учётом того, что конкуренты не сидят, сложа руки. AMD предлагает собственный SDK (Stream Computing), да и Intel объявила о решении (Ct), хотя оно ещё не доступно. Грядёт война стандартов, и на рынке явно не будет места для трёх конкурентов, пока другой игрок, например, Microsoft, не выйдет с предложением общего API, что, конечно, облегчит жизнь разработчикам.

Поэтому у nVidia есть немало трудностей на пути утверждения CUDA. Хотя технологически перед нами, без сомнения, успешное решение, ещё остаётся убедить разработчиков в его перспективах - и это будет сделать нелегко. Впрочем, судя по многим недавним объявлениям и новостям по поводу API, будущее выглядит отнюдь не печальным.

Технология CUDA

Владимир Фролов, [email protected]

Аннотация

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

1. Введение

Развитие вычислительных технологий последние десятки лет шло быстрыми темпами. Настолько быстрыми, что уже сейчас разработчики процессоров практически подошли к так называемому «кремниевому тупику». Безудержный рост тактовой частоты стал невозможен в силу целого ряда серьезных технологических причин.

Отчасти поэтому все производители современных вычислительных систем идут в сторону увеличения числа процессоров и ядер, а не увеличивают частоту одного процессора. Количество ядер центрального процессора (CPU) в передовых системах сейчас уже равняется 8.

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

Однако если посмотреть в сторону графических процессоров GPU (Graphics Processing Unit), то там по пути параллелизма пошли гораздо раньше. В сегодняшних видеокартах, например в GF8800GTX, число процессоров может достигать 128. Производительность подобных систем при умелом их программировании может быть весьма значительной (рис. 1).

Рис. 1. Количество операций с плавающей точкой для CPU и GPU

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

Рисунки 2,3. Игры Doom и Wolfenstein 3D

Со времени создания компанией 3Dfx первых видеокарт Voodoo, (1996 г.) и вплоть до 2001 года в GPU был реализован только фиксированный набор операций над входными данными.

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

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

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

Сегодня появилась тенденция нетрадиционного использования видеокарт для решения задач в областях квантовой механики, искусственного интеллекта, физических расчетов, криптографии, физически корректной визуализации, реконструкции по фотографиям, распознавания и.т.п. Эти задачи неудобно решать в рамках графических API (DirectX, OpenGL), так как эти API создавались совсем для других применений.

Развитие программирования общего назначения на GPU (General Programming on GPU, GPGPU) логически привело к возникновению технологий, нацеленных на более широкий круг задач, чем растеризация. В результате компанией Nvidia была создана технология Compute Unified Device Architecture (или сокращенно CUDA), а конкурирующей компанией ATI - технология STREAM.

Следует заметить, что на момент написания этой статьи, технология STREAM сильно отставала в развитии от CUDA, и поэтому здесь она рассматриваться не будет. Мы сосредоточимся на CUDA - технологии GPGPU, позволяющей писать программы на подмножестве языка C++.

2. Принципиальная разница между CPU и GPU

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

2.1. Возможности

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

Для GPU это не так. Как вы узнаете, прочитав эту статью, в CUDA имеется целых 6 видов памяти. Читать можно из любой ячейки, доступной физически, но вот записывать – не во все ячейки. Причина заключается в том, что GPU в любом случае представляет собой специфическое устройство, предназначенное для конкретных целей. Это ограничение введено ради увеличения скорости работы определенных алгоритмов и снижения стоимости оборудования.

2.2. Быстродействие памяти

Извечная проблема большинства вычислительных систем заключена в том, что память работает медленнее процессора. Производители CPU решают ее путем введения кэшей. Наиболее часто используемые участки памяти помещается в сверхоперативную или кэш-память, работающую на частоте процессора. Это позволяет сэкономить время при обращении к наиболее часто используемым данным и загрузить процессор собственно вычислениями.

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

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

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

3. Ядро CUDA

3.1. Потоковая модель

Вычислительная архитектура CUDA основана на концепции одна команда на множество данных (Single Instruction Multiple Data , SIMD) и понятии мультипроцессора .

Концепция SIMD подразумевает, что одна инструкция позволяет одновременно обработать множество данных. Например, команда addps в процессоре Pentium 3 и в более новых моделях Pentium позволяет складывать одновременно 4 числа с плавающей точкой одинарной точности.

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

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

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

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

Фактически, та часть вашей программы, которая работает на CPU - это хост, а ваша видеокарта - устройство. Логически устройство можно представить как набор мультипроцессоров (рис. 4) плюс драйвер CUDA.

Рис. 4. Устройство

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

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

Рис. 5. Организация потоков

На рис. 5. ядро обозначено как Kernel. Все потоки, выполняющие это ядро, объединяются в блоки (Block), а блоки, в свою очередь, объединяются в сетку (Grid).

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

В общем случае индексы представляют собой трехмерные векторы. Для каждого потока будут известны: индекс потока внутри блока threadIdx и индекс блока внутри сетки blockIdx. При запуске все потоки будут отличаться только этими индексами. Фактически, именно через эти индексы программист осуществляет управление, определяя, какая именно часть его данных обрабатывается в каждом потоке.

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

Блок задач (потоков) выполняется на мультипроцессоре частями, или пулами, называемыми warp. Размер warp на текущий момент в видеокартах с поддержкой CUDA равен 32 потокам. Задачи внутри пула warp исполняются в SIMD стиле, т.е. во всех потоках внутри warp одновременно может выполняться только одна инструкция .

Здесь следует сделать одну оговорку. В архитектурах, современных на момент написания этой статьи, количество процессоров внутри одного мультипроцессора равно 8, а не 32. Из этого следует, что не весь warp исполняется одновременно, он разбивается на 4 части, которые выполняются последовательно (т.к. процессоры скалярные).

Но, во-первых, разработчики CUDA не регламентируют жестко размер warp. В своих работах они упоминают параметр warp size, а не число 32. Во-вторых, с логической точки зрения именно warp является тем минимальным объединением потоков, про который можно говорить, что все потоки внутри него выполняются одновременно - и при этом никаких допущений относительно остальной системы сделано не будет .

3.1.1. Ветвления

Сразу же возникает вопрос: если в один и тот же момент времени все потоки внутри warp исполняют одну и ту же инструкцию, то как быть с ветвлениями? Ведь если в коде программы встречается ветвление, то инструкции будут уже разные. Здесь применяется стандартное для SIMD программирования решение (рис 6).

Рис. 6. Организация ветвления в SIMD

Пусть имеется следующий код:

if(cond) B;

В случае SISD (Single Instruction Single Data) мы выполняем оператор A, проверяем условие, затем выполняем операторы B и D (если условие истинно).

Пусть теперь у нас есть 10 потоков, исполняющихся в стиле SIMD. Во всех 10 потоках мы выполняем оператор A, затем проверяем условие cond и оказывается, что в 9 из 10 потоках оно истинно, а в одном потоке - ложно.

Понятно, что мы не можем запустить 9 потоков для выполнения оператора B, а один оставшийся- для выполнения оператора C, потому что одновременно во всех потоках может исполняться только одна инструкция. В этом случае нужно поступить так: сначала «убиваем» отколовшийся поток так, чтобы он не портил ничьи данные, и выполняем 9 оставшихся потоков. Затем «убиваем» 9 потоков, выполнивших оператор B, и проходим один поток с оператором C. После этого потоки опять объединяются и выполняют оператор D все одновременно .

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

Однако не все так плохо, как может показаться на первый взгляд. К очень большому плюсу технологии можно отнести то, что эти фокусы выполняются динамически драйвером CUDA и для программиста они совершенно прозрачны. В то же время, имея дело с SSE командами современных CPU (именно в случае попытки выполнения 4 копий алгоритма одновременно), программист сам должен заботиться о деталях: объединять данные по четверкам, не забывать о выравнивании, и вообще писать на низком уровне, фактически как на ассемблере .

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

3.1.2. Взаимодействие между потоками

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

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

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

3.2. Память

В CUDA выделяют шесть видов памяти (рис. 7). Это регистры, локальная, глобальная, разделяемая, константная и текстурная память.

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

Рис. 7. Виды памяти в CUDA

3.2.0. Регистры

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

При обычном разделении в 64 потока на блок получается всего 128 регистров (существуют некие объективные критерии, но 64 подходит в среднем для многих задач). Реально, 128 регистров nvcc никогда не выделит. Обычно он не дает больше 40, а остальные переменные попадпют в локальную память. Так происходит потому что на одном мультипроцессоре может исполняться несколько блоков. Компилятор старается максимизировать число одновременно работающих блоков. Для большей большей эффективности надо стараться занимать меньше чем 32 регистра. Тогда теоретически может быть запущено 4 блока (8 warp-ов, если 64 треда в одном блоке) на одном мультипроцессоре. Однако здесь еще следует учитывать объем разделяемой памяти, занимаемой потоками, так как если один блок занимает всю разделяемую память, два таких блока не могут выполняться на мультипроцессоре одновременно .

3.2.1. Локальная память

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

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

3.2.2. Глобальная память

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

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

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

Переменные, объявленные с квалификатором __global__, размещаются в глобальной памяти. Глобальную память также можно выделить динамически, вызвав функцию cudaMalloc(void* mem, int size) на хосте. Из устройства эту функцию вызывать нельзя. Отсюда следует, что распределением памяти должна заниматься программа-хост, работающая на CPU. Данные с хоста можно отправлять в устройство вызовом функции cudaMemcpy:

cudaMemcpy(void* gpu_mem, void* cpu_mem, int size, cudaMemcpyHostToDevice);

Точно таким же образом можно проделать и обратную процедуру:

cudaMemcpy(void* cpu_mem, void* gpu_mem, int size, cudaMemcpyDeviceToHost);

Этот вызов тоже осуществляется с хоста.

При работе с глобальной памятью важно соблюдать правило коалесинга (coalescing). Основная идея в том, что треды должны обращаться к последоваетльным ячейкам памяти, причем 4,8 или 16 байтовым. При этом, самый первый тред должен обращаться по адресу, выровненному на границу соответственно 4,8 или 16 байт. Адреса, возвращаемые cudaMalloc выровнены как минимум по границе 256 байт.

3.2.3. Разделяемая память

Разделяемая память - это некэшируемая, но быстрая память. Ее и рекомендуется использовать как управляемый кэш. На один мультипроцессор доступно всего 16KB разделяемой памяти. Разделив это число на количество задач в блоке, получим максимальное количество разделяемой памяти, доступной на один поток (если планируется использовать ее независимо во всех потоках).

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

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

Переменные, объявленные с квалификатором __shared__, размещаются в разделяемой памяти.

Shared__ float mem_shared;

Следует еще раз подчеркнуть, что разделяемая память для блока одна. Поэтому если нужно использовать ее просто как управляемый кэш, следует обращаться к разным элементам массива, например, так:

float x = mem_shared;

Где threadIdx.x - индекс x потока внутри блока.

3.2.4. Константная память

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

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

#define N 100

Constant__ int gpu_buffer[N];

void host_function()

int cpu_buffer[N];

cudaMemcpyToSymbol(gpu_buffer, cpu_buffer, sizeof(int )*N);

// __global__ означает, что device_kernel - ядро, которое может быть запущено на GPU

Global__ void device_kernel()

int a = gpu_buffer;

int b = gpu_buffer + gpu_buffer;

// gpu_buffer = a; ОШИБКА! константная память доступна только для чтения

Так как для константной памяти используется кэш, доступ к ней в общем случае довольно быстрый. Единственный, но очень большой недостаток константной памяти заключается в том, что ее размер составляет всего 64 Kбайт (на все устройство). Из этого следует, что в контекстной памяти имеет смысл хранить лишь небольшое количество часто используемых данных.

3.2.5. Текстурная память

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

Название текстурной памяти (и, к сожалению, функциональность) унаследовано от понятий «текстура» и «текстурирование». Текстурирование - это процесс наложения текстуры (просто картинки) на полигон в процессе растеризации. Текстурная память оптимизирована под выборку 2D данных и имеет следующие возможности:

быстрая выборка значений фиксированного размера (байт, слово, двойное или учетверенное слово) из одномерного или двухмерного массива;

нормализованная адресация числами типа float в интервале . Затем можно их выбирать, используя нормализованную адресацию. Результирующим значением будетет слово типа float4, отображенное в интервал ;

CudaMalloc((void**) &gpu_memory, N*sizeof (uint4 )); // выделим память в GPU

// настройка параемтров текстуры texture

Texture.addressMode = cudaAddressModeWrap; // режим Wrap

Texture.addressMode = cudaAddressModeWrap;

Texture.filterMode = cudaFilterModePoint; // ближайшеезначение

Texture.normalized = false; // не использовать нормализованную адресацию

CudaBindTexture (0, texture , gpu _ memory , N ) // отныне эта память будет считаться текстурной


CudaMemcpy (gpu _ memory , cpu _ buffer , N * sizeof (uint 4), cudaMemcpyHostToDevice ); // копируем данные на GPU

// __global__ означает, что device_kernel - ядро, которое нужно распараллелить

Global__ void device_kernel()

uint4 a = tex1Dfetch(texture,0); // можно выбирать данные только таким способом!

uint4 b = tex1Dfetch(texture,1);

int c = a.x * b.y;

...

3.3. Простой пример

В качестве простого примера предлагается рассмотреть программу cppIntegration из CUDA SDK. Она демонстрирует приемы работы с CUDA, а также использование nvcc (специальный компилятор подмножества С++ от Nvidia) в сочетании с MS Visual Studio, что сильно упрощает разработку программ на CUDA.

4.1. Правильно проводите разбиение вашей задачи

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

4.2. Выбор типа памяти

Помещайте свои данные в текстурную или константную память, если все задачи одного блока обращаются к одному и тому же участку памяти или к близко расположенным участкам. Двухмерные данные могут быть эффективно обработаны при помощи функций text2Dfetch и text2D. Текстурная память специально оптимизирована под двухмерную выборку.

Используйте глобальную память в сочетании с разделяемой памятью, если все задачи обращаются бессистемно к разным, далеко расположенным друг от друга участкам памяти (с сильно различными адресами или координатами, если это 2D/3D данные).

глобальная память => разделяемая память

Syncthreads();

Обработать данные в разделяемой памяти

Syncthreads();

глобальная память <= разделяемая память

4.3. Включите счетчики памяти

Флаг компилятора --ptxas-options=-v позволяет точно сказать, сколько и какой памяти (регистров, разделяемой, локальной, константной) вы используете. Если компилятор использует локальную память, вы точно знаете об этом. Анализ данных о количестве и типах используемой памяти может сильно помочь вам при оптимизации программы.

4.4. Старайтесь минимизировать использование регистров и разделяемой памяти

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

4.5. Разделяемая память вместо локальной

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

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

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

Хотя автоматическое размещение компилятором таких переменных в локальной памяти может показаться удобным, на самом деле это не так. Непросто будет найти узкое место при последующих модификациях программы, если переменная начнет использоваться чаще. Компилятор может перенести такую переменную в регистровую память, а может и не перенести. Если же модификатор __global__ будет указан явно, программист скорее обратит на это внимание.

4.6. Разворачивание циклов

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

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

int a[N]; int summ;

for (int i=0;i

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

template

class ArraySumm

Device__ static T exec(const T* arr) { return arr + ArraySumm(arr+1); }

template

class ArraySumm<0,T>

Device__ static T exec(const T* arr) { return 0; }

for (int i=0;i

summ+= ArraySumm<4,int>::exec(a);

Следует отметить одну интересную особенность компилятора nvcc. Компилятор всегда будет встраивать функции типа __device__ по умолчанию (чтобы это отменить, существует специальная директива __noinline__) .

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

4.7. Выравнивание данных и выборка по 16 байт

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

Если структура занимает 8 байт или меньше, можно выравнивать ее по 8 байт. Но в этом случае можно выбрать сразу две переменные за один раз, объединив две 8-байтовые переменные в структуру с помощью union или приведения указателей. Приведением следует пользоваться осторожно, так как компилятор может поместить данные в локальную память, а не в регистры.

4.8. Конфликты банков разделяемой памяти

Разделяемая память организована в виде 16 (всего-то!) банков памяти с шагом в 4 байта. Во время выполнения пула потоков warp на мультипроцессоре, он делится на две половинки (если warp-size = 32) по 16 потоков, которые осуществляют доступ к разделяемой памяти по очереди.

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

Оптимальные шаги - 4, 12, 28, ..., 2^n-4 байт (рис. 8).



Рис. 8. Оптимальные шаги.

Не оптимальные шаги – 1, 8, 16, 32, ..., 2^n байт (рис. 9).



Рис. 9. Неоптимальные шаги

4.9. Минимизация перемещений данных Host <=> Device

Старайтесь как можно реже передавать промежуточные результаты на host для обработки с помощью CPU. Реализуйте если не весь алгоритм, то, по крайней мере, его основную часть на GPU, оставляя CPU лишь управляющие задачи.

5. CPU/GPU переносимая математическая библиотека

Автором этой статьи написана переносимая библиотека MGML_MATH для работы с простыми пространственными объектами, код которой работоспособен как на устройстве, так и на хосте.

Библиотека MGML_MATH может быть использована как каркас для написания CPU/GPU переносимых (или гибридных) систем расчета физических, графических или других пространственных задач. Основное ее достоинство в том, что один и тот же код может использоваться как на CPU, так и на GPU, и при этом во главу требований, предъявляемых к библиотеке, ставится скорость.

6 . Литература

Крис Касперски. Техника оптимизации программ. Эффективное использование памяти. - Спб.: БХВ-Петербург, 2003. - 464 с.: ил.

CUDA Programming Guide 1.1 (http://developer.download.nvidia.com/compute/cuda/1_1/NVIDIA_CUDA_Programming_Guide_1.1.pdf )

CUDA Programming Guide 1.1. page 14-15

CUDA Programming Guide 1.1. page 48