Вычисления с GPU-ускорением на Python

❯ Для чего нужен GPU?

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

Разница между ЦП и GPU. Источник.

❯ PyCuda

Инструмент CUDA от компании NVIDIA, написанный на Python, предлагает API для работы с драйверами и среду выполнения, в которой упрощается ускоренная обработка задач с поддержкой GPU и с применением различных инструментариев и библиотек. Язык Python широко известен как наиболее популярный язык для научных вычислений, инженерии, анализа данных, а также для глубокого обучения. Python – интерпретируемый язык, поэтому его часто критикуют за относительно невысокую производительность. В этой области следует отметить библиотеку PyCuda, обеспечивающую прямолинейный и Python-подобный способ обращаться к API CUDA от NVIDIA, предназначенному для параллельных вычислений. В этом отношении у PyCuda есть ряд ключевых преимуществ, в частности:

• Полнота: PyCuda предоставляет доступ к целому спектру возможностей, заключённому в API CUDA, что позволит использовать эти возможности на полную мощность.

• Автоматизированная проверка ошибок: любые ошибки CUDA, возникающие в процессе выполнения, автоматически преобразуются в исключения Python. Так становится удобнее и обрабатывать ошибки, и отлаживать код.

• Скорость: в основе PyCuda лежит уровень, реализованный на C++, благодаря чему вы можете пользоваться всеми вышеупомянутыми достоинствами, не жертвуя производительностью.

Более подробно можно изучить эту тему по полезной документации, выложенной здесь.

❯ Базовая терминология CUDA

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

Архитектура грида GPU.

• Ядро: функция, параллельно выполняемая множеством потоков в GPU. Это фундаментальная концепция программирования GPU. В ядре представлен код, выполняемый на устройстве GPU.

• Грид: коллекция блоков потоков. Представляет общую компоновку потоков, выполняющих функцию ядра на GPU.

• Поток: мельчайшая единица выполнения, в которой решается конкретная задача в рамках ядра. Представляет отдельный экземпляр кода, выполняемого на GPU.

❯ Выполняем первую программу на GPU

Перед началом работы убедитесь, что у вас NVIDIA GPU. Проверить, какой GPU стоит на вашей машине, можно по адресу https://www.pcmag.com/how-to/what-graphics-card-do-i-have

Если у вас на машине уже установлен Python, то можно выполнить следующую команду для установки PyCuda:

pip install pycuda

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

import pycuda.driver as cuda
import pycuda.autoinit
import numpy

from pycuda.compiler import SourceModule


mod = SourceModule("""
  __global__ void doublify(float *a)
  {
    int idx = threadIdx.x + threadIdx.y*4;
    a[idx] *= 2;
  }
  """)

# Инициализация массива
a = numpy.random.randn(4,4).astype(numpy.float32)

# Выделение памяти на устройстве (GPU)
a_gpu = cuda.mem_alloc(a.nbytes)

# Перенос данных с хоста (CPU) на устройство (GPU)
cuda.memcpy_htod(a_gpu, a)

# Получение ссылки на нашу функцию
func = mod.get_function("doublify")

# Вызов функции с размером блока (4,4)
func(a_gpu, block=(4,4,1))

# Выборка данных из GPU
a_doubled = numpy.empty_like(a)
cuda.memcpy_dtoh(a_doubled, a_gpu)

# Дальнейшая обработка нового массива
a_doubled = ...

На первом шаге в этом коде мы импортируем PyCuda и инициализируем CUDA при помощи pycuda.autoinit. Далее код компилируется при помощи SourceModule, затем проверяется на наличие каких-либо ошибок, а если ошибки не найдены – загружается на устройство. Наконец мы применяем класс pycuda.driver.Function, при помощи которого генерируем ссылку на нашу функцию и вызываем её, предоставив требуемые аргументы. В следующих разделах будет подробнее пошагово объяснено, что требуется сделать для выполнения работы ядра на GPU.

❯ Код хоста и устройства

В программе CUDA решение о том, где именно будет выполняться функция – на ЦП или на GPU — принимается в зависимости от её сигнатуры. Таким образом, это могут решать как разработчики, так и сама программа.

Код, предназначенный для выполнения на ЦП, называется кодом хоста, а код для выполнения на GPU — кодом устройства. Как было сказано выше, код этих двух типов различается по сигнатуре функций. Если функция относится к коду устройства, то в самом её начале присутствует ключевое слово __global__ или __device__ , тогда как в именах функций из кода хоста такие квалификаторы отсутствуют. Иногда в функциях с хоста также встречается и опциональное ключевое слово __host__.

❯ Как в коде устройства работает индексирование?

В CUDA индексирование массивов/векторов может выполняться в одном, двух или трёх измерениях, в зависимости от размерности массива. Рассмотрим, как устроено индексирование массивов в CUDA для случаев 1D, 2D и 3D:

1D:

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

• Поэтому индексу можно обращаться к соответствующему элементу в одномерном массиве. Возвращаясь к предыдущему примеру (случай 1D):

mod = SourceModule("""
  __global__ void doublify(float *a)
  {
    int idx = threadIdx.x;
    a[idx] *= 2;
  }
  """)

2D:

• В двумерном массиве потоки упорядочены в виде плоской сетки (грида), где как блокам потоков, так и самим потокам присваиваются ID одновременно по осям x и y.

• Чтобы вычислить индекс при работе с двумерным массивом, обычно требуется умножить индекс строки (threadIdx.y или blockDim.y * blockIdx.y + threadIdx.y) на ширину массива и приплюсовать индекс столбца (threadIdx.x или blockDim.x * blockIdx.x + threadIdx.x).

• По этому индексу можно обращаться к желаемому элементу в двумерном массиве. Возвращаясь к предыдущему примеру (случай 2D):

mod = SourceModule("""
  __global__ void doublify(float a[4][4])
  {
    int i = blockDim.x * blockIdx.x + threadIdx.x;
    int j = blockDim.y * blockIdx.y + threadIdx.y;
    a[i][j] *= 2;
  }
  """)

3D:

• Аналогично, в случае с трёхмерным массивом, потоки упорядочиваются в виде 3D-сетки, а блоки потоков и отдельные потоки получают ID по осям x, y и z.

• Чтобы вычислить индекс при работе с трёхмерным массивом, требуется умножить индекс глубины (threadIdx.z или blockDim.z * blockIdx.z + threadIdx.z) на произведение высоты и ширины массива.

• Затем требуется умножить индекс строки (threadIdx.y или blockDim.y * blockIdx.y + threadIdx.y) на ширину массива и приплюсовать индекс столбца (threadIdx.x или blockDim.x * blockIdx.x + threadIdx.x).

• По полученному в результате индексу можно обращаться к желаемому элементу в 3D-массиве. Возвращаясь к предыдущему примеру (случай 3D):

mod = SourceModule("""
  __global__ void doublify(float a[4][4][4])
  {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    int j = blockIdx.y * blockDim.y + threadIdx.y;
    int k = blockIdx.z * blockDim.z + threadIdx.z;

    a[i][j][k] *= 2;
  }
  """)

Опираясь на эти схемы индексирования, CUDA обеспечивает эффективный параллельный доступ к элементам массива через различные потоки в ядре GPU. Подробнее об индексировании потоков рассказано здесь.

❯ Распространённые ошибки при программировании с помощью PyCuda

Недопустимое обращение к памяти: в таком случае может всплыть сообщение об ошибке: “RuntimeError: CUDA error: an illegal memory access was encountere”. Данная ошибка — CUDA-специфичный вариант “Index out of range”. Решить эту проблему можно следующим образом:

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

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

Ошибка «Нехватка памяти»: вам может встретиться сообщение об ошибке следующего содержания: “RuntimeError: CUDA out of memory.” . Подобное иногда происходит при работе с крупными датасетами или сложными моделями, требующими от GPU много памяти. Чтобы решить эту проблему, можно предпринять следующее:

• Уменьшитьразмер данных: например, можно уменьшить размер обрабатываемых данных. В частности, попробуйте прореживание (downsampling) данных или выберите для анализа меньшее подмножество данных. Уменьшив размер данных, вы высвободите часть памяти GPU и избежите подобных ошибок.

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

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

❯ Дополнение: шаги (strides)

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

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

При работе с вычислениями на базе GPU разбираться в шагах становится принципиально важно. И вот почему:

При программировании для CUDA именно расстановка шагов играет решающую роль для эффективного доступа к данным и объединения памяти (memory coalescing). Под объединением понимается обращение к смежным участкам памяти так, как если бы они располагались непрерывным блоком. При параллельных вычислениях такой подход позволяет существенно увеличить пропускную способность памяти и общую производительность системы.

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

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

❯ Заключение

Разработка программ, которые предполагается выполнять на GPU порой даётся сложнее, чем обычное программирование на C++ или Python, поскольку здесь применяется иная модель программирования и требуется глубоко понимать то железо, на котором работает программа. Тем не менее, на многих высокоуровневых языках и, в частности, на Python, предлагаются библиотеки, при помощи которых можно программировать для GPU, не погружаясь при этом в низкоуровневые детали кода. Правда, рекомендую хорошо изучить основы параллельных вычислений и архитектуру GPU, а лишь потом пробовать свои силы в таком программировании.

Ссылки:

• Документация по Pycuda: https://documen.tician.de/pycuda/index.html

• Сравнение кода хоста и кода устройств: https://dev.to/zenulabidin/an-overview-of-cuda-part-2-host-and-device-code-69d

• Индексирование потоков в Cuda: https://anuradha-15.medium.com/cuda-thread-indexing-fb9910cba084

• Руководство по программированию в CUDA: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html

Известно, насколько сложно программировать для GPU в силу уникальной архитектуры этих процессоров, которая, к тому же, постоянно развивается. Подробнее эта тема рассмотрена здесь: https://dl.acm.org/doi/10.1145/3611643.3616365.