Карта с технологией cuda. Nvidia CUDA? неграфические вычисления на графических процессорах

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

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

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

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

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

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

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

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

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

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

Проблемы CPU

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

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

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

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

Развитие GPU

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

…

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

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

…

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

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

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

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

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

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

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

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

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

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

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

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

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

Что такое CUDA

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

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

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

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

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

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

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

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

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

К примеру:

GeForce

GTX 460 — 336 CUDA ядер

GTX 580 — 512 CUDA ядер

8800GTX — 128 CUDA ядер

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

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

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

Приступим.

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

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

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

CUDA и язык C:

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

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

Как было сказано, спецификаторы функций определяют, как и откуда буду вызываться функции. Всего в 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 (по крайней мере, на моей видеокарте).

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

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

Переходим непосредственно к написанию кода:

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

// Функция сложения двух векторов
__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), где

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

Возвращает:

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

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

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

Возвращает:

cudaSuccess – при удачном копировании
cudaErrorInvalidValue – неверные параметры аргумента (например, размер копирования отрицателен)
cudaErrorInvalidDevicePointer – неверный указатель памяти в видеокарте
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), где

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

Возвращает:

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

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

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

Возвращает:

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

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

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

Возвращает:

cudaSuccess – в случае успеха
cudaErrorInitializationError – ошибка инициализации
cudaErrorPriorLaunchFailure – ошибка при предыдущем асинхронном запуске функции
cudaErrorInvalidValue – неверное значение
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: Задавайте вопросы.

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

Приступим.

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

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

CUDA и язык C:

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

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

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

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

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

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

CUDA host API:

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

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

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

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

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

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

Переходим непосредственно к написанию кода:

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

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

Пишем код, которые отвечает за 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 .

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

Возвращает:

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

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

Возвращает:

cudaSuccess – при удачном копировании
cudaErrorInvalidValue – неверные параметры аргумента (например, размер копирования отрицателен)
cudaErrorInvalidDevicePointer – неверный указатель памяти в видеокарте
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 .

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

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

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

Возвращает:

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

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

Возвращает:

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

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

Возвращает:

cudaSuccess – в случае успеха
cudaErrorInitializationError – ошибка инициализации
cudaErrorPriorLaunchFailure – ошибка при предыдущем асинхронном запуске функции
cudaErrorInvalidValue – неверное значение
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 .

Заключение

Теги:

CUDA
gpgpu
nvidia

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

В развитии современных процессоров намечается тенденция к постепенному увеличению количества ядер, что повышает их возможности в параллельных вычислениях. Однако уже давно имеются GPU, значительно превосходящие центральные процессоры в данном отношении. И эти возможности графических процессоров уже взяты на заметку некоторыми компаниями. Первые попытки использовать графические ускорители для нецелевых вычислений предпринимались еще с конца 90-х годов. Но только появление шейдеров стало толчком к развитию абсолютно новой технологии, и в 2003 году появилось понятие GPGPU (General-purpose graphics processing units). Немаловажную роль в развитии данной инициативы сыграл BrookGPU, который является специальным расширением для языка C. До появления BrookGPU программисты могли работать с GPU лишь через API Direct3D или OpenGL. Brook позволил разработчикам работать с привычной средой, а уже сам компилятор с помощью специальных библиотек реализовал взаимодействие с GPU на низком уровне.

Такой прогресс не мог не привлечь внимания лидеров данной индустрии - AMD и NVIDIA, которые занялись разработкой собственных программных платформ для неграфических вычислений на своих видеокартах. Никто лучше разработчиков GPU не знает в совершенстве все нюансы и особенности своих продуктов, что позволяет этим же компаниям максимально эффективно оптимизировать программный комплекс для конкретных аппаратных решений. На данный момент NVIDIA развивает платформу CUDA (Compute Unified Device Architecture), у AMD подобная технология именуется CTM (Close To Metal) или AMD Stream Computing. Мы рассмотрим некоторые возможности CUDA и на практике оценим вычислительные возможности графического чипа G92 видеокарты GeForce 8800 GT.

Но прежде рассмотрим некоторые нюансы выполнения расчетов при помощи графических процессоров. Основное преимущество их заключается в том, что графический чип изначально проектируется под выполнение множества потоков, а каждое ядро обычного CPU выполняет поток последовательных инструкций. Любой современный GPU является мультипроцессором, состоящим из нескольких вычислительных кластеров, с множеством ALU в каждом. Самый мощный современный чип GT200 состоит из 10 таких кластеров, на каждый из которых приходится 24 потоковых процессора. У тестируемой видеокарты GeForce 8800 GT на базе чипа G92 семь больших вычислительных блока по 16 потоковых процессоров. CPU используют SIMD блоки SSE для векторных вычислений (single instruction multiple data - одна инструкция выполняется над многочисленными данными), что требует трансформации данных в 4х векторы. GPU скалярно обрабатывает потоки, т.е. одна инструкция применяется над несколькими потоками (SIMT - single instruction multiple threads). Это избавляет разработчиков от преобразования данных в векторы, и допускает произвольные ветвления в потоках. Каждый вычислительный блок GPU имеет прямой доступ к памяти. Да и пропускная способность видеопамяти выше, благодаря использованию нескольких раздельных контроллеров памяти (на топовом G200 это 8 каналов по 64-бит) и высоких рабочих частот.

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

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

Но как бы там ни было, есть огромный потенциал графических процессоров, который можно использовать, и который требует специфического подхода к разработке программных продуктов. Все это реализовано в аппаратно-программной среде CUDA, которая состоит из нескольких программных уровней - высокоуровневый CUDA Runtime API и низкоуровневый CUDA Driver API.

CUDA использует для программирования стандартный язык C, что является одним из основных ее преимуществ для разработчиков. Изначально CUDA включает библиотеки BLAS (базовый пакет программ линейной алгебры) и FFT (расчёт преобразований Фурье). Также CUDA имеет возможность взаимодействия с графическими API OpenGL или DirectX, возможность разработки на низком уровне, характеризуется оптимизированным распределением потоков данных между CPU и GPU. Вычисления CUDA выполняются одновременно с графическими, в отличие от аналогичной платформы AMD, где для расчетов на GPU вообще запускается специальная виртуальная машина. Но такое «сожительство» чревато и возникновением ошибок в случае создания большой нагрузки графическим API при одновременной работе CUDA - ведь графические операции имеют все же более высокий приоритет. Платформа совместима с 32- и 64-битными операционными системами Windows XP, Windows Vista, MacOS X и различными версиями Linux. Платформа открытая и на сайте, кроме специальных драйверов для видеокарты, можно загрузить программные пакеты CUDA Toolkit, CUDA Developer SDK, включающие компилятор, отладчик, стандартные библиотеки и документацию.

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

В качестве практических испытаний новой технологии на видеокарте MSI NX8800GT-T2D256E-OC мы воспользуемся программой TMPGEnc. Данный продукт является коммерческим (полная версия стоит $100), но к видеокартам MSI он поставляется в качестве бонуса в trial-версии сроком на 30 дней. Скачать данную версию можно и с сайта разработчика, но для установки TMPGEnc 4.0 XPress MSI Special Edition необходим оригинальный диск с драйверами от карты MSI - без него программа не инсталлируется.

Для отображения максимально полной информации о вычислительных возможностях в CUDA и сравнения с другими видеоадаптерами можно использовать специальную утилиту CUDA-Z. Вот какую информацию она выдает о нашей видеокарте GeForce 8800GT:

Относительно референсных моделей наш экземпляр работает на более высоких частотах: растровый домен на 63 МГц выше номинала, а шейдерные блоки быстрее на 174 МГц, память - на 100 МГц.

Мы сравним скорость конвертации одного и того же HD-видео при расчетах только с помощью CPU и при дополнительной активации CUDA в программе TMPGEnc на следующей конфигурации:

Процессор: Pentium Dual-Core E5200 2,5 ГГц;
Материнская плата: Gigabyte P35-S3;
Память: 2х1GB GoodRam PC6400 (5-5-5-18-2T)
Видеокарта: MSI NX8800GT-T2D256E-OC;
Жесткий диск: 320GB WD3200AAKS;
Блок питания: CoolerMaster eXtreme Power 500-PCAP;
Операционная система: Windows XP SP2;
TMPGEnc 4.0 XPress 4.6.3.268;
Драйвера видеокарты: ForceWare 180.60.

Для тестов процессор разгонялся до 3 ГГц (в конфигурации 11,5x261 МГц) и до 4 ГГц (11,5x348 МГц) при частоте оперативной памяти 835 МГц в первом и втором случае. Видеоролик в разрешении Full HD 1920х1080 продолжительностью одну минуту двадцать секунд. Для создания дополнительной нагрузки включался фильтр шумоподавления, настройки которого оставлены по умолчанию.

Кодирование осуществлялось с помощью кодека DivX 6.8.4. В настройках качества этого кодека все значения оставлены по умолчанию, multithreading включен.

Поддержка многопоточности в TMPGEnc изначально включена во вкладке настроек CPU/GPU. В этом же разделе активируется и CUDA.

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

По итогам проведенных тестов получены следующие данные:

При частоте процессора 4 ГГц с активацией CUDA мы выиграли всего пару секунд (или 2%), что не особо впечатляет. А вот на более низкой частоте прирост от активации данной технологии позволяет сэкономить уже около 13% времени, что будет довольно ощутимо при обработке больших файлов. Но все равно результаты не столь впечатляющие, как ожидалось.

В программе TMPGEnc есть индикатор загрузки CPU и CUDA, в данной тестовой конфигурации он показывал загрузку центрального процессора примерно на 20%, а графического ядра на оставшиеся 80%. В итоге у нас те же 100%, что и при конвертации без CUDA и разницы по времени вообще может и не быть (но она все-таки есть). Небольшой объем памяти в 256 MB так же не является сдерживающим фактором. Судя по показаниям RivaTuner, в процессе работы использовалось не более 154 MB видеопамяти.

Выводы

Программа TMPGEnc является одной из тех, кто вводит технологию CUDA в массы. Использование GPU в данной программе позволяет ускорить процесс обработки видео и значительно разгрузить центральный процессор, что позволит пользователю комфортно заниматься и другими задачами в это же время. В нашем конкретном примере видеокарта GeForce 8800GT 256MB незначительно улучшила временные показатели при конвертации видео на базе разогнанного процессора Pentium Dual-Core E5200. Но отчетливо видно, что при снижении частоты увеличивается прирост от активации CUDA, на слабых процессорах прирост от ее использования будет намного больше. На фоне такой зависимости вполне логично предположить что и при увеличении нагрузки (например, использование очень большого количества дополнительных видео-фильтров) результаты системы с CUDA будут выделяется более значимой дельтой разницы затраченного времени на процесс кодирования. Также не стоит забывать, что и G92 на данный момент не самый мощный чип, и более современные видеокарты обеспечат значительно более высокую производительность в подобных приложениях. Однако в процессе работы приложения GPU загружен не полностью и, вероятно, распределение нагрузки зависит от каждой конфигурации отдельно, а именно от связки процессор/видеокарта, что в итоге может дать и больший (или меньший) прирост в процентном соотношении от активации CUDA. В любом случае, тем, кто работает с большими объемами видеоданных, такая технология все равно позволит значительно сэкономить свое время.

Правда, CUDA еще не обрела повсеместную популярность, качество программного обеспечения, работающего с этой технологией, требует доработок. В рассмотренной нами программе TMPGEnc 4.0 XPress данная технология не всегда работала. Один и тот же ролик можно было перекодировать несколько раз, а потом вдруг, при следующем запуске, загрузка CUDA уже была равна 0%. И это явление носило совершенно случайный характер на абсолютно разных операционных системах. Также рассмотренная программа отказывалась использовать CUDA при кодировании в формат XviD, но с популярным кодеком DivX никаких проблем не было.

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

Благодарим следующие компании за предоставленное тестовое оборудование: