Cuda что это такое в видеокарте
Сложность вычислительных заданий требует резкого увеличения ресурсов и скорости компьютеров. Наиболее перспективным направлением повышения скорости решения задач является внедрение идей параллелизма в работу вычислительных систем.
Сегодня спроектированы и испытаны сотни различных компьютеров, которые используют в своей архитектуре тот или иной вид параллельной обработки данных. Основная сложность при проектировании параллельных программ – обеспечение правильной последовательности взаимодействия между разными вычислительными процессами, а также координация ресурсов, которые разделяются между ними.
CUDA – это программно-аппаратная архитектура параллельных вычислений, позволяющая существенно увеличить вычислительную продуктивность благодаря использованию графических процессоров NVIDIA.
При использовании данной технологии необходимо знать следующие понятия:
- устройство (device) – сама видеокарта, графический процессор (GPU) – выполняет команды центрального процессора;
- хост (host) – центральный процессор (CPU) – запускает различные задания, выделяет память, etc.;
- ядро (kernel) – функция (задание), которая будет выполняться на GPU.
CUDA позволяет программистам реализовывать на специальном упрощенном диалекте языка C алгоритмы, которые используются в графических процессорах NVIDIA, и включать специальные функции в текст программы на C.
"Архитектура CUDA позволяет разработчику на свое усмотрение организовывать доступ к набору инструкций GPU и управлять его памятью."
Эта технология поддерживает несколько языков программирования. Среди них Java, Python и некоторые другие.
Рассмотрим, как происходит запуск программы на графическом процессоре:
- Хост выделяет необходимое количество памяти на устройстве.
- Хост копирует данные из своей памяти в память устройства.
- Хост запускает ядро на устройстве.
- Устройство исполняет это ядро.
- Хост копирует результаты из памяти устройства в свою память.
На рисунке изображены все перечисленные шаги запуска программы, кроме первого (источник).
Взаимодействие CPU и GPU
Как видно из рисунка, центральный процессор взаимодействует с графическим через CUDA Runtime API, CUDA Driver API и CUDA Libraries. Runtime и Driver API отличаются уровнем абстракции. Грубо говоря, первый вариант более высокого уровня в плане программирования, более абстрактный, а второй – напротив, более низкого (уровень драйвера).
В целом Runtime API является абстрактной оберткой Driver API. Во время программирования вы можете использовать любой из представленных вариантов. Из личного опыта: при использовании Driver API нужно написать немного «лишнего» кода + данный вариант сложнее.
Также необходимо понять одну важную вещь, которая впоследствии сэкономит вам время и нервы:
"Если отношение времени, потраченного на работу ядер, окажется меньше времени, потраченного на выделение памяти и запуск этих ядер, вы получите нулевую эффективность от использования GPU."
Давайте разберем написанное подробнее. Чтобы запустить некоторые задачи на GPU, необходимо потратить «немного» времени на выделение памяти, копирование результата, etc., поэтому не нужно выполнять на графическом процессоре легкие задания, которые на деле занимают буквально миллисекунды. Зачем выполнять на GPU то, с чем легко, а главное, быстрее справится центральный процессор?
У вас возникнет вопрос: «Тогда зачем вообще использовать GPU, если при этом приходится тратить драгоценное время на выделение памяти и другие ненужные вещи?». Это заблуждение, и со временем вы поймете, что CUDA – действительно мощная технология. Дальше разберемся, почему это так.
Архитектура GPU построена несколько иначе, нежели CPU. Поскольку графические процессоры сперва использовались только для графических расчетов, которые допускают независимую параллельную обработку данных, то GPU и предназначены именно для параллельных вычислений. Он спроектирован таким образом, чтобы выполнять огромное количество потоков (элементарных параллельных процессов).
Архитектура CPU и GPU
Как видно из картинки – в GPU есть много простых арифметически-логических устройств (АЛП), которые объединены в несколько групп и обладают общей памятью. Это помогает повысить продуктивность в вычислительных заданиях, но немного усложняет программирование.
«Для достижения лучшего ускорения необходимо продумывать стратегии доступа к памяти и учитывать аппаратные особенности.»
GPU ориентирован на выполнение программ с большим объемом данных и расчетов и представляет собой массив потоковых процессоров (Streaming Processor Array), что состоит из кластеров текстурных процессоров (Texture Processor Clusters, TPC). TPC в свою очередь состоит из набора мультипроцессоров (SM – Streaming Multi-processor), в каждом из которых несколько потоковых процессоров (SP – Streaming Processors) или ядер (в современных процессорах количество ядер превышает 1024).
Набор ядер каждого мультипроцессора работает по принципу SIMD (но с некоторым отличием) – реализация, которая позволяет группе процессоров, работающих параллельно, работать с различными данными, но при этом все они в любой момент времени должны выполнять одинаковую команду. Говоря проще, несколько потоков выполняют одно и то же задание.
Мультипроцессоры, SM
В результате GPU фактически стал устройством, которое реализует потоковую вычислительную модель (stream computing model): есть потоки входящих и исходящих данных, что состоят из одинаковых элементов, которые могут быть обработаны независимо друг от друга.
Продолжаем разбираться с CUDA. Каждая видеокарта обладает так называемыми compute capabilities – количественными характеристиками скорости выполнения определенных операций на графическом процессоре. Данное число показывает, насколько быстро видеокарта будет выполнять свою работу.
В NVIDIA эту характеристику обозначают Compute Capability Version. В таблице приведены некоторые видеокарты и соответствующие им вычислительные возможности:
Полный перечень можно посмотреть здесь. Compute Capability Version описывает множество параметров, среди которых: количество потоков на блок, максимальное количество блоков и потоков, размер warp, а также многое другое.
CUDA использует большое количество отдельных потоков для расчетов. Все они группируются в иерархию – grid / block / thread.
Верхний уровень – grid – отвечает ядру и объединяет все потоки, которые выполняет данное ядро. Grid – одномерный или двумерный массив блоков (block). Каждый блок (block) представляет собой полностью независимый набор скоординированных между собой потоков. Потоки из разных блоков не могут взаимодействовать.
Мы упоминали об отличии от SIMD-архитектуры. Есть такое понятие, как warp – группа из 32 потоков (в зависимости от архитектуры GPU, но почти всегда 32). Только потоки в рамках одной группы (warp) могут физически выполняться одновременно. Потоки разных варпов могут находиться на разных стадиях выполнения программы. Такой метод обработки данных обозначается термином SIMT (Single Instruction – Multiple Theads). Управление работой варпов выполняется на аппаратном уровне.
Выше уже было написано, что не стоит выполнять на GPU слишком простые задания. Чтобы понять, следует определить два термина:
- Задержка – это преимущественно время ожидания между запросом на какой-либо ресурс и получением доступа к данному ресурсу.
- Пропускная способность – количество операций, которые выполняются за единицу времени.
Таким образом, главный вопрос состоит в следующем: почему графический процессор иногда «тупит»? Объясняем на простом примере.
У нас есть 2 автомобиля:
- легковой фургон – скорость 120 км/ч, способен вместить 9 человек;
- автобус – скорость 90 км/ч, способен вместить 30 человек.
Если одна операция – это передвижение одного человека на определенное расстояние (пусть будет 1 км), то задержка (время, за которое один человек пройдет 1 км) для первого авто составит 3600/120 = 30 сек, а пропускная способность – 9/30 = 0,3. Для автобуса – 3600/90 = 40 сек и 30/40 = 0,75.
CPU – это фургон, а GPU – автобус: у него большая задержка, но также и большая пропускная способность. Если для вашего задания задержка каждой конкретной операции не так важна, как количество этих самых операций в секунду, то стоит рассмотреть использование GPU.
Отличительными чертами GPU в сравнении с CPU являются:
- архитектура, максимально нацеленная на увеличение скорости расчета текстур и сложных графических объектов;
- предельная мощность типичного GPU намного больше, чем у CPU;
- благодаря специализированной конвейерной архитектуре GPU более эффективен в обработке графической информации, нежели центральный процессор.
Главный минус CUDA в том, что данная технология поддерживается только видеокартами NVIDIA без каких-либо альтернатив.
Графический процессор не всегда может дать ускорение при выполнении определенных алгоритмов. Поэтому перед использованием GPU для вычислений стоит хорошо подумать, а нужен ли он в данном случае. Вы можете использовать видеокарту для сложных вычислений: работа с графикой или изображениями, инженерные расчеты, криптографические задачи (майнинг), и т. д., но не используйте GPU для решения простых задач (разумеется, вы можете, но тогда эффективность будет равняться нулю).
Помните о задаче с фургоном и автобусом, а также не забывайте, что использование графического процессора гораздо вероятнее замедлит программу, нежели ускорит ее.
Внутренняя модель nVidia GPU – ключевой момент в понимании GPGPU с использованием CUDA. В этот раз я постараюсь наиболее детально рассказать о программном устройстве GPUs. Я расскажу о ключевых моментах компилятора CUDA, интерфейсе CUDA runtime API, ну, и в заключение, приведу пример использования CUDA для несложных математических вычислений.
CUDA host API:
Перед тем, как приступить к непосредственному использованию CUDA для вычислений, необходимо ознакомиться с так называемым CUDA host API, который является связующим звеном между CPU и GPU. CUDA host API в свою очередь можно разделить на низкоуровневое API под названием CUDA driver API, который предоставляет доступ к драйверу пользовательского режима CUDA, и высокоуровневое API – CUDA runtime API. В своих примерах я буду использовать CUDA runtime API.
- Device Management – включает функции для общего управления GPU (получение инфор-мации о возможностях GPU, переключение между GPU при работе SLI-режиме и т.д.).
- Thread Management – управление нитями.
- Stream Management – управление потоками.
- Event Management – функция создания и управления event’ами.
- Execution Control – функции запуска и исполнения ядра CUDA.
- Memory Management – функции управлению памятью GPU.
- Texture Reference Manager – работа с объектами текстур через CUDA.
- OpenGL Interoperability – функции по взаимодействию с OpenGL API.
- Direct3D 9 Interoperability – функции по взаимодействию с Direct3D 9 API.
- Direct3D 10 Interoperability – функции по взаимодействию с Direct3D 10 API.
- Error Handling – функции обработки ошибок.
Понимаем работу GPU:
Как было сказано, нить – непосредственный исполнитель вычислений. Каким же тогда образом происходит распараллеливание вычислений между нитями? Рассмотрим работу отдельно взятого блока.
Задача. Требуется вычислить сумму двух векторов размерностью N элементов.
Нам известна максимальные размеры нашего блока: 512*512*64 нитей. Так как вектор у нас одномерный, то пока ограничимся использованием x-измерения нашего блока, то есть задействуем только одну полосу нитей из блока (рис. 3).
Рис. 3. Наша полоса нитей из используемого блока.
- Получить данные для расчетов.
- Скопировать эти данные в GPU память.
- Произвести вычисление в GPU через функцию ядра.
- Скопировать вычисленные данные из GPU памяти в ОЗУ.
- Посмотреть результаты.
- Высвободить используемые ресурсы.
Первым делом напишем функцию ядра, которая и будет осуществлять сложение векторов:
// Функция сложения двух векторов
__global__ void addVector( float * left, float * right, float * result)
//Получаем id текущей нити.
int idx = threadIdx.x;
//Расчитываем результат.
result[idx] = left[idx] + right[idx];
>
* This source code was highlighted with Source Code Highlighter .
Таким образом, распараллеливание будет выполнено автоматически при запуске ядра. В этой функции так же используется встроенная переменная threadIdx и её поле x, которая позволяет задать соответствие между расчетом элемента вектора и нитью в блоке. Делаем расчет каждого элемента вектора в отдельной нити.
Пишем код, которые отвечает за 1 и 2 пункт в программе:
//Инициализируем значения векторов
for ( int i = 0; i < SIZE; i++)
vec1[i] = i;
vec2[i] = i;
>
//Указатели на память видеокарте
float * devVec1;
float * devVec2;
float * devVec3;
//Копируем данные в память видеокарты
cudaMemcpy(devVec1, vec1, sizeof ( float ) * SIZE, cudaMemcpyHostToDevice);
cudaMemcpy(devVec2, vec2, sizeof ( float ) * SIZE, cudaMemcpyHostToDevice);
…
>
* This source code was highlighted with Source Code Highlighter .
- devPtr – указатель, в который записывается адрес выделенной памяти,
- count – размер выделяемой памяти в байтах.
- cudaSuccess – при удачном выделении памяти
- cudaErrorMemoryAllocation – при ошибке выделения памяти
- dst – указатель, содержащий адрес места-назначения копирования,
- src – указатель, содержащий адрес источника копирования,
- count – размер копируемого ресурса в байтах,
- cudaMemcpyKind – перечисление, указывающее направление копирования (может быть cudaMemcpyHostToDevice, cudaMemcpyDeviceToHost, cudaMemcpyHostToHost, cudaMemcpyDeviceToDevice).
- cudaSuccess – при удачном копировании
- cudaErrorInvalidValue – неверные параметры аргумента (например, размер копирования отрицателен)
- cudaErrorInvalidDevicePointer – неверный указатель памяти в видеокарте
- cudaErrorInvalidMemcpyDirection – неверное направление (например, перепутан источник и место-назначение копирования)
…
dim3 gridSize = dim3(1, 1, 1); //Размер используемого грида
dim3 blockSize = dim3(SIZE, 1, 1); //Размер используемого блока
//Выполняем вызов функции ядра
addVector>>(devVec1, devVec2, devVec3);
…
* This source code was highlighted with Source Code Highlighter .
В нашем случае определять размер грида и блока необязательно, так как используем всего один блок и одно измерение в блоке, поэтому код выше можно записать:
* This source code was highlighted with Source Code Highlighter .
Теперь нам остаеться скопировать результат расчета из видеопамяти в память хоста. Но у функций ядра при этом есть особенность – асинхронное исполнение, то есть, если после вызова ядра начал работать следующий участок кода, то это ещё не значит, что GPU выполнил расчеты. Для завершения работы заданной функции ядра необходимо использовать средства синхронизации, например event’ы. Поэтому, перед копированием результатов на хост выполняем синхронизацию нитей GPU через event.
Код после вызова ядра:
//Выполняем вызов функции ядра
addVector>>(devVec1, devVec2, devVec3);
//Хендл event'а
cudaEvent_t syncEvent;
cudaEventCreate(&syncEvent); //Создаем event
cudaEventRecord(syncEvent, 0); //Записываем event
cudaEventSynchronize(syncEvent); //Синхронизируем event
//Только теперь получаем результат расчета
cudaMemcpy(vec3, devVec3, sizeof ( float ) * SIZE, cudaMemcpyDeviceToHost);
* This source code was highlighted with Source Code Highlighter .
Рассмотрим более подробно функции из Event Managment API.
- *event – указатель для записи хендла event’а.
- cudaSuccess – в случае успеха
- cudaErrorInitializationError – ошибка инициализации
- cudaErrorPriorLaunchFailure – ошибка при предыдущем асинхронном запуске функции
- cudaErrorInvalidValue – неверное значение
- cudaErrorMemoryAllocation – ошибка выделения памяти
- event – хендл хаписываемого event’а,
- stream – номер потока, в котором записываем (в нашем случае это основной нулевой по-ток).
- cudaSuccess – в случае успеха
- cudaErrorInvalidValue – неверное значение
- cudaErrorInitializationError – ошибка инициализации
- cudaErrorPriorLaunchFailure – ошибка при предыдущем асинхронном запуске функции
- cudaErrorInvalidResourceHandle – неверный хендл event’а
- event – хендл event’а, прохождение которого ожидается.
- cudaSuccess – в случае успеха
- cudaErrorInitializationError – ошибка инициализации
- cudaErrorPriorLaunchFailure – ошибка при предыдущем асинхронном запуске функции
- cudaErrorInvalidValue – неверное значение
- cudaErrorInvalidResourceHandle – неверный хендл event’а
Рис. 4. Синхронизация работы основоной и GPU прграмм.
На рисунке 4 блок «Ожидание прохождения Event’а» и есть вызов функции cudaEventSynchronize.
Ну и в заключении выводим результат на экран и чистим выделенные ресурсы.
cudaFree(devVec1);
cudaFree(devVec2);
cudaFree(devVec3);
delete[] vec1; vec1 = 0;
delete[] vec2; vec2 = 0;
delete[] vec3; vec3 = 0;
* This source code was highlighted with Source Code Highlighter .
Думаю, что описывать функции высвобождения ресурсов нет необходимости. Разве что, можно напомнить, что они так же возвращают значения cudaError_t, если есть необходимость проверки их работы.
CUDA и язык C:
- Спецификаторы функций, которые показывают, как и откуда буду выполняться функции.
- Спецификаторы переменных, которые служат для указания типа используемой памяти GPU.
- Спецификаторы запуска ядра GPU.
- Встроенные переменные для идентификации нитей, блоков и др. параметров при исполнении кода в ядре GPU .
- Дополнительные типы переменных.
- __host__ — выполнятся на CPU, вызывается с CPU (в принципе его можно и не указывать).
- __global__ — выполняется на GPU, вызывается с CPU.
- __device__ — выполняется на GPU, вызывается с GPU.
- gridSize – размерность сетки блоков (dim3), выделенную для расчетов,
- blockSize – размер блока (dim3), выделенного для расчетов,
- sharedMemSize – размер дополнительной памяти, выделяемой при запуске ядра,
- cudaStream – переменная cudaStream_t, задающая поток, в котором будет произведен вызов.
- gridDim – размерность грида, имеет тип dim3. Позволяет узнать размер гридa, выделенного при текущем вызове ядра.
- blockDim – размерность блока, так же имеет тип dim3. Позволяет узнать размер блока, выделенного при текущем вызове ядра.
- blockIdx – индекс текущего блока в вычислении на GPU, имеет тип uint3.
- threadIdx – индекс текущей нити в вычислении на GPU, имеет тип uint3.
- warpSize – размер warp’а, имеет тип int (сам еще не пробовал использовать).
Дополнительные типы переменных и их спецификаторы будут рассмотрены непосредственно в примерах работы с памятью.
Заключение
Надеюсь, что этот материал поможет вам понять, как функционирует GPU. Я описал самые главные моменты, которые необходимо знать для работы с CUDA. Попробуйте сами написать сложение двух матриц, но не забывайте об аппаратных ограничениях видеокарты.
P.S.: Получилось не очень кратко. Надеюсь, что не утомил. Если нужен весь исходный код, то могу выслать на почту.
P.S.S: Задавайте вопросы.
На Хабре было уже немало хороших статей по CUDA — раз, два и другие. Однако, поиск комбинации «CUDA scan» выдал всего 2 статьи никак не связанные с, собственно, алгоритмом scan на GPU — а это один из самых базовых алгоритмов. Поэтому, вдохновившись только что просмотренным курсом на Udacity — Intro to Parallel Programming, я и решился написать более полную серию статей о CUDA. Сразу скажу, что серия будет основываться именно на этом курсе, и если у вас есть время — намного полезнее будет пройти его.
Содержание
На данный момент планируются следующие статьи:
Часть 1: Введение.
Часть 2: Аппаратное обеспечение GPU и шаблоны параллельной коммуникации.
Часть 3: Фундаментальные алгоритмы GPU: свертка (reduce), сканирование (scan) и гистограмма (histogram).
Часть 4: Фундаментальные алгоритмы GPU: уплотнение (compact), сегментированное сканирование (segmented scan), сортировка. Практическое применение некоторых алгоритмов.
Часть 5: Оптимизация GPU программ.
Часть 6: Примеры параллелизации последовательных алгоритмов.
Часть 7: Дополнительные темы параллельного программирования, динамический параллелизм.
Задержка vs пропускная способность
Первый вопрос, который должен задать каждый перед применением GPU для решения своих задач — а для каких целей хорош GPU, когда стоит его применять? Для ответа нужно определить 2 понятия:
Задержка (latency) — время, затрачиваемое на выполнение одной инструкции/операции.
Пропускная способность — количество инструкций/операций, выполняемых за единицу времени.
Простой пример: имеем легковой автомобиль со скоростью 90 км/ч и вместимостью 4 человека, и автобус со скоростью 60 км/ч и вместимостью 20 человек. Если за операцию принять перемещение 1 человека на 1 километр, то задержка легкового автомобиля — 3600/90=40с — за столько секунд 1 человек преодолеет расстояние в 1 километр, пропускная способность автомобиля — 4/40=0.1 операций/секунду; задержка автобуса — 3600/60=60с, пропускная способность автобуса — 20/60=0.3(3) операций/секунду.
Так вот, CPU — это автомобиль, GPU — автобус: он имеет большую задержку но также и большую пропускную способность. Если для вашей задачи задержка каждой конкретной операции не настолько важна как количество этих операций в секунду — стоит рассмотреть применение GPU.
Базовые понятия и термины CUDA
- Устройство (device) — GPU. Выполняет роль «подчиненного» — делает только то, что ему говорит CPU.
- Хост (host) — CPU. Выполняет управляющую роль — запускает задачи на устройстве, выделяет память на устройстве, перемещает память на/с устройства. И да, использование CUDA предполагает, что как устройство так и хост имеют свою отдельную память.
- Ядро (kernel) — задача, запускаемая хостом на устройстве.
Основные этапы CUDA-программы
- Хост выделяет нужное количество памяти на устройстве.
- Хост копирует данные из своей памяти в память устройства.
- Хост стартует выполнение определенных ядер на устройстве.
- Устройство выполняет ядра.
- Хост копирует результаты из памяти устройства в свою память.
- Внутри ядер вы имеете возможность узнать «идентификатор» или, проще говоря, позицию потока, который сейчас выполняется — используя эту позицию мы добиваемся того, что одно и то же ядро будет работать с разными данными в зависимости от потока, в котором оно запущено. Кстати, такая организация параллельных вычислений называется SIMD (Single Instruction Multiple Data) — когда несколько процессоров выполняют одновременно одну и ту же операцию но на разных данных.
- В некоторых случаях в коде ядра необходимо использовать различные способы синхронизации.
- Сначала задаются размеры так называемой сетки (grid), в 3D координатах: grid_x, grid_y, grid_z. В результате, сетка будет состоять из grid_x*grid_y*grid_z блоков.
- Потом задаются размеры блока в 3D координатах: block_x, block_y, block_z. В результате, блок будет состоять из block_x*block_y*block_z потоков. Итого, имеем grid_x*grid_y*grid_z*block_x*block_y*block_z потоков. Важное замечание — максимальное количество потоков в одном блоке ограничено и зависит от модели GPU — типичны значения 512 (более старые модели) и 1024 (более новые модели).
- Внутри ядра доступны переменные threadIdx и blockIdx с полями x, y, z — они содержат 3D координаты потока в блоке и блока в сетке соответственно. Также доступны переменные blockDim и gridDim с теми же полями — размеры блока и сетки соответственно.
Пишем первую программу на CUDA
Тут все довольно очевидно — читаем файл с изображением, подготавливаем указатели на цветное и в оттенках серого изображение, запускаем вариант
с OpenMP и вариант с CUDA, замеряем время. Функция prepareImagePointers имеет следующий вид:
Я пошел на небольшую хитрость: дело в том, что мы выполняем очень мало работы на каждый пиксел изображения — то-есть при варианте с CUDA встает упомянутая выше проблема соотношения времени выполнения полезных операций к времени выделения памяти и копирования данных, и в результате общее время CUDA варианта будет больше OpenMP варианта, а мы же хотим показать что CUDA быстрее:) Поэтому для CUDA будет измеряться только время, потраченное на выполнение собственно конвертации изображения — без учета операций с памятью. В свое оправдание скажу, что для большого класса задач время полезной работы будет все-таки доминировать, и CUDA будет быстрее даже с учетом операций с памятью.
Далее напишем код для OpenMP варианта:
Все довольно прямолинейно — мы всего лишь добавили директиву omp parallel for к однопоточному коду — в этом вся красота и мощь OpenMP. Я пробовал поиграться с параметром schedule, но получалось только хуже, чем без него.
Наконец, переходим к CUDA. Тут распишем более детально. Сначала нужно выделить память под входные данные, переместить их с CPU на GPU и выделить память под выходные данные:
Стоит обратить внимание на стандарт именования переменных в CUDA — данные на CPU начинаются с h_ (host), данные да GPU — с d_ (device). checkCudaErrors — макрос, взят с github-репозитория Udacity курса. Имеет следующий вид:
cudaMalloc — аналог malloc для GPU, cudaMemcpy — аналог memcpy, имеет дополнительный параметр в виде enum-а, который указывает тип копирования: cudaMemcpyHostToDevice, cudaMemcpyDeviceToHost, cudaMemcpyDeviceToDevice.
Далее необходимо задать размеры сетки и блока и вызвать ядро, не забыв измерить время:
Здесь мы вычисляем координаты y и x обрабатываемого пиксела, используя ранее описанные переменные threadIdx, blockIdx и blockDim, ну и выполняем конвертацию. Обратите внимание на проверку if (x>=numCols || y>=numRows) — так как размеры изображения не обязательно будут делится нацело на размеры блоков, некоторые блоки могут «выходить за рамки» изображения — поэтому необходима эта проверка. Также, функция ядра должна помечаться спецификатором __global__ .
Последний шаг — cкопировать результат назад с GPU на CPU и освободить выделенную память:
Кстати, CUDA позволяет использовать C++ компилятор для host-кода — так что запросто можно написать обертки для автоматического освобождения памяти.
Итак, запускаем, измеряем (размер входного изображения — 10,109 × 4,542):
Конфигурация машины, на которой проводились тесты:
Процессор: Intel® Core(TM) i7-3615QM CPU @ 2.30GHz.
GPU: NVIDIA GeForce GT 650M, 1024 MB, 900 MHz.
RAM: DD3, 2x4GB, 1600 MHz.
OS: OS X 10.9.5.
Компилятор: g++ (GCC) 4.9.2 20141029.
CUDA компилятор: Cuda compilation tools, release 6.0, V6.0.1.
Поддерживаемая версия OpenMP: OpenMP 4.0.
Получилось как-то не очень впечатляюще:) А проблема все та же — слишком мало работы выполняется над каждым пикселом — мы запускаем тысячи потоков, каждый из которых отрабатывает практически моментально. В случае с CPU такой проблемы не возникает — OpenMP запустит сравнительно малое количество потоков (8 в моем случае) и разделит работу между ними поровну — таким образом процессоры будет занят практически на все 100%, в то время как с GPU мы, по сути, не используем всю его мощь. Решение довольно очевидное — обрабатывать несколько пикселов в ядре. Новое, оптимизированное, ядро будет выглядеть следующим образом:
Здесь не все так просто как с предыдущим ядром. Если разобраться, теперь каждый поток будет обрабатывать elemsPerThread пикселов, причем не подряд, а с расстоянием в WARP_SIZE между ними. Что такое WARP_SIZE, почему оно равно 32, и зачем обрабатывать пиксели пободным образом, будет более детально рассказано в следующих частях, сейчас только скажу что этим мы добиваемся более эффективной работы с памятью. Каждый поток теперь обрабатывает elemsPerThread пикселов с расстоянием в WARP_SIZE между ними, поэтому x-координата первого пиксела для этого потока исходя из его позиции в блоке теперь рассчитывается по несколько более сложной формуле чем раньше.
Запускается это ядро следующим образом:
Количество блоков по x-координате теперь рассчитывается как numCols / (threadNum*elemsPerThread) + 1 вместо numCols / threadNum + 1. В остальном все осталось так же.
Запускаем:
Получили прирост по скорости в 2.76 раза (опять же, не учитывая время на операции с памятью) — для такой простой проблемы это довольно неплохо. Да-да, эта задача слишком простая — с ней достаточно хорошо справляется и CPU. Как видно из второго теста, простая реализация на GPU может даже проигрывать по скорости реализации на CPU.
На сегодня все, в следующей части рассмотрим аппаратное обеспечение GPU и основные шаблоны параллельной коммуникации.
Весь исходный код доступен на bitbucket.
Многие видели моё введение в современные технологии высокопроизводительных вычислений и оценки производительности, теперь я продолжу тему более подробным рассказом о технологии CUDA.
Для тех кто не смотрел предыдущие серии: CUDA позволяет писать и запускать на видеокартах nVidia(8xxx и выше) программы написанные на С++ со специальными расширениями. На правильных задачах достигается значительное превосходство по производительности на $ по сравнению с обычными CPU.
Достижимая производительность — 1 трлн и выше операций в секунду на GTX295.
NB: Статья — краткое введение, покрыть все ньюансы программирования под CUDA в одной статье вряд ли возможно :-)
CUDA работает на видеокартых начиная с 8400GS и выше. Разные видеокарты имеют разые возможности. В целом, если вы видите что в видеокарте например 128 SP(Streaming Processor) — это значит что там 8 SIMD MP (multiprocessor), каждый из которых делает одновременно 16 операций. На один MP есть 16кб shared memory, 8192 штуки 4-хбайтных регистров (В картах серии GTX2xx значения больше). Также есть 64кб констант общие для всех MP, они кешируются, при непопадании в кеш — достаточно большая задержка (400-600 тактов). Есть глобальная память видеокарты, доступ туда не кешируется, и текстуры (кешируется, кеш оптимизирован для 2D выборок). Для использования нескольких видеокарт нужно во первый отключать SLI в дровах, а во вторых — на каждую видеокарту запускать по потоку, и вызывать cudaSetDevice().
Самый быстрый способ научиться программировать на CUDA — это взять какой-нибуть пример из SDK, запустить его, и затем модифицировать, пока работает(собственно так я и делал, когда писал свой BarsWF ) :-)
Для начала идем на http://www.nvidia.com/object/cuda_get.html и качаем SDK и Toolkit под вашу операционную систему нужной битности. (к сожалению например 32-х битный SDK и 64-хбитный toolkit мешать нельзя). Полезно обновить драйвер видеокарты до последней версии (т.к. CUDA быстро развивается, всегда полезно иметь последние дрова, и вам и пользователям ваших программ).Сдесь я буду рассматривать разработку под Windows в Visual Studio (2005, недавно с 2008 тоже стало можно).
Для примера возьмем пример SDK Mandelbrot. Самое важное — это .cu файл, обратите внимание на его Custom Build Rule:
$(CUDA_BIN_PATH)\nvcc.exe -ccbin "$(VCInstallDir)bin" -c -DWIN32 -D_CONSOLE -D_MBCS -Xcompiler /EHsc,/W3,/nologo,/Wp64,/O2,/Zi,/MT -I"$(CUDA_INC_PATH)" -I./ -I../../common/inc -I"$(DXSDK_DIR)\Include" -o $(ConfigurationName)\Mandelbrot_sm10.obj Mandelbrot_sm10.cu
Его вы можете использовать во всех своих проектах, только вместо "../../common/inc " можно указать абсолютный путь (или переменную окружения).
nvcc — это и есть великий и ужасный компилатор CUDA. На выходе он генерирует объектный файл, в котором уже включена откомпилированная программа для видеокарты.
Обратите внимение на описание интерфейса в Mandelbrot_kernel.h — тут руками приходится описывать kernel-ы которые мы собираемся вызывать из основной С++ программы (впрочем их обычно не много, так что это не страшно).
После того как вам удалось запустить пример SDK, можно рассмотреть, чем же CUDA программа отличается от обычной.
NB: Если вы добавите параметр -keep то после компиляции сможете найти много занимательных промежуточных файлов.
Перед функциями в .cu файле могут стоять следующие «модификаторы»:
__device__ — это означает, что функция выполняется только на видеокарте. Из программы выполняющейся на обычном процессоре(host) её вызвать нельзя.
__global__ — Эта функция — начало вашего вычислительного ядра. Выполняется на видеокарте, но запускается только с хоста.
__host__ — Выполняется и запускается только с хоста (т.е. обычная функция C++). Если вы перед функцией укажите например __host__ и __device__ — то будут скомпилированы 2 версии функции (не все комбинации допустимы).
__device__ — означает что переменная находится в глобальной памяти видеокарты (т.е. которой там 512-1024Мб и выше). Очень медленная память по меркам вычислений на видеокарте(хоть и быстрее памяти центрального процессора в несколько раз), рекомендуется использовать как можно реже. В этой памяти данные сохраняются между вызовами разных вычислительных ядер. Данные сюда можно записывать и читать из host-части с помощью
cudaMemcpy(device_variable, host_variable, size, cudaMemcpyHostToDevice); //cudaMemcpyDeviceToHost - в обратную сторону
__constant__ — задает переменную в константной памяти. Следует обратить внимание, что значения для констант нужно загружать функцией cudaMemcpyToSymbol. Константы доступны из всех тредов, скорость работы сравнима с регистрами(когда в кеш попадает).
__shared__ — задает переменную в общей памяти блока тредов (т.е. и значение будет общее на всех). Тут нужно подходить с осторожностью — компилятор агрессивно оптимизирует доступ сюда(можно придушить модификатором volatile), можно получать race condition, нужно использовать __syncthreads(); чтобы данные гарантированно записались. Shared memory разделена на банки, и когда 2 потока одновременно пытаются обратиться к одному банку, возникает bank conflict и падает скорость.
Все локальные переменные которые вы определеили в ядре (__device__) — в регистрах, самая высокая скорость доступа.
Основая идея CUDA в том, что для решения вашей задачи вы запускаете тысячи и тысячи потоков, поэтому не стоит пугаться того что тут будет дальше написано :-)
Допустим, надо сделать какую-то операцию над картинкой 200x200. Картинка разбивается на куски 10x10, и на каждый пиксел такого кусочка запускаем по потоку. Выглядить это будет так:
dim3 threads(10, 10);//размер квардатика, 10*10
dim3 grid(20, 20);//сколько квадратиков нужно чтобы покрыть все изображение
your_kernel>>(image, 200,200);//Эта строка запустит 40'000 потоков (не одновременно, одновременно работать будет 200-2000 потоков примерно).
В отличии от Brook+ от AMD, где мы сразу определяем какому потоку над какими данными работать, в CUDA все не так: передаваеиые kernel-у параметры одинаковые для всех потоков, и поток должен сам получить данные для себя, чтобы сделать это, потоку нужно вычислить, в каком месте изображения он находится. В этом помогают магические переменные blockDim, blockIdx.
const int ix = blockDim.x * blockIdx.x + threadIdx.x;
const int iy = blockDim.y * blockIdx.y + threadIdx.y;
В ix и iy — координаты, с помощью которых можно получить исходные данные из массива image, и записать результат работы.
- Как можно меньше используйте __global__ память.
- При работе с __shared__ памятью избегайте конфлктов банков (впрочем многие задачи могут быть решены без shared памяти).
- Как можно меньше ветвлений в коде, где разные потоки идут по разным путям. Такой код не выполняется параллельно.
- Используйте как можно меньше памяти. Чем меньше памяти вы используете, тем агрессивнее компилятор и железо смогут запускать ваш kernel (например он может взять 100 тредов, и используя в 100 больше регистров запустить одновременно на одном MP, радикально уменьшая задержки)
Надеюсь это введение поможет людям, решившим попробовать программирование для видеокарт. Если есть проблемы/вопросы — буду рад помочь. Ну а в переди нас ждет введение в Brook+ и SIMD x86
На современных видеокартах можно делать параллельные математические вычисления, не связанные с 3D-графикой. Наиболее известные программные интерфейсы Direct3D (DirectX) или OpenGL, ответственные за вывод графики, в данном случае не подходят. На сегодня есть два основных типа интерфейсов для доступа к таким вычислениям, в том числе векторным, на GPU – CUDA и OpenCL. Стоит также уточнить, что в отличие от многоядерных CPU, ядра видеокарт менее универсальны и более просты. Из-за этого, их может быть на видеокарте огромное количество. nVidia GeForce RTX 2070 имеет 2560 CUDA ядер.
CUDA - аббревиатура запатентована фирмой nVidia и расшифровывается как Compute Unified Device Architecture. Это закрытая программно-аппаратная архитектура. CUDA состоит из самих вычислительных ядер и «родного» API - программного интерфейса. К дополнительным преимуществам можно отнести полную аппаратную поддержку целочисленных и побитовых операций и некоторые особенности в работе с памятью, что также ускоряет вычисления.
OpenCL – идеологически схожий открытый проект, который разрабатывается некоммерческим консорциумом «Khronos Group». В него входят такие крупные компании, как Intel, AMD, Apple и даже nVidia. Хотя она этого и не афиширует. У AMD Radeon доступ к ресурсам вычислений завязан на OpenCL. В отличие от CUDA не является «родным» и оптимизированным под конкретную архитектуру, хотя и более универсальный. Поэтому видеокарты от AMD могут проигрывать в среднем от 10% до 60% в тестах на такие параллельные математические вычисления. С выходом всё более новых драйверов и версий самого стандарта, ситуация улучшается. nVidia также поддерживает OpenCL, но «поверх» своего API CUDA, и в тестах этого стандарта будет еще менее производительной.
Количество исполнительных ядер в моделях при схожей цене у nVidia и AMD не будет сильно отличаться. А ядра у этих двух брендов объединены в потоковые мультипроцессоры. При схожей архитектуре и характеристиках, за счет CUDA AMD будет проигрывать в параллельных вычислениях на GPU.
Распараллеливание операций также возможно и на CPU. Даже одновременное использование мощностей видеокарты и центрального процессора. Но с той лишь разницей, что при наличии дискретной видеокарты, CPU станет узким местом, а общая производительность упадет.
Популярные рендер движки: V-Ray, Arnold, Redshift, Octane, RenderMan и другие менее известные, использование видеокарт от AMD через OpenCL на данный момент не поддерживаются. То же касается и Adobe Substance 3D Painter-a, как стандарта для текстурирования и запекания текстурных карт. Также всё не очень хорошо и с программами для симуляций, VFX, композа и видеомонтажа. Marmoset Toolbag также не поддерживает OpenCL на данный момент. Единственный пакет 3D-моделирования с поддержкой OpenCL рендера – Blender со встроенным рендер движком Cycles, EEVEE OpenCL не поддерживает.
Думаю, что в ближайшем будущем ситуация должна измениться, а стандарт OpenCL будет более популярным и получит поддержку основными программными пакетами в области графики и видео. Кроме общего развития стандартов, на это может повлиять кризис полупроводников и цены на рынке видеокарт среднего и топ-сегмента, которые в несколько раз выше рекомендованной и не доступны среднему пользователю. В начале января на выставке CES 2022, компания Intel представила свою линейку видеокарт на новой архитектуре Arc Alchemist. Но скорее всего в продажу они поступят не ранее третьего квартала 2022 года. Наиболее вероятным, будет использования Intel-ом OpenCL, так как внедрить новый стандарт в индустрию сей час крайне сложно и дорого.
В итоге всё сводится к программно-аппаратной реализации доступа и поддержке производителями софта. Вывод только один - на данный момент альтернативы CUDA нет.
Пишите своё мнение в комментариях, задавайте вопросы и подписывайтесь на канал. Если вам понравилась статья, поддержите лайком. Спасибо.
Вычислительная модель GPU:
-
Верхний уровень ядра GPU состоит из блоков, которые группируются в сетку или грид (grid) размерностью N1 * N2 * N3. Это можно изобразить следующим образом:
Рис. 1. Вычислительное устройство GPU.
При использовании GPU вы можете задействовать грид необходимого размера и сконфигурировать блоки под нужды вашей задачи.
Читайте также: