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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

Состав Nvidia CUDA

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

Выводы

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

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

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

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

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

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

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

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

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

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

Ядра 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 .

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

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

Технология 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.

Технология CUDA

Владимир Фролов, [email protected]

Аннотация

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

1. Введение

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

Отчасти поэтому все производители современных вычислительных систем идут в сторону увеличения числа процессоров и ядер, а не увеличивают частоту одного процессора. Количество ядер центрального процессора (CPU) в передовых системах сейчас уже равняется 8.

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

Однако если посмотреть в сторону графических процессоров GPU (Graphics Processing Unit), то там по пути параллелизма пошли гораздо раньше. В сегодняшних видеокартах, например в GF8800GTX, число процессоров может достигать 128. Производительность подобных систем при умелом их программировании может быть весьма значительной (рис. 1).

Рис. 1. Количество операций с плавающей точкой для CPU и GPU

Когда первые видеокарты только появились в продаже, они представляли собой достаточно простые (по сравнению с центральным процессором) узкоспециализированные устройства, предназначенные для того чтобы снять с процессора нагрузку по визуализации двухмерных данных. С развитием игровой индустрии и появлением таких трехмерных игр как Doom (рис. 2) и Wolfenstein 3D (рис. 3) возникла необходимость в 3D визуализации.

Рисунки 2,3. Игры Doom и Wolfenstein 3D

Со времени создания компанией 3Dfx первых видеокарт Voodoo, (1996 г.) и вплоть до 2001 года в GPU был реализован только фиксированный набор операций над входными данными.

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

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

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

Сегодня появилась тенденция нетрадиционного использования видеокарт для решения задач в областях квантовой механики, искусственного интеллекта, физических расчетов, криптографии, физически корректной визуализации, реконструкции по фотографиям, распознавания и.т.п. Эти задачи неудобно решать в рамках графических API (DirectX, OpenGL), так как эти API создавались совсем для других применений.

Развитие программирования общего назначения на GPU (General Programming on GPU, GPGPU) логически привело к возникновению технологий, нацеленных на более широкий круг задач, чем растеризация. В результате компанией Nvidia была создана технология Compute Unified Device Architecture (или сокращенно CUDA), а конкурирующей компанией ATI - технология STREAM.

Следует заметить, что на момент написания этой статьи, технология STREAM сильно отставала в развитии от CUDA, и поэтому здесь она рассматриваться не будет. Мы сосредоточимся на CUDA - технологии GPGPU, позволяющей писать программы на подмножестве языка C++.

2. Принципиальная разница между CPU и GPU

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

2.1. Возможности

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

Для GPU это не так. Как вы узнаете, прочитав эту статью, в CUDA имеется целых 6 видов памяти. Читать можно из любой ячейки, доступной физически, но вот записывать – не во все ячейки. Причина заключается в том, что GPU в любом случае представляет собой специфическое устройство, предназначенное для конкретных целей. Это ограничение введено ради увеличения скорости работы определенных алгоритмов и снижения стоимости оборудования.

2.2. Быстродействие памяти

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

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

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

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

3. Ядро CUDA

3.1. Потоковая модель

Вычислительная архитектура CUDA основана на концепции одна команда на множество данных (Single Instruction Multiple Data , SIMD) и понятии мультипроцессора .

Концепция SIMD подразумевает, что одна инструкция позволяет одновременно обработать множество данных. Например, команда addps в процессоре Pentium 3 и в более новых моделях Pentium позволяет складывать одновременно 4 числа с плавающей точкой одинарной точности.

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

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

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

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

Фактически, та часть вашей программы, которая работает на CPU - это хост, а ваша видеокарта - устройство. Логически устройство можно представить как набор мультипроцессоров (рис. 4) плюс драйвер CUDA.

Рис. 4. Устройство

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

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

Рис. 5. Организация потоков

На рис. 5. ядро обозначено как Kernel. Все потоки, выполняющие это ядро, объединяются в блоки (Block), а блоки, в свою очередь, объединяются в сетку (Grid).

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

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

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

Блок задач (потоков) выполняется на мультипроцессоре частями, или пулами, называемыми warp. Размер warp на текущий момент в видеокартах с поддержкой CUDA равен 32 потокам. Задачи внутри пула warp исполняются в SIMD стиле, т.е. во всех потоках внутри warp одновременно может выполняться только одна инструкция .

Здесь следует сделать одну оговорку. В архитектурах, современных на момент написания этой статьи, количество процессоров внутри одного мультипроцессора равно 8, а не 32. Из этого следует, что не весь warp исполняется одновременно, он разбивается на 4 части, которые выполняются последовательно (т.к. процессоры скалярные).

Но, во-первых, разработчики CUDA не регламентируют жестко размер warp. В своих работах они упоминают параметр warp size, а не число 32. Во-вторых, с логической точки зрения именно warp является тем минимальным объединением потоков, про который можно говорить, что все потоки внутри него выполняются одновременно - и при этом никаких допущений относительно остальной системы сделано не будет .

3.1.1. Ветвления

Сразу же возникает вопрос: если в один и тот же момент времени все потоки внутри warp исполняют одну и ту же инструкцию, то как быть с ветвлениями? Ведь если в коде программы встречается ветвление, то инструкции будут уже разные. Здесь применяется стандартное для SIMD программирования решение (рис 6).

Рис. 6. Организация ветвления в SIMD

Пусть имеется следующий код:

if(cond) B;

В случае SISD (Single Instruction Single Data) мы выполняем оператор A, проверяем условие, затем выполняем операторы B и D (если условие истинно).

Пусть теперь у нас есть 10 потоков, исполняющихся в стиле SIMD. Во всех 10 потоках мы выполняем оператор A, затем проверяем условие cond и оказывается, что в 9 из 10 потоках оно истинно, а в одном потоке - ложно.

Понятно, что мы не можем запустить 9 потоков для выполнения оператора B, а один оставшийся- для выполнения оператора C, потому что одновременно во всех потоках может исполняться только одна инструкция. В этом случае нужно поступить так: сначала «убиваем» отколовшийся поток так, чтобы он не портил ничьи данные, и выполняем 9 оставшихся потоков. Затем «убиваем» 9 потоков, выполнивших оператор B, и проходим один поток с оператором C. После этого потоки опять объединяются и выполняют оператор D все одновременно .

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

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

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

3.1.2. Взаимодействие между потоками

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

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

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

3.2. Память

В CUDA выделяют шесть видов памяти (рис. 7). Это регистры, локальная, глобальная, разделяемая, константная и текстурная память.

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

Рис. 7. Виды памяти в CUDA

3.2.0. Регистры

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

При обычном разделении в 64 потока на блок получается всего 128 регистров (существуют некие объективные критерии, но 64 подходит в среднем для многих задач). Реально, 128 регистров nvcc никогда не выделит. Обычно он не дает больше 40, а остальные переменные попадпют в локальную память. Так происходит потому что на одном мультипроцессоре может исполняться несколько блоков. Компилятор старается максимизировать число одновременно работающих блоков. Для большей большей эффективности надо стараться занимать меньше чем 32 регистра. Тогда теоретически может быть запущено 4 блока (8 warp-ов, если 64 треда в одном блоке) на одном мультипроцессоре. Однако здесь еще следует учитывать объем разделяемой памяти, занимаемой потоками, так как если один блок занимает всю разделяемую память, два таких блока не могут выполняться на мультипроцессоре одновременно .

3.2.1. Локальная память

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

Физически локальная память является аналогом глобальной памяти, и работает с той же скоростью. На момент написания статьи не было никаких механизмов, позволяющих явно запретить компилятору использование локальной памяти для конкретных переменных. Так как проконтролировать локальную память довольно трудно, лучше не использовать ее вовсе (см. раздел 4 «Рекомендации по оптимизации»).

3.2.2. Глобальная память

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

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

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

Переменные, объявленные с квалификатором __global__, размещаются в глобальной памяти. Глобальную память также можно выделить динамически, вызвав функцию cudaMalloc(void* mem, int size) на хосте. Из устройства эту функцию вызывать нельзя. Отсюда следует, что распределением памяти должна заниматься программа-хост, работающая на CPU. Данные с хоста можно отправлять в устройство вызовом функции cudaMemcpy:

cudaMemcpy(void* gpu_mem, void* cpu_mem, int size, cudaMemcpyHostToDevice);

Точно таким же образом можно проделать и обратную процедуру:

cudaMemcpy(void* cpu_mem, void* gpu_mem, int size, cudaMemcpyDeviceToHost);

Этот вызов тоже осуществляется с хоста.

При работе с глобальной памятью важно соблюдать правило коалесинга (coalescing). Основная идея в том, что треды должны обращаться к последоваетльным ячейкам памяти, причем 4,8 или 16 байтовым. При этом, самый первый тред должен обращаться по адресу, выровненному на границу соответственно 4,8 или 16 байт. Адреса, возвращаемые cudaMalloc выровнены как минимум по границе 256 байт.

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

Разделяемая память - это некэшируемая, но быстрая память. Ее и рекомендуется использовать как управляемый кэш. На один мультипроцессор доступно всего 16KB разделяемой памяти. Разделив это число на количество задач в блоке, получим максимальное количество разделяемой памяти, доступной на один поток (если планируется использовать ее независимо во всех потоках).

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

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

Переменные, объявленные с квалификатором __shared__, размещаются в разделяемой памяти.

Shared__ float mem_shared;

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

float x = mem_shared;

Где threadIdx.x - индекс x потока внутри блока.

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

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

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

#define N 100

Constant__ int gpu_buffer[N];

void host_function()

int cpu_buffer[N];

cudaMemcpyToSymbol(gpu_buffer, cpu_buffer, sizeof(int )*N);

// __global__ означает, что device_kernel - ядро, которое может быть запущено на GPU

Global__ void device_kernel()

int a = gpu_buffer;

int b = gpu_buffer + gpu_buffer;

// gpu_buffer = a; ОШИБКА! константная память доступна только для чтения

Так как для константной памяти используется кэш, доступ к ней в общем случае довольно быстрый. Единственный, но очень большой недостаток константной памяти заключается в том, что ее размер составляет всего 64 Kбайт (на все устройство). Из этого следует, что в контекстной памяти имеет смысл хранить лишь небольшое количество часто используемых данных.

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

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

Название текстурной памяти (и, к сожалению, функциональность) унаследовано от понятий «текстура» и «текстурирование». Текстурирование - это процесс наложения текстуры (просто картинки) на полигон в процессе растеризации. Текстурная память оптимизирована под выборку 2D данных и имеет следующие возможности:

быстрая выборка значений фиксированного размера (байт, слово, двойное или учетверенное слово) из одномерного или двухмерного массива;

нормализованная адресация числами типа float в интервале . Затем можно их выбирать, используя нормализованную адресацию. Результирующим значением будетет слово типа float4, отображенное в интервал ;

CudaMalloc((void**) &gpu_memory, N*sizeof (uint4 )); // выделим память в GPU

// настройка параемтров текстуры texture

Texture.addressMode = cudaAddressModeWrap; // режим Wrap

Texture.addressMode = cudaAddressModeWrap;

Texture.filterMode = cudaFilterModePoint; // ближайшеезначение

Texture.normalized = false; // не использовать нормализованную адресацию

CudaBindTexture (0, texture , gpu _ memory , N ) // отныне эта память будет считаться текстурной


CudaMemcpy (gpu _ memory , cpu _ buffer , N * sizeof (uint 4), cudaMemcpyHostToDevice ); // копируем данные на GPU

// __global__ означает, что device_kernel - ядро, которое нужно распараллелить

Global__ void device_kernel()

uint4 a = tex1Dfetch(texture,0); // можно выбирать данные только таким способом!

uint4 b = tex1Dfetch(texture,1);

int c = a.x * b.y;

...

3.3. Простой пример

В качестве простого примера предлагается рассмотреть программу cppIntegration из CUDA SDK. Она демонстрирует приемы работы с CUDA, а также использование nvcc (специальный компилятор подмножества С++ от Nvidia) в сочетании с MS Visual Studio, что сильно упрощает разработку программ на CUDA.

4.1. Правильно проводите разбиение вашей задачи

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

4.2. Выбор типа памяти

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

Используйте глобальную память в сочетании с разделяемой памятью, если все задачи обращаются бессистемно к разным, далеко расположенным друг от друга участкам памяти (с сильно различными адресами или координатами, если это 2D/3D данные).

глобальная память => разделяемая память

Syncthreads();

Обработать данные в разделяемой памяти

Syncthreads();

глобальная память <= разделяемая память

4.3. Включите счетчики памяти

Флаг компилятора --ptxas-options=-v позволяет точно сказать, сколько и какой памяти (регистров, разделяемой, локальной, константной) вы используете. Если компилятор использует локальную память, вы точно знаете об этом. Анализ данных о количестве и типах используемой памяти может сильно помочь вам при оптимизации программы.

4.4. Старайтесь минимизировать использование регистров и разделяемой памяти

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

4.5. Разделяемая память вместо локальной

Если компилятор Nvidia по какой-то причине расположил данные в локальной памяти (обычно это заметно по очень сильному падению производительности в местах, где ничего ресурсоемкого нет), выясните, какие именно данные попали в локальную память, и поместите их в разделяемую память (shared memory).

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

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

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

4.6. Разворачивание циклов

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

Вот как можно развернуть цикл нахождения суммы массива (например, целочисленного):

int a[N]; int summ;

for (int i=0;i

Разумеется, циклы можно развернуть и вручную (как показано выше), но это малопроизводительный труд. Гораздо лучше использовать шаблоны С++ в сочетание со встраиваемыми функциями.

template

class ArraySumm

Device__ static T exec(const T* arr) { return arr + ArraySumm(arr+1); }

template

class ArraySumm<0,T>

Device__ static T exec(const T* arr) { return 0; }

for (int i=0;i

summ+= ArraySumm<4,int>::exec(a);

Следует отметить одну интересную особенность компилятора nvcc. Компилятор всегда будет встраивать функции типа __device__ по умолчанию (чтобы это отменить, существует специальная директива __noinline__) .

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

4.7. Выравнивание данных и выборка по 16 байт

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

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

4.8. Конфликты банков разделяемой памяти

Разделяемая память организована в виде 16 (всего-то!) банков памяти с шагом в 4 байта. Во время выполнения пула потоков warp на мультипроцессоре, он делится на две половинки (если warp-size = 32) по 16 потоков, которые осуществляют доступ к разделяемой памяти по очереди.

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

Оптимальные шаги - 4, 12, 28, ..., 2^n-4 байт (рис. 8).



Рис. 8. Оптимальные шаги.

Не оптимальные шаги – 1, 8, 16, 32, ..., 2^n байт (рис. 9).



Рис. 9. Неоптимальные шаги

4.9. Минимизация перемещений данных Host <=> Device

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

5. CPU/GPU переносимая математическая библиотека

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

Библиотека MGML_MATH может быть использована как каркас для написания CPU/GPU переносимых (или гибридных) систем расчета физических, графических или других пространственных задач. Основное ее достоинство в том, что один и тот же код может использоваться как на CPU, так и на GPU, и при этом во главу требований, предъявляемых к библиотеке, ставится скорость.

6 . Литература

Крис Касперски. Техника оптимизации программ. Эффективное использование памяти. - Спб.: БХВ-Петербург, 2003. - 464 с.: ил.

CUDA Programming Guide 1.1 (http://developer.download.nvidia.com/compute/cuda/1_1/NVIDIA_CUDA_Programming_Guide_1.1.pdf )

CUDA Programming Guide 1.1. page 14-15

CUDA Programming Guide 1.1. page 48

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




Поделись статьей:
Похожие статьи

17 апреля 2015
Обзор ноутбука Lenovo ThinkPad X1 Carbon (2018): лёгкий, удобный, мощный

17 апреля 2015
Стереофония Устройство для приема передач на расстоянии

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

17 апреля 2015
Как подключиться к компьютеру по локальной сети?