Что такое NVIDIA CUDA? Nvidia CUDA? неграфические вычисления на графических процессорах карта с технологией cuda
И предназначен для трансляции host-кода (главного, управляющего кода) и device-кода (аппаратного кода) (файлов с расширением.cu) в объектные файлы, пригодные в процессе сборки конечной программы или библиотеки в любой среде программирования, например в NetBeans .
В архитектуре CUDA используется модель памяти грид , кластерное моделирование потоков и SIMD -инструкции. Применима не только для высокопроизводительных графических вычислений, но и для различных научных вычислений с использованием видеокарт nVidia. Ученые и исследователи широко используют CUDA в различных областях, включая астрофизику , вычислительную биологию и химию, моделирование динамики жидкостей, электромагнитных взаимодействий, компьютерную томографию, сейсмический анализ и многое другое. В CUDA имеется возможность подключения к приложениям, использующим OpenGL и Direct3D . CUDA - кроссплатформенное программное обеспечение для таких операционных систем как Linux , Mac OS X и Windows .
22 марта 2010 года nVidia выпустила CUDA Toolkit 3.0, который содержал поддержку OpenCL .
Оборудование
Платформа CUDA Впервые появились на рынке с выходом чипа NVIDIA восьмого поколения G80 и стала присутствовать во всех последующих сериях графических чипов, которые используются в семействах ускорителей GeForce , Quadro и NVidia Tesla .
Первая серия оборудования, поддерживающая CUDA SDK, G8x, имела 32-битный векторный процессор одинарной точности , использующий CUDA SDK как API (CUDA поддерживает тип double языка Си, однако сейчас его точность понижена до 32-битного с плавающей запятой). Более поздние процессоры GT200 имеют поддержку 64-битной точности (только для SFU), но производительность значительно хуже, чем для 32-битной точности (из-за того, что SFU всего два на каждый потоковый мультипроцессор, а скалярных процессоров - восемь). Графический процессор организует аппаратную многопоточность, что позволяет задействовать все ресурсы графического процессора. Таким образом, открывается перспектива переложить функции физического ускорителя на графический ускоритель (пример реализации - nVidia PhysX). Также открываются широкие возможности использования графического оборудования компьютера для выполнения сложных неграфических вычислений: например, в вычислительной биологии и в иных отраслях науки.
Преимущества
По сравнению с традиционным подходом к организации вычислений общего назначения посредством возможностей графических API, у архитектуры CUDA отмечают следующие преимущества в этой области:
Ограничения
- Все функции, выполнимые на устройстве, не поддерживают рекурсии (в версии CUDA Toolkit 3.1 поддерживает указатели и рекурсию) и имеют некоторые другие ограничения
Поддерживаемые GPU и графические ускорители
Перечень устройств от производителя оборудования Nvidia с заявленной полной поддержкой технологии CUDA приведён на официальном сайте Nvidia: CUDA-Enabled GPU Products (англ.) .
Фактически же, в настоящее время на рынке аппаратных средств для ПК поддержку технологии CUDA обеспечивают следующие периферийные устройства :
| Версия спецификации | GPU | Видеокарты |
|---|---|---|
| 1.0 | G80, G92, G92b, G94, G94b | GeForce 8800GTX/Ultra, 9400GT, 9600GT, 9800GT, Tesla C/D/S870, FX4/5600, 360M, GT 420 |
| 1.1 | G86, G84, G98, G96, G96b, G94, G94b, G92, G92b | GeForce 8400GS/GT, 8600GT/GTS, 8800GT/GTS, 9600 GSO, 9800GTX/GX2, GTS 250, GT 120/30/40, FX 4/570, 3/580, 17/18/3700, 4700x2, 1xxM, 32/370M, 3/5/770M, 16/17/27/28/36/37/3800M, NVS420/50 |
| 1.2 | GT218, GT216, GT215 | GeForce 210, GT 220/40, FX380 LP, 1800M, 370/380M, NVS 2/3100M |
| 1.3 | GT200, GT200b | GeForce GTX 260, GTX 275, GTX 280, GTX 285, GTX 295, Tesla C/M1060, S1070, Quadro CX, FX 3/4/5800 |
| 2.0 | GF100, GF110 | GeForce (GF100) GTX 465, GTX 470, GTX 480, Tesla C2050, C2070, S/M2050/70, Quadro Plex 7000, Quadro 4000, 5000, 6000, GeForce (GF110) GTX 560 TI 448, GTX570, GTX580, GTX590 |
| 2.1 | GF104, GF114, GF116, GF108, GF106 | GeForce 610M, GT 430, GT 440, GTS 450, GTX 460, GTX 550 Ti, GTX 560, GTX 560 Ti, 500M, Quadro 600, 2000 |
| 3.0 | GK104, GK106, GK107 | GeForce GTX 690, GTX 680, GTX 670, GTX 660 Ti, GTX 660, GTX 650 Ti, GTX 650, GT 640, GeForce GTX 680MX, GeForce GTX 680M, GeForce GTX 675MX, GeForce GTX 670MX, GTX 660M, GeForce GT 650M, GeForce GT 645M, GeForce GT 640M |
| 3.5 | GK110 |
|
|
|
|
|
- Модели Tesla C1060, Tesla S1070, Tesla C2050/C2070, Tesla M2050/M2070, Tesla S2050 позволяют производить вычисления на GPU с двойной точностью.
Особенности и спецификации различных версий
| Feature support (unlisted features are supported for all compute capabilities) |
Compute capability (version) | ||||
|---|---|---|---|---|---|
| 1.0 | 1.1 | 1.2 | 1.3 | 2.x | |
32-bit words in global memory |
Нет | Да | |||
floating point values in global memory |
|||||
| Integer atomic functions operating on 32-bit words in shared memory |
Нет | Да | |||
| atomicExch() operating on 32-bit floating point values in shared memory |
|||||
| Integer atomic functions operating on 64-bit words in global memory |
|||||
| Warp vote functions | |||||
| Double-precision floating-point operations | Нет | Да | |||
| Atomic functions operating on 64-bit integer values in shared memory |
Нет | Да | |||
| Floating-point atomic addition operating on 32-bit words in global and shared memory |
|||||
| _ballot() | |||||
| _threadfence_system() | |||||
| _syncthreads_count(), _syncthreads_and(), _syncthreads_or() |
|||||
| Surface functions | |||||
| 3D grid of thread block | |||||
| Technical specifications | Compute capability (version) | ||||
|---|---|---|---|---|---|
| 1.0 | 1.1 | 1.2 | 1.3 | 2.x | |
| Maximum dimensionality of grid of thread blocks | 2 | 3 | |||
| Maximum x-, y-, or z-dimension of a grid of thread blocks | 65535 | ||||
| Maximum dimensionality of thread block | 3 | ||||
| Maximum x- or y-dimension of a block | 512 | 1024 | |||
| Maximum z-dimension of a block | 64 | ||||
| Maximum number of threads per block | 512 | 1024 | |||
| Warp size | 32 | ||||
| Maximum number of resident blocks per multiprocessor | 8 | ||||
| Maximum number of resident warps per multiprocessor | 24 | 32 | 48 | ||
| Maximum number of resident threads per multiprocessor | 768 | 1024 | 1536 | ||
| Number of 32-bit registers per multiprocessor | 8 K | 16 K | 32 K | ||
| Maximum amount of shared memory per multiprocessor | 16 KB | 48 KB | |||
| Number of shared memory banks | 16 | 32 | |||
| Amount of local memory per thread | 16 KB | 512 KB | |||
| Constant memory size | 64 KB | ||||
| Cache working set per multiprocessor for constant memory | 8 KB | ||||
| Cache working set per multiprocessor for texture memory | Device dependent, between 6 KB and 8 KB | ||||
| Maximum width for 1D texture |
8192 | 32768 | |||
| Maximum width for 1D texture reference bound to linear memory |
2 27 | ||||
| Maximum width and number of layers for a 1D layered texture reference |
8192 x 512 | 16384 x 2048 | |||
| Maximum width and height for 2D texture reference bound to linear memory or a CUDA array |
65536 x 32768 | 65536 x 65535 | |||
| Maximum width, height, and number of layers for a 2D layered texture reference |
8192 x 8192 x 512 | 16384 x 16384 x 2048 | |||
| Maximum width, height and depth for a 3D texture reference bound to linear memory or a CUDA array |
2048 x 2048 x 2048 | ||||
| Maximum number of textures that can be bound to a kernel |
128 | ||||
| Maximum width for a 1D surface reference bound to a CUDA array |
Not supported |
8192 | |||
| Maximum width and height for a 2D surface reference bound to a CUDA array |
8192 x 8192 | ||||
| Maximum number of surfaces that can be bound to a kernel |
8 | ||||
| Maximum number of instructions per kernel |
2 million | ||||
Пример
CudaArray* cu_array; texture< float , 2 > tex; // Allocate array cudaMalloc( & cu_array, cudaCreateChannelDesc< float> () , width, height ) ; // Copy image data to array cudaMemcpy( cu_array, image, width* height, cudaMemcpyHostToDevice) ; // Bind the array to the texture cudaBindTexture( tex, cu_array) ; // Run kernel dim3 blockDim(16 , 16 , 1 ) ; dim3 gridDim(width / blockDim.x , height / blockDim.y , 1 ) ; kernel<<< gridDim, blockDim, 0 >>> (d_odata, width, height) ; cudaUnbindTexture(tex) ; __global__ void kernel(float * odata, int height, int width) { unsigned int x = blockIdx.x * blockDim.x + threadIdx.x ; unsigned int y = blockIdx.y * blockDim.y + threadIdx.y ; float c = texfetch(tex, x, y) ; odata[ y* width+ x] = c; }
Import pycuda.driver as drv import numpy drv.init () dev = drv.Device (0 ) ctx = dev.make_context () mod = drv.SourceModule (""" __global__ void multiply_them(float *dest, float *a, float *b) { const int i = threadIdx.x; dest[i] = a[i] * b[i]; } """ ) multiply_them = mod.get_function ("multiply_them" ) a = numpy.random .randn (400 ) .astype (numpy.float32 ) b = numpy.random .randn (400 ) .astype (numpy.float32 ) dest = numpy.zeros_like (a) multiply_them( drv.Out (dest) , drv.In (a) , drv.In (b) , block= (400 , 1 , 1 ) ) print dest-a*b
CUDA как предмет в вузах
По состоянию на декабрь 2009 года, программная модель CUDA преподается в 269 университетах по всему миру. В России обучающие курсы по CUDA читаются в Санкт-Петербургском политехническом университете , Ярославском государственном университете им. П. Г. Демидова , Московском , Нижегородском , Санкт-Петербургском , Тверском , Казанском , Новосибирском , Новосибирском государственном техническом университете Омском и Пермском государственных университетах, Международном университете природы общества и человека «Дубна» , Ивановском государственном энергетическом университете , Белгородский государственный университет , МГТУ им. Баумана , РХТУ им. Менделеева , Межрегиональном суперкомпьютерном центре РАН, . Кроме того, в декабре 2009 года было объявлено о начале работы первого в России научно-образовательного центра «Параллельные вычисления», расположенного в городе Дубна , в задачи которого входят обучение и консультации по решению сложных вычислительных задач на GPU.
На Украине курсы по CUDA читаются в Киевском институте системного анализа.
Ссылки
Официальные ресурсы
- CUDA Zone (рус.) - официальный сайт CUDA
- CUDA GPU Computing (англ.) - официальные веб-форумы, посвящённые вычислениям CUDA
Неофициальные ресурсы
Tom"s Hardware- Дмитрий Чеканов. nVidia CUDA: вычисления на видеокарте или смерть CPU? . Tom"s Hardware (22 июня 2008 г.). Архивировано
- Дмитрий Чеканов. nVidia CUDA: тесты приложений на GPU для массового рынка . Tom"s Hardware (19 мая 2009 г.). Архивировано из первоисточника 4 марта 2012. Проверено 19 мая 2009.
- Алексей Берилло. NVIDIA CUDA - неграфические вычисления на графических процессорах. Часть 1 . iXBT.com (23 сентября 2008 г.). Архивировано из первоисточника 4 марта 2012. Проверено 20 января 2009.
- Алексей Берилло. NVIDIA CUDA - неграфические вычисления на графических процессорах. Часть 2 . iXBT.com (22 октября 2008 г.). - Примеры внедрения NVIDIA CUDA. Архивировано из первоисточника 4 марта 2012. Проверено 20 января 2009.
- Боресков Алексей Викторович. Основы CUDA (20 января 2009 г.). Архивировано из первоисточника 4 марта 2012. Проверено 20 января 2009.
- Владимир Фролов. Введение в технологию CUDA . Сетевой журнал «Компьютерная графика и мультимедиа» (19 декабря 2008 г.). Архивировано из первоисточника 4 марта 2012. Проверено 28 октября 2009.
- Игорь Осколков. NVIDIA CUDA – доступный билет в мир больших вычислений . Компьютерра (30 апреля 2009 г.). Проверено 3 мая 2009.
- Владимир Фролов. Введение в технологию CUDA (1 августа 2009 г.). Архивировано из первоисточника 4 марта 2012. Проверено 3 апреля 2010.
- GPGPU.ru . Использование видеокарт для вычислений
- . Центр Параллельных Вычислений
Примечания
См. также
| Nvidia | ||||||||
|---|---|---|---|---|---|---|---|---|
| Графические процессоры |
| |||||||
Согласно Дарвинской теории эволюции, первая человекообразная обезьяна (если
быть точным – homo antecessor, человек-предшественник) превратилась впоследствии
в нас. Многотонные вычислительные центры с тысячью и больше радиоламп,
занимающие целые комнаты, сменились полукилограммовыми ноутами, которые, кстати,
не уступят в производительности первым. Допотопные печатные машинки превратились
в печатающие что угодно и на чем угодно (даже на теле человека)
многофункциональные устройства. Процессорные гиганты вдруг вздумали замуровать
графическое ядро в «камень». А видеокарты стали не только показывать картинку с
приемлемым FPS и качеством графики, но и производить всевозможные вычисления. Да
еще как производить! О технологии многопоточных вычислений средствами GPU, и пойдет речь.
Почему GPU?
Интересно, почему всю вычислительную мощь решили переложить на графический
адаптер? Как видно, процессоры еще в моде, да и вряд ли уступят свое теплое
местечко. Но у GPU есть пара козырей в рукаве вместе с джокером, да и рукавов
хватает. Современный центральный процессор заточен под получение максимальной
производительности при обработке целочисленных данных и данных с плавающей
запятой, особо не заботясь при этом о параллельной обработке информации. В то же
время архитектура видеокарты позволяет быстро и без проблем «распараллелить»
обработку данных. С одной стороны, идет обсчет полигонов (за счет 3D-конвейера),
с другой – пиксельная обработка текстур. Видно, что происходит «слаженная
разбивка» нагрузки в ядре карты. Кроме того, работа памяти и видеопроцессора
оптимальнее, чем связка «ОЗУ-кэш-процессор». В тот момент, когда единица данных
в видеокарте начинает обрабатываться одним потоковым процессором GPU, другая
единица параллельно загружается в другой, и, в принципе, легко можно достичь
загруженности графического процессора, сравнимой с пропускной способностью шины,
однако для этого загрузка конвейеров должна осуществляться единообразно, без
всяких условных переходов и ветвлений. Центральный же процессор в силу своей
универсальности требует для своих процессорных нужд кэш, заполненный
информацией.
Ученые мужи задумались насчет работы GPU в параллельных вычислениях и
математике и вывели теорию, что многие научные расчеты во многом схожи с
обработкой 3D-графики. Многие эксперты считают, что основополагающим фактором в
развитии GPGPU
(General Purpose computation on GPU – универсальные
расчеты средствами видеокарты
) стало появление в 2003 году проекта Brook GPU.
Создателям проекта из Стэндфордского университета предстояло решить непростую
проблему: аппаратно и программно заставить графический адаптер производить
разноплановые вычисления. И у них это получилось. Используя универсальный язык C,
американские ученые заставили работать GPU как процессор, с поправкой на
параллельную обработку. После Brook появился целый ряд проектов по VGA-расчетам,
таких как библиотека Accelerator, библиотека Brahma, система
метапрограммирования GPU++ и другие.
CUDA!
Предчувствие перспективности разработки заставило AMD
и NVIDIA
вцепиться в Brook GPU, как питбуль. Если опустить маркетинговую политику, то,
реализовав все правильно, можно закрепиться не только в графическом секторе
рынка, но и в вычислительном (посмотри на специальные вычислительные карты и
серверы Tesla
с сотнями мультипроцессоров), потеснив привычные всем CPU.
Естественно, «повелители FPS» разошлись у камня преткновения каждый по своей
тропе, но основной принцип остался неизменным – производить вычисления
средствами GPU. И сейчас мы подробнее рассмотрим технологию «зеленых» – CUDA
(Compute Unified Device Architecture
).
Работа нашей «героини» заключается в обеспечении API, причем сразу двух.
Первый – высокоуровневый, CUDA Runtime, представляет собой функции, которые
разбиваются на более простые уровни и передаются нижнему API – CUDA Driver. Так
что фраза «высокоуровневый» применима к процессу с натяжкой. Вся соль находится
именно в драйвере, и добыть ее помогут библиотеки, любезно созданные
разработчиками NVIDIA
: CUBLAS (средства для математических расчетов) и
FFT (расчет посредством алгоритма Фурье). Ну что ж, перейдем к практической
части материала.
Терминология CUDA
NVIDIA
оперирует весьма своеобразными определениями для CUDA API. Они
отличаются от определений, применяемых для работы с центральным процессором.
Поток (thread)
– набор данных, который необходимо обработать (не
требует больших ресурсов при обработке).
Варп (warp)
– группа из 32 потоков. Данные обрабатываются только
варпами, следовательно варп – это минимальный объем данных.
Блок (block)
– совокупность потоков (от 64 до 512) или совокупность
варпов (от 2 до 16).
Сетка (grid)
– это совокупность блоков. Такое разделение данных
применяется исключительно для повышения производительности. Так, если число
мультипроцессоров велико, то блоки будут выполняться параллельно. Если же с
картой не повезло (разработчики рекомендуют для сложных расчетов использовать
адаптер не ниже уровня GeForce 8800 GTS 320 Мб), то блоки данных обработаются
последовательно.
Также NVIDIA вводит такие понятия, как ядро (kernel)
, хост (host)
и девайс (device)
.
Работаем!
Для полноценной работы с CUDA нужно:
1. Знать строение шейдерных ядер GPU, так как суть программирования
заключается в равномерном распределении нагрузки между ними.
2. Уметь программировать в среде C, с учетом некоторых аспектов.
Разработчики NVIDIA
раскрыли «внутренности» видеокарты несколько
иначе, чем мы привыкли видеть. Так что волей-неволей придется изучать все
тонкости архитектуры. Разберем строение «камня» G80 легендарной GeForce 8800
GTX
.
Шейдерное ядро состоит из восьми TPC (Texture Processor Cluster) – кластеров
текстурных процессоров (так, у GeForce GTX 280
– 15 ядер, у 8800 GTS
их шесть, у 8600
– четыре и т.д.). Те, в свою очередь, состоят из двух
потоковых мультипроцессоров (streaming multiprocessor – далее SM). SM (их всего
16) состоит из front end (решает задачи чтения и декодирования инструкций) и
back end (конечный вывод инструкций) конвейеров, а также восьми scalar SP (shader
processor) и двумя SFU (суперфункциональные блоки). За каждый такт (единицу
времени) front end выбирает варп и обрабатывает его. Чтобы все потоки варпа
(напомню, их 32 штуки) обработались, требуется 32/8 = 4 такта в конце конвейера.
Каждый мультипроцессор обладает так называемой общей памятью (shared memory).
Ее размер составляет 16 килобайт и предоставляет программисту полную свободу
действий. Распределяй как хочешь:). Shared memory обеспечивает связь потоков в
одном блоке и не предназначена для работы с пиксельными шейдерами.
Также SM могут обращаться к GDDR. Для этого им «пришили» по 8 килобайт
кэш-памяти, хранящих все самое главное для работы (например, вычислительные
константы).
Мультипроцессор имеет 8192 регистра. Число активных блоков не может быть
больше восьми, а число варпов – не больше 768/32 = 24. Из этого видно, что G80
может обработать максимум 32*16*24 = 12288 потоков за единицу времени. Нельзя не
учитывать эти цифры при оптимизации программы в дальнейшем (на одной чашу весов
– размер блока, на другой – количество потоков). Баланс параметров может сыграть
важную роль в дальнейшем, поэтому NVIDIA
рекомендует использовать блоки
со 128 или 256 потоками. Блок из 512 потоков неэффективен, так как обладает
повышенными задержками. Учитывая все тонкости строения GPU видеокарты плюс
неплохие навыки в программировании, можно создать весьма производительное
средство для параллельных вычислений. Кстати, о программировании...
Программирование
Для «творчества» вместе с CUDA требуется видеокарта GeForce не ниже
восьмой серии
. С
официального сайта нужно скачать три программных пакета: драйвер с
поддержкой CUDA (для каждой ОС – свой), непосредственно пакет CUDA SDK (вторая
бета-версия) и дополнительные библиотеки (CUDA toolkit). Технология поддерживает
операционные системы Windows (XP и Vista), Linux и Mac OS X. Для изучения я
выбрал Vista Ultimate Edition x64 (забегая вперед, скажу, что система вела себя
просто превосходно). В момент написания этих строк актуальным для работы был
драйвер ForceWare 177.35. В качестве набора инструментов использовался
программный пакет Borland C++ 6 Builder (хотя подойдет любая среда, работающая с
языком C).
Человеку, знающему язык, будет легко освоиться в новой среде. Требуется лишь
запомнить основные параметры. Ключевое слово _global_ (ставится перед функцией)
показывает, что функция относится к kernel (ядру). Ее будет вызывать центральный
процессор, а вся работа произойдет на GPU. Вызов _global_ требует более
конкретных деталей, а именно размер сетки, размер блока и какое ядро будет
применено. Например, строчка _global_ void saxpy_parallel<<
размер сетки, а Y – размер блока, задает эти параметры.
Символ _device_ означает, что функцию вызовет графическое ядро, оно же
выполнит все инструкции. Эта функция располагается в памяти мультипроцессора,
следовательно, получить ее адрес невозможно. Префикс _host_ означает, что вызов
и обработка пройдут только при участии CPU. Надо учитывать, что _global_ и
_device_ не могут вызывать друг друга и не могут вызывать самих себя.
Также язык для CUDA имеет ряд функций для работы с видеопамятью: cudafree
(освобождение памяти между GDDR и RAM), cudamemcpy и cudamemcpy2D (копирование
памяти между GDDR и RAM) и cudamalloc (выделение памяти).
Все программные коды проходят компиляцию со стороны CUDA API. Сначала берется
код, предназначенный исключительно для центрального процессора, и подвергается
стандартной компиляции, а другой код, предназначенный для графического адаптера,
переписывается в промежуточный язык PTX (сильно напоминает ассемблер) для
выявления возможных ошибок. После всех этих «плясок» происходит окончательный
перевод (трансляция) команд в понятный для GPU/CPU язык.
Набор для изучения
Практически все аспекты программирования описаны в документации, идущей
вместе с драйвером и двумя приложениями, а также на сайте разработчиков. Размера
статьи не хватит, чтобы описать их (заинтересованный читатель должен приложить
малую толику стараний и изучить материал самостоятельно).
Специально для новичков разработан CUDA SDK Browser. Любой желающий может
ощутить силу параллельных вычислений на своей шкуре (лучшая проверка на
стабильность – работа примеров без артефактов и вылетов). Приложение имеет
большой ряд показательных мини-программок (61 «тест»). К каждому опыту имеется
подробная документация программного кода плюс PDF-файлы. Сразу видно, что люди,
присутствующие со своими творениями в браузере, занимаются серьезной работой.
Тут же можно сравнить скорости работы процессора и видеокарты при обработке
данных. Например, сканирование многомерных массивов видеокартой GeForce 8800
GT
512 Мб с блоком с 256 потоками производит за 0.17109 миллисекунды.
Технология не распознает SLI-тандемы, так что если у тебя дуэт или трио,
отключай функцию «спаривания» перед работой, иначе CUDA увидит только один
девайс. Двуядерный AMD Athlon 64 X2
(частота ядра 3000 МГц) тот же опыт
проходит за 2.761528 миллисекунды. Получается, что G92 более чем в 16 раз
быстрее «камня» AMD
! Как видишь, далеко не экстремальная система в
тандеме с нелюбимой в массах операционной системой показывает неплохие
результаты.
Помимо браузера существует ряд полезных обществу программ. Adobe
адаптировала свои продукты к новой технологии. Теперь Photoshop CS4 в полной
мере использует ресурсы графических адаптеров (необходимо скачать специальный
плагин). Такими программами, как Badaboom media converter и RapiHD можно
произвести декодирование видео в формат MPEG-2. Для обработки звука неплохо
подойдет бесплатная утилита Accelero. Количество софта, заточенного под CUDA API,
несомненно, будет расти.
А в это время…
А пока ты читаешь сей материал, трудяги из процессорных концернов
разрабатывают свои технологии по внедрению GPU в CPU. Со стороны AMD
все
понятно: у них есть большущий опыт, приобретенный вместе с ATI
.
Творение «микродевайсеров», Fusion, будет состоять из нескольких ядер под
кодовым названием Bulldozer и видеочипа RV710 (Kong). Их взаимосвязь будет
осуществляться за счет улучшенной шины HyperTransport. В зависимости от
количества ядер и их частотных характеристик AMD планирует создать целую ценовую
иерархию «камней». Также планируется производить процессоры как для ноутбуков (Falcon),
так и для мультимедийных гаджетов (Bobcat). Причем именно применение технологии
в портативных устройствах будет первоначальной задачей для канадцев. С развитием
параллельных вычислений применение таких «камней» должно быть весьма популярно.
Intel
немножко отстает по времени со своей Larrabee. Продукты AMD
,
если ничего не случится, появятся на прилавках магазинов в конце 2009 – начале
2010 года. А решение противника выйдет на свет божий только почти через два
года.
Larrabee будет насчитывать большое количество (читай – сотни) ядер. Вначале
же выйдут продукты, рассчитанные на 8 – 64 ядера. Они очень сходны с Pentium, но
довольно сильно переработаны. Каждое ядро имеет 256 килобайт кэша второго уровня
(со временем его размер увеличится). Взаимосвязь будет осуществляться за счет
1024-битной двунаправленной кольцевой шины. Интел говорит, что их «дитя» будет
отлично работать с DirectX и Open GL API (для «яблочников»), поэтому никаких
программных вмешательств не потребуется.
А к чему я все это тебе поведал? Очевидно, что Larrabee и Fusion не вытеснят
обычные, стационарные процессоры с рынка, так же, как не вытеснят с рынка
видеокарты. Для геймеров и экстремалов пределом мечтаний по-прежнему останется
многоядерный CPU и тандем из нескольких топовых VGA. Но то, что даже
процессорные компании переходят на параллельные вычисления по принципам,
аналогичным GPGPU, говорит уже о многом. В частности о том, что такая
технология, как CUDA, имеет право на существование и, по всей видимости, будет
весьма популярна.
Небольшое резюме
Параллельные вычисления средствами видеокарты – всего лишь хороший инструмент
в руках трудолюбивого программиста. Вряд ли процессорам во главе с законом Мура
придет конец. Компании NVIDIA
предстоит пройти еще длинный путь по
продвижению в массы своего API (то же можно сказать и о детище ATI/AMD
).
Какой он будет, покажет будущее. Так что CUDA will be back:).
P.S. Начинающим программистам и заинтересовавшимся людям рекомендую посетить
следующие «виртуальные заведения»:
официальный сайт NVIDIA и сайт
GPGPU.com . Вся
предоставленная информация – на английском языке, но, спасибо хотя бы, что не на
китайском. Так что дерзай! Надеюсь, что автор хоть немного помог тебе в
захватывающих начинаниях познания CUDA!
На протяжении десятилетий действовал закон Мура, который гласит, что каждые два года количество транзисторов на кристалле будет удваиваться. Однако это было в далеком 1965 году, а последние 5 лет стала бурно развиваться идея физической многоядерности в процессорах потребительского класса: в 2005 году Intel представила Pentium D, а AMD – Athlon X2. Тогда приложений, использующих 2 ядра, можно было пересчитать по пальцам одной руки. Однако следующее поколение процессоров Intel, совершившее революцию, имело именно 2 физических ядра. Более того, в январе 2007 года появилась серия Quad, тогда же и сам Мур признался, что вскоре его закон перестанет действовать.
Что же сейчас? Двухядерные процессоры даже в бюджетных офисных системах, а 4 физических ядра стало нормой и это всего за 2-3 года. Частота процессоров не наращивается, а улучшается архитектура, увеличивается количество физических и виртуальных ядер. Однако идея использования видеоадаптеров, наделенных десятками, а то и сотнями вычислительных «блоков» витала давно.
И хотя перспективы вычислений силами GPU огромны, наиболее популярное решение – Nvidia CUDA бесплатно, имеет множество документаций и в целом весьма несложное в реализации, приложений, использующих эту технологию не так много. В основном это всевозможные специализированные расчеты, до которых рядовому пользователю в большинстве случаев нет дела. Но есть и программы, рассчитанные на массового пользователя, о них мы и поговорим в данной статье.
Для начала немного о самой технологии и с чем ее едят. Т.к. при написании статьи я ориентируюсь на широкий круг читателей, то и объяснить постараюсь доступным языком без сложных терминов и несколько вкратце.
CUDA (англ. Compute Unified Device Architecture) - программно-аппаратная архитектура, позволяющая производить вычисления с использованием графических процессоров NVIDIA, поддерживающих технологию GPGPU (произвольных вычислений на видеокартах). Архитектура CUDA впервые появились на рынке с выходом чипа NVIDIA восьмого поколения - G80 и присутствует во всех последующих сериях графических чипов, которые используются в семействах ускорителей GeForce, Quadro и Tesla. (с) Wikipedia.org
Входящие потоки обрабатываются независимо друг от друга, т.е. параллельно.
При этом существует разделение на 3 уровня:
Grid – ядро. Содержит одно/двух/трехмерный массив блоков.
Block – содержит в себе множество потоков (thread). Потоки разных блоков между собой взаимодействовать не могут. Для чего нужно было вводить блоки? Каждый блок по сути отвечает за свою подзадачу. Например, большое изображение (которое является матрицей) можно разбить на несколько более мелких частей (матриц) и параллельно работать с каждой частью изображения.
Thread – поток. Потоки внутри одного блока могут взаимодействовать либо через общую (shared) память, которая, кстати, куда быстрее глобальной (global) памяти, либо через средства синхронизации потоков.
Warp – это объединение взаимодействующих между собой потоков, для всех современных GPU размер Warp’а равен 32. Далее идет half-warp , являющийся половинкой warp’a, т.к. обращение к памяти обычно идет раздельно для первой и второй половины warp’a.
Как можно заметить, данная архитектура отлично подходит для распараллеливания задач. И хотя программирование ведется на языке Си с некоторыми ограничениями, на деле не все так просто, т.к. не все можно распараллелить. Нет же и стандартных функций для генерации случайных чисел (или инициализации), все это приходится реализовывать отдельно. И хотя готовых вариантов имеется в достаточном количестве, радости все это не приносит. Возможность использования рекурсии появилась сравнительно недавно.
Для наглядности была написана небольшая консольная (для минимизации кода) программа, производящая операции с двумя массивами типа float, т.е. с нецелочисленными значениями. По указанным выше причинам инициализация (заполнение массива различными произвольными значениями) производилось силами CPU. Далее с соответствующими элементами из каждого массива производилось 25 всевозможных операций, промежуточные результаты записывались в третий массив. Менялся размер массива, результаты следующие:
Всего было проведено 4 теста:
1024 элемента в каждом массиве:
Наглядно видно, что при таком малом количестве элементов толку от параллельных вычислений немного, т.к. сами вычисления проходят куда быстрее, чем их подготовка.
4096 элементов в каждом массиве:
И вот уже видно, что видеокарта в 3 раза быстрее производит операции над массивами, чем процессор. Более того, время выполнения данного теста на видеокарте не увеличилось (незначительное уменьшение времени можно сослать на погрешность).
Теперь 12288 элементов в каждом массиве:
Отрыв видеокарты увеличился еще в 2 раза. Опять же стоит обратить внимание, что время выполнения на видеокарте увеличилось
незначительно, а вот на процессоре более чем в 3 раза, т.е. пропорционально усложнению задачи.
И последний тест – 36864 элемента в каждом массиве:
В данном случае ускорение достигает внушительных значений – почти в 22 раза быстрее на видеокарте. И опять же время выполнения на видеокарте возросло незначительно, а на процессоре – положенные 3 раза, что опять же пропорционально усложнению задачи.
Если же и дальше усложнять вычисления, то видеокарта выигрывает все больше и больше. Хоть и пример несколько утрированный, но в целом ситуацию показывает наглядно. Но как упоминалось выше, не все можно распараллелить. Например, вычисление числа Пи. Существуют лишь примеры, написанные посредством метода Monte Carlo, но точность вычислений составляет 7 знаков после запятой, т.е. обычный float. Для того, чтобы увеличить точность вычислений необходима длинная арифметика, а вот тут то и наступают проблемы, т.к. эффективно это реализовать очень и очень сложно. В интернете найти примеров, использующих CUDA и рассчитывающих число Пи до 1 миллиона знаков после запятой мне не удалось. Были предприняты попытки написать такое приложение, но самый простой и эффективный метод расчета числа Пи – это алгоритм Брента - Саламина или формула Гаусса. В известном SuperPI скорее всего (судя по скорости работы и количеству итераций) используется формула Гаусса. И, судя по
тому, что SuperPI однопоточный, отсутствию примеров под CUDA и провалу моих попыток, эффективно распараллелить подсчет Pi невозможно.
Кстати, можно заметить, как в процессе выполнения вычислений повышается нагрузка на GPU, а так же происходит выделение памяти.
Теперь же перейдем к более практической пользе от CUDA, а именно существующие на данный момент программы, использующие данную технологию. В большинстве своем это всевозможные аудио/видео конвертеры и редакторы.
В тестировании использовалось 3 различных видеофайла:
- *История создания фильма Аватар - 1920x1080, MPEG4, h.264.
- *Серия "Lie to me" - 1280х720, MPEG4, h.264.
- *Серия "В Филадельфии всегда солнечно" - 624х464, xvid.
Контейнер и размер первых двух файлов был.mkv и 1,55 гб, а последнего - .avi и 272 мб.
Начнем с весьма нашумевшего и популярного продукта – Badaboom . Использовалась версия – 1.2.1.74 . Стоимость программы составляет $29.90 .
Интерфейс программы простой и наглядный – слева выбираем исходный файл или диск, а справа – необходимое устройство, для которого будем кодировать. Есть и пользовательский режим, в котором вручную задаются параметры, он и использовался.
Для начала рассмотрим насколько быстро и качественно кодируется видео «в себя же», т.е. в то же разрешение и примерно тот же размер. Скорость будем измерять в fps, а не в затраченном времени – так удобнее и сравнивать, и подсчитывать сколько будет сжиматься видео произвольной продолжительности. Т.к. сегодня мы рассматриваем технологию «зеленых», то и графики будут соответствующие -)
Скорость кодирования напрямую зависит от качества, это очевидно. Стоит отметить, что легкое разрешение (назовем его традиционно – SD) не проблема для Badaboom – скорость кодирования в 5,5 раз превысила исходный (24 fps) фреймрейт видео. Да и даже тяжелый 1080p видеоролик программа преобразует в реальном времени. Стоит отметить, что качество итогового видео очень близко к исходному видеоматериалу, т.е. кодирует Badaboom весьма и весьма качественно.
Но обычно перегоняют видео в более низкое разрешение, посмотрим как обстоят дела в этом режиме. При снижении разрешения снижался и битрейт видео. Он составлял 9500 кбит/с для 1080p выходного файла, 4100 кбит/с для 720 p и 2400 кбит/с для 720х404. Выбор сделан исходя из разумного соотношения размер/качество.
Комментарии излишни. Если делать из 720p рип до обычного SD качества, то на перекодирование фильма длительностью 2 часа уйдет около 30 минут. И при этом загрузка процессора будет незначительной, можно заниматься своими делами не ощущая дискомфорта.
А что если перегнать видео в формат для мобильного устройства? Для этого выберем профиль iPhone (битрейт 1 мбит/с, 480х320) и посмотрим на скорость кодирования:
Нужно ли что-то говорить? Двухчасовой фильм в обычном качестве для iPhone перекодируется менее чем за 15 минут. С HD качеством сложнее, но все равно весьма быстро. Главное, что качество выходного видеоматериала остается на довольно высоком уровне при просмотре на дисплее телефона.
В целом впечатления от Badaboom положительные, скорость работы радует, интерфейс простой и понятный. Всевозможные баги ранних версий (пользовался еще бетой в 2008-ом году) вылечены. Кроме одного – путь к исходному файлу, а так же к папке, в которую сохраняется готовое видео, не должен содержать русских букв. Но на фоне достоинств программы этот недостаток незначителен.
Следующим на очереди у нас будет Super LoiLoScope . За обычную его версию просят 3 280 рублей , а за touch версию, поддерживающую сенсорное управление в Windows 7, просят аж 4 440 рублей . Попробуем разобраться за что разработчик хочет таких денег и зачем видеоредактору поддержка multitouch. Использовалась последняя версия – 1.8.3.3 .
Описать интерфейс программы словами довольно сложно, поэтому я решил снять небольшой видеоролик. Сразу скажу, что, как и все видеоконвертеры под CUDA, ускорение средствами GPU поддерживается только для вывода видео в MPEG4 с кодеком h.264.
Во время кодирования загрузка процессора составляет 100%, однако дискомфорта это не вызывает. Браузер и другие не тяжелые приложения не тормозят.
Теперь перейдем к производительности. Для начала все тоже самое, что и с Badaboom – перекодирование видео в аналогичное по качеству.
Результаты куда лучше, чем у Badaboom. Качество так же на высоте, разницу с оригиналом можно заметить только сравнивая попарно кадры под лупой.
Ого, а вот тут LoiloScope обходит Badaboom в 2,5 раза. При этом можно запросто параллельно резать и кодировать другое видео, читать новости и даже смотреть кино, причем даже FullHD проигрываются без проблем, хоть загрузка процессора и максимальна.
Теперь же попробуем сделать видео для мобильного устройства, профиль назовем так же, как он назывался в Badaboom – iPhone (480x320, 1 мбит/с):
Никакой ошибки нет. Все перепроверялось несколько раз, каждый раз результат был аналогичным. Скорее всего, это происходит по той простой причине, что SD файл записан с другим кодеком и в другом контейнере. При перекодировании видео сначала декодируется, разбивается на матрицы определенного размера, сжимается. ASP декодер, использующийся в случае с xvid, медленнее, чем AVC (для h.264) при параллельном декодировании. Однако и 192 fps – это в 8 раз быстрее, чем скорость исходного видео, серия длительностью 23 минуты сжимается менее чем за 4 минуты. Ситуация повторялась и с другими файлами, пережатыми в xvid/DivX.
LoiloScope оставил о себе только приятные впечатления – интерфейс, несмотря на свою необычность, удобный и функциональный, а скорость работы выше всяких похвал. Несколько расстраивает относительно бедный функционал, но в зачастую при простом монтаже требуется лишь немного подкорректировать цвета, сделать плавные переходы, наложить текст, а с этим LoiloScope отлично справляется. Несколько пугает и цена – более $100 за обычную версию нормально для зарубежья, но нам такие цифры пока кажутся несколько дикими. Хотя, признаюсь, что если бы я, например, часто снимал и монтировал домашнее видео, то возможно и задумался над покупкой. Заодно, кстати, проверил возможность редактирования HD (а точнее AVCHD) контента прямо из видеокамеры без предварительного конвертирования в другой формат, у LoiloScope никаких проблем с файлами типа.mts не выявлено.
В развитии современных процессоров намечается тенденция к постепенному увеличению количества ядер, что повышает их возможности в параллельных вычислениях. Однако уже давно имеются 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.
Кодирование осуществлялось с помощью кодека 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 это один из этапов в процессе данного объединения.
Благодарим следующие компании за предоставленное тестовое оборудование:
Новая технология — как вновь возникший эволюционный вид. Странное создание, непохожее на многочисленных старожилов. Местами неуклюжее, местами смешное. И поначалу его новые качества кажутся ну никак не подходящими для этого обжитого и стабильного мира.
Однако проходит немного времени, и оказывается, что новичок бегает быстрее, прыгает выше и вообще сильнее. И мух он лопает больше его соседей-ретроградов. И вот тогда эти самые соседи начинают понимать, что ссориться с этим бывшим неуклюжим не стоит. Лучше с ним дружить, а еще лучше организовать симбиоз. Глядишь, и мух перепадет побольше.
Технология GPGPU (General-Purpose Graphics Processing Units — графический процессор общего назначения) долгое время существовала только в теоретических выкладках мозговитых академиков. А как иначе? Предложить кардинально изменить сложившийся за десятилетия вычислительный процесс, доверив расчет его параллельных веток видеокарте, — на это только теоретики и способны.
Логотип технологии CUDA напоминает о том, что выросла она в недрах
3D-графики.
Но долго пылиться на страницах университетских журналов технология GPGPU не собиралась. Распушив перья своих лучших качеств, она привлекла к себе внимание производителей. Так на свет появилась CUDA — реализация GPGPU на графических процессорах GeForce производства компании nVidia.
Благодаря CUDA технологии GPGPU стали мейнстримом. И ныне только самый недальновидный и покрытый толстым слоем лени разработчик систем программирования не заявляет о поддержке своим продуктом CUDA. IT-издания почли за честь изложить подробности технологии в многочисленных пухлых научно-популярных статьях, а конкуренты срочно уселись за лекала и кросскомпиляторы, чтобы разработать нечто подобное.
Публичное признание — это мечта не только начинающих старлеток, но и вновь зародившихся технологий. И CUDA повезло. Она на слуху, о ней говорят и пишут.
Вот только пишут так, словно продолжают обсуждать GPGPU в толстых научных журналах. Забрасывают читателя грудой терминов типа «grid», «SIMD», «warp», «хост», «текстурная и константная память». Погружают его по самую маковку в схемы организации графических процессоров nVidia, ведут извилистыми тропами параллельных алгоритмов и (самый сильный ход) показывают длинные листинги кода на языке Си. В результате получается, что на входе статьи мы имеем свежего и горящего желанием понять CUDA читателя, а на выходе — того же читателя, но с распухшей головой, заполненной кашей из фактов, схем, кода, алгоритмов и терминов.
А между тем цель любой технологии — сделать нашу жизнь проще. И CUDA прекрасно с этим справляется. Результаты ее работы — именно это убедит любого скептика лучше сотни схем и алгоритмов.
Далеко не везде
CUDA поддерживается высокопроизводительными суперкомпьютерами
nVidia Tesla.
И все же прежде, чем взглянуть на результаты трудов CUDA на поприще облегчения жизни рядового пользователя, стоит уяснить все ее ограничения. Точно как с джинном: любое желание, но одно. У CUDA тоже есть свои ахиллесовы пятки. Одна из них — ограничения платформ, на которых она может трудиться.
Перечень видеокарт производства nVidia, поддерживающих CUDA, представлен в специальном списке, именуемом CUDA Enabled Products. Список весьма внушительный, но легко классифицируемый. В поддержке CUDA не отказывают:
Модели nVidia GeForce 8-й, 9-й, 100-й, 200-й и 400-й серий с минимумом 256 мегабайт видеопамяти на борту. Поддержка распространяется как на карты для настольных систем, так и на мобильные решения.
Подавляющее большинство настольных и мобильных видеокарт nVidia Quadro.
Все решения нетбучного ряда nvidia ION.
Высокопроизводительные HPC (High Performance Computing) и суперкомпьютерные решения nVidia Tesla, используемые как для персональных вычислений, так и для организации масштабируемых кластерных систем.
Поэтому, прежде чем применять программные продукты на базе CUDA, стоит свериться с этим списком избранных.
Кроме самой видеокарты, для поддержки CUDA требуется соответствующий драйвер. Именно он является связующим звеном между центральным и графическим процессором, выполняя роль своеобразного программного интерфейса для доступа кода и данных программы к многоядерной сокровищнице GPU. Чтобы наверняка не ошибиться, nVidia рекомендует посетить страничку драйверов и получить наиболее свежую версию.
...но сам процесс
Как работает CUDA? Как объяснить сложный процесс параллельных вычислений на особой аппаратной архитектуре GPU так, чтобы не погрузить читателя в пучину специфических терминов?
Можно попытаться это сделать, представив, как центральный процессор выполняет программу в симбиозе с процессором графическим.
Архитектурно центральный процессор (CPU) и его графический собрат (GPU) устроены по-разному. Если проводить аналогию с миром автопрома, то CPU — универсал, из тех, которые называют «сарай». Выглядит легковым авто, но при этом (с точки зрения разработчиков) «и швец, и жнец, и на дуде игрец». Выполняет роль маленького грузовика, автобуса и гипертрофированного хечбэка одновременно. Универсал, короче. Цилиндров-ядер у него немного, но они «тянут» практически любые задачи, а внушительная кэш-память способна разместить кучу данных.
А вот GPU — это спорткар. Функция одна: доставить пилота на финиш как можно быстрее. Поэтому никакой большой памяти-багажника, никаких лишних посадочных мест. Зато цилиндров-ядер в сотни раз больше, чем у CPU.
Благодаря CUDA разработчикам программ GPGPU не требуется вникать в сложности программи-
рования под такие графические движки, как DirectX и OpenGL
В отличие от центрального процессора, способного решать любую задачу, в том числе и графическую, но с усредненной производительностью, графический процессор адаптирован на высокоскоростное решение одной задачи: превращение куч полигонов на входе в кучу пикселов на выходе. Причем задачу эту можно решать параллельно на сотнях относительно простых вычислительных ядер в составе GPU.
Так какой же может быть тандем из универсала и спорткара? Работа CUDA происходит примерно так: программа выполняется на CPU до тех пор, пока в ней появляется участок кода, который можно выполнить параллельно. Тогда, вместо того, чтобы он медленно выполнялся на двух (да пусть даже и восьми) ядрах самого крутого CPU, его передают на сотни ядер GPU. При этом время выполнения этого участка сокращается в разы, а значит, сокращается и время выполнения всей программы.
Технологически для программиста ничего не меняется. Код CUDA-программ пишется на языке Си. Точнее, на особом его диалекте «С with streams» (Си с потоками). Разработанное в Стэнфорде, это расширение языка Си получило название Brook. В качестве интерфейса, передающего Brook-код на GPU, выступает драйвер видеокарты, поддерживающей CUDA. Он организует весь процесс обработки этого участка программы так, что для программиста GPU выглядит как сопроцессор CPU. Очень похоже на использование математического сопроцессора на заре персональных компьютеров. С появлением Brook, видеокарт с поддержкой CUDA и драйверов для них любой программист стал способен в своих программах обращаться к GPU. А ведь раньше этим шаманством владел узкий круг избранных, годами оттачивающих технику программирования под графические движки DirectX или OpenGL.
В бочку этого пафосного меда — дифирамбов CUDA — стоит положить ложку дегтя, то бишь ограничений. Далеко не любая задача, которую нужно запрограммировать, подходит для решения с помощью CUDA. Добиться ускорения решения рутинных офисных задач не получится, а вот доверить CUDA обсчет поведения тысячи однотипных бойцов в World of Warcraft — пожалуйста. Но это задача, высосанная из пальца. Рассмотрим же примеры того, что CUDA уже очень эффективно решает.
Труды праведные
CUDA — весьма прагматичная технология. Реализовав ее поддержку в своих видеокартах, компания nVidia весьма справедливо рассчитывала на то, что знамя CUDA будет подхвачено множеством энтузиастов как в университетской среде, так и в коммерции. Так и случилось. Проекты на базе CUDA живут и приносят пользу.
NVIDIA PhysX
Рекламируя очередной игровой шедевр, производители частенько напирают на его 3D-реалистичность. Но каким бы реальным ни был игровой 3D-мир, если элементарные законы физики, такие как тяготение, трение, гидродинамика, будут реализованы неправильно, фальшь почувствуется моментально.
Одна из возможностей физического движка NVIDIA PhysX — реалистичная работа с тканями.
Реализовать алгоритмы компьютерной симуляции базовых физических законов — дело очень трудоемкое. Наиболее известными компаниями на этом поприще являются ирландская компания Havok с ее межплатформенным физическим Havok Physics и калифорнийская Ageia — прародитель первого в мире физического процессора (PPU — Physics Processing Unit) и соответствующего физического движка PhysX. Первая из них, хотя и приобретена компанией Intel, активно трудится сейчас на поприще оптимизации движка Havok для видеокарт ATI и процессоров AMD. А вот Ageia с ее движком PhysX стала частью nVidia. При этом nVidia решила достаточно сложную задачу адаптации PhysX под технологию CUDA.
Возможным это стало благодаря статистике. Статистически было доказано, что, какой бы сложный рендеринг ни выполнял GPU, часть его ядер все равно простаивает. Именно на этих ядрах и работает движок PhysX.
Благодаря CUDA львиная доля вычислений, связанных с физикой игрового мира, стала выполняться на видеокарте. Освободившаяся мощь центрального процессора была брошена на решение других задач геймплея. Результат не заставил себя ждать. По оценкам экспертов, прирост производительности игрового процесса с PhysX, работающем, на CUDA возрос минимум на порядок. Выросло и правдоподобие реализации физических законов. CUDA берет на себя рутинный расчет реализации трения, тяготения и прочих привычных нам вещей для многомерных объектов. Теперь не только герои и их техника идеально вписываются в законы привычного нам физического мира, но и пыль, туман, взрывная волна, пламя и вода.
CUDA-версия пакета сжатия текстур NVIDIA Texture Tools 2
Нравятся реалистичные объекты в современных играх? Стоит сказать спасибо разработчикам текстур. Но чем больше реальности в текстуре, тем больше ее объем. Тем больше она занимает драгоценной памяти. Чтобы этого избежать, текстуры предварительно сжимают и динамически распаковывают по мере надобности. А сжатие и распаковка — это сплошные вычисления. Для работы с текстурами nVidia выпустила пакет NVIDIA Texture Tools. Он поддерживает эффективное сжатие и распаковку текстур стандарта DirectX (так называемый ВЧЕ-формат). Вторая версия этого пакета может похвастаться поддержкой алгоритмов сжатия BC4 и BC5, реализованных в технологии DirectX 11. Но главное то, что в NVIDIA Texture Tools 2 реализована поддержка CUDA. По оценке nVidia, это дает 12-кратный прирост производительности в задачах сжатия и распаковки текстур. А это значит, что фреймы игрового процесса будут грузиться быстрее и радовать игрока своей реалистичностью.
Пакет NVIDIA Texture Tools 2 заточен под работу с CUDA. Прирост производительности при сжатии и распаковке текстур налицо.
Использование CUDA позволяет существенно повысить эффективность видеослежки.
Обработка видеопотока в реальном времени
Как ни крути, а нынешний мир, с точки зрения соглядатайства, куда ближе к миру оруэлловского Большого Брата, чем кажется. Пристальные взгляды видеокамер ощущают на себе и водители авто, и посетители общественных мест.
Полноводные реки видеоинформации стекаются в центры ее обработки и... наталкиваются на узкое звено — человека. Именно он в большинстве случаев — последняя инстанция, следящая за видеомиром. Причем инстанция не самая эффективная. Моргает, отвлекается и норовит уснуть.
Благодаря CUDA появилась возможность реализации алгоритмов одновременного слежения за множеством объектов в видеопотоке. При этом процесс происходит в реальном масштабе времени, а видео является полноценным 30 fps. По сравнению с реализацией такого алгоритма на современных многоядерных CPU CUDA дает двух-, трехкратный прирост производительности, а это, согласитесь, немало.
Конвертирование видео, фильтрация аудио
Видеоконвертер Badaboom — первая ласточка, использующая CUDA для ускорения конвертирования.
Приятно посмотреть новинку видеопроката в FullHD-качестве и на большом экране. Но большой экран не возьмешь с собой в дорогу, а видеокодек FullHD будет икать на маломощном процессоре мобильного гаджета. На помощь приходит конвертирование. Но большинство тех, кто с ним сталкивался на практике, сетуют на длительное время конвертации. Оно и понятно, процесс рутинный, пригодный к распараллеливанию, и его выполнение на CPU не очень оптимально.
А вот CUDA с ним справляется на ура. Первая ласточка — конвертер Badaboom от компании Elevental. Разработчики Badaboom, выбрав CUDA, не просчитались. Тесты показывают, что стандартный полуторачасовый фильм на нем конвертируется в формат iPhone/iPod Touch менее чем за двадцать минут. И это при том, что при использовании только CPU этот процесс занимает больше часа.
Помогает CUDA и профессиональным меломанам. Любой из них полцарства отдаст за эффективный FIR-кроссовер — набор фильтров, разделяющих звуковой спектр на несколько полос. Процесс этот весьма трудоемкий и при большом объеме аудиоматериала заставляет звукорежиссера сходить на несколько часов «покурить». Реализация FIR-кроссовера на базе CUDA ускоряет его работу в сотни раз.
CUDA Future
Сделав технологию GPGPU реальностью, CUDA не собирается почивать на лаврах. Как это происходит повсеместно, в CUDA работает принцип рефлексии: теперь не только архитектура видеопроцессоров nVidia влияет на развитие версий CUDA SDK, а и сама технология CUDA заставляет nVidia пересматривать архитектуру своих чипов. Пример такой рефлексии — платформа nVidia ION. Ее вторая версия специально оптимизирована для решения CUDA-задач. А это означает, что даже в относительно недорогих аппаратных решениях потребители получат всю мощь и блестящие возможности CUDA.
