CUDA

Open book.svg Авторство
С. А. Гончаров
Согласовано: 23.01.2019
CUDA
CUDA.png
Разработчики: NVIDIA Corporation
Выпущена: 23 June 2007 года; 12 years ago (2007-06-23)
Постоянный выпуск: 10.0 / 18 September 2018 года; 18 months ago (2018-09-18)
Состояние разработки: активное
Операционная система: Windows XP и выше, OS X, Linux
Локализация: Английский язык
Тип ПО: GPGPU
Веб-сайт developer.nvidia.com/cuda-zone

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

Обзор платформы

Направление вычислений эволюционирует от «централизованной обработки данных» на центральном процессоре до «совместной обработки» на CPU и GPU. Для реализации новой вычислительной парадигмы компания NVIDIA изобрела архитектуру параллельных вычислений CUDA, на данный момент представленную в графических процессорах GeForce, ION, Quadro и Tesla и обеспечивающую необходимую базу разработчикам ПО. Платформа параллельных вычислений CUDA® обеспечивает набор расширений для языков C и C++, позволяющих выражать как параллелизм данных, так и параллелизм задач на уровне мелких и крупных структурных единиц (рисунок 1). Программист может выбрать средства разработки: языки высокого уровня, такие как C, C++, Fortran или же открытые стандарты, такие как директивы OpenACC. Платформа параллельных вычислений CUDA используется на сегодняшний день в тысячах GPU-ускоренных приложений и тысячах опубликованных научных статьях.

Рисунок 1 – Сравнение кодов стандартного и параллельного языка С

Этапы запуска программы на GPU

Рассмотрим, как происходит запуск программы на графическом процессоре:

  1. Хост[1] выделяет необходимое количество памяти на устройстве.
  2. Хост копирует данные из своей памяти в память устройства.
  3. Хост запускает ядро на устройстве.
  4. Устройство исполняет это ядро.
  5. Хост копирует результаты из памяти устройства в свою память.

На рисунке 2 изображены все перечисленные шаги запуска программы, кроме первого.

Рисунок 2 – Шаги запуска программы

Как видно из рисунка 3, центральный процессор взаимодействует с графическим через CUDA Runtime API, CUDA Driver API и CUDA Libraries. Runtime и Driver API отличаются уровнем абстракции. Первый вариант более высокого уровня в плане программирования, более абстрактный, а второй – напротив, более низкого (уровень драйвера).

Рисунок 3 – Взаимодействие CPU и GPU

В целом Runtime API[2] является абстрактной оберткой Driver API. Во время программирования вы можете использовать любой из представленных вариантов[Источник 2].

Аппаратная часть

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

Рисунок 4 – Архитектура CPU и GPU

Как видно из рисунка 4 – в GPU есть много простых арифметически-логических устройств (АЛП), которые объединены в несколько групп и обладают общей памятью. Это помогает повысить продуктивность в вычислительных заданиях, но немного усложняет программирование.

GPU ориентирован на выполнение программ с большим объемом данных и расчетов и представляет собой массив потоковых процессоров (Streaming Processor Array [3]), что состоит из кластеров текстурных процессоров (Texture Processor Clusters, TPC). TPC в свою очередь состоит из набора мультипроцессоров (SM – Streaming Multi-processor, рисунок 5), в каждом из которых несколько потоковых процессоров (SP – Streaming Processors) или ядер (в современных процессорах количество ядер превышает 1024).

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

Рисунок 5 – Мультипроцессоры, SM

В результате GPU фактически стал устройством, которое реализует потоковую вычислительную модель (stream computing model): есть потоки входящих и исходящих данных, что состоят из одинаковых элементов, которые могут быть обработаны независимо друг от друга (рисунок 6).

Рисунок 6 – Потоковая вычислительная модель

Потоки, блоки и сетки

CUDA использует большое количество отдельных потоков для расчетов. Все они группируются в иерархию – grid / block / thread (рисунок 7).

Рисунок 7 – Структура блоков

Верхний уровень – grid – отвечает ядру и объединяет все потоки, которые выполняет данное ядро. Grid – одномерный или двумерный массив блоков (block). Каждый блок (block) представляет собой полностью независимый набор скоординированных между собой потоков. Потоки из разных блоков не могут взаимодействовать.

Отличие от SIMD-архитектуры: существует понятие warp – группа из 32 потоков (в зависимости от архитектуры GPU, но почти всегда 32). Только потоки в рамках одной группы (warp) могут физически выполняться одновременно. Потоки разных варпов могут находиться на разных стадиях выполнения программы. Такой метод обработки данных обозначается термином SIMT (Single Instruction – Multiple Theads). Управление работой варпов выполняется на аппаратном уровне.

CUDA и язык C

Сама технология CUDA (компилятор nvcc.exe) вводит ряд дополнительных расширений для языка C, которые необходимы для написания кода для GPU:

  • Спецификаторы функций, которые показывают, как и откуда буду выполняться функции.
  • Спецификаторы переменных, которые служат для указания типа используемой памяти GPU.
  • Спецификаторы запуска ядра GPU.
  • Встроенные переменные для идентификации нитей, блоков и др. параметров при исполнении кода в ядре GPU .
  • Дополнительные типы переменных.

Как было сказано, спецификаторы функций определяют, как и откуда буду вызываться функции. Всего в CUDA 3 таких спецификатора:

  • __host__ — выполнятся на CPU, вызывается с CPU (в принципе его можно и не указывать).
  • __global__ — выполняется на GPU, вызывается с CPU.
  • __device__ — выполняется на GPU, вызывается с GPU.

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

myKernelFunc<<<gridSize, blockSize, sharedMemSize, cudaStream>>>(float* param1,float* param2), где * gridSize – размерность сетки блоков (dim3), выделенную для расчетов,

  • blockSize – размер блока (dim3), выделенного для расчетов,
  • sharedMemSize – размер дополнительной памяти, выделяемой при запуске ядра,
  • cudaStream – переменная cudaStream_t, задающая поток, в котором будет произведен вызов.

myKernelFunc – функция ядра (спецификатор __global__). Некоторые переменные при вызове ядра можно опускать, например sharedMemSize и cudaStream[Источник 3].

Так же стоит упомянуть о встроенных переменных:

  • gridDim – размерность грида, имеет тип dim3. Позволяет узнать размер гридa, выделенного при текущем вызове ядра.
  • blockDim – размерность блока, так же имеет тип dim3. Позволяет узнать размер блока, выделенного при текущем вызове ядра.
  • blockIdx – индекс текущего блока в вычислении на GPU, имеет тип uint3.
  • threadIdx – индекс текущей нити в вычислении на GPU, имеет тип uint3.
  • warpSize – размер warp’а, имеет тип int (сам еще не пробовал использовать).

gridDim и blockDim - переменные, которые мы передаем при запуске ядра GPU.

CUDA host API

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 – функции обработки ошибок.

Пример

Этот пример кода на C ++ загружает текстуру из изображения в массив на GPU:

texture<float, 2, cudaReadModeElementType> tex;

void foo()
{
  cudaArray* cu_array;

  // Allocate array
  cudaChannelFormatDesc description = cudaCreateChannelDesc<float>();
  cudaMallocArray(&cu_array, &description, width, height);

  // Copy image data to array
  cudaMemcpyToArray(cu_array, image, width*height*sizeof(float), cudaMemcpyHostToDevice);

  // Set texture parameters (default)
  tex.addressMode[0] = cudaAddressModeClamp;
  tex.addressMode[1] = cudaAddressModeClamp;
  tex.filterMode = cudaFilterModePoint;
  tex.normalized = false; // do not normalize coordinates

  // Bind the array to the texture
  cudaBindTextureToArray(tex, cu_array);

  // Run kernel
  dim3 blockDim(16, 16, 1);
  dim3 gridDim((width + blockDim.x - 1)/ blockDim.x, (height + blockDim.y - 1) / blockDim.y, 1);
  kernel<<< gridDim, blockDim, 0 >>>(d_data, height, width);

  // Unbind the array from the texture
  cudaUnbindTexture(tex);
} //end foo()

__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;
   if (x < width && y < height) {
      float c = tex2D(tex, x, y);
      odata[y*width+x] = c;
   }
}

Ниже приведен пример кода, написанный на Python, который вычисляет произведение двух массивов на GPU:

import pycuda.compiler as comp
import pycuda.driver as drv
import numpy
import pycuda.autoinit

mod = comp.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

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

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

Рисунок 8 – Модель программирования GPGPU

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

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

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

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

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

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

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

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

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

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

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


Примечания

Источники

  1. Что такое CUDA? // Nvidia. [2018]. Дата обновления: 01.12.2018. URL: https://www.nvidia.ru/object/cuda-parallel-computing-ru.html (дата обращения 01.12.2018)
  2. Знакомство с программно-аппаратной архитектурой CUDA // Proglib. [2018]. Дата обновления: 01.12.2018. URL: https://proglib.io/p/cuda/ (дата обращения 01.12.2018)
  3. CUDA: Как работает GPU // Habr. [2018]. Дата обновления: 01.12.2018. URL: https://habr.com/post/54707/ (дата обращения 01.12.2018)
  4. Nvidia CUDA? Неграфические вычисления на графических процессорах // Ixbt. [2018]. Дата обновления: 01.12.2018. URL: https://www.ixbt.com/video3/cuda-1.shtml (дата обращения 01.12.2018)

Ссылки

  • Nvidia [Электронный ресурс]: Nvidia CUDA Zone/ Дата обращения: 01.12.2018. Режим доступа: https://developer.nvidia.com/cuda-zone
  • Nvworld [Электронный ресурс]: Параллельные вычислительные процессоры NVIDIA: настоящее и будущее/ Дата обращения: 01.12.2018. Режим доступа: https://nvworld.ru/articles/cuda-parallel/