Я расскажу о ключевых моментах компилятора 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

Владимир Фролов, [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

На протяжении десятилетий действовал закон Мура, который гласит, что каждые два года количество транзисторов на кристалле будет удваиваться. Однако это было в далеком 1965 году, а последние 5 лет стала бурно развиваться идея физической многоядерности в процессорах потребительского класса: в 2005 году Intel представила Pentium D, а AMD – Athlon X2. Тогда приложений, использующих 2 ядра, можно было пересчитать по пальцам одной руки. Однако следующее поколение процессоров Intel, совершившее революцию, имело именно 2 физических ядра. Более того, в январе 2007 года появилась серия Quad, тогда же и сам Мур признался, что вскоре его закон перестанет действовать.

Что же сейчас? Двухядерные процессоры даже в бюджетных офисных системах, а 4 физических ядра стало нормой и это всего за 2-3 года. Частота процессоров не наращивается, а улучшается архитектура, увеличивается количество физических и виртуальных ядер. Однако идея использования видеоадаптеров, наделенных десятками, а то и сотнями вычислительных «блоков» витала давно.

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

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

CUDA (англ. Compute Unified Device Architecture) - программно-аппаратная архитектура, позволяющая производить вычисления с использованием графических процессоров NVIDIA, поддерживающих технологию GPGPU (произвольных вычислений на видеокартах). Архитектура CUDA впервые появились на рынке с выходом чипа NVIDIA восьмого поколения - G80 и присутствует во всех последующих сериях графических чипов, которые используются в семействах ускорителей GeForce, Quadro и Tesla. (с) Wikipedia.org

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



При этом существует разделение на 3 уровня:



Grid – ядро. Содержит одно/двух/трехмерный массив блоков.

Block – содержит в себе множество потоков (thread). Потоки разных блоков между собой взаимодействовать не могут. Для чего нужно было вводить блоки? Каждый блок по сути отвечает за свою подзадачу. Например, большое изображение (которое является матрицей) можно разбить на несколько более мелких частей (матриц) и параллельно работать с каждой частью изображения.

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

Warp – это объединение взаимодействующих между собой потоков, для всех современных GPU размер Warp’а равен 32. Далее идет half-warp , являющийся половинкой warp’a, т.к. обращение к памяти обычно идет раздельно для первой и второй половины warp’a.

Как можно заметить, данная архитектура отлично подходит для распараллеливания задач. И хотя программирование ведется на языке Си с некоторыми ограничениями, на деле не все так просто, т.к. не все можно распараллелить. Нет же и стандартных функций для генерации случайных чисел (или инициализации), все это приходится реализовывать отдельно. И хотя готовых вариантов имеется в достаточном количестве, радости все это не приносит. Возможность использования рекурсии появилась сравнительно недавно.

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

Всего было проведено 4 теста:

1024 элемента в каждом массиве:



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

4096 элементов в каждом массиве:



И вот уже видно, что видеокарта в 3 раза быстрее производит операции над массивами, чем процессор. Более того, время выполнения данного теста на видеокарте не увеличилось (незначительное уменьшение времени можно сослать на погрешность).

Теперь 12288 элементов в каждом массиве:



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

И последний тест – 36864 элемента в каждом массиве:



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

Если же и дальше усложнять вычисления, то видеокарта выигрывает все больше и больше. Хоть и пример несколько утрированный, но в целом ситуацию показывает наглядно. Но как упоминалось выше, не все можно распараллелить. Например, вычисление числа Пи. Существуют лишь примеры, написанные посредством метода Monte Carlo, но точность вычислений составляет 7 знаков после запятой, т.е. обычный float. Для того, чтобы увеличить точность вычислений необходима длинная арифметика, а вот тут то и наступают проблемы, т.к. эффективно это реализовать очень и очень сложно. В интернете найти примеров, использующих CUDA и рассчитывающих число Пи до 1 миллиона знаков после запятой мне не удалось. Были предприняты попытки написать такое приложение, но самый простой и эффективный метод расчета числа Пи – это алгоритм Брента - Саламина или формула Гаусса. В известном SuperPI скорее всего (судя по скорости работы и количеству итераций) используется формула Гаусса. И, судя по
тому, что SuperPI однопоточный, отсутствию примеров под CUDA и провалу моих попыток, эффективно распараллелить подсчет Pi невозможно.

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



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

В тестировании использовалось 3 различных видеофайла:

*История создания фильма Аватар - 1920x1080, MPEG4, h.264.
*Серия "Lie to me" - 1280х720, MPEG4, h.264.
*Серия "В Филадельфии всегда солнечно" - 624х464, xvid.

Контейнер и размер первых двух файлов был.mkv и 1,55 гб, а последнего - .avi и 272 мб.

Начнем с весьма нашумевшего и популярного продукта – Badaboom . Использовалась версия – 1.2.1.74 . Стоимость программы составляет $29.90 .



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

Для начала рассмотрим насколько быстро и качественно кодируется видео «в себя же», т.е. в то же разрешение и примерно тот же размер. Скорость будем измерять в fps, а не в затраченном времени – так удобнее и сравнивать, и подсчитывать сколько будет сжиматься видео произвольной продолжительности. Т.к. сегодня мы рассматриваем технологию «зеленых», то и графики будут соответствующие -)



Скорость кодирования напрямую зависит от качества, это очевидно. Стоит отметить, что легкое разрешение (назовем его традиционно – SD) не проблема для Badaboom – скорость кодирования в 5,5 раз превысила исходный (24 fps) фреймрейт видео. Да и даже тяжелый 1080p видеоролик программа преобразует в реальном времени. Стоит отметить, что качество итогового видео очень близко к исходному видеоматериалу, т.е. кодирует Badaboom весьма и весьма качественно.

Но обычно перегоняют видео в более низкое разрешение, посмотрим как обстоят дела в этом режиме. При снижении разрешения снижался и битрейт видео. Он составлял 9500 кбит/с для 1080p выходного файла, 4100 кбит/с для 720 p и 2400 кбит/с для 720х404. Выбор сделан исходя из разумного соотношения размер/качество.



Комментарии излишни. Если делать из 720p рип до обычного SD качества, то на перекодирование фильма длительностью 2 часа уйдет около 30 минут. И при этом загрузка процессора будет незначительной, можно заниматься своими делами не ощущая дискомфорта.

А что если перегнать видео в формат для мобильного устройства? Для этого выберем профиль iPhone (битрейт 1 мбит/с, 480х320) и посмотрим на скорость кодирования:



Нужно ли что-то говорить? Двухчасовой фильм в обычном качестве для iPhone перекодируется менее чем за 15 минут. С HD качеством сложнее, но все равно весьма быстро. Главное, что качество выходного видеоматериала остается на довольно высоком уровне при просмотре на дисплее телефона.

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

Следующим на очереди у нас будет Super LoiLoScope . За обычную его версию просят 3 280 рублей , а за touch версию, поддерживающую сенсорное управление в Windows 7, просят аж 4 440 рублей . Попробуем разобраться за что разработчик хочет таких денег и зачем видеоредактору поддержка multitouch. Использовалась последняя версия – 1.8.3.3 .

Описать интерфейс программы словами довольно сложно, поэтому я решил снять небольшой видеоролик. Сразу скажу, что, как и все видеоконвертеры под CUDA, ускорение средствами GPU поддерживается только для вывода видео в MPEG4 с кодеком h.264.

Во время кодирования загрузка процессора составляет 100%, однако дискомфорта это не вызывает. Браузер и другие не тяжелые приложения не тормозят.

Теперь перейдем к производительности. Для начала все тоже самое, что и с Badaboom – перекодирование видео в аналогичное по качеству.



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



Ого, а вот тут LoiloScope обходит Badaboom в 2,5 раза. При этом можно запросто параллельно резать и кодировать другое видео, читать новости и даже смотреть кино, причем даже FullHD проигрываются без проблем, хоть загрузка процессора и максимальна.

Теперь же попробуем сделать видео для мобильного устройства, профиль назовем так же, как он назывался в Badaboom – iPhone (480x320, 1 мбит/с):



Никакой ошибки нет. Все перепроверялось несколько раз, каждый раз результат был аналогичным. Скорее всего, это происходит по той простой причине, что SD файл записан с другим кодеком и в другом контейнере. При перекодировании видео сначала декодируется, разбивается на матрицы определенного размера, сжимается. ASP декодер, использующийся в случае с xvid, медленнее, чем AVC (для h.264) при параллельном декодировании. Однако и 192 fps – это в 8 раз быстрее, чем скорость исходного видео, серия длительностью 23 минуты сжимается менее чем за 4 минуты. Ситуация повторялась и с другими файлами, пережатыми в xvid/DivX.

LoiloScope оставил о себе только приятные впечатления – интерфейс, несмотря на свою необычность, удобный и функциональный, а скорость работы выше всяких похвал. Несколько расстраивает относительно бедный функционал, но в зачастую при простом монтаже требуется лишь немного подкорректировать цвета, сделать плавные переходы, наложить текст, а с этим LoiloScope отлично справляется. Несколько пугает и цена – более $100 за обычную версию нормально для зарубежья, но нам такие цифры пока кажутся несколько дикими. Хотя, признаюсь, что если бы я, например, часто снимал и монтировал домашнее видео, то возможно и задумался над покупкой. Заодно, кстати, проверил возможность редактирования HD (а точнее AVCHD) контента прямо из видеокамеры без предварительного конвертирования в другой формат, у LoiloScope никаких проблем с файлами типа.mts не выявлено.

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

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

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

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

Итак, разберемся с терминологией CUDA:


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

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

1. Хост выделяет нужное количество памяти на устройстве.
2. Хост копирует данные из своей памяти в память устройства.
3. Хост стартует выполнение определенных ядер на устройстве.
4. Устройство выполняет ядра.
5. Хост копирует результаты из памяти устройства в свою память.
Естественно, для наибольшей эффективности использования GPU нужно чтобы соотношение времени, потраченного на работу ядер, к времени, потраченному на выделение памяти и перемещение данных, было как можно больше.

Ядра

Рассмотрим более детально процесс написания кода для ядер и их запуска. Важный принцип - ядра пишутся как (практически) обычные последовательные программы - то-есть вы не увидите создания и запуска потоков в коде самих ядер. Вместо этого, для организации параллельных вычислений GPU запустит большое количество копий одного и того же ядра в разных потоках - а точнее, вы сами говорите сколько потоков запустить. И да, возвращаясь к вопросу эффективности использования GPU - чем больше потоков вы запускаете (при условии что все они будут выполнять полезную работу) - тем лучше.
Код для ядер отличается от обычного последовательного кода в таких моментах:
1. Внутри ядер вы имеете возможность узнать «идентификатор» или, проще говоря, позицию потока, который сейчас выполняется - используя эту позицию мы добиваемся того, что одно и то же ядро будет работать с разными данными в зависимости от потока, в котором оно запущено. Кстати, такая организация параллельных вычислений называется SIMD (Single Instruction Multiple Data) - когда несколько процессоров выполняют одновременно одну и ту же операцию но на разных данных.
2. В некоторых случаях в коде ядра необходимо использовать различные способы синхронизации.
Каким же образом мы задаем количество потоков, в которых будет запущено ядро? Поскольку GPU это все таки Graphics Processing Unit, то это, естественно, повлияло на модель CUDA, а именно на способ задания количества потоков:
• Сначала задаются размеры так называемой сетки (grid), в 3D координатах: grid_x, grid_y, grid_z . В результате, сетка будет состоять из grid_x*grid_y*grid_z блоков.
• Потом задаются размеры блока в 3D координатах: block_x, block_y, block_z . В результате, блок будет состоять из block_x*block_y*block_z потоков. Итого, имеем grid_x*grid_y*grid_z*block_x*block_y*block_z потоков. Важное замечание - максимальное количество потоков в одном блоке ограничено и зависит от модели GPU - типичны значения 512 (более старые модели) и 1024 (более новые модели).
• Внутри ядра доступны переменные threadIdx и blockIdx с полями x, y, z - они содержат 3D координаты потока в блоке и блока в сетке соответственно. Также доступны переменные blockDim и gridDim с теми же полями - размеры блока и сетки соответственно.
Как видите, данный способ запуска потоков действительно подходит для обработки 2D и 3D изображений: например, если нужно определенным образом обработать каждый пиксел 2D либо 3D изображения, то после выбора размеров блока (в зависимости от размеров картинки, способа обработки и модели GPU) размеры сетки выбираются такими, чтобы было покрыто все изображение, возможно, с избытком - если размеры изображения не делятся нацело на размеры блока.

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

Довольно теории, время писать код. Инструкции по установке и конфигурации CUDA для разных ОС - docs.nvidia.com/cuda/index.html . Также, для простоты работы с файлами изображений будем использовать OpenCV , а для сравнения производительности CPU и GPU - OpenMP .
Задачу поставим довольно простую: конвертация цветного изображения в оттенки серого . Для этого, яркость пиксела pix в серой шкале считается по формуле: Y = 0.299*pix.R + 0.587*pix.G + 0.114*pix.B .
Сначала напишем скелет программы:


main.cpp

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



Тут все довольно очевидно - читаем файл с изображением, подготавливаем указатели на цветное и в оттенках серого изображение, запускаем вариант
с OpenMP и вариант с CUDA, замеряем время. Функция prepareImagePointers имеет следующий вид:


prepareImagePointers

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



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


openMP.hpp

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



Все довольно прямолинейно - мы всего лишь добавили директиву omp parallel for к однопоточному коду - в этом вся красота и мощь OpenMP. Я пробовал поиграться с параметром schedule , но получалось только хуже, чем без него.
Наконец, переходим к CUDA. Тут распишем более детально. Сначала нужно выделить память под входные данные, переместить их с CPU на GPU и выделить память под выходные данные:


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

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



Стоит обратить внимание на стандарт именования переменных в CUDA - данные на CPU начинаются с h_ (h ost), данные да GPU - с d_ (d evice). checkCudaErrors - макрос, взят с github-репозитория Udacity курса. Имеет следующий вид:


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

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



cudaMalloc - аналог malloc для GPU, cudaMemcpy - аналог memcpy , имеет дополнительный параметр в виде enum-а, который указывает тип копирования: cudaMemcpyHostToDevice, cudaMemcpyDeviceToHost, cudaMemcpyDeviceToDevice.
Далее необходимо задать размеры сетки и блока и вызвать ядро, не забыв измерить время:


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

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



Обратите внимание на формат вызова ядра - kernel_name<<>> . Код самого ядра также не очень сложный:


rgba_to_grayscale_simple

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



Здесь мы вычисляем координаты y и x обрабатываемого пиксела, используя ранее описанные переменные threadIdx , blockIdx и blockDim , ну и выполняем конвертацию. Обратите внимание на проверку if (x>=numCols || y>=numRows) - так как размеры изображения не обязательно будут делится нацело на размеры блоков, некоторые блоки могут «выходить за рамки» изображения - поэтому необходима эта проверка. Также, функция ядра должна помечаться спецификатором __global__ .
Последний шаг - cкопировать результат назад с GPU на CPU и освободить выделенную память:


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

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



Кстати, CUDA позволяет использовать C++ компилятор для host-кода - так что запросто можно написать обертки для автоматического освобождения памяти.
Итак, запускаем, измеряем (размер входного изображения - 10,109 × 4,542):
OpenMP time (ms):45 CUDA time simple (ms): 43.1941
Конфигурация машины, на которой проводились тесты:


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

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



Получилось как-то не очень впечатляюще:) А проблема все та же - слишком мало работы выполняется над каждым пикселом - мы запускаем тысячи потоков, каждый из которых отрабатывает практически моментально. В случае с CPU такой проблемы не возникает - OpenMP запустит сравнительно малое количество потоков (8 в моем случае) и разделит работу между ними поровну - таким образом процессоры будет занят практически на все 100%, в то время как с GPU мы, по сути, не используем всю его мощь. Решение довольно очевидное - обрабатывать несколько пикселов в ядре. Новое, оптимизированное, ядро будет выглядеть следующим образом:


rgba_to_grayscale_optimized

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


Здесь не все так просто как с предыдущим ядром. Если разобраться, теперь каждый поток будет обрабатывать elemsPerThread пикселов, причем не подряд, а с расстоянием в WARP_SIZE между ними. Что такое WARP_SIZE, почему оно равно 32, и зачем обрабатывать пиксели пободным образом, будет более детально рассказано в следующих частях, сейчас только скажу что этим мы добиваемся более эффективной работы с памятью. Каждый поток теперь обрабатывает elemsPerThread пикселов с расстоянием в WARP_SIZE между ними, поэтому x-координата первого пиксела для этого потока исходя из его позиции в блоке теперь рассчитывается по несколько более сложной формуле чем раньше.
Запускается это ядро следующим образом:


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

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



Количество блоков по x-координате теперь рассчитывается как numCols / (threadNum*elemsPerThread) + 1 вместо numCols / threadNum + 1 . В остальном все осталось так же.
Запускаем:
OpenMP time (ms):44 CUDA time simple (ms): 53.1625 CUDA time optimized (ms): 15.9273
Получили прирост по скорости в 2.76 раза (опять же, не учитывая время на операции с памятью) - для такой простой проблемы это довольно неплохо. Да-да, эта задача слишком простая - с ней достаточно хорошо справляется и CPU. Как видно из второго теста, простая реализация на GPU может даже проигрывать по скорости реализации на CPU.
На сегодня все, в следующей части рассмотрим аппаратное обеспечение GPU и основные шаблоны параллельной коммуникации.
Весь исходный код доступен на bitbucket .

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

В развитии современных процессоров намечается тенденция к постепенному увеличению количества ядер, что повышает их возможности в параллельных вычислениях. Однако уже давно имеются 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 это один из этапов в процессе данного объединения.


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