• Гаджеты
• Android
• Windows
• Ios
• Блог

Куда ведет cuda: практическое применение технологии gpgpu - лучшее оборудование. CUDA мы катимся: технология NVIDIA CUDA Поддерживаемые GPU и графические ускорители

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

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

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

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

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



К примеру:

GeForce

GTX 460 — 336 CUDA ядер

GTX 580 — 512 CUDA ядер

8800GTX — 128 CUDA ядер

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

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

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

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

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

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

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

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



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



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

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

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

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

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

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

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

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



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

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



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

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



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

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



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

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

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



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

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

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

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

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



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

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



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

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



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

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



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

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

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

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

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

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



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



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

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



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

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

И предназначен для трансляции host-кода (главного, управляющего кода) и device-кода (аппаратного кода) (файлов с расширением.cu) в объектные файлы, пригодные в процессе сборки конечной программы или библиотеки в любой среде программирования, например в NetBeans .

В архитектуре CUDA используется модель памяти грид , кластерное моделирование потоков и SIMD -инструкции. Применима не только для высокопроизводительных графических вычислений, но и для различных научных вычислений с использованием видеокарт nVidia. Ученые и исследователи широко используют CUDA в различных областях, включая астрофизику , вычислительную биологию и химию, моделирование динамики жидкостей, электромагнитных взаимодействий, компьютерную томографию, сейсмический анализ и многое другое. В CUDA имеется возможность подключения к приложениям, использующим OpenGL и Direct3D . CUDA - кроссплатформенное программное обеспечение для таких операционных систем как Linux , Mac OS X и Windows .

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

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

Платформа CUDA Впервые появились на рынке с выходом чипа NVIDIA восьмого поколения G80 и стала присутствовать во всех последующих сериях графических чипов, которые используются в семействах ускорителей GeForce , Quadro и NVidia Tesla .

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

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

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

Ограничения

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

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

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

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

Версия спецификации	GPU	Видеокарты
1.0	G80, G92, G92b, G94, G94b	GeForce 8800GTX/Ultra, 9400GT, 9600GT, 9800GT, Tesla C/D/S870, FX4/5600, 360M, GT 420
1.1	G86, G84, G98, G96, G96b, G94, G94b, G92, G92b	GeForce 8400GS/GT, 8600GT/GTS, 8800GT/GTS, 9600 GSO, 9800GTX/GX2, GTS 250, GT 120/30/40, FX 4/570, 3/580, 17/18/3700, 4700x2, 1xxM, 32/370M, 3/5/770M, 16/17/27/28/36/37/3800M, NVS420/50
1.2	GT218, GT216, GT215	GeForce 210, GT 220/40, FX380 LP, 1800M, 370/380M, NVS 2/3100M
1.3	GT200, GT200b	GeForce GTX 260, GTX 275, GTX 280, GTX 285, GTX 295, Tesla C/M1060, S1070, Quadro CX, FX 3/4/5800
2.0	GF100, GF110	GeForce (GF100) GTX 465, GTX 470, GTX 480, Tesla C2050, C2070, S/M2050/70, Quadro Plex 7000, Quadro 4000, 5000, 6000, GeForce (GF110) GTX 560 TI 448, GTX570, GTX580, GTX590
2.1	GF104, GF114, GF116, GF108, GF106	GeForce 610M, GT 430, GT 440, GTS 450, GTX 460, GTX 550 Ti, GTX 560, GTX 560 Ti, 500M, Quadro 600, 2000
3.0	GK104, GK106, GK107	GeForce GTX 690, GTX 680, GTX 670, GTX 660 Ti, GTX 660, GTX 650 Ti, GTX 650, GT 640, GeForce GTX 680MX, GeForce GTX 680M, GeForce GTX 675MX, GeForce GTX 670MX, GTX 660M, GeForce GT 650M, GeForce GT 645M, GeForce GT 640M
3.5	GK110

Nvidia GeForce для настольных компьютеров GeForce GTX 590 GeForce GTX 580 GeForce GTX 570 GeForce GTX 560 Ti GeForce GTX 560 GeForce GTX 550 Ti GeForce GTX 520 GeForce GTX 480 GeForce GTX 470 GeForce GTX 465 GeForce GTX 460 GeForce GTS 450 GeForce GTX 295 GeForce GTX 285 GeForce GTX 280 GeForce GTX 275 GeForce GTX 260 GeForce GTS 250 GeForce GT 240 GeForce GT 220 GeForce 210 GeForce GTS 150 GeForce GT 130 GeForce GT 120 GeForce G100 GeForce 9800 GX2 GeForce 9800 GTX+ GeForce 9800 GTX GeForce 9800 GT GeForce 9600 GSO GeForce 9600 GT GeForce 9500 GT GeForce 9400 GT GeForce 9400 mGPU GeForce 9300 mGPU GeForce 8800 GTS 512 GeForce 8800 GT GeForce 8600 GTS GeForce 8600 GT GeForce 8500 GT GeForce 8400 GS

Nvidia GeForce для мобильных компьютеров GeForce GTX 580M GeForce GTX 570M GeForce GTX 560M GeForce GT 555M GeForce GT 540M GeForce GT 525M GeForce GT 520M GeForce GTX 485M GeForce GTX 480M GeForce GTX 470M GeForce GTX 460M GeForce GT 445M GeForce GT 435M GeForce GT 425M GeForce GT 420M GeForce GT 415M GeForce GTX 285M GeForce GTX 280M GeForce GTX 260M GeForce GTS 360M GeForce GTS 350M GeForce GTS 160M GeForce GTS 150M GeForce GT 335M GeForce GT 330M GeForce GT 325M GeForce GT 240M GeForce GT 130M GeForce G210M GeForce G110M GeForce G105M GeForce 310M GeForce 305M GeForce 9800M GTX GeForce 9800M GT GeForce 9800M GTS GeForce 9700M GTS GeForce 9700M GT GeForce 9650M GS GeForce 9600M GT GeForce 9600M GS GeForce 9500M GS GeForce 9500M G GeForce 9300M GS GeForce 9300M G GeForce 9200M GS GeForce 9100M G GeForce 8800M GTS GeForce 8700M GT GeForce 8600M GT GeForce 8600M GS GeForce 8400M GT GeForce 8400M GS

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

Nvidia Quadro для настольных компьютеров Quadro 6000 Quadro 5000 Quadro 4000 Quadro 2000 Quadro 600 Quadro FX 5800 Quadro FX 5600 Quadro FX 4800 Quadro FX 4700 X2 Quadro FX 4600 Quadro FX 3700 Quadro FX 1700 Quadro FX 570 Quadro FX 470 Quadro FX 380 Low Profile Quadro FX 370 Quadro FX 370 Low Profile Quadro CX Quadro NVS 450 Quadro NVS 420 Quadro NVS 290 Quadro Plex 2100 D4 Quadro Plex 2200 D2 Quadro Plex 2100 S4 Quadro Plex 1000 Model IV

Nvidia Quadro для мобильных компьютеров Quadro 5010M Quadro 5000M Quadro 4000M Quadro 3000M Quadro 2000M Quadro 1000M Quadro FX 3800M Quadro FX 3700M Quadro FX 3600M Quadro FX 2800M Quadro FX 2700M Quadro FX 1800M Quadro FX 1700M Quadro FX 1600M Quadro FX 880M Quadro FX 770M Quadro FX 570M Quadro FX 380M Quadro FX 370M Quadro FX 360M Quadro NVS 5100M Quadro NVS 4200M Quadro NVS 3100M Quadro NVS 2100M Quadro NVS 320M Quadro NVS 160M Quadro NVS 150M Quadro NVS 140M Quadro NVS 135M Quadro NVS 130M

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

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

Feature support (unlisted features are supported for all compute capabilities)	Compute capability (version)
	1.0	1.1	1.2	1.3	2.x
32-bit words in global memory	Нет	Да
floating point values in global memory
Integer atomic functions operating on 32-bit words in shared memory	Нет		Да
atomicExch() operating on 32-bit floating point values in shared memory
Integer atomic functions operating on 64-bit words in global memory
Warp vote functions
Double-precision floating-point operations	Нет			Да
Atomic functions operating on 64-bit integer values in shared memory	Нет				Да
Floating-point atomic addition operating on 32-bit words in global and shared memory
_ballot()
_threadfence_system()
_syncthreads_count(), _syncthreads_and(), _syncthreads_or()
Surface functions
3D grid of thread block

Technical specifications	Compute capability (version)
	1.0	1.1	1.2	1.3	2.x
Maximum dimensionality of grid of thread blocks	2				3
Maximum x-, y-, or z-dimension of a grid of thread blocks	65535
Maximum dimensionality of thread block	3
Maximum x- or y-dimension of a block	512				1024
Maximum z-dimension of a block	64
Maximum number of threads per block	512				1024
Warp size	32
Maximum number of resident blocks per multiprocessor	8
Maximum number of resident warps per multiprocessor	24		32		48
Maximum number of resident threads per multiprocessor	768		1024		1536
Number of 32-bit registers per multiprocessor	8 K		16 K		32 K
Maximum amount of shared memory per multiprocessor	16 KB				48 KB
Number of shared memory banks	16				32
Amount of local memory per thread	16 KB				512 KB
Constant memory size	64 KB
Cache working set per multiprocessor for constant memory	8 KB
Cache working set per multiprocessor for texture memory	Device dependent, between 6 KB and 8 KB
Maximum width for 1D texture	8192				32768
Maximum width for 1D texture reference bound to linear memory	2 27
Maximum width and number of layers for a 1D layered texture reference	8192 x 512				16384 x 2048
Maximum width and height for 2D texture reference bound to linear memory or a CUDA array	65536 x 32768				65536 x 65535
Maximum width, height, and number of layers for a 2D layered texture reference	8192 x 8192 x 512				16384 x 16384 x 2048
Maximum width, height and depth for a 3D texture reference bound to linear memory or a CUDA array	2048 x 2048 x 2048
Maximum number of textures that can be bound to a kernel	128
Maximum width for a 1D surface reference bound to a CUDA array	Not supported				8192
Maximum width and height for a 2D surface reference bound to a CUDA array					8192 x 8192
Maximum number of surfaces that can be bound to a kernel					8
Maximum number of instructions per kernel	2 million

Пример

CudaArray* cu_array; texture< float , 2 > tex; // Allocate array cudaMalloc( & cu_array, cudaCreateChannelDesc< float> () , width, height ) ; // Copy image data to array cudaMemcpy( cu_array, image, width* height, cudaMemcpyHostToDevice) ; // Bind the array to the texture cudaBindTexture( tex, cu_array) ; // Run kernel dim3 blockDim(16 , 16 , 1 ) ; dim3 gridDim(width / blockDim.x , height / blockDim.y , 1 ) ; kernel<<< gridDim, blockDim, 0 >>> (d_odata, width, height) ; cudaUnbindTexture(tex) ; __global__ void kernel(float * odata, int height, int width) { unsigned int x = blockIdx.x * blockDim.x + threadIdx.x ; unsigned int y = blockIdx.y * blockDim.y + threadIdx.y ; float c = texfetch(tex, x, y) ; odata[ y* width+ x] = c; }

Import pycuda.driver as drv import numpy drv.init () dev = drv.Device (0 ) ctx = dev.make_context () mod = drv.SourceModule (""" __global__ void multiply_them(float *dest, float *a, float *b) { const int i = threadIdx.x; dest[i] = a[i] * b[i]; } """ ) multiply_them = mod.get_function ("multiply_them" ) a = numpy.random .randn (400 ) .astype (numpy.float32 ) b = numpy.random .randn (400 ) .astype (numpy.float32 ) dest = numpy.zeros_like (a) multiply_them( drv.Out (dest) , drv.In (a) , drv.In (b) , block= (400 , 1 , 1 ) ) print dest-a*b

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

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

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

Ссылки

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

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

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

Tom"s Hardware
• Дмитрий Чеканов. nVidia CUDA: вычисления на видеокарте или смерть CPU? . Tom"s Hardware (22 июня 2008 г.). Архивировано
• Дмитрий Чеканов. nVidia CUDA: тесты приложений на GPU для массового рынка . Tom"s Hardware (19 мая 2009 г.). Архивировано из первоисточника 4 марта 2012. Проверено 19 мая 2009.
iXBT.com
• Алексей Берилло. NVIDIA CUDA - неграфические вычисления на графических процессорах. Часть 1 . iXBT.com (23 сентября 2008 г.). Архивировано из первоисточника 4 марта 2012. Проверено 20 января 2009.
• Алексей Берилло. NVIDIA CUDA - неграфические вычисления на графических процессорах. Часть 2 . iXBT.com (22 октября 2008 г.). - Примеры внедрения NVIDIA CUDA. Архивировано из первоисточника 4 марта 2012. Проверено 20 января 2009.
Другие ресурсы
• Боресков Алексей Викторович. Основы CUDA (20 января 2009 г.). Архивировано из первоисточника 4 марта 2012. Проверено 20 января 2009.
• Владимир Фролов. Введение в технологию CUDA . Сетевой журнал «Компьютерная графика и мультимедиа» (19 декабря 2008 г.). Архивировано из первоисточника 4 марта 2012. Проверено 28 октября 2009.
• Игорь Осколков. NVIDIA CUDA – доступный билет в мир больших вычислений . Компьютерра (30 апреля 2009 г.). Проверено 3 мая 2009.
• Владимир Фролов. Введение в технологию CUDA (1 августа 2009 г.). Архивировано из первоисточника 4 марта 2012. Проверено 3 апреля 2010.
• GPGPU.ru . Использование видеокарт для вычислений
• . Центр Параллельных Вычислений

Примечания

См. также

Nvidia
Графические процессоры	Ранние NV1 NV2 Семейство RIVA TNT TNT2 Семейство GeForce GeForce 256 GeForce 2 GeForce 3 GeForce 4 GeForce FX GeForce 6 GeForce 7 GeForce 8 GeForce 9 GeForce 100 GeForce 200 GeForce 400 GeForce 500 GeForce 600

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