Знакомство с программно-аппаратной архитектурой CUDA


Оглавление (нажмите, чтобы открыть):

Технология CUDA

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

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

Кстати, заметим. что современные локомотивы вычислительной техники — самые мощные суперкомпьютеры, такие как TITANIUM, построены на «детских самокатах» — видеокартах игровых компьютеров, имеющих графические процессоры с тысячами параллельно работающих ядер.

Программная архитектура CUDA

Первоначальная версия CUDA SDK была представлена 15 февраля 2007 года. В основе CUDA API лежит язык Си с некоторыми ограничениями. Для успешной трансляции кода на этом языке, в состав CUDA SDK входит собственный Си-компилятор командной строки nvcc компании Nvidia. Компилятор nvcc создан на основе открытого компилятора Open64 и предназначен для трансляции host-кода (главного, управляющего кода) и device-кода (аппаратного кода) (файлов с расширением .cu) в объектные файлы, пригодные в процессе сборки конечной программы или библиотеки в любой среде программирования, например в NetBeans.

Использует grid-модель памяти, кластерное моделирование потоков и SIMD инструкции. Применим в основном для высокопроизводительных графических вычислений и разработок NVIDIA-совместимого графического API. Включена возможность подключения к приложениям, использующим OpenGL и Microsoft Direct3D v9. Создан в версиях для Linux, Mac OS X, Windows.

22 марта 2010 года nVidia выпустила CUDA Toolkit 3.0, который содержал поддержку OpenCL.

Оборудование CUDA

Первая серия оборудования, поддерживающая CUDA SDK, G8x, имела 32-битный векторный процессор одинарной точности, использующий CUDA SDK как API (CUDA поддерживает тип double языка Си, однако сейчас его точность понижена до 32-битного с плавающей запятой). Более поздние процессоры GT200 имеют поддержку 64-битной точности (только для SFU), но производительность значительно хуже, чем для 32-битной точности (из-за того что SFU всего 2 на каждый потоковый мультипроцессор, а скалярных процессоров 8). Графический процессор организует аппаратную многопоточность, что позволяет задействовать все ресурсы графического процессора. Таким образом, открывается перспектива переложить функции физического ускорителя на графический ускоритель (пример реализации — nVidia PhysX). Также открываются широкие возможности использования графического оборудования компьютера для выполнения сложных неграфических вычислений: например, в вычислительной биологии и в иных отраслях науки.

Преимущества CUDA

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

  • Интерфейс программирования приложений CUDA (CUDA API) основан на стандартном языке программирования Си с некоторыми ограничениями. По мнению разработчиков, это должно упростить и сгладить процесс изучения архитектуры CUDA
  • Разделяемая между потоками память (shared memory) размером в 16 Кб может быть использована под организованный пользователем кэш с более широкой полосой пропускания, чем при выборке из обычных текстур
  • Более эффективные транзакции между памятью центрального процессора и видеопамятью
  • Полная аппаратная поддержка целочисленных и побитовых операций

Ограничения CUDA

Все функции, выполнимые на устройстве, не поддерживают рекурсии (в версии CUDA Toolkit 3.1 поддерживает указатели и рекурсию) и имеют некоторые другие ограничения

Архитектуру CUDA поддерживает и развивает только производитель NVidia

Поддерживаемые GPU и графические ускорители CUDA

Перечень устройств от производителя оборудования Nvidia с заявленной полной поддержкой технологии CUDA приведён на официальном сайте Nvidia: CUDA-Enabled GPU Products (англ.).

Фактически же, в настоящее время на рынке аппаратных средств для ПК поддержку технологии CUDA обеспечивают следующие периферийные устройства:

версия CUDA

Графические процессоры (GPU)

G86, G84, G98, G96, G96b, G94, G94b, G92, G92b

GT218, GT216, GT215

GT200, GT200b

GF100, GF110

GF108, GF106, GF104, GF114, GF116

Nvidia GeForce для десктопов

GeForce GTX 590, GTX 580, GTX 570, GTX 560 Ti, GTX 560, GTX 550 Ti, GTX 520, GTX 480, GTX 470, GTX 465, GTX 460, GTS 450, GTX 295, GTX 285, GTX 280, GTX 275, GTX 260, GTS 250, GT 240, GT 220, GT 210, GTS 150, GT 130, GT 120, G100, 9800 GX2, 9800 GTX+, 9800 GTX, 9800 GT, 9600 GSO, 9600 GT, 9500 GT, 9400 GT, 9400 mGPU, 9300 mGPU, 8800 GTS 512, 8800 GT, 8600 GTS, 8600 GT, 8500 GT, 8400 GS

Nvidia GeForce для мобильных устройств

GeForce GTX 580M, GTX 570M, GTX 560M, GT 555M, GT 540M, GT 525M, GT 520M, GTX 485M, GTX 480M, GTX 470M, GTX 460M, GT 445M, GT 435M, GT 425M, GT 420M, GT 415M, GTX 285M, GTX 280M, GTX 260M, GTS 360M, GTS 350M, GTS 160M, GTS 150M, GT 335M, GT 330M, GT 325M, GT 240M, GT 130M, G110M, G105M, 310M, 305M, 9800M GTX, 9800M GT, 9800M GTS, 9700M GTS, 9700M GT, 9650M GS, 9600M GT, 9600M GS, 9500M GS, 9500M G, 9300M GS, 9300M G, 9200M GS, 9100M G, 8800M GTS, 8700M GT, 8600M GT, 8600M GS, 8400M GT, 8400M GS

Tesla C2050/C2070, Tesla M2050/M2070/M2090, Tesla S2050, Tesla S1070, Tesla M1060, Tesla C1060, Tesla C870, Tesla D870, Tesla S870

Nvidia Quadro для десктопов

Quadro 6000, Quadro 5000, Quadro 4000, Quadro 2000, Quadro 600, Quadro FX 5800, Quadro FX 5600, Quadro FX 4800, Quadro FX 4700 X2, Quadro FX 4600, Quadro FX 3700, Quadro FX 1700, Quadro FX 570, Quadro FX 470, Quadro FX 380 Low Profile, Quadro FX 370, Quadro FX 370 Low Profile, Quadro CX, Quadro NVS 450, Quadro NVS 420, Quadro NVS 290, Quadro Plex 2100 D4, Quadro Plex 2200 D2, Quadro Plex 2100 S4, Quadro Plex 1000 Model IV

Nvidia Quadro для мобильных устройств

Quadro 5010M, Quadro 5000M, Quadro 4000M, Quadro 3000M, Quadro 2000M, Quadro 1000M, Quadro FX 3800M, Quadro FX 3700M, Quadro FX 3600M, Quadro FX 2800M, Quadro FX 2700M, Quadro FX 1800M, Quadro FX 1700M, Quadro FX 880M, Quadro FX 770M, Quadro FX 570M, Quadro FX 380M, Quadro FX 370M, Quadro FX 360M, Quadro NVS 5100M, Quadro NVS 3100M, Quadro NVS 2100M, Quadro NVS 320M, Quadro NVS 160M, Quadro NVS 150M, Quadro NVS 140M, Quadro NVS 135M, Quadro NVS 130M

Модели Tesla C1060, Tesla S1070, Tesla C2050/C2070, Tesla M2050/M2070, Tesla S2050 позволяют производить вычисления на GPU с двойной точностью.

Особенности и спецификации различных версий

Пример программирования CUDA

Этот пример кода на C загрузки текстур из изображения в массив на GPU:

Пример программы на языке Python, перемножающий матрицы средствами GPU. Взаимодействие идёт с использованием PyCUDA

CUDA как предмет в вузах

По состоянию на декабрь 2009 года, программная модель CUDA преподается в 269 университетах по всему миру. В России обучающие курсы по CUDA читаются в Санкт-Петербургском политехническом университете, Ярославском государственном университете им. П.Г. Демидова, Московском, Нижегородском, Санкт-Петербургском, Тверском, Казанском, Новосибирском, Новосибирском государственном техническом университете Омском и Пермском государственных университетах, Международном университете природы общества и человека “Дубна”, Объединённом институте ядерных исследований, Московском институте электронной техники, Ивановском государственном энергетическом университете, БГТУ им. В. Г. Шухова, МГТУ им. Баумана, РХТУ им.Менделеева, Российском научном центре “Курчатовский институт”, Межрегиональном суперкомпьютерном центре РАН, Таганрогском технологическом институте (ТТИ ЮФУ). Кроме того, в декабре 2009 года было объявлено о начале работы первого в России научно-образовательного центра “Параллельные вычисления”, расположенного в городе Дубна, в задачи которого входят обучение и консультации по решению сложных вычислительных задач на GPU.

На Украине курсы по CUDA читаются в Киевском институте системного анализа.

Официальные ресурсы CUDA

CUDA Zone (рус.) — официальный сайт CUDA

CUDA GPU Computing (англ.) — официальные веб-форумы, посвящённые вычислениям CUDA

Технология CUDA

Введение

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

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

ИСТОРИЯ

Направление вычислений эволюционирует от “централизованной обработки данных” на центральном процессоре до “совместной обработки” на CPU и GPU. Для реализации новой вычислительной парадигмы компания NVIDIA изобрела архитектуру параллельных вычислений CUDA, на данный момент представленную в графических процессорах GeForce, ION, Quadro и Tesla и обеспечивающую необходимую базу разработчикам ПО.

Говоря о потребительском рынке, стоит отметить, что почти все основные приложения для работы с видео уже оборудованы, либо будут оснащены поддержкой CUDA-ускорения, включая продукты от Elemental Technologies, MotionDSP и LoiLo.

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

На финансовом рынке компании Numerix и CompatibL анонсировали поддержку CUDA в новом приложении анализа риска контрагентов и достигли ускорения работы в 18 раз. Numerix используется почти 400 финансовыми институтами.

Показателем роста применения CUDA является также рост использования графических процессоров Tesla в GPU вычислениях. На данный момент более 700 GPU кластеров установлены по всему миру в компаниях из списка Fortune 500, таких как Schlumberger и Chevron в энергетическом секторе, а также BNP Paribas в секторе банковских услуг.

Благодаря ожидаемым в ближайшем будущем системам Microsoft Windows 7 и Apple Snow Leopard, вычисления на GPU займут свои позиции в секторе массовых решений. В этих новых операционных системах GPU предстанет не только графическим процессором, но также и универсальным процессором для параллельных вычислений, работающим с любым приложением.

Состав NV >CUDA включает два API: высокого уровня (CUDA Runtime API) и низкого (CUDA Driver API), хотя в одной программе одновременное использование обоих невозможно, нужно использовать или один или другой. Высокоуровневый работает “сверху” низкоуровневого, все вызовы runtime транслируются в простые инструкции, обрабатываемые низкоуровневым Driver API. Но даже “высокоуровневый” API предполагает знания об устройстве и работе видеочипов NVIDIA, слишком высокого уровня абстракции там нет.

Есть и ещё один уровень, даже более высокий — две библиотеки:

CUBLAS — CUDA вариант BLAS (Basic Linear Algebra Subprograms), предназначенный для вычислений задач линейной алгебры и использующий прямой доступ к ресурсам G PU;

CUFFT — CUDA вариант библиотеки Fast Fourier Transform для расчёта быстрого преобразования Фурье, широко используемого при обработке сигналов. Поддерживаются следующие типы преобразований: complex-complex (C2C), real-complex (R2C) и complex-real (C2R).

Рассмотрим эти библиотеки подробнее. CUBLAS — это переведённые на язык CUDA стандартные алгоритмы линейной алгебры, на данный момент поддерживается только определённый набор основных функций CUBLAS. Библиотеку очень легко использовать: нужно создать матрицу и векторные объекты в памяти видеокарты, заполнить их данными, вызвать требуемые функции CUBLAS, и загрузить результаты из видеопамяти обратно в системную. CUBLAS содержит специальные функции для создания и уничтожения объектов в памяти GPU, а также для чтения и записи данных в эту память. Поддерживаемые функции BLAS: уровни 1, 2 и 3 для действительных чисел, уровень 1 CGEMM для комплексных. Уровень 1 — это векторно-векторные операции, уровень 2 — векторно-матричные операции, уровень 3 — матрично-матричные операции.

CUFFT — CUDA вариант функции быстрого преобразования Фурье — широко используемой и очень важной при анализе сигналов, фильтрации и т.п. CUFFT предоставляет простой интерфейс для эффективного вычисления FFT на видеочипах производства NVIDIA без необходимости в разработке собственного варианта FFT для GPU. CUDA вариант FFT поддерживает 1D, 2D, и 3D преобразования комплексных и действительных данных, пакетное исполнение для нескольких 1D трансформаций в параллели, размеры 2D и 3D трансформаций могут быть в пределах [2, 16384], для 1D поддерживается размер до 8 миллионов элементов.

Среда программирования

В состав CUDA входят runtime библиотеки:

общая часть, предоставляющая встроенные векторные типы и подмножества вызовов RTL, поддерживаемые на CPU и GPU;

CPU-компонента, для управления одним или несколькими GPU;


GPU-компонента, предоставляющая специфические функции для GPU.

Основной процесс приложения CUDA работает на универсальном процессоре (host), он запускает несколько копий процессов kernel на видеокарте. Код для CPU делает следующее: инициализирует GPU, распределяет память на видеокарте и системе, копирует константы в память видеокарты, запускает несколько копий процессов kernel на видеокарте, копирует полученный результат из видеопамяти, освобождает память и завершает работу.

В качестве примера для понимания приведем CPU код для сложения векторов, представленный в CUDA:

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

Программы на CUDA могут взаимодействовать с графическими API: для рендеринга данных, сгенерированных в программе, для считывания результатов рендеринга и их обработки средствами CUDA (например, при реализации фильтров постобработки). Для этого ресурсы графических API могут быть отображены (с получением адреса ресурса) в пространство глобальной памяти CUDA. Поддерживаются следующие типы ресурсов графических API: Buffer Objects (PBO / VBO) в OpenGL, вершинные буферы и текстуры (2D, 3D и кубические карты) D irect3D9.

Стадии компиляции CUDA-приложения:

Файлы исходного кода на CUDA C компилируются при помощи программы NVCC, которая является оболочкой над другими инструментами, и вызывает их: cudacc, g++, cl и др. NVCC генерирует: код для центрального процессора, который компилируется вместе с остальными частями приложения, написанными на чистом Си, и объектный код PTX для видеочипа. Исполнимые файлы с кодом на CUDA в обязательном порядке требуют наличия библиотек CUDA runtime library (cudart) и CUDA core library (cuda).

Практические выводы

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

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

Определение функций

Перед функциями в .cu файле могут стоять следующие “модификаторы”:

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

__global__ — Эта функция — начало вашего вычислительного ядра. Выполняется на видеокарте, но запускается только с хоста.

__host__ — Выполняется и запускается только с хоста (т.е. обычная функция C++). Если вы перед функцией укажите например __host__ и __device__ — то будут скомпилированы 2 версии функции (не все комбинации допустимы).

__device__ — означает что переменная находится в глобальной памяти видеокарты (т.е. которой там 512-1024Мб и выше). Очень медленная память по меркам вычислений на видеокарте(хоть и быстрее памяти центрального процессора в несколько раз), рекомендуется использовать как можно реже. В этой памяти данные сохраняются между вызовами разных вычислительных ядер. Данные сюда можно записывать и читать из host-части с помощью

cudaMemcpy(device_variable, host_variable, size, cudaMemcpyHostToDevice); //cudaMemcpyDeviceToHost — в обратную сторону

__constant__ — задает переменную в константной памяти. Следует обратить внимание, что значения для констант нужно загружать функцией cudaMemcpyToSymbol. Константы доступны из всех тредов, скорость работы сравнима с регистрами(когда в кеш попадает).

__shared__ — задает переменную в общей памяти блока тредов (т.е. и значение будет общее на всех). Тут нужно подходить с осторожностью — компилятор агрессивно оптимизирует доступ сюда(можно придушить модификатором volatile), можно получать race condition, нужно использовать __syncthreads(); чтобы данные гарантированно записались. Shared memory разделена на банки, и когда 2 потока одновременно пытаются обратиться к одному банку, возникает bank conflict и падает скорость.

Все локальные переменные которые вы определеили в ядре (__device__) — в регистрах, самая высокая скорость доступа.

Как поток узнает над чем ему работать

Допустим, надо сделать какую-то операцию над картинкой 200×200. Картинка разбивается на куски 10×10, и на каждый пиксел такого кусочка запускаем по потоку. Выглядить это будет так:

dim3 th reads(10, 10);//размер квардатика, 10*10

dim3 grid(20, 20);//сколько квадратиков нужно чтобы покрыть все изображение

your_kernel >>(image, 200,200);//Эта строка запустит 40’000 потоков (не одновременно, одновременно работать будет 200-2000 потоков примерно).

В отличии от Brook+ от AMD, где мы сразу определяем какому потоку над какими данными работать, в CUDA все не так: передаваеиые kernel-у параметры одинаковые для всех потоков, и поток должен сам получить данные для себя, чтобы сделать это, потоку нужно вычислить, в каком месте изображения он находится. В этом помогают магические переменные blockDim, blockIdx.

const int ix = blockDim.x * blockIdx.x + threadIdx.x;

const int iy = blockDim.y * blockIdx.y + threadIdx.y;

В ix и iy — координаты, с помощью которых можно получить исходные данные из массива image, и записать результат работы.

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

•Как можно меньше используйте __global__ память.

•При работе с __shared__ памятью избегайте конфлктов банков (впрочем многие задачи могут быть решены без shared памяти).

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

•Используйте как можно меньше памяти. Чем меньше памяти вы используете, тем агрессивнее компилятор и железо смогут запускать ваш kernel (например он может взять 100 тредов, и используя в 100 больше регистров запустить одновременно на одном MP, радикально уменьшая задержки)

Примеры Программ на CUDA C++

A[n][n] – матрица размерности n x n;

b[n] – вектор, состоящий из n элементов.

c[n] – вектор из n элементов.

//Последовательный алгоритм умножения матрицы на вектор

Теперь рассмотрим решение этой задачи на видеокарте. Следующий код иллюстрирует пример вызова функции CUDA:

// обычные массивы в оперативной памяти

h_a = new float[Size*Size];

h_b = new float[Size];

h_c = new float[Size];

// указатели на массивы в видеопамяти

// копирование из оперативной памяти в видеопамять

CUDA_SAFE_CALL(cudaMemcpy(d_a, h_a, sizeof(float)*Size*Size,

CUDA_SAFE_CALL(cudaMemcpy(d_b, h_b, sizeof(float)*Size,

// установка количества блоков

dim3 grid((Size+255)/256, 1, 1);

// установка количества потоков в блоке

dim3 threads(256, 1, 1);

MatrVectMul >> (d_c, d_a, d_b,Size);

// копирование из видеопамяти в оперативную память

CUDA_SAFE_CALL(cudaMemcpy(h_c, d_c, sizeof(float)*Size,

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

1. инициализация CUDA

2. выделение видеопамяти для хранения данных программы

3. копирование необходимых для работы функции данных из оперативной памяти в видеопамять

4. вызов функции CUDA

5. копирование возвращаемых данных из видеопамяти в оперативную

6. освобождение видеопамяти

Пример функции, исполнимой на видеокарте

extern «C» __global__ void MatrVectMul(float *d_c, float *d_a, float *d_b, int Size)

int i = blockIdx.x*blockDim.x+threadIdx.x;

threadIdx.x – идентификатор потока в блоке по координате x,

blockIdx.x – идентификатор блока в гриде по координате x,

blockDim.x – количество потоков в одном блоке.

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


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

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

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

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

Замер времени в CUDA

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

unsigned int timer;

… // код, время исполнения которого необходимо замерить

Знакомство с программно-аппаратной архитектурой CUDA

Устройства для превращения персональных компьютеров в маленькие суперкомпьютеры известны довольно давно. Ещё в 80-х годах прошлого века на рынке предлагались так называемые транспьютеры, которые вставлялись в распространенные тогда слоты расширения ISA. Первое время их производительность в соответствующих задачах впечатляла, но затем рост быстродействия универсальных процессоров ускорился, они усилили свои позиции в параллельных вычислениях, и смысла в транспьютерах не осталось. Хотя подобные устройства существуют и сейчас — это разнообразные специализированные ускорители. Но зачастую сфера их применения узка и особого распространения такие ускорители не получили.

Но в последнее время эстафета параллельных вычислений перешла к массовому рынку, так или иначе связанному с трёхмерными играми. Универсальные устройства с многоядерными процессорами для параллельных векторных вычислений, используемых в 3D-графике, достигают высокой пиковой производительности, которая универсальным процессорам не под силу. Конечно, максимальная скорость достигается лишь в ряде удобных задач и имеет некоторые ограничения, но такие устройства уже начали довольно широко применять в сферах, для которых они изначально и не предназначались. Отличным примером такого параллельного процессора является процессор Cell, разработанный альянсом Sony-Toshiba-IBM и применяемый в игровой приставке Sony PlayStation 3, а также и все современные видеокарты от лидеров рынка — компаний Nvidia и AMD.

Cell мы сегодня трогать не будем, хоть он и появился раньше и является универсальным процессором с дополнительными векторными возможностями, речь сегодня не о нём. Для 3D видеоускорителей ещё несколько лет назад появились первые технологии неграфических расчётов общего назначения GPGPU (General-Purpose computation on GPUs). Ведь современные видеочипы содержат сотни математических исполнительных блоков, и эта мощь может использоваться для значительного ускорения множества вычислительно интенсивных приложений. И нынешние поколения GPU обладают достаточно гибкой архитектурой, что вместе с высокоуровневыми языками программирования и программно-аппаратными архитектурами, подобными рассматриваемой в этой статье, раскрывает эти возможности и делает их значительно более доступными.

На создание GPCPU разработчиков побудило появление достаточно быстрых и гибких шейдерных программ, которые способны исполнять современные видеочипы. Разработчики задумали сделать так, чтобы GPU рассчитывали не только изображение в 3D приложениях, но и применялись в других параллельных расчётах. В GPGPU для этого использовались графические API: OpenGL и Direct3D, когда данные к видеочипу передавались в виде текстур, а расчётные программы загружались в виде шейдеров. Недостатками такого метода является сравнительно высокая сложность программирования, низкая скорость обмена данными между CPU и GPU и другие ограничения, о которых мы поговорим далее.

Вычисления на GPU развивались и развиваются очень быстро. И в дальнейшем, два основных производителя видеочипов, Nvidia и AMD, разработали и анонсировали соответствующие платформы под названием CUDA (Compute Unified Device Architecture) и CTM (Close To Metal или AMD Stream Computing), соответственно. В отличие от предыдущих моделей программирования GPU, эти были выполнены с учётом прямого доступа к аппаратным возможностям видеокарт. Платформы не совместимы между собой, CUDA — это расширение языка программирования C, а CTM — виртуальная машина, исполняющая ассемблерный код. Зато обе платформы ликвидировали некоторые из важных ограничений предыдущих моделей GPGPU, использующих традиционный графический конвейер и соответствующие интерфейсы Direct3D или OpenGL.

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

Именно поэтому компания Nvidia выпустила платформу CUDA — C-подобный язык программирования со своим компилятором и библиотеками для вычислений на GPU. Конечно же, написание оптимального кода для видеочипов совсем не такое простое и эта задача нуждается в длительной ручной работе, но CUDA как раз и раскрывает все возможности и даёт программисту больший контроль над аппаратными возможностями GPU. Важно, что поддержка Nvidia CUDA есть у чипов G8x, G9x и GT2xx, применяемых в видеокартах Geforce серий 8, 9 и 200, которые очень широко распространены. В настоящее время выпущена финальная версия CUDA 2.0, в которой появились некоторые новые возможности, например, поддержка расчётов с двойной точностью. CUDA доступна на 32-битных и 64-битных операционных системах Linux, Windows и MacOS X.

Разница между CPU и GPU в параллельных расчётах

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

Специализированные векторные возможности (SSE2 и SSE3) для четырехкомпонентных (одинарная точность вычислений с плавающей точкой) и двухкомпонентных (двойная точность) векторов появились в универсальных процессорах из-за возросших требований графических приложений, в первую очередь. Именно поэтому для определённых задач применение GPU выгоднее, ведь они изначально сделаны для них.

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

Самое главное — эти несколько ядер мультипроцессора в GPU являются SIMD (одиночный поток команд, множество потоков данных) ядрами. И эти ядра исполняют одни и те же инструкции одновременно, такой стиль программирования является обычным для графических алгоритмов и многих научных задач, но требует специфического программирования. Зато такой подход позволяет увеличить количество исполнительных блоков за счёт их упрощения.

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

Разработчики CPU стараются добиться выполнения как можно большего числа инструкций параллельно, для увеличения производительности. Для этого, начиная с процессоров Intel Pentium, появилось суперскалярное выполнение, обеспечивающее выполнение двух инструкций за такт, а Pentium Pro отличился внеочередным выполнением инструкций. Но у параллельного выполнения последовательного потока инструкций есть определённые базовые ограничения и увеличением количества исполнительных блоков кратного увеличения скорости не добиться.

У видеочипов работа простая и распараллеленная изначально. Видеочип принимает на входе группу полигонов, проводит все необходимые операции, и на выходе выдаёт пиксели. Обработка полигонов и пикселей независима, их можно обрабатывать параллельно, отдельно друг от друга. Поэтому, из-за изначально параллельной организации работы в GPU используется большое количество исполнительных блоков, которые легко загрузить, в отличие от последовательного потока инструкций для CPU. Кроме того, современные GPU также могут исполнять больше одной инструкции за такт (dual issue). Так, архитектура Tesla в некоторых условиях запускает на исполнение операции MAD+MUL или MAD+SFU одновременно.

GPU отличается от CPU ещё и по принципам доступа к памяти. В GPU он связанный и легко предсказуемый — если из памяти читается тексель текстуры, то через некоторое время придёт время и для соседних текселей. Да и при записи то же — пиксель записывается во фреймбуфер, и через несколько тактов будет записываться расположенный рядом с ним. Поэтому организация памяти отличается от той, что используется в CPU. И видеочипу, в отличие от универсальных процессоров, просто не нужна кэш-память большого размера, а для текстур требуются лишь несколько (до 128-256 в нынешних GPU) килобайт.

Да и сама по себе работа с памятью у GPU и CPU несколько отличается. Так, не все центральные процессоры имеют встроенные контроллеры памяти, а у всех GPU обычно есть по несколько контроллеров, вплоть до восьми 64-битных каналов в чипе Nvidia GT200. Кроме того, на видеокартах применяется более быстрая память, и в результате видеочипам доступна в разы большая пропускная способность памяти, что также весьма важно для параллельных расчётов, оперирующих с огромными потоками данных.

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

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

Есть множество различий и в поддержке многопоточности. CPU исполняет 1-2 потока вычислений на одно процессорное ядро, а видеочипы могут поддерживать до 1024 потоков на каждый мультипроцессор, которых в чипе несколько штук. И если переключение с одного потока на другой для CPU стоит сотни тактов, то GPU переключает несколько потоков за один такт.

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

Вкратце можно сказать, что в отличие от современных универсальных CPU, видеочипы предназначены для параллельных вычислений с большим количеством арифметических операций. И значительно большее число транзисторов GPU работает по прямому назначению — обработке массивов данных, а не управляет исполнением (flow control) немногочисленных последовательных вычислительных потоков. Это схема того, сколько места в CPU и GPU занимает разнообразная логика:

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

Выполнение расчётов на GPU показывает отличные результаты в алгоритмах, использующих параллельную обработку данных. То есть, когда одну и ту же последовательность математических операций применяют к большому объёму данных. При этом лучшие результаты достигаются, если отношение числа арифметических инструкций к числу обращений к памяти достаточно велико. Это предъявляет меньшие требования к управлению исполнением (flow control), а высокая плотность математики и большой объём данных отменяет необходимость в больших кэшах, как на CPU.

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

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

Первые попытки применения расчётов на GPU

Видеочипы в параллельных математических расчётах пытались использовать довольно давно. Самые первые попытки такого применения были крайне примитивными и ограничивались использованием некоторых аппаратных функций, таких, как растеризация и Z-буферизация. Но в нынешнем веке, с появлением шейдеров, начали ускорять вычисления матриц. В 2003 году на SIGGRAPH отдельная секция была выделена под вычисления на GPU, и она получила название GPGPU (General-Purpose computation on GPU) — универсальные вычисления на GPU).

Наиболее известен BrookGPU — компилятор потокового языка программирования Brook, созданный для выполнения неграфических вычислений на GPU. До его появления разработчики, использующие возможности видеочипов для вычислений, выбирали один из двух распространённых API: Direct3D или OpenGL. Это серьёзно ограничивало применение GPU, ведь в 3D графике используются шейдеры и текстуры, о которых специалисты по параллельному программированию знать не обязаны, они используют потоки и ядра. Brook смог помочь в облегчении их задачи. Эти потоковые расширения к языку C, разработанные в Стэндфордском университете, скрывали от программистов трёхмерный API, и представляли видеочип в виде параллельного сопроцессора. Компилятор обрабатывал файл .br с кодом C++ и расширениями, производя код, привязанный к библиотеке с поддержкой DirectX, OpenGL или x86.

Естественно, у Brook было множество недостатков, на которых мы останавливались, и о которых ещё подробнее поговорим далее. Но даже просто его появление вызвало значительный прилив внимания тех же Nvidia и ATI к инициативе вычислений на GPU, так как развитие этих возможностей серьёзно изменило рынок в дальнейшем, открыв целый новый его сектор — параллельные вычислители на основе видеочипов.

В дальнейшем, некоторые исследователи из проекта Brook влились в команду разработчиков Nvidia, чтобы представить программно-аппаратную стратегию параллельных вычислений, открыв новую долю рынка. И главным преимуществом этой инициативы Nvidia стало то, что разработчики отлично знают все возможности своих GPU до мелочей, и в использовании графического API нет необходимости, а работать с аппаратным обеспечением можно напрямую при помощи драйвера. Результатом усилий этой команды стала Nvidia CUDA (Compute Unified Device Architecture) — новая программно-аппаратная архитектура для параллельных вычислений на Nvidia GPU, которой посвящена эта статья.

Области применения параллельных расчётов на GPU

Чтобы понять, какие преимущества приносит перенос расчётов на видеочипы, приведём усреднённые цифры, полученные исследователями по всему миру. В среднем, при переносе вычислений на GPU, во многих задачах достигается ускорение в 5-30 раз, по сравнению с быстрыми универсальными процессорами. Самые большие цифры (порядка 100-кратного ускорения и даже более!) достигаются на коде, который не очень хорошо подходит для расчётов при помощи блоков SSE, но вполне удобен для GPU.

Это лишь некоторые примеры ускорений синтетического кода на GPU против SSE-векторизованного кода на CPU (по данным Nvidia):

  • Флуоресцентная микроскопия: 12x;
  • Молекулярная динамика (non-bonded force calc): 8-16x;
  • Электростатика (прямое и многоуровневое суммирование Кулона): 40-120x и 7x.

А это табличка, которую очень любит Nvidia, показывая её на всех презентациях, на которой мы подробнее остановимся во второй части статьи, посвящённой конкретным примерам практических применений CUDA вычислений:

Как видите, цифры весьма привлекательные, особенно впечатляют 100-150-кратные приросты. В следующей статье, посвящённой CUDA, мы подробно разберём некоторые из этих цифр. А сейчас перечислим основные приложения, в которых сейчас применяются вычисления на GPU: анализ и обработка изображений и сигналов, симуляция физики, вычислительная математика, вычислительная биология, финансовые расчёты, базы данных, динамика газов и жидкостей, криптография, адаптивная лучевая терапия, астрономия, обработка звука, биоинформатика, биологические симуляции, компьютерное зрение, анализ данных (data mining), цифровое кино и телевидение, электромагнитные симуляции, геоинформационные системы, военные применения, горное планирование, молекулярная динамика, магнитно-резонансная томография (MRI), нейросети, океанографические исследования, физика частиц, симуляция свёртывания молекул белка, квантовая химия, трассировка лучей, визуализация, радары, гидродинамическое моделирование (reservoir simulation), искусственный интеллект, анализ спутниковых данных, сейсмическая разведка, хирургия, ультразвук, видеоконференции.

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

Возможности Nvidia CUDA

Технология CUDA — это программно-аппаратная вычислительная архитектура Nvidia, основанная на расширении языка Си, которая даёт возможность организации доступа к набору инструкций графического ускорителя и управления его памятью при организации параллельных вычислений. CUDA помогает реализовывать алгоритмы, выполнимые на графических процессорах видеоускорителей Geforce восьмого поколения и старше (серии Geforce 8, Geforce 9, Geforce 200), а также Quadro и Tesla.

Хотя трудоёмкость программирования GPU при помощи CUDA довольно велика, она ниже, чем с ранними GPGPU решениями. Такие программы требуют разбиения приложения между несколькими мультипроцессорами подобно MPI программированию, но без разделения данных, которые хранятся в общей видеопамяти. И так как CUDA программирование для каждого мультипроцессора подобно OpenMP программированию, оно требует хорошего понимания организации памяти. Но, конечно же, сложность разработки и переноса на CUDA сильно зависит от приложения.

Набор для разработчиков содержит множество примеров кода и хорошо документирован. Процесс обучения потребует около двух-четырёх недель для тех, кто уже знаком с OpenMP и MPI. В основе API лежит расширенный язык Си, а для трансляции кода с этого языка в состав CUDA SDK входит компилятор командной строки nvcc, созданный на основе открытого компилятора Open64.

Перечислим основные характеристики CUDA:

  • унифицированное программно-аппаратное решение для параллельных вычислений на видеочипах Nvidia;
  • большой набор поддерживаемых решений, от мобильных до мультичиповых
  • стандартный язык программирования Си;
  • стандартные библиотеки численного анализа FFT (быстрое преобразование Фурье) и BLAS (линейная алгебра);
  • оптимизированный обмен данными между CPU и GPU;
  • взаимодействие с графическими API OpenGL и DirectX;
  • поддержка 32- и 64-битных операционных систем: Windows XP, Windows Vista, Linux и MacOS X;
  • возможность разработки на низком уровне.

Касательно поддержки операционных систем нужно добавить, что официально поддерживаются все основные дистрибутивы Linux (Red Hat Enterprise Linux 3.x/4.x/5.x, SUSE Linux 10.x), но, судя по данным энтузиастов, CUDA прекрасно работает и на других сборках: Fedora Core, Ubuntu, Gentoo и др.

Среда разработки CUDA (CUDA Toolkit) включает:

  • компилятор nvcc;
  • библиотеки FFT и BLAS;
  • профилировщик;
  • отладчик gdb для GPU;
  • CUDA runtime драйвер в комплекте стандартных драйверов Nvidia
  • руководство по программированию;
  • CUDA Developer SDK (исходный код, утилиты и документация).

В примерах исходного кода: параллельная битонная сортировка (bitonic sort), транспонирование матриц, параллельное префиксное суммирование больших массивов, свёртка изображений, дискретное вейвлет-преобразование, пример взаимодействия с OpenGL и Direct3D, использование библиотек CUBLAS и CUFFT, вычисление цены опциона (формула Блэка-Шоулза, биномиальная модель, метод Монте-Карло), параллельный генератор случайных чисел Mersenne Twister, вычисление гистограммы большого массива, шумоподавление, фильтр Собеля (нахождение границ).

Преимущества и ограничения CUDA

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

Чтобы перенести вычисления на GPU в рамках такой модели, нужен специальный подход. Даже поэлементное сложение двух векторов потребует отрисовки фигуры на экране или во внеэкранный буфер. Фигура растеризуется, цвет каждого пикселя вычисляется по заданной программе (пиксельному шейдеру). Программа считывает входные данные из текстур для каждого пикселя, складывает их и записывает в выходной буфер. И все эти многочисленные операции нужны для того, что в обычном языке программирования записывается одним оператором!

Поэтому, применение GPGPU для вычислений общего назначения имеет ограничение в виде слишком большой сложности обучения разработчиков. Да и других ограничений достаточно, ведь пиксельный шейдер — это всего лишь формула зависимости итогового цвета пикселя от его координаты, а язык пиксельных шейдеров — язык записи этих формул с Си-подобным синтаксисом. Ранние методы GPGPU являются хитрым трюком, позволяющим использовать мощность GPU, но без всякого удобства. Данные там представлены изображениями (текстурами), а алгоритм — процессом растеризации. Нужно особо отметить и весьма специфичную модель памяти и исполнения.

Программно-аппаратная архитектура для вычислений на GPU компании Nvidia отличается от предыдущих моделей GPGPU тем, что позволяет писать программы для GPU на настоящем языке Си со стандартным синтаксисом, указателями и необходимостью в минимуме расширений для доступа к вычислительным ресурсам видеочипов. CUDA не зависит от графических API, и обладает некоторыми особенностями, предназначенными специально для вычислений общего назначения.

Преимущества CUDA перед традиционным подходом к GPGPU вычислениям:

  • интерфейс программирования приложений CUDA основан на стандартном языке программирования Си с расширениями, что упрощает процесс изучения и внедрения архитектуры CUDA;
  • CUDA обеспечивает доступ к разделяемой между потоками памяти размером в 16 Кб на мультипроцессор, которая может быть использована для организации кэша с широкой полосой пропускания, по сравнению с текстурными выборками;
  • более эффективная передача данных между системной и видеопамятью
  • отсутствие необходимости в графических API с избыточностью и накладными расходами;
  • линейная адресация памяти, и gather и scatter, возможность записи по произвольным адресам;
  • аппаратная поддержка целочисленных и битовых операций.

Основные ограничения CUDA:

  • отсутствие поддержки рекурсии для выполняемых функций;
  • минимальная ширина блока в 32 потока;
  • закрытая архитектура CUDA, принадлежащая Nvidia.


Слабыми местами программирования при помощи предыдущих методов GPGPU является то, что эти методы не используют блоки исполнения вершинных шейдеров в предыдущих неунифицированных архитектурах, данные хранятся в текстурах, а выводятся во внеэкранный буфер, а многопроходные алгоритмы используют пиксельные шейдерные блоки. В ограничения GPGPU можно включить: недостаточно эффективное использование аппаратных возможностей, ограничения полосой пропускания памяти, отсутствие операции scatter (только gather), обязательное использование графического API.

Основные преимущества CUDA по сравнению с предыдущими методами GPGPU вытекают из того, что эта архитектура спроектирована для эффективного использования неграфических вычислений на GPU и использует язык программирования C, не требуя переноса алгоритмов в удобный для концепции графического конвейера вид. CUDA предлагает новый путь вычислений на GPU, не использующий графические API, предлагающий произвольный доступ к памяти (scatter или gather). Такая архитектура лишена недостатков GPGPU и использует все исполнительные блоки, а также расширяет возможности за счёт целочисленной математики и операций битового сдвига.

Кроме того, CUDA открывает некоторые аппаратные возможности, недоступные из графических API, такие как разделяемая память. Это память небольшого объёма (16 килобайт на мультипроцессор), к которой имеют доступ блоки потоков. Она позволяет кэшировать наиболее часто используемые данные и может обеспечить более высокую скорость, по сравнению с использованием текстурных выборок для этой задачи. Что, в свою очередь, снижает чувствительность к пропускной способности параллельных алгоритмов во многих приложениях. Например, это полезно для линейной алгебры, быстрого преобразования Фурье и фильтров обработки изображений.

Удобнее в CUDA и доступ к памяти. Программный код в графических API выводит данные в виде 32-х значений с плавающей точкой одинарной точности (RGBA значения одновременно в восемь render target) в заранее предопределённые области, а CUDA поддерживает scatter запись — неограниченное число записей по любому адресу. Такие преимущества делают возможным выполнение на GPU некоторых алгоритмов, которые невозможно эффективно реализовать при помощи методов GPGPU, основанных на графических API.

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

История развития CUDA

Разработка CUDA была анонсирована вместе с чипом G80 в ноябре 2006, а релиз публичной бета-версии CUDA SDK состоялся в феврале 2007 года. Версия 1.0 вышла в июне 2007 года под запуск в продажу решений Tesla, основанных на чипе G80, и предназначенных для рынка высокопроизводительных вычислений. Затем, в конце года вышла бета-версия CUDA 1.1, которая, несмотря на малозначительное увеличение номера версии, ввела довольно много нового.

Из появившегося в CUDA 1.1 можно отметить включение CUDA-функциональности в обычные видеодрайверы Nvidia. Это означало, что в требованиях к любой CUDA программе достаточно было указать видеокарту серии Geforce 8 и выше, а также минимальную версию драйверов 169.xx. Это очень важно для разработчиков, при соблюдении этих условий CUDA программы будут работать у любого пользователя. Также было добавлено асинхронное выполнение вместе с копированием данных (только для чипов G84, G86, G92 и выше), асинхронная пересылка данных в видеопамять, атомарные операции доступа к памяти, поддержка 64-битных версий Windows и возможность мультичиповой работы CUDA в режиме SLI.

На данный момент актуальной является версия для решений на основе GT200 — CUDA 2.0, вышедшая вместе с линейкой Geforce GTX 200. Бета-версия была выпущена ещё весной 2008 года. Во второй версии появились: поддержка вычислений двойной точности (аппаратная поддержка только у GT200), наконец-то поддерживается Windows Vista (32 и 64-битные версии) и Mac OS X, добавлены средства отладки и профилирования, поддерживаются 3D текстуры, оптимизированная пересылка данных.

Что касается вычислений с двойной точностью, то их скорость на текущем аппаратном поколении ниже одинарной точности в несколько раз. Причины рассмотрены в нашей базовой статье по Geforce GTX 280. Реализация в GT200 этой поддержки заключается в том, блоки FP32 не используются для получения результата в четыре раза меньшем темпе, для поддержки FP64 вычислений в Nvidia решили сделать выделенные вычислительные блоки. И в GT200 их в десять раз меньше, чем блоков FP32 (по одному блоку двойной точности на каждый мультипроцессор).

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

Ещё одна полезная особенность CUDA 2.0 не имеет отношения к GPU, как ни странно. Просто теперь можно компилировать код CUDA в высокоэффективный многопоточный SSE код для быстрого исполнения на центральном процессоре. То есть, теперь эта возможность годится не только для отладки, но и реального использования на системах без видеокарты Nvidia. Ведь использование CUDA в обычном коде сдерживается тем, что видеокарты Nvidia хоть и самые популярные среди выделенных видеорешений, но имеются не во всех системах. И до версии 2.0 в таких случаях пришлось бы делать два разных кода: для CUDA и отдельно для CPU. А теперь можно выполнять любую CUDA программу на CPU с высокой эффективностью, пусть и с меньшей скоростью, чем на видеочипах.

Решения с поддержкой Nvidia CUDA

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

Актуальный список поддерживающих CUDA продуктов можно получить на вебсайте Nvidia. На момент написания статьи расчёты CUDA поддерживали все продукты серий Geforce 200, Geforce 9 и Geforce 8, в том числе и мобильные продукты, начиная с Geforce 8400M, а также и чипсеты Geforce 8100, 8200 и 8300. Также поддержкой CUDA обладают современные продукты Quadro и все Tesla: S1070, C1060, C870, D870 и S870.

Особо отметим, что вместе с новыми видеокартами Geforce GTX 260 и 280, были анонсированы и соответствующие решения для высокопроизводительных вычислений: Tesla C1060 и S1070 (представленные на фото выше), которые будут доступны для приобретения осенью этого года. GPU в них применён тот же — GT200, в C1060 он один, в S1070 — четыре. Зато, в отличие от игровых решений, в них используется по четыре гигабайта памяти на каждый чип. Из минусов разве что меньшая частота памяти и ПСП, чем у игровых карт, обеспечивающая по 102 гигабайт/с на чип.

Состав Nvidia CUDA

CUDA включает два API: высокого уровня (CUDA Runtime API) и низкого (CUDA Driver API), хотя в одной программе одновременное использование обоих невозможно, нужно использовать или один или другой. Высокоуровневый работает «сверху» низкоуровневого, все вызовы runtime транслируются в простые инструкции, обрабатываемые низкоуровневым Driver API. Но даже «высокоуровневый» API предполагает знания об устройстве и работе видеочипов Nvidia, слишком высокого уровня абстракции там нет.

Есть и ещё один уровень, даже более высокий — две библиотеки:

CUBLAS — CUDA вариант BLAS (Basic Linear Algebra Subprograms), предназначенный для вычислений задач линейной алгебры и использующий прямой доступ к ресурсам GPU;

CUFFT — CUDA вариант библиотеки Fast Fourier Transform для расчёта быстрого преобразования Фурье, широко используемого при обработке сигналов. Поддерживаются следующие типы преобразований: complex-complex (C2C), real-complex (R2C) и complex-real (C2R).

Рассмотрим эти библиотеки подробнее. CUBLAS — это переведённые на язык CUDA стандартные алгоритмы линейной алгебры, на данный момент поддерживается только определённый набор основных функций CUBLAS. Библиотеку очень легко использовать: нужно создать матрицу и векторные объекты в памяти видеокарты, заполнить их данными, вызвать требуемые функции CUBLAS, и загрузить результаты из видеопамяти обратно в системную. CUBLAS содержит специальные функции для создания и уничтожения объектов в памяти GPU, а также для чтения и записи данных в эту память. Поддерживаемые функции BLAS: уровни 1, 2 и 3 для действительных чисел, уровень 1 CGEMM для комплексных. Уровень 1 — это векторно-векторные операции, уровень 2 — векторно-матричные операции, уровень 3 — матрично-матричные операции.

CUFFT — CUDA вариант функции быстрого преобразования Фурье — широко используемой и очень важной при анализе сигналов, фильтрации и т.п. CUFFT предоставляет простой интерфейс для эффективного вычисления FFT на видеочипах производства Nvidia без необходимости в разработке собственного варианта FFT для GPU. CUDA вариант FFT поддерживает 1D, 2D, и 3D преобразования комплексных и действительных данных, пакетное исполнение для нескольких 1D трансформаций в параллели, размеры 2D и 3D трансформаций могут быть в пределах [2, 16384], для 1D поддерживается размер до 8 миллионов элементов.

Основы создания программ на CUDA

Для понимания дальнейшего текста следует разбираться в базовых архитектурных особенностях видеочипов Nvidia. GPU состоит из нескольких кластеров текстурных блоков (Texture Processing Cluster). Каждый кластер состоит из укрупнённого блока текстурных выборок и двух-трех потоковых мультипроцессоров, каждый из которых состоит из восьми вычислительных устройств и двух суперфункциональных блоков. Все инструкции выполняются по принципу SIMD, когда одна инструкция применяется ко всем потокам в warp (термин из текстильной промышленности, в CUDA это группа из 32 потоков — минимальный объём данных, обрабатываемых мультипроцессорами). Этот способ выполнения назвали SIMT (single instruction multiple threads — одна инструкция и много потоков).

Каждый из мультипроцессоров имеет определённые ресурсы. Так, есть специальная разделяемая память объемом 16 килобайт на мультипроцессор. Но это не кэш, так как программист может использовать её для любых нужд, подобно Local Store в SPU процессоров Cell. Эта разделяемая память позволяет обмениваться информацией между потоками одного блока. Важно, что все потоки одного блока всегда выполняются одним и тем же мультипроцессором. А потоки из разных блоков обмениваться данными не могут, и нужно помнить это ограничение. Разделяемая память часто бывает полезной, кроме тех случаев, когда несколько потоков обращаются к одному банку памяти. Мультипроцессоры могут обращаться и к видеопамяти, но с большими задержками и худшей пропускной способностью. Для ускорения доступа и снижения частоты обращения к видеопамяти, у мультипроцессоров есть по 8 килобайт кэша на константы и текстурные данные.

Мультипроцессор использует 8192-16384 (для G8x/G9x и GT2xx, соответственно) регистра, общие для всех потоков всех блоков, выполняемых на нём. Максимальное число блоков на один мультипроцессор для G8x/G9x равно восьми, а число warp — 24 (768 потоков на один мультипроцессор). Всего топовые видеокарты серий Geforce 8 и 9 могут обрабатывать до 12288 потоков единовременно. Geforce GTX 280 на основе GT200 предлагает до 1024 потоков на мультипроцессор, в нём есть 10 кластеров по три мультипроцессора, обрабатывающих до 30720 потоков. Знание этих ограничений позволяет оптимизировать алгоритмы под доступные ресурсы.

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

Видеочип получает ядро и создает копии для каждого элемента данных. Эти копии называются потоками (thread). Поток содержит счётчик, регистры и состояние. Для больших объёмов данных, таких как обработка изображений, запускаются миллионы потоков. Потоки выполняются группами по 32 штуки, называемыми warp’ы. Warp’ам назначается исполнение на определенных потоковых мультипроцессорах. Каждый мультипроцессор состоит из восьми ядер — потоковых процессоров, которые выполняют одну инструкцию MAD за один такт. Для исполнения одного 32-поточного warp’а требуется четыре такта работы мультипроцессора (речь о частоте shader domain, которая равна 1.5 ГГц и выше).

Мультипроцессор не является традиционным многоядерным процессором, он отлично приспособлен для многопоточности, поддерживая до 32 warp’ов единовременно. Каждый такт аппаратное обеспечение выбирает, какой из warp’ов исполнять, и переключается от одного к другому без потерь в тактах. Если проводить аналогию с центральным процессором, это похоже на одновременное исполнение 32 программ и переключение между ними каждый такт без потерь на переключение контекста. Реально ядра CPU поддерживают единовременное выполнение одной программы и переключаются на другие с задержкой в сотни тактов.

Модель программирования CUDA

Повторимся, что CUDA использует параллельную модель вычислений, когда каждый из SIMD процессоров выполняет ту же инструкцию над разными элементами данных параллельно. GPU является вычислительным устройством, сопроцессором (device) для центрального процессора (host), обладающим собственной памятью и обрабатывающим параллельно большое количество потоков. Ядром (kernel) называется функция для GPU, исполняемая потоками (аналогия из 3D графики — шейдер).

Мы говорили выше, что видеочип отличается от CPU тем, что может обрабатывать одновременно десятки тысяч потоков, что обычно для графики, которая хорошо распараллеливается. Каждый поток скалярен, не требует упаковки данных в 4-компонентные векторы, что удобнее для большинства задач. Количество логических потоков и блоков потоков превосходит количество физических исполнительных устройств, что даёт хорошую масштабируемость для всего модельного ряда решений компании.

Модель программирования в CUDA предполагает группирование потоков. Потоки объединяются в блоки потоков (thread block) — одномерные или двумерные сетки потоков, взаимодействующих между собой при помощи разделяемой памяти и точек синхронизации. Программа (ядро, kernel) исполняется над сеткой (grid) блоков потоков (thread blocks), см. рисунок ниже. Одновременно исполняется одна сетка. Каждый блок может быть одно-, двух- или трехмерным по форме, и может состоять из 512 потоков на текущем аппаратном обеспечении.

Блоки потоков выполняются в виде небольших групп, называемых варп (warp), размер которых — 32 потока. Это минимальный объём данных, которые могут обрабатываться в мультипроцессорах. И так как это не всегда удобно, CUDA позволяет работать и с блоками, содержащими от 64 до 512 потоков.

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

Модель памяти CUDA

Модель памяти в CUDA отличается возможностью побайтной адресации, поддержкой как gather, так и scatter. Доступно довольно большое количество регистров на каждый потоковый процессор, до 1024 штук. Доступ к ним очень быстрый, хранить в них можно 32-битные целые или числа с плавающей точкой.

Каждый поток имеет доступ к следующим типам памяти:

Глобальная память — самый большой объём памяти, доступный для всех мультипроцессоров на видеочипе, размер составляет от 256 мегабайт до 1.5 гигабайт на текущих решениях (и до 4 Гбайт на Tesla). Обладает высокой пропускной способностью, более 100 гигабайт/с для топовых решений Nvidia, но очень большими задержками в несколько сот тактов. Не кэшируется, поддерживает обобщённые инструкции load и store, и обычные указатели на память.

Локальная память — это небольшой объём памяти, к которому имеет доступ только один потоковый процессор. Она относительно медленная — такая же, как и глобальная.

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

Память констант — область памяти объемом 64 килобайта (то же — для нынешних GPU), доступная только для чтения всеми мультипроцессорами. Она кэшируется по 8 килобайт на каждый мультипроцессор. Довольно медленная — задержка в несколько сот тактов при отсутствии нужных данных в кэше.

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

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

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

Типичный, но не обязательный шаблон решения задач:

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

Среда программирования

В состав CUDA входят runtime библиотеки:

  • общая часть, предоставляющая встроенные векторные типы и подмножества вызовов RTL, поддерживаемые на CPU и GPU;
  • CPU-компонента, для управления одним или несколькими GPU;
  • GPU-компонента, предоставляющая специфические функции для GPU.

Основной процесс приложения CUDA работает на универсальном процессоре (host), он запускает несколько копий процессов kernel на видеокарте. Код для CPU делает следующее: инициализирует GPU, распределяет память на видеокарте и системе, копирует константы в память видеокарты, запускает несколько копий процессов kernel на видеокарте, копирует полученный результат из видеопамяти, освобождает память и завершает работу.

В качестве примера для понимания приведем CPU код для сложения векторов, представленный в CUDA:

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

Программы на CUDA могут взаимодействовать с графическими API: для рендеринга данных, сгенерированных в программе, для считывания результатов рендеринга и их обработки средствами CUDA (например, при реализации фильтров постобработки). Для этого ресурсы графических API могут быть отображены (с получением адреса ресурса) в пространство глобальной памяти CUDA. Поддерживаются следующие типы ресурсов графических API: Buffer Objects (PBO / VBO) в OpenGL, вершинные буферы и текстуры (2D, 3D и кубические карты) Direct3D9.

Стадии компиляции CUDA-приложения:

Файлы исходного кода на CUDA C компилируются при помощи программы NVCC, которая является оболочкой над другими инструментами, и вызывает их: cudacc, g++, cl и др. NVCC генерирует: код для центрального процессора, который компилируется вместе с остальными частями приложения, написанными на чистом Си, и объектный код PTX для видеочипа. Исполнимые файлы с кодом на CUDA в обязательном порядке требуют наличия библиотек CUDA runtime library (cudart) и CUDA core library (cuda).

Оптимизация программ на CUDA

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

В общих чертах, при оптимизации программы CUDA нужно постараться добиться оптимального баланса между размером и количеством блоков. Большее количество потоков в блоке снизит влияние задержек памяти, но снизит и доступное число регистров. Кроме того, блок из 512 потоков неэффективен, сама Nvidia рекомендует использовать блоки по 128 или 256 потоков, как компромиссное значение для достижения оптимальных задержек и количества регистров.

Среди основных моментов оптимизации программ CUDA: как можно более активное использование разделяемой памяти, так как она значительно быстрее глобальной видеопамяти видеокарты; операции чтения и записи из глобальной памяти должны быть объединены (coalesced) по возможности. Для этого нужно использовать специальные типы данных для чтения и записи сразу по 32/64/128 бита данных одной операцией. Если операции чтения трудно объединить, можно попробовать использовать текстурные выборки.

Выводы

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

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

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

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

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

Но, конечно, GPU не заменят CPU. В их нынешнем виде они и не предназначены для этого. Сейчас что видеочипы движутся постепенно в сторону CPU, становясь всё более универсальными (расчёты с плавающей точкой одинарной и двойной точности, целочисленные вычисления), так и CPU становятся всё более «параллельными», обзаводясь большим количеством ядер, технологиями многопоточности, не говоря про появление блоков SIMD и проектов гетерогенных процессоров. Скорее всего, GPU и CPU в будущем просто сольются. Известно, что многие компании, в том числе Intel и AMD работают над подобными проектами. И неважно, будут ли GPU поглощены CPU, или наоборот.

В статье мы в основном говорили о преимуществах CUDA. Но есть и ложечка дёгтя. Один из немногочисленных недостатков CUDA — слабая переносимость. Эта архитектура работает только на видеочипах этой компании, да ещё и не на всех, а начиная с серии Geforce 8 и 9 и соответствующих Quadro и Tesla. Да, таких решений в мире очень много, Nvidia приводит цифру в 90 миллионов CUDA-совместимых видеочипов. Это просто отлично, но ведь конкуренты предлагают свои решения, отличные от CUDA. Так, у AMD есть Stream Computing, у Intel в будущем будет Ct.


Которая из технологий победит, станет распространённой и проживёт дольше остальных — покажет только время. Но у CUDA есть неплохие шансы, так как по сравнению с Stream Computing, например, она представляет более развитую и удобную для использования среду программирования на обычном языке Си. Возможно, в определении поможет третья сторона, выпустив некое общее решение. К примеру, в следующем обновлении DirectX под версией 11, компанией Microsoft обещаны вычислительные шейдеры, которые и могут стать неким усреднённым решением, устраивающим всех, или почти всех.

Судя по предварительным данным, этот новый тип шейдеров заимствует многое из модели CUDA. И программируя в этой среде уже сейчас, можно получить преимущества сразу и необходимые навыки для будущего. С точки зрения высокопроизводительных вычислений, у DirectX также есть явный недостаток в виде плохой переносимости, так как этот API ограничен платформой Windows. Впрочем, разрабатывается и ещё один стандарт — открытая мультиплатформенная инициатива OpenCL, которая поддерживается большинством компаний, среди которых Nvidia, AMD, Intel, IBM и многие другие.

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

Знакомство с NV >

В очередной статье о технологиях корпорации NVIDIA, мы поговорим о NVIDIA CUDA — технологии параллельных вычислений с помощью GPU. В основе статьи лежит инетервью с ведущим специалистом по технологиям NVIDIA — Юрием Уральским. Юрий достаточно подробно рассказал о CUDA и перспективах применения этой технологии для увеличения произвоидтельности в различных областях. Однако мы сделали упор на применении CUDA в CG и особенно для визуализации, а так же вычислений сложных эффектов с применением GPU как основного инструмента ускорения вычислений.

На этот раз, вводный ролик к статье будет обладать легким оттенком PR, почему, вы узнаете в конце статьи — в разделе Making Of.

Все смотрели программу MythBusters (Разрушители мифов), и все, знают Джеми и Адама? На прошлом NVIDIA NVISION 08 эти ребята в самой доступной и наглядной форме показали различия между CPU и GPU вычислениям.

Видео-введение

В данном видео-введении автор статьи познакомит Вас с темой статьи.

Введение

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

Здесь вступает в игру не так давно получившее массовое признание, направление GPGPU — параллельные вычисления на графических процессорах.

В представленном вашему вниманию обзоре-интервью мы продолжим знакомство с технологиями корпорации NVIDIA и поговорим о NVIDIA CUDA. Но мы сделаем упор на применение этого инструмента в нашей с вами области — компьютерной графике (CG). Статья представляет собой интервью с ведущим специалистом по технологиям NVIDIA – Юрием Уральским. Некоторые вопросы дополнены комментарием автора и более детальным разъяснением. Также, имеется видео-версия этого интервью, с дополнительными видео-материалами. В конце статьи даны ссылки на все использованные в статье материалы и программное обеспечение.

Не обойдем стороной и новые технологии, использующие NVIDIA CUDA для самых разных задач. Рассмотрим NVIDIA PhysX с точки зрения его реализации, познакомимся с новыми системами визуализации от mental images – разработчика легендарного mental ray, и многое другое.

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

Интервью с Юрием Уральским (NVIDIA Corp.)

Когда я этим летом ездил в Москву на CG EVENT 2009, мне довелось посетить Российское представительство корпорации NVIDIA в Москве. Тема посещения, разумеется, была — технология NVIDIA CUDA и применение её в CG. С вопросами и за разъяснениями я обратился к Юрию Уральскому — ведущему специалисту по технологиям NVIDIA.

Ведущий специалист по технологиям NVIDIA – Юрий Уральский.

Интервью представлено в виде диалога с Юрием, как в текстовом, так и в видео варианте.

Первая часть видео-версии интервью с Юрием Уральским.

Dimson3d | Добрый день Юрий, спасибо, что согласились дать интервью для on-line журнала render.ru. Расскажите, пожалуйста, о NVIDIA CUDA, как появилось направление GPGPU и что стало предпосылкой к созданию технологии CUDA?

Стоит начать с вопроса — Что такое CUDA? CUDA — платформа, а не какой-то язык программирования, не какая-то среда программирования, а платформа, которая позволяет использовать мощности графических процессоров (GPU) для вычислений общего назначения. Почему это стало интересно? Примерно в 2003 году сформировалось сообщество энтузиастов, которые стали использовать GPU для запуска произвольного кода. Тогда еще не было CUDA, были стандартные графические API. Но примерно в то время GPU стали достаточно программируемыми. То есть появилось понятие графического шейдера, Open GL и Direct X фактически позволял вам сконфигурировать графический pipeline визуализации. К примеру, расчет цвета пикселя, указать какие текстуры используются и по каким правилам смешиваются, каким образом получается пиксель. Примерно в 2003 году появилась программируемость, разработчики получили возможность написать произвольную программу, которая выполнялась на каждом пикселе и вычисляла, какой ни будь алгоритм. Вместо того, чтобы фиксированным образом конфигурировать графический pipeline, появилась возможность описывать действия, которые производятся над каждой вершиной или пикселем в виде программы. При этом интерфейс GPU оставался сугубо графическим. Т.е. все равно нужно было использовать Direct X или Open GL. Можно так сказать, что графический конвейер остался тем же самым, но мы добавили программируемость в определенные места этого конвейера. Т.е. в конфигурируемом конвейере появилась возможность заменять участки на программы. При этом энтузиасты поняли, что эту возможность можно использовать не только для графики, но и для любых вычислений. Если представить, что мы вычисляем не цвета пикселей, а какие то значения, решаем уравнения. При этом соответственно ваш экран или массив пикселей, который вы считаете, становится выходным массивом данных, а ваша «текстура», которую вы подаете — входным массивом данных. Это то, что называется GPGPU.

Естественно, программирование, таким образом, было сопряжено с большими трудностями, потому что разработчикам нужно было изучать API, понимать, как работает графический ускоритель, и для многих это оставалось недоступным. Если человек исследователь, в какой-то предметной области несвязанной никак с графикой, ему соответственно требовалось время и ресурсы чтобы изучить саму программную модель того же Direct X. Но, тем не менее, люди этим занимались. Почему это было интересно? Потому что для большого класса задач оказалось, что производительность, которую они получают, используя GPU, значительно превосходила CPU.

Примерно в то время мы за этим очень внимательно наблюдали и попытались проанализировать все проблемы, с которыми разработчики сталкивались, и мы решили создать свою программную модель, которая все эти проблемы решала. То есть модель-платформа, которая позволяла запускать произвольный код на GPU. Собственно это идея, которая лежит в основе CUDA. Сама CUDA на рынке появилась в конце 2006 года. С приходом архитектуры G80. Когда появились карты GeForce 8800 — первые наши продукты, которые поддерживали CUDA.

Dimson3d | Архитектура у G80 была существенно переработана?

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

Архитектура чипа G80.

Еще раз подчеркну, что CUDA — не какой то конкретный язык программирования. Люди зачастую ассоциируют с языком C и в принципе изначально мы такие заявления делали — что CUDA — язык программирования. Теперь мы пытаемся от этого уйти, то есть CUDA – именно платформа. Набор средств, которые позволяют вам использовать возможности GPU для вычислений. На данный момент у нас есть компилятор с языка С, но это не значит, что в будущем будет только он. У нас есть планы по поддержке других языков, будет Open CL, DX Compute, все это языки программирования, программные среды поддерживают одну и туже базовую модель вычислений, и туже базовую архитектуру. Вот эта архитектура называется CUDA.

Dimson3d | Скажите, а какие различия существуют, между разработкой для CPU и GPU?

Наверное, в первую очередь, стоит отметить, какие архитектурные отличия существуют между CPU и GPU. Потому что из этого будут следовать отличия программирования. То есть GPU по сути своей — массивно параллельный процессор. Традиционная модель вычислений на CPU подразумевает последовательную модель исполнения. То есть у вас есть программа набор инструкций и эти инструкции вы выполняете одну за другой последовательно. А GPU подходит к проблеме совсем с другой стороны. Мы изначально предполагаем что, в наборе данных задач, которые мы обрабатываем, присутствует очень высокая степень параллелизма. То есть если мы говорим о графике, то совершенно естественно у вас большое количество пикселей на экране, большое количество вершин, треугольников и с этим большим массивом данных вам необходимо провести действия, но которые сходны. Одна последовательность действий, одна программа, которая исполняется сразу в массивно-параллельном режиме. В соответствии со свойством таких систем, мы изначально имели дело с очень массивно-параллельными задачами — коими являются графические задачи. Исходя из этого, архитектура GPU сильно отличается от архитектуры CPU. Мы имеем дело с большим массивом вычислительных модулей. Каждый, из которых работает в какой-то степени независимо, но все они выполняют одну последовательность действий одну программу над большим массивом данных. CPU в отличии от этой модели, использует другой подход. Вместо того, что бы использовать очень много маленьких процессоров, которые все делают, они строят достаточно увесистые ядра. Большие размеры кэша, обеспечивая достаточно высокую производительность при исполнении последовательного кода, но при этом очень трудно вместить большое количество ядер на кристалл. Допустим, сейчас Intel говорит о 4 — 8 ядрах на кристалле, в то время как у нас это море процессоров, массив процессоров исчисляется сотнями. Наш последний GPU например Tesla 10-ой серии — 240 процессоров, которые работают параллельно. Кардинальные отличия между архитектурами – то, что мы тратим площадь кристалла на сами вычислительные модули и делаем их очень много. В то время как центральный процессор, тратит площади кристалла, допустим, на кэш и достаточно небольшое количество относительно мощных вычислительных ядер. В соответствии с этим, программные модели тоже кардинально отличаются. Собственно поэтому мы и вынуждены были создать CUDA потому что традиционные языки программирования такие — как традиционный С не совсем подходят для выражения этого параллелизма. Нам пришлось взять стандартный С, расширить его мета-конструкциями для того, что бы программист мог выражать параллельные задачи и создать свой компилятор. В этом, наверное, и состоит главное отличие — масштаб параллелизма, к которому мы идем.

Пример кода на стандартном языке C, и пример кода с дополнениями для параллелизма.

Изначально CUDA как модель поддерживала гетерогенные вычисления — вы строите свою программу в виде секций. Есть секции, которые массивно параллельны — секции, которые будут исполняться на GPU. Есть секции, которые исполняются последовательно на CPU. С самой первой версии CUDA это было реализовано и поддерживалось. В одной программе вы могли смешивать параллельные и последовательные участки кода.

Дополнение. Параллелизм и новая архитектура NVIDIA Fermi

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

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

За счет увеличения количества вычислительных ядер (ALU) достигается производительность системы в расчетах.

На GTC 2009 была презентована новая архитектура — NVIDIA Fermi. О ней я и хочу поговорить в этом дополнении. Начнем с того, что теперь в новых чипах более 3х миллиардов транзисторов это конечно стало возможным благодаря переходу на 40 нм техпроцесс. Так же на долю вычислений предоставлено 512 так называемых CUDA ядер. На заметку в предыдущем чипе (G200) было 240 ядер, а в G80 их всего 128.

Архитектура Fermi. Теперь, как вы видите, используется достаточно емкий кэш 2-го уровня.

Теперь чип состоит из 16-ти потоковых мультипроцессоров, которые содержат по 32 шейдерных ядра. Что в совокупности и дает нам 512 CUDA ядер. Блоки расположены вокруг общей кэш-памяти второго уровня. Каждый из блоков состоит из планировщика и организатора, исполнительных модулей и файлов регистров и кэш-памяти первого уровня.

Детальная архитектура мультипроцессора.

Другое достаточно заманчивое решение, которое реализовано в Fermi – поддержка коррекции ошибок памяти (ЕСС). Что конечно скажется на производительности в лучшую сторону. А если учитывать то, что данная архитектура больше направлена и сориентирована на вычисления, то это позволит использовать продукты следующего поколения Tesla или Quadro в еще более сложных задачах и требующих продолжительных вычислений с большими объемами данных.

Одно из ядер мультипроцессора.

Так же следует отметить поддержку максимального объема памяти — 6 Гб, чип оснащается шестью 64-разрядными контроллерами памяти GDDR5, что дает 384-битную шину памяти. Для ускорения математических вычислений и выполнения других функций предназначена технология NVIDIA Parallel DataCache. В завершении скажу, что также Fermi поддерживает аппаратно такие средства программирования как C, C++, FORTRAN, и многие другие функции — такие как Open CL и Direct Compute.

GTC 2009 Презентация NVIDIA Fermi.

Вторая часть видео-версии интервью с Юрием Уральским.

Dimson3d | В демонстрационных материалах вы часто показываете, диаграмму производительности GPU по сравнению с CPU. Есть ли предел на данный момент производительности GPU? Центральные процессоры все равно упираются в определенный потолок.

Предел? (улыбаясь, переспросил Юрий) Хороший вопрос. Наверное, предел есть, просто мы живем в физическом мире. Вопрос можно поставить так: «Где этот предел находится?» — насколько этот потолок большой. Сейчас на самом деле пределом как я представляю, является так называемый power wall. Мы подходим к тому моменту, когда высокопроизводительным кластерам невозможно дальше наращивать производительность из-за того, что количество энергии, которое нужно для подпитывания машин слишком велико. В этом смысле как раз GPU, и вообще параллельные машины предлагают более эффективное решение. Один параметр, которым мы обычно хвастаемся — Performance per Watt. Это процент производительности на Ватт затраченной энергии или сколько GFlop потребляет энергии. Почему это так? Если вернемся к архитектуре возникает такая ситуация благодаря тому, что мы используем площадь кристалла на сами вычислители. Мы стараемся очень эффективно использовать доступную площадь и доступный бюджет при потребляемой энергии. Вместо того, чтобы кэшировать данные, наш механизм исполнения базируется на совсем других предположениях. Традиционные процессоры вынуждены строить большие массивные кэши и если программа использует набор данных, который не помещается в кэш, производительность процессоров сильно падает. Соответственно, чем выше мы хотим получить финальную производительность тем, больше нам нужно строить кэш, а кэш — структура достаточно не эффективная в плане потребляемой мощности и кэш сам по себе это не вычислительный модуль. Те транзисторы, которые вы тратите на кэш это просто занятая площадь, не выполняющая вычислений.

График производительности GPU по сравнению с CPU.

Dark | Как обстоят дела с коммуникационными расходами?

Совершенно верно, коммуникационные расходы — здесь ключевым моментом является то, что мы не строим большие массивные кэши, а используем площадь кристалла и энергетические ресурсы для добавления новых параллельных модулей — позволяет нам находиться на кривой роста производительности. Наращивать параллелизм, наращивать количество «маленьких вычислителей» значительно проще, чем увеличивать тактовую частоту. Центральные процессоры не могут позволить себе такой рост именно из-за того, что не могут наращивать параллелизм такими же темпами. Поскольку они вынуждены хорошо исполнять традиционные приложения — OS, прикладные программы (Word, Excel к примеру) и они просто не могут себе позволить уменьшить вычислительное ядро. Они не могут деградировать производительность всех этих приложений.

Dark | А есть ли какой то предел количеству процессоров, после которого производительность просто падает?

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

Dimson3d | В линейке вашей продукции для вычислений с помощью CUDA, представлены решения GPU GeForce, Quadro FX и Tesla. Все эти графические процессоры и системы могут использоваться для параллельных расчетов. Расскажите, пожалуйста, о NVIDIA Tesla — её основных отличиях, что позволяет её выделить среди других.

В первую очередь я скажу, что GeForce, Quadro и Tesla — просто названия продуктов. Это продукты, которые ориентированы на тот или иной рынок, все они используют одну и туже базовую архитектуру. Базовая архитектура, заложенная, в эти продукты поддерживает NVIDIA CUDA и в принципе используют одну и туже технологию. Скажем, различия начинаются на более высоком уровне, в плане того, что мы пытаемся ориентировать продукт на конкретный рынок. GeForce – продукты для Consumer направления — геймеров, использования в стандартных компьютерах. Quadro имеет ориентацию на более профессиональный рынок — визуализация, CAD, рабочие станции (workstations). Tesla – продукт, который ориентирован на вычисления в кластерах, если вы хотите построить кластер на основе GPU, то вы будете использовать Tesla. Если говорить о различиях между этими продуктами, то Tesla в частности отличается тем, что чипы проходят наиболее жесткое тестирование. Качество памяти, которая устанавливается на плату значительно выше, чем у GeForce. Поскольку мы сами выпускаем Tesla, мы даем гарантию, что она будет работать у вас 24 часа в сутки 7 дней в неделю. Это зависит от способа применения данного конкретного продукта. GeForce — скорей всего вы включите компьютер, сделаете все необходимое, поиграете и выключите его, и сможет ли он проработать целую неделю с постоянной нагрузкой 24/7? Если у вас сервер, вы проводите, какие то сложные научные расчеты, то вам важно, что бы он был доступен постоянно. Это продукты, которые ориентированы на разные способы применения.

Дополнение. NVIDIA Tesla краткое описание

Технические характеристики решений Tesla в desktop и server вариантах.

NVIDIA Tesla представлена в виде серверного решения и desktop варианта. Для разработчиков, и вообще для специалистов, кому не требуется большая производительность подойдет идеально подойдет desktop вариант, так же такое решение прекрасно подходит для создания персональных высокопроизводительных рабочих станций нацеленных на научные исследования и вычисления. Кстати, вот можно использовать такую комбинацию GPU Quadro или GeForce и NVIDIA Tesla. GPU выполняет визуализацию, а Tesla не обременяясь расчетами картинки выполняет только поставленные задачи по вычислениям. Плюс, их можно использовать совместно, непосредственно для вычислений.

С технической стороны данные решения выглядят так. Модель Tesla С1060 для настольных систем представлена в виде графического ускорителя, такого же, как у нас в системных блоках, но он не обладает портами для подключения мониторов. Внутри Tesla C1060 содержит чип T10 (GT200). У Tesla C1060 на борту 4 Gb GDDR3 графической памяти с интерфейсом в 512-bit. Потребление энергии данным устройством составляет 160W.

Но это решение для рабочих станций, где вы можете работать и выполнять расчеты непосредственно на рабочем месте. Но что делать, когда доходит дело до вычислений в огромных масштабах, в серверных залах с возможностью масштабирования? Для этого разработана NVIDIA Tesla S1070. Данное решение представлено в виде 1U корпуса, который устанавливается в стойку над или под сервером (см. иллюстрацию ниже). Данное решение представлено 4мя чипами T10 (GT200), каждому из них отведено по 4 Gb GDDR3 памяти, при этом в сумме получаем 16 Gb памяти. Но это решение к системам подключается через специальный HIC (host interface card) и кабель. В сами сервера x86 или рабочие станции в шину PCI-E x16 или x8.

Dimson3d | Мне как специалисту по визуализации, интересно, можно ли применить CUDA непосредственно в этой области. Возможна ли полная реализация системы визуализации? Хочется заметить, что многие пользователи жалуются — «Вот использовать эту мощность для визуализации финальных сцен, прикрутить к V-Ray или mental ray». Будет ли реализована новая система визуализации, которая будет полностью использовать потенциал графических ускорителей, или же можно будет перекладывать отдельные задачи, модули на GPU. Расчеты трассировки лучей или GI и т.п.

Да. Естественно мы считаем, что CUDA должна и может использоваться для финальной визуализации. К сожалению, пока о каких-то готовых продуктах, которые используют CUDA для визуализации говорить рано. Мы в этом направлении работаем можно отметить существование библиотеки NVRT демонстрация, которой была на SIGGRAPH 2007, она уже успешно применяется для визуализации. Многие компании из индустрии её используют. Но это незаконченный финальный продукт — именно библиотека, которая позволяет вам использовать CUDA GPU для решения задач трассировки лучей. Это базовый уровень. Предполагается, что на основе этой библиотеки вы будете строить свои приложения — визуализация с использованием Raytracing, или какие либо другие задачи. Это оптимизированная реализация трассировки лучей в общем смысле не обязательно для рендеринга и оптимизированная для наших процессоров GPU. Внутри компании, мы так же достаточно плотно занимаемся вопросами реализации, например альтернативных конвейеров рендеринга на GPU. У нас есть проект реализации Rays на CUDA. Так же проект из Microsoft Research где, они реализовали RenderMan с помощью CUDA и получили очень хорошее ускорение. Но пока это находится на стадии разработки и экспериментов.

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

Еще можно наверное сказать относительно применения CUDA вообще в задачах для симуляции. Есть такой plug-in к 3ds Max — Ray Fire он не использует CUDA на прямую, а использует PhysX, а PhusX в свою очередь базируется на CUDA. Это как пример, как CUDA может помочь в решении задач не связанных на прямую с визуализацией.

Дополнение. OptiX, iray, RealityServer

Слайд с презентации NVIDIA на SIGGRAPH 2009 на котором представлена информация о интерактивном рейтрейсере — NVIDIA OptiX.

Вот мы и добрались до самого вкусного и интересного раздела дополнений. Хочется сразу вернуться на 1,5 года назад, когда я предрекал некоторым своим друзьям о скором будущем, когда Raytracing можно будет рассчитывать за минуты и причем в очень сложных сценах. Конечно, тогда все смеялись над таким заявлением. Но теперь смеяться не над чем. И так, OptiX — интерактивный визуализатор трассировки лучей, базирующийся на NVIDIA CUDA, и который может быть встроен в любой софт. Дело только за разработчиками. OptiX был представлен на SIGGRAPH 2009 на стенде NVIDIA.

На данный момент, OptiX на финальной стадии разработки, но то, что сделано уже сейчас впечатляет. Давайте посмотрим на то, что мы получаем и, что можно рассчитывать с помощью OptiX’a:

  • Программируемые — поверхности, шейдеры, камеры;
  • Не лимитированы только расчетом графики — лучи могут содержать любые пользовательские данные;
  • Возможность масштабирования оборудования;
  • Ambient Occlusion;
  • Reflections;
  • Refractions;
  • Photon Mapping GI;


Ключевые особенности NVIDIA OptiX.

При этом, главное преимущество OptiX практически полная открытость для расширения и надстроек. Вы можете создавать свои собственные шейдеры, ИС, виртуальные камеры и многое другое. После чего, используя GPU NVIDIA Quadro FX создавать свои проекты с меньшими затратами времени на ожидание и исправления. Хотя пока на данный момент все это требует серьезной доработки в плане дружественности с пользователем, но в течение ближайшего года следует ожидать кардинальных изменений в нашем с вами программном обеспечении.

NVIDIA OptiX является частью AXE — Application Acceleration Engines, набор модулей для разработчиков в который помимо OptiX включены еще такие модули как SceniX и CompleX, но о них я расскажу в одной из следующих статей.

Слайд с презентации NVIDIA на SIGGRAPH 2009 на котором, представлена информация о новой системе визуализации mental images iray.

Наконец, для пользователей системы визуализации mental ray появилась возможность использовать все потенциалы GPU для визуализации самых сложных и насыщенных различными сложными поверхностями сцены. На SIGGRAPH 2009 был презентован новый продукт от mental images – iray. iray базируясь на архитектуре NVIDIA CUDA, позволяет использовать для визуализации все возможности GPU. iray является как компонента mental ray, так и RealityServer. Плюс ко всему, mental ray 3.8 с iray в течение ближайшего года будет интегрирован во все программы, создаваемые в OEM партнерстве с mental images. К примеру, следующие 3ds Max или Maya, а также SI, возможно будут обладать возможностью расчета финальной картинки с помощью mental ray и iray.

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

Схема использования GPU приложениями использующими код на CUDA. Для примера показана схема для OptiX или iray. Учтите что OptiX — интерактивная система, а iray это компонент mental ray 3.8 или RealityServer. Как интегрированные в DCC приложения, так и работающих независимо от них.

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

Но, а на какую графическую подсистему рассчитывается применение технологии OptiX и iray? Все как всегда сводится к вопросу о производительности и стабильности, и здесь как всегда ответ сам приходит на ум — Quadro FX и Tesla. Преимущество Tesla состоит в том, что она занимается исключительно расчетами — в таком случае можно использовать её как мощнейший сопроцессор. В случае с Quadro и Tesla мы получаем стабильность и гарантию того что все будет выполнено как задумано художником и не будет такого момента как перегоревшая в самый неподходящий момент GPU или сбои в работе оборудования. Хотя это можно и приписать к минусу — ведь далеко не все пользователи позволят себе покупку более дорогой Quadro и тем более Tesla. Однако подождем выхода конечных продуктов, непосредственно в приложениях для DCC и посмотрим, как будут реализованы эти технологии.

Ах, да совсем забыл сказать, что на данный момент для OptiX рекомендуется применение GPU NVIDIA Quadro FX 3800 и выше. Хотя я с достаточно большим успехом запустил демонстрационные примеры и на своей Quadro FX 1800. Все дело упирается в количество потоковых процессоров. У NVIDIA Quadro FX 1800 их 64, а у Quadro FX 3800 их уже 192 (как и у Quadro FX 4800\CX). И чем больше потоковых процессоров — тем лучше. Хм, раньше мы пытались под «кипение кристалла ЦП» заставить визуализироваться сложные сцены, теперь будем пытаться заставить «кипеть» кристалл графического ускорителя 8).

Это видео демонстрирует возможнсоти iray в RealityServer 2.4.

Но не одним iray мы едины. Учитывая, что есть и другие рендереры — V-Ray, Brazil и Final Render. У Chaosgroup (разработчик V-Ray) есть так же решение, разработанное с учетом того, что оно будет использовать GPU — V-Ray RT. Хотя в первом релизе на данный момент нет прямой поддержки GPU, и данное решение встраивается в 3ds Max в качестве Active Shade во Viewport. Однако хочется сказать, что разработчики уделяют этому огромную долю внимания — ведь получить визуализатор, который будет использваоть GPU для увеличения скорости рендера. Это очень большой шаг вперед. Но стоит другой вопрос — как быть, если у потенциального клиента и покупателя используются GPU не NVIDIA? Тут ведь нельзя просто взять и написать программный код, оптимизированный только под одну архитектуру и производителя. Есть отличный ответ — использовать Open CL (Open Computing Language). Данный язык не привязан к конкретной продукции — будь то NVIDIA или ATI. И разработчики в Chaosgroup как раз присматриваются именно к использованию Open CL, хотя в демонстрации, которую вы увидите ниже используется и NVIDIA CUDA.

GPU accelerated rendering of VRay part 1 — 20x speedup!

GPU accelerated rendering of VRay part 2 — 20x speedup!

Но вернемся к iray. Ниже даны примеры визуализации интерьера с помощью технологии iray. Достоверно я не могу сказать, сколько времени заняла визуализация этих изображений. Но на видео-демонстрации нового RealityServer 3.0 картинка визуализировалась неимоверно быстро — учитывая, что рассчитывались GI, Reflections. Refractions. Буквально несколько секунд, и готово.

Пример визуализации интерьера офисного помещения с помощью iray.

Что хочется добавить в заключение и повториться в этом разделе дополнений. А вот что. Главное отличие iray от того же V-Ray RT и OptiX — в первую очередь он не является интерактивным средством визуализации. iray предназначается для увеличения скорости вычислений с помощью GPU и множества процессоров, которые находятся внутри графического чипа. Поэтому чем больше GPU будет у вас в системе и чем больше вычислительных ядер в каждом GPU (64 — 240 и выше), тем лучше. При этом если имеется такой продукт как NVIDIA Tesla, то мы получим высокопроизводительное решение для визуализации самых разнообразных сцен и моделей за достаточно короткие сроки.

Если применять RealityServer и новые решения Tesla RS, здесь появляется новый термин доселе не применявшийся в области визуализации — «облачные вычисления», об этом мы поговорим отдельно в статье посвященной iray и RealityServer. Весной 2010 года.

Видео презентации интерактивного Raytracing на модели Bugatti Verona.

Nvidia: Raytracing Bugatti Veyron

Это видео демонстрирует, как работает интерактивный трассировщик луча на примере модели Bugatti Veyron.

Рубрика: Введение в программирование на CUDA

Занятие 11. Программирование на CUDA. Часть 7. Множество Мандельброта

Реализация фрактала Множество Мандельброта

В данном примере будет продемонстрировано построение множества Мандельброта на CUDA C с использованием OpenCV (Open Source Computer Vision Library), которая из себя представляет библиотеку алгоритмов компьютерного зрения. В программе будут применяться следующие модули:

Занятие 10. Программирование на CUDA. Часть 6.

В этом примере будет реализована битонная сортировка. Битонная сортировка – это параллельный алгоритм сортировки, основанный на битонной последовательности. Такой последовательностью называют последовательность, которая монотонно не убывает, а затем монотонно не возрастает. Пример такой последовательности – <8, 10, 16, 12, 4, -2, -8>.

Занятие 9. Программирование на CUDA C. Часть 5.

В данном примере будет показано, как с помощью GPU можно посчитать квадрат числа N . Также, как и в предыдущих примерах особое внимание будет уделено специфике CUDA C при написании программ.

Занятие 8. Программирование на CUDA C. Часть 4.

Сложение векторов и сравнение с количеством тредов в памяти

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

Занятие 7. Программирование на CUDA C. Часть 3.

Увеличение значение у элементов в матрице на единицу

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

Занятие 6. Программирование на CUDA C. Часть 2.

Скалярное произведение с использованием shared-памяти

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

Занятие 5. Программирование на CUDA C. Часть 1.

На данном занятии читатель познакомится с основами практики программирования на CUDA C – будет подробно описано написание кода, его суть и исполнение. В качестве примеров будут разобраны самые популярные и полезные задачи, среди которых можно выделить:

Занятие 4. Первая программа «Hello, CUDA!»

Традиционно каждый программист, изучая новый язык программирования или новую технологию, начинает свой путь с написания программы “Hello, World!” Также поступим и мы, изучая архитектуру CUDA, но, изменив слово “Мир” на “CUDA” — это будет оригинальнее.

Занятие 3. Установка и настройка рабочего пространства CUDA

Приступая к работе, нужно убедиться, что на вашем рабочем месте, будь то компьютер или ноутбук, присутствует дискретная видеокарта NVIDIA с чипом восьмого поколения G80 (NVIDIA GeForce 8).

Занятие 2. Память в CUDA

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

Найдите нас

Адрес
г. Владивосток, кампус ДВФУ (остров Русский),
лабораторный корпус, L449, L451.

О сайте

Сайт работает на серверах университета ДВФУ, г. Владивосток.

Рубрика: Введение в программирование на CUDA

Занятие 11. Программирование на CUDA. Часть 7. Множество Мандельброта

Реализация фрактала Множество Мандельброта

В данном примере будет продемонстрировано построение множества Мандельброта на CUDA C с использованием OpenCV (Open Source Computer Vision Library), которая из себя представляет библиотеку алгоритмов компьютерного зрения. В программе будут применяться следующие модули:

Занятие 10. Программирование на CUDA. Часть 6.

В этом примере будет реализована битонная сортировка. Битонная сортировка – это параллельный алгоритм сортировки, основанный на битонной последовательности. Такой последовательностью называют последовательность, которая монотонно не убывает, а затем монотонно не возрастает. Пример такой последовательности – <8, 10, 16, 12, 4, -2, -8>.

Занятие 9. Программирование на CUDA C. Часть 5.

В данном примере будет показано, как с помощью GPU можно посчитать квадрат числа N . Также, как и в предыдущих примерах особое внимание будет уделено специфике CUDA C при написании программ.

Занятие 8. Программирование на CUDA C. Часть 4.

Сложение векторов и сравнение с количеством тредов в памяти

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

Занятие 7. Программирование на CUDA C. Часть 3.

Увеличение значение у элементов в матрице на единицу

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

Занятие 6. Программирование на CUDA C. Часть 2.

Скалярное произведение с использованием shared-памяти

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

Занятие 5. Программирование на CUDA C. Часть 1.

На данном занятии читатель познакомится с основами практики программирования на CUDA C – будет подробно описано написание кода, его суть и исполнение. В качестве примеров будут разобраны самые популярные и полезные задачи, среди которых можно выделить:

Занятие 4. Первая программа «Hello, CUDA!»


Традиционно каждый программист, изучая новый язык программирования или новую технологию, начинает свой путь с написания программы “Hello, World!” Также поступим и мы, изучая архитектуру CUDA, но, изменив слово “Мир” на “CUDA” — это будет оригинальнее.

Занятие 3. Установка и настройка рабочего пространства CUDA

Приступая к работе, нужно убедиться, что на вашем рабочем месте, будь то компьютер или ноутбук, присутствует дискретная видеокарта NVIDIA с чипом восьмого поколения G80 (NVIDIA GeForce 8).

Занятие 2. Память в CUDA

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

Найдите нас

Адрес
г. Владивосток, кампус ДВФУ (остров Русский),
лабораторный корпус, L449, L451.

О сайте

Сайт работает на серверах университета ДВФУ, г. Владивосток.

Знакомство с технологией CUDA

В развитии современных процессоров намечается тенденция к постепенному увеличению количества ядер, что повышает их возможности в параллельных вычислениях. Однако уже давно имеются 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,5×261 МГц) и до 4 ГГц (11,5×348 МГц) при частоте оперативной памяти 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 это один из этапов в процессе данного объединения.


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

  • MSI за видеокарту NX8800GT-T2D256E-OC;
  • WD за жесткий диск WD3200AAKS.

Особенности архитектуры и программирования графических ускорителей

Программирование › Архитектура ЭВМ и систем › Особенности архитектуры и программирования графических ускорителей

В этой теме 0 ответов, 1 участник, последнее обновление Васильев Владимир Сергеевич 2 года/лет, 4 мес. назад.

Развитие современных графических ускорителей или Graphic Processor Unit (в дальнейшем GPU) привело к появлению графических ускорителей с программируемым конвейером. Если раньше программист графических приложений мог использовать только фиксированную функциональность GPU, то теперь 2 стадии обработки информации в конвейере GPU стали программируемыми. Более точно программируемой стала обработка вершин, что позволило создавать, в частности, более разнообразные модели «освещения», а так же обработка фрагментов. Столь коренное изменение архитектуры GPU привело к возможности использования GPU не только для целей компьютерной графики, но и для задач, которые раннее решались на CPU, таким образом, появилась Технология Вычислений Общего Назначения на Графическом Ускорителе или General Purpose computation on Graphic Processor Unit (в дальнейшем просто GPGPU).

Идея применения специализированных арифметических ускорителей при построении суперкомпьютерных систем за последнее 5-10 лет стала довольно популярной благодаря возможности существенного повышения производительности при сохранении уровня энергопотребления и количества вычислительных узлов. Если проследить за эволюцией списка самых быстрых суперкомпьютеров мира Top500, то легко заметить, что переход к гетерогенным архитектурам не раз позволял соответствующим вычислительным системам занимать первые места с существенным отрывом от «классических» кластеров.

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

Применить технологии, отработанные при создании мощных игровых видеокарт для параллельных вычислений, впервые удалось NVIDIA, не остался в стороне и AMD со своим ускорителем FireStream. Использование графических ускорителей (GPGPU) позволяет получить значительную вычислительную мощность в десятки раз дешевле по деньгам и потребляемой мощности. В нынешнем списке TOP500три из первой пятерки суперкомпьютеров используют ускорители NVIDIA TESLA. GPGPU — единственная доступная по деньгам возможность получить «персональный суперкомпьютер» терафлопсной мощности в обычном настольном корпусе. Однако не все так просто, программирование для графических ускорителей не самое простое занятие. Также возникают вопросы, что выбрать — дорогую специализированную Tesla или топовую графическую карту, которая быстрее и дешевле? В любом случае, альтернативы множеству «легких» ядер пока не просматривается, а значит, приходится программировать все более параллельно. Сейчас выбор ПО, предназначенного для параллельной обработки данных с использованием не только лишь CPU весьма и весьма велик.

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

  • Развитие графических ускорителей;
  • Преимущества GPGPU;
  • Архитектура GPGPU;
  • Технологии программирования GPGPU: CUDA, OpenCL, Direct X.

1 Общие сведения о графических ускорителях

Современная видеокарта состоит из следующих частей:

  • графический процессор (Graphics processin gunit – графическое процессорное устройство) — занимается расчётами выводимого изображения, освобождая от этой обязанности центральный процессор, производит расчёты для обработки команд трёхмерной графики.
    видеоконтроллер — отвечает за формирование изображения в видеопамяти, даёт команды RAMDAC на формирование сигналов развёртки для монитора и осуществляет обработку запросов центрального процессора. Кроме этого, обычно присутствуют контроллер внешней шины данных (например, PCI или AGP), контроллер внутренней шины данных и контроллер видеопамяти.
  • видеопамять — выполняет роль кадрового буфера, в котором хранится изображение, генерируемое и постоянно изменяемое графическим процессором и выводимое на экран монитора (или нескольких мониторов). цифро-аналоговый преобразователь (ЦАП, RAMDAC – Random Access Memory Digital-to-Analog Converter) — служит для преобразования изображения, формируемого видеоконтроллером, в уровни интенсивности цвета, подаваемые на аналоговый монитор.
  • видео-ПЗУ (Video ROM) — постоянное запоминающее устройство, в которое записаны видео-BIOS, экранные шрифты, служебные таблицы и т. п. ПЗУ не используется видеоконтроллером напрямую — к нему обращается только центральный процессор.
  • система охлаждения — предназначена для сохранения температурного режима видеопроцессора и видеопамяти в допустимых пределах.

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

Для примера приведем некоторые числа:
Если у Intel 3.0 GHz Pentium 4 пиковая вычислительная мощность оценивается как 12 GFLOPs, то у видеокарты ATI Radeon X1800XT 120 GFLOPs, для этого же оборудования пиковая скорость обмена данными между процессором и соответствующей памятью составляет соответственно 5.96 GB/s у CPU против 42 GB/s у GPU1.

На рисунке 1 показана производительность видеокарт разных поколений с разрешением 1920х1080:

Но GPU не только быстры и мощны, но и что важно, их мощности растут очень быстрыми темпами.

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

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

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

  • операции над интерполированными значениями;
  • доступ к текстурам;
  • наложение текстур;
  • создание эффекта дымки;
  • наложение цветов.

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

2 Особенности архитектура GPU

2.1 Конвейер GPU

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

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

На вход подается набор вершин, поступающий из приложения затем вершины, обрабатываются в вершинных процессорах, классифицируемых как MIMD, программа для вершинных процессоров называется вершинным шейдером. После этого результат работы вершинного шейдера поступает на сборку примитивов, таких как полигоны или линии, далее следует тесты видимости отсечения и прочие стандартные операции компьютерной графики, после них данные поступают на растеризацию. Здесь объемное изображение проецируется на плоский экран, и полученная картинка масштабируется соответствии с параметрами окна приложения. Результатом этой операции является картинка, текстура представляющая из себя прямоугольный массив так называемых пикселей. Пиксель может быть представлен 1-4 числами, например, в формате RGBA, те red, green, blue, alpha (коэфицент прозрачности).

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

Схематичное изображение работы GPU-конвейера:

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

2.2 Введение в модель программирования на GPU

Программы, работающие с трёхмерной графикой и видео (игры, GIS, CAD, CAM и др.), используют шейдеры для определения параметров геометрических объектов или изображения, для изменения изображения (для создания эффектов сдвига, отражения, преломления, затемнения с учётом заданных параметров поглощения и рассеяния света, для наложения текстур на геометрические объекты и др.).


Для решения проблемы в видеокарты стали добавлять (аппаратно) алгоритмы, востребованные разработчиками. Вскоре стало ясно, что реализовать все алгоритмы невозможно и нецелесообразно; решили дать разработчикам доступ к видеокарте — позволить собирать блоки графического процессора в произвольные конвейеры, реализующие разные алгоритмы. Программы, предназначенные для выполнения на процессорах видеокарты, получили название «шейдеры». Были разработаны специальные языки для составления шейдеров. Теперь в видеокарты загружались не только данные о геометрических объектах («геометрия»), текстуры и другие данные, необходимые для рисования (формировании изображения), но и инструкции для GPU.

Разделяют два вида устройств – то которое управляет общей логикой – host, и то которое умеет быстро выполнить некоторый набор инструкций над большим объемом данных – device.

В роли хоста обычно выступает центральный процессор (CPU – например i5/i7).

В роли вычислительного устройства – видеокарта (GPU – GTX690/HD7970). Видеокарта содержит Compute Units – процессорные ядра. Неразбериху вводят и производители NVidia называет свои Streaming Multiprocessor unit или SMX , а ATI – SIMD Engine или Vector Processor. В современных игровых видеокартах – их 8-32.

Процессорные ядра могут исполнять несколько потоков за счет того, что в каждом содержится несколько (8-16) потоковых процессоров (Stream Cores или Stream Processor). Для карт NVidia – вычисления производятся непосредственно на потоковых процессорах, но ATI ввели еще один уровень абстракции – каждый потоковый процессор, состоит из processing elements – PE (иногда называемых ALU – arithmetic and logic unit) – и вычисления происходят на них.

Архитектура модели GPU:

2.3 Модель памяти GPU

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

  1. Регистровая;
  2. Разделяемая;
  3. Локальная;
  4. Глобальная;
  5. Константная;
  6. Текстурная

GPU так же как CPU имеет свои собственные регистры и кэш, для ускорения доступа к данным во время вычислений. Помимо этого, GPU имеет свою собственную основную память (графическая память), поэтому для того, чтобы программист мог выполнять вычисления на GPU он должен предварительно передать данные из оперативной памяти CPU в память GPU. Эта операция традиционно является достаточно дорогостоящей с точки зрения производительности ввиду относительно не высокой скорости передачи данных между памятью приложения и памятью видеокарты, хотя современны шины PCI Express и специальные чипы на материнской плате (такие как NV3 и NV4) заметно ускорили этот процесс.

В отличие от памяти CPU память GPU имеет некоторые серьезные ограничения, и доступ к ней может быть осуществлен только посредством некоторых абстракций графического программного интерфейса. Каждая такая абстракция может считаться как тип потока со своим собственным набором правил доступа. Таких потоков непосредственно доступных для программиста может быть выделено 4 – потоки вершин потоки фрагментов потоки буфера кадров потоки текстур.

Модель памяти в CUDA отличается возможностью побайтной адресации, поддержкой как gather, так и scatter. Доступно довольно большое количество регистров на каждый потоковый процессор, до 1024 штук. Доступ к ним очень быстрый, хранить в них можно 32-битные целые или числа с плавающей точкой.

2.4 Архитектура графических адаптеров Nvidia

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

Первой видеокартой на основе графического процессора GK104 стала GeForce GTX 680. Она заменила переставший выпускаться GTX 580 на базе GF110. С одной стороны, принцип наименования видеокарт NVIDIA вроде бы не изменился, топовая модель получила изменение первой цифры индекса. С другой — даже судя по кодовому имени чипа GK104, он изначально вряд ли планировался в роли именно топового решения.

Вполне вероятно, изначально это должна была быть GTX 670 (приставка Ti по вкусу) или что-то в этом роде, но потом решили повременить с настоящим топовым чипом, раз у TSMC с новым техпроцессом до сих пор дела неидеальны, и выпустить в качестве верхней модели менее мощный чип. Впрочем, мы не раз говорили о том, что наименование видеокарт всегда является маркетинговым решением, которое не особенно влияет на технические характеристики.
Правда, оно сильно влияет на розничную цену решений. По всем внешним признакам (сложность чипа, сложность печатной платы, энергопотребление, да и себестоимость, скорее всего), GTX 680 больше похожа на решение из верхнего среднего диапазона. На это же намекают и кодовые имена чипов: GF104 — GK104. Впрочем, после выхода видеоплат AMD по высоким ценам, которые не слишком сильно обогнали GeForce GTX 580, у NVIDIA появился большой соблазн поднять GK104 в верхний ценовой диапазон, заработав самим и давая заработать своим партнёрам по выпуску видеокарт.

Все вычислительные ядра на графическом адаптере объединены в независимые блоки TPC (Texture process cluster) количество которых зависит как от версии чипа (G80 – максимум 8, G200 – максимум 10), так и просто от конкретной видеокарты внутри чипа (GeForce 220GT – 2, GeForce 275 – 10). Так же, как от видеокарты к видеокарте может меняться количество TPC, так же может меняться и количество DRAM партиций и соот ветственно общий объем оперативной памяти. DRAM партиции имеют кэш второго уровня и объединены между собой коммуникационной сетью, в которую, так же подключены все TPC. Любая видеокарта подключается к CPU через мост, который может быть, как интегрирован в CPU (Intel Core i7), так и дискретным (Intel Core 2 Duo). От чипа к чипу (от G80 до G200) менялись детали в TPC, а общая архитектура оставалась одинаковой. В новом чипе Fermi произошли изменения и в общей архитектуре, поэтому о нем речь пойдет отдельно.

Архитектура чипа G80

Любой из 128 потоковых процессоров G80 представляет собой обычное вычислительное устройство, способное работать с данными в формате с плавающей запятой. Следовательно, он не только способен обрабатывать шейдеры любого типа – вершинные, пиксельные или геометрические, но и использоваться для просчета физической модели или других расчетов, в рамках концепции Compute Unified Device Architecture (CUDA), причем, независимо от других процессоров. Иными словами, одна часть GeForce 8800 может заниматься какими-либо расчетами, а другая, к примеру, визуализацией их результатов, благо, поточная архитектура позволяет использовать результаты, полученные на выходе одного из процессоров в качестве источника данных для другого.

Создавая GeForce 8800, NVIDIA заботилась не только о производительности, но и о качестве изображения – свойстве, в котором предыдущие разработки компании зачастую уступали изделиям ATI Technologies. Изменения и улучшения коснулись как полноэкранного сглаживания, так и анизотропной фильтрации.

Основные составляющие TPC:

  • TEX – логика работы с текстурами, содержит в себе участки конвейера, предназначенные для обработки особой текстурной памяти о которой речь пойдет позже.
  • SM – потоковый мультипроцессор, самостоятельный вычислительный модуль, именно на нем осуществляется выполнение блока. В архитектуре чипа G80 в одном TPC находится 2 SM.

Основные составляющие SM:

  • SP – потоковый процессор, непосредственно вычислительный модуль, способен совершать арифметические операции с целочисленными операндами и с операндами с плавающей точкой (одинарная точность). Не является самостоятельной единицей, управляется SM.
  • SFU – модуль сложных математических функций. Проводит вычисления сложных математических функций (exp, sqr, log). Использует вычислительные мощности SP.
  • Регистровый файл – единый банк регистров, на каждом SM имеется 32Кб. Самый быстрый тип памяти на графическом адаптере.
  • Разделяемая память – специальный тип памяти, предназначенный для совместного использования данных тредов из одного блока. На каждом SM – 16Кб разделяемой памяти.
  • Кэш констант – место кэширования особого типа памяти (константной). Об особенностях применения речь пойдет позже.
  • Кэш инструкций, блок выборки инструкций – управляющая система SM. Не играет роли при программировании.

Итого, чип G80 имеет максимально 128 вычислительных модулей (SP) способных выполнять вычисления с целыми числами и числами с плавающей точкой с одинарной точностью. Такого функционала было недостаточно для многих научных задач, требовалась двойная точность.

3 Особенности программирования GPGPU (OpenCL)

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

Модель платформы (platform model) дает высокоуровневое описание гетерогенной системы. Центральным элементом данной модели выступает понятие хоста (host) — первичного устройства, которое управляет OpenCL-вычислениями и осуществляет все взаимодействия с пользователем. Хост всегда представлен в единственном экземпляре, в то время как OpenCL-устройства (devices), на которых выполняются OpenCL-инструкции могут быть представлены во множественном числе. OpenCL-устройством может быть CPU, GPU, DSP или любой другой процессор в системе, поддерживающийся установленными в системе OpenCL-драйверами. OpenCL-устройства логически делятся моделью на вычислительные модули (compute units), которые в свою очередь делятся на обрабатывающие элементы (processing elements). Вычисления на OpenCL-устройствах в действительности происходят на обрабатывающих элементах. На рисунке схематически изображена OpenCL-платформа из 3-х устройств:

OpenCL первоначально был разработан в компании Apple Inc. Apple внесла предложения по разработке спецификации в комитет Khronos. Вскоре компания AMD решила поддержать разработку OpenCL (и DirectX 11), который должен заменить фреймворк Close to Metal.

Программная модель OpenCL

Ядро (kernel) — функция, исполняемая устройством. Имеет в описании спецификацию __kernel.

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

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

Команда (command) — операция OpenCL, предназначенная для исполнения (исполнение ядра на устройстве, манипуляции с памятью и т. д.)

Объекты OpenCL

Объект (object) — абстрактное представление ресурса, управляемого OpenCL API (объект ядра, памяти и т. д.)
Дескриптор (handle) — непрозрачный тип, ссылающийся на объект, выделяемый OpenCL. Любая операция с объектом выполняется через дескриптор.

Очередь команд (command-queue) — объект, содержащий команды для исполнения на устройстве.

Объект ядра (kernel object) — хранит отдельную функцию ядра программы вместе со значениями аргументов.

Объект события (event object) — хранит состояние команды. Предназначен для синхронизации.

Объект буфера (buffer object) — последовательный набор байт. Доступен из ядра через указатель и из управляющего узла при помощи вызовов API.

Объект памяти (memory object) — ссылается на область глобальной памяти.

Устройства OpenCL

Рабочий элемент (work-item) — набор параллельно исполняемых ядер на устройстве, вызванных при помощи команды.

Рабочая группа (work-group) — набор взаимодействующих рабочих элементов, исполняющихся на одном устройстве.

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

Обрабатывающий элемент (processing element) — виртуальный скалярный процессор. Рабочий элемент может выполняться на одном или нескольких обрабатывающих элементах.

Вычислительный узел (compute unit) — исполняет одну рабочую группу. Устройство может состоять из одного или нескольких вычислительных элементов.

Контекст

Контекст (context) — среда, в которой выполняются ядра, а также область определения синхронизации и управления памятью. Включает:

  • набор устройств;
  • память, доступную устройствам;
  • свойства памяти;
  • одну или несколько очередей команд.

Объект программы (program object) — включает:

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

Пример программирования GPU на языке MC#

Базовая структура программы на языке MC#, предназначенной для исполнения на графическом процессоре, состоит из:

  1. описания конфигурации графического процессора, в котором, в частности, задается количество (параллельных) потоков, запускаемых на GPU, и параметры их объединения в блоки и решетку;
  2. gpu-функции (метода), которая будет исполняться в рамках одного вычислительного потока на GPU.

Gpu-метод в программе задается путем приписывания его определению модификатора gpu синтаксически располагающегося на месте типа возвращаемого значения (см. ниже пример gpu-функции vecadd). Определение конфигурации GPU производится с помощью создания объекта класса GpuConfig и задания его параметров.

Ниже представлен полный текст простой программы на языке MC#, предназначенной для сложения двух векторов целых чисел с использованием GPU. В этой программе, исходные векторы ? и ?, а также результирующий вектор ? имеют длину ?. Данное число служит размером (одномерного) блока потоков, запускаемых на GPU – соответственно, ?-ый поток выполняет сложение ?-ых компонентов ?[?] и ?[?] исходных векторов. Также в данной программе предполагается, что длина векторов ? не превосходит размера блока потоков, допускаемого для конкретного графического устройства, на котором предполагается исполнение данной программы.

На примере этой программы, отметим некоторые ключевые особенности MC#-программ, предназначенных для исполнения на GPU, которые будут детализированы в последующих разделах:

  1. Для исполнения на GPU, требуется использование библиотеки GPU.NET, включенной в состав системы программирования MC#.
    Ссылка на эту библиотеку задается с помощью оператора using :
    using GpuDotNet.Cuda;
  2. Основными методами, устанавливающими параметры конфигурации графического процессора, являются:
    • SetDeviceNumber (номер графического устройства);
    • SetGridSize (размеры решетки потоков);
    • SetBlockSize (размеры блоков потоков).

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

Программно-аппаратная архитектура параллельных вычислений CUDA

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

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


Параллельные вычисления с CUDA

Направление вычислений эволюционирует от «централизованной обработки данных» на центральном процессоре до «совместной обработки» на CPU(центральное процессорное устройство) и GPU. Для реализации новой вычислительной парадигмы компания NVIDIA изобрела архитектуру параллельных вычислений CUDA, на данный момент представленную в графических процессорах GeForce, ION, Quadro и Tesla и обеспечивающую необходимую базу разработчикам ПО.

Говоря о потребительском рынке, стоит отметить, что почти все основные приложения для работы с видео уже оборудованы, либо будут оснащены поддержкой CUDA-ускорения, включая продукты от Elemental Technologies, MotionDSP и LoiLo.

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

На финансовом рынке компании Numerix и CompatibL анонсировали поддержку CUDA в новом приложении анализа риска контрагентов и достигли ускорения работы в 18 раз. Numerix используется почти 400 финансовыми институтами.

Показателем роста применения CUDA является также рост использования графических процессоров Tesla в GPU вычислениях. На данный момент более 700 GPU кластеров установлены по всему миру в компаниях из списка Fortune 500, таких как Schlumberger и Chevron в энергетическом секторе, а также BNP Paribas в секторе банковских услуг.

Благодаря относительно недавно выпущенным системам Microsoft Windows 7 и Apple Snow Leopard, вычисления на GPU займут свои позиции в секторе массовых решений. В этих новых операционных системах GPU предстанет не только графическим процессором, но также и универсальным процессором для параллельных вычислений, работающим с любым приложением.

CUDA в электроэнергетике

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

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

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

Было проведено тестирование трех базовых алгоритмов и разработанного гибридного подхода для построения прогноза энергопотребления. Для проведения исследования были рассмотрены показатели энергопотребления в Костромской области за 2005 год.

В качестве критерия оптимальности построенного прогноза был использован индекс Тейла:

Индекс Тейла U(X, Y) измеряет несовпадение временных рядов X(t) и Y(t), и чем ближе он к нулю, тем ближе сравниваемые ряды. По результатам испытаний данный показатель при прогнозировании с помощью эволюционного моделирования составил 0,0436, фильтра Винера — 0,0329, нейронной сети — 0,0322, гибридного алгоритма — 0,0239. Анализ полученных данных показывает, что разработка гибридного алгоритма позволила решить поставленную задачу — повышение точности построения прогноза энергопотребления.

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

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

Технология параллельного программирования CUDA

вторник, 3 февраля 2015 г.

— Русская Народная Сказка

1. Введение

В этой статье мы поговорим об одной из популярных технологий высокопроизводительных вычислений использующей GPU (graphics processing units). Изначально этот класс устройств разрабатывался для обработки графики. Техника совершенствовалась, GPU наращивали производительность и в какой-то момент оказалось, что GPU можно успешно использовать не только для задач компьютерной графики но и как математический сопроцессор для CPU машины, получая при этом существенный прирост в производительности.

Этот класс технологий получил название GPGPU (General-Purpose Graphics Processing Units) — использование графического процессора для общего назначения. Популярный представитель этого класса — CUDA (Compute Unified Device Architecture), она была впервые представлена компанией NVIDIA в 2007 году [1]. CUDA может использоваться на GPU производства NVIDIA, таких как GeForce, Quadro, Tesla. Последняя линейка устройств (Tesla) видеовыхода вообще не имеют и предназначены исключительно для высокопроизводительных вычислений.

CUDA работает только с устройствами производства NVIDIA, но это не беда, помимо CUDA существуют и другие аналогичные технологии, например OpenCL и AMD FireStream, но их описание выходит за рамки этой статьи.

2. Аппаратная чаcть

Архитектура GPU построена иначе, чем у универсальных CPU [2], и в ней изначально заложена определенная специализация. Задачи компьютерной графики предполагают независимую параллельную обработку данных и GPU изначально предназначен для параллельных вычислений. Он спроектирован так, чтобы выполнять большое количество тредов (элементарных параллельных процессов). GPU содержит много относительно простых арифметико-логических устройств (Рис.2), которые объединены в группы, и реализует модель параллельного вычислителя над общей памятью, но с некоторыми особенностями.

GPU ориентирован на выполнение программ с большим объемом вычислений (распараллеливание по данным типа SIMD). Память GPU имеет иерархическую структуру и оптимизирована под максимальную пропускную способность. Вместо системы кэшей CPU и сложных арифметико-логических схем (АЛУ), GPU имеет много упрощённых АЛУ, имеющих общую память. Это помогает повысить производительность в вычислительных задачах но несколько усложняет программирование. Для достижения наилучшего ускорения необходимо продумывать стратегии доступа к памяти и учитывать аппаратные особенности.

GPU представляет собой массив потоковых процессоров (Streaming Processor Array), состоящий из кластеров текстурных процессоров (Texture Processor Clusters, TPC). TPC состоит из набора мультипроцессоров (SM – Streaming Multi-processor), каждый из которых содержит несколько потоковых процессоров (SP – Streaming Processors) или ядер. Набор ядер каждого мультипроцессора работает по принципу SIMD (одиночный поток команд, множество потоков данных). На рис.3 представлена общая аппаратная схема работы мультипроцессоров (SM) GPU.

В сравнении с универсальным CPU, конструкция GPU налагает ряд дополнительных ограничений для программирования, они зависят от конкретной модели. Для NVIDIA Quadro FX1700 это следующие особенности: GPU не поддерживает рекурсию и вычисления с двойной (double) точностью (возможна только одинарная точность — float).

3. Программная часть

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

Существуют привязки CUDA для разных языков программирования, полный список можно посмотреть на сайте NVIDIA. Здесь для реализации тестовых примеров мы будем использовать CUDA C — адаптированный для программирования GPU диалект языка C++.

Пакет CUDA-разработчика для Linux содержит компилятор nvcc , который генерирует код для работы с GPU.

Каждое устройство CUDA имеет специальное свойство — уровень вычислительных возможностей (compute capability), которое определяет набор доступных устройству возможностей из всего функционала CUDA. Для NV >’-arch sm_11′ .

4. Программирование на CUDA C

На листинге 2 представлена простейшая CUDA C программа, она не делает ничего. Ядро, которое описано функцией kern , выполняется на GPU, эта функция имеет спецификатор __global__ , который говорит компилятору, что функция вызывается с CPU и выполняется на GPU. Для функций есть и другие спецификаторы(см.таб.1).

спецификатор выполняется вызывается
__host__ host host
__global__ device host
__device__ device device

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

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

Запуск ядра порождает процессы (треды) на GPU в соответствии с заданным параметрами (рис.4). Всё множество процессов, порождаемых запуском ядра, в терминологии CUDA называеться грид (gr >блоков (block), блок из тредов (thread). Тред это элементарный параллельный процесс. Треды в блоках и блоки в гриде могут быть представлены в виде одно-, двух- или трёхмерной решетки.

Для FX1700 максимальный размер грида 65535 x 65535 блоков. Максимально возможные значения индексов номера треда в блоке 512 x 512 x 64, но при этом количество тредов в блоке не должно превышать 512.

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

При конфигурировании топологии процессов необходимо учитывать аппаратные особенности. Здесь следует ввести понятие варпа. Варп (warp) это группа тредов, размер варпа для FX1700 — 32 треда. Полуварп (half-warp) — половина тредов варпа. Все треды одного варпа выполняются одновременно и синхронно (SIMD) на своём мультипроцессоре. При доступе треда к основной памяти GPU её части могут кэшироваться в локальной памяти данного мультипроцессора. Если все данные, которые требуются тредам полуварпа будут находиться в этом кэше то это может повысить производительность.

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

4.1. Основная (глобальная) память

4.2. О компоновке тредов и блоков

На листинге 4 представлен пример следующей организации тредов. Размер блока 8x8x8, что как раз равно 512 тредов на один блок, размер грида 8х32 блока, таким образом общее количество параллельных процессов 131072=8x8x8x8x32.

При этом, адресация выделенной памяти — линейная и сложный номер треда пересчитывается в индекс ячейки масива памяти.

Для того, что бы задать трех- или двухмерную решетку тредов и/или блоков необходимо определить переменную размеров типа dim3 , и передать её в качестве параметра (в угловых скобках) при запуске ядра. Кроме dim3 CUDA содержит и другие векторные типы char2, int3, float3 и др.

Для определения номера треда и размеров блока и грида в CUDA есть следующие системные переменные (таблица 2).

тип имя назначение
dim3 gridDim размер грида
dim3 blockDim размер блока
uint3 blockIdx номер блока в гриде
uint3 threadIdx номер треда в блоке
int warpSize количество тредов в варпе

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

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

На листинге 6 представлен пример работы с разделяемой памятью, которая объявляется внутри ядра с модификатором __shared__ . Каждый тред записывает в «свою» ячейку разделяемой памяти данные, затем изменяет данные в «чужой» ячейке, после этого результат переписывается в основную память GPU, а затем и в память CPU.

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

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

Рассмотрим простейший пример работы с текстурной памятью (листинг 7).

4.6. Атомарные операции

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

4.7. Параллельный запуск нескольких ядер на GPU

По умолчанию в CUDA-программе есть только одна исполняемая очередь или нуль-стрим. В представленной выше программе мы с помощью cudaStreamCreate() создаём ещё два стрима stream0, stream1 и запускаем два ядра во вновь созданные стримы, добавив соответствующий параметр при запуске ядра. Далее мы ждём пока оба ядра завершат работу, используя cudaStreamSynchronize() , переписываем результат в память CPU и закрываем оба отработанных стрима с помощью cudaStreamDestroy() .

4.8. Асинхронное копирование данных и блокированная память

Перед тем как асинхронно копировать данные из памяти хоста на девайс и обратно необходимо заблокировать эту область памяти на хосте. Блокировка области памяти (page-locked) означает, что ОС хоста запрещается перемещать данные из этой области физической памяти, т.е. к этой области памяти нельзя применять свопинг и т.п. Этот метод нужно использовать осторожно, потому как мы лишаемся всех преимуществ виртуальной памяти, и должны следить за тем, что бы у нас оставалось достаточно физической памяти для работы.

На листинге 10 приведен пример использования асинхронного копирования. В начале проверяем с помощью cudaGetDeviceProperties() поддерживает ли устройство overlap, выделяем блокированную память на хосте с помощью cudaHostAlloc() , выделяем память на GPU и запускаем асинхронное копирование данных на GPU cudaMemcpyAsync() , запускаем ядра в созданные стримы и асинхронное копирование результата с GPU, после этого ожидаем завершения работы стримов с помощью cudaStreamSynchronize() . Когда стримы завершатся — печатаем результат и освобождаем ресурсы — cudaFreeHost() .

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

4.9. Использование памяти CPU

В начале мы проверяем поддерживает ли устройство отображение памяти. Далее включаем эту функциональность с помощью cudaSetDeviceFlags(cudaDeviceMapHost) . Выделяем блокированную память CPU с помощью cudaHostAlloc() , при этом для входного буфера можно указать флаг cudaHostAllocWriteCombined , который устанавливает особый режим кэширования этой области памяти для CPU, это ускорят чтение этих данных для GPU, но при этом скорость работы с ними CPU существенно падает. Отображаем выделенные буферы памяти CPU в память GPU с помощью cudaHostGetDevicePointer() и запускаем ядро. Перед тем как печатать результат необходимо вызвать cudaThreadSynchronize() , что бы буферы памяти GPU и CPU синхронизировались. В конце работы освобождаем ресурсы.

4.10. Операции с матрицами

На листинге 12 представлен пример использования cuBLAS, эта программа генерирует две матрицы чисел одинакового размера, выполняет с ними действия — сложение, транспонирование и умножение. В начале мы выделяем память для данных обычным способом, далее инициализируем cuBLAS с помощью cublasCreate() , копируем данные host на device с помощью cublasSetMatrix() . После этого выполняем операции с матрицами: сложение — cublasSgeam() , умножение — cublasSgemm() , домножение на скаляр — cublasSscal() . Каждая из этих функций cuBLAS самостоятельно формирует и запускает на device нужное количество процессов, при этом каждая такая функция может выполнять сразу несколько операций, например — транспонирование и умножение. После завершения вычислений результат копируем обратно на host с помощью cublasGetMatrix() , печатаем результат и освобождаем ресурсы — cublasDestroy() .

4.11. Оценка затраченного на вычисления времени

На листинге 13 приведена чать кода, которая замеряет время выполнения программы. В начале создаём временные метки start, stop с помощью cudaEventCreate() , далее фиксируем время старта с помощью cudaEventRecord() , выполняем вычисления, фиксируем время завершения. Поскольку некоторые операции могут выполняться асинхронно Вызываем cudaEventSynchronize() , что бы метка записалась корректно. Далее вычисляем разницу во времени с помощью cudaEventElapsedTime() и печатаем результат.

Мастер Йода рекомендует:  Интервью для программистов
Добавить комментарий