CUDA мы катимся: технология NVIDIA CUDA. Смотреть что такое "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 обеспечивают следующие периферийные устройства :

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

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

Пример

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

Неофициальные ресурсы

• Дмитрий Чеканов. 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 . Использование видеокарт для вычислений
• . Центр Параллельных Вычислений

Примечания

См. также

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

Анализ

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

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

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

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

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

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

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

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

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

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

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

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

Согласно Дарвинской теории эволюции, первая человекообразная обезьяна (если
быть точным – 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<<>>, где X –
размер сетки, а 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!

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

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

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

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

К примеру:

GeForce

GTX 460 — 336 CUDA ядер

GTX 580 — 512 CUDA ядер

8800GTX — 128 CUDA ядер

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

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

На протяжении десятилетий действовал закон Мура, который гласит, что каждые два года количество транзисторов на кристалле будет удваиваться. Однако это было в далеком 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 не выявлено.