Глава 3: Параллельные модели программирования в дизайне GPU

Графические процессорные устройства (GPU) эволюционировали от фиксированных графических ускорителей до высокопараллельных, программируемых вычислительных двигателей, способных ускорять широкий спектр приложений. Для того, чтобы программисты могли эффективно использовать массивный параллелизм в GPU, были разработаны несколько параллельных моделей программирования и API, таких как NVIDIA CUDA, OpenCL и DirectCompute. Эти модели программирования предоставляют абстракции, которые позволяют программистам выражать параллелизм в своих приложениях, скрывая при этом низкоуровневые детали аппаратного обеспечения GPU.

В этой главе мы исследуем ключевые концепции и принципы параллельных моделей программирования для GPU, сосредоточившись на модели выполнения SIMT (Single Instruction, Multiple Thread), модели программирования CUDA и API, а также на фреймворке OpenCL. Мы также обсудим техники отображения алгоритмов на архитектуры GPU для достижения высокой производительности и эффективности.

Модель выполнения SIMT (Single Instruction, Multiple Thread)

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

Ядра и иерархия потоков

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

• Сетка представляет всю пространство задачи и состоит из одного или нескольких блоков.
• Блок - это группа потоков, которые могут сотрудничать и синхронизироваться друг с другом через разделяемую память и барьеры. Потоки внутри блока выполняются на одном ядре GPU (называемом мультипроцессором потоков).Вот перевод на русский язык с сохранением оригинального кода:

Иерархия потоков в модели SIMT (Single Instruction, Multiple Thread)

• Каждый поток имеет уникальный идентификатор (ID) внутри своего блока и сетки, который может использоваться для вычисления адресов памяти и принятия решений о потоке управления.

Эта иерархическая организация позволяет программистам выражать как параллелизм данных (когда одна и та же операция применяется к нескольким элементам данных), так и параллелизм задач (когда различные задачи выполняются параллельно).

На рисунке 3.1 показана иерархия потоков в модели SIMT.

            Сетка
       ________________
      /   /   /   /   /
     /   /   /   /   / 
    /   /   /   /   /
   /   /   /   /   /
  /__/__/__/__/__/
    |   |   |   |
    |   | Блок |
    |   |   |   |
  Поток Поток ...

Рисунок 3.1: Иерархия потоков в модели SIMT.

Выполнение SIMT

В модели SIMT каждый поток выполняет одну и ту же инструкцию, но работает с различными данными. Однако, в отличие от SIMD (Single Instruction, Multiple Data), где все вычислительные элементы выполняются синхронно, SIMT позволяет потокам иметь независимые пути выполнения и расходиться на инструкциях ветвления.

Когда варп (группа из 32 потоков в GPU NVIDIA или 64 потоков в GPU AMD) сталкивается с инструкцией ветвления, аппаратное обеспечение GPU оценивает условие ветвления для каждого потока в варпе. Если все потоки выбирают один и тот же путь (сходятся), варп продолжает выполнение в обычном режиме. Однако, если некоторые потоки выбирают разные пути (расходятся), варп разделяется на два или более подварпа, каждый из которых следует по разному пути. Аппаратное обеспечение GPU последовательно выполняет расходящиеся пути, маскируя неактивные потоки в каждом подварпе. Когда все пути завершены, подварпы снова сходятся и продолжают выполнение синхронно.

На рисунке 3.2 показано выполнение SIMT с расходящимся потоком управления.

         Варп
   ________________
  /   /   /   /   /
 /   /   /   /   / 
/   /   /   /   /
   |   |   |
   | Ветвление |
   |   |   |
  / \ / \ / \
 /   X   \   \
/   / \   \   \
   /   \   \
  /     \   \
 /       \   \
/         \   \
           \
```Реконвергенция

Рисунок 3.2: Выполнение SIMT с расходящимся управляющим потоком.

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

Иерархия памяти

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

• Глобальная память: Самое большое, но самое медленное пространство памяти, доступное для всех потоков в ядре. Глобальная память, как правило, реализуется с использованием высокоскоростной памяти GDDR или HBM.
• Разделяемая память: Быстрое, встроенное в чип пространство памяти, общее для всех потоков в блоке. Разделяемая память используется для межпоточной связи и обмена данными внутри блока.
• Постоянная память: Только для чтения пространство памяти, используемое для трансляции данных только для чтения во все потоки.
• Текстурная память: Только для чтения пространство памяти, оптимизированное для пространственной локальности и доступа через текстурные кэши. Текстурная память чаще используется в графических рабочих нагрузках.
• Локальная память: Частное пространство памяти для каждого потока, используемое для переполнения регистров и больших структур данных. Локальная память, как правило, отображается на глобальную память.

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

На рисунке 3.3 показана иерархия памяти GPU.

|            |
|   Общая    |
|   Память   |
 ____________
      |
 ____________ 
|            |
|   Локальная |
|   Память   |
 ____________

Рисунок 3.3: Иерархия памяти GPU.

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

CUDA (Compute Unified Device Architecture) - это платформа параллельных вычислений и модель программирования, разработанная NVIDIA для общих вычислений на GPU. CUDA предоставляет набор расширений к стандартным языкам программирования, таким как C, C++ и Fortran, которые позволяют программистам выражать параллелизм и использовать вычислительную мощность GPU NVIDIA.

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

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

CUDA вводит несколько ключевых абстракций для облегчения параллельного программирования:

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

CUDA также предоставляет встроенные переменные (например, threadIdx, blockIdx, blockDim, gridDim), которые позволяют потокам идентифицировать себя и вычислять адреса памяти на основе их положения в иерархии потоков.

Рисунок 3.4 иллюстрирует модель программирования CUDA.

Сетка

/ / / / / / / / / / / / / / / / / / / / /////__/ | | | | | | Блок | | | | | Поток Поток ...

Рисунок 3.4: Модель программирования CUDA.

### Иерархия памяти CUDAВот перевод на русский язык с сохранением оригинального кода:

CUDA предоставляет программисту иерархию памяти GPU, позволяя явно управлять размещением и перемещением данных. Основные пространства памяти в CUDA:

- Глобальная память: Доступна всем потокам в ядре и сохраняется между запусками ядра. Глобальная память имеет самую высокую задержку и, как правило, используется для больших структур данных.
- Разделяемая память: Быстрая, встроенная в чип память, общая для всех потоков в блоке. Разделяемая память используется для межпоточной связи и обмена данными внутри блока.
- Константная память: Пространство только для чтения, используемое для трансляции данных только для чтения во все потоки. Константная память кэшируется и обеспечивает низкую задержку доступа.
- Текстурная память: Пространство только для чтения, оптимизированное для локальности в пространстве и доступа через текстурные кэши. Текстурная память чаще используется в графических рабочих нагрузках.
- Локальная память: Частное пространство памяти для каждого потока, используемое для переполнения регистров и больших структур данных. Локальная память, как правило, отображается на глобальную память.

Программисты могут выделять и передавать данные между памятью хоста (CPU) и устройства (GPU) с помощью API-интерфейсов среды выполнения CUDA, таких как cudaMalloc, cudaMemcpy и cudaFree.

На рисунке 3.5 показана иерархия памяти CUDA.

| | | Global | | Memory |

|

| | | Constant | | Memory |

|

| | | Texture | | Memory |

| |

| | | Shared | | Memory |

|

| | | Local | | Memory |

Рисунок 3.5: Иерархия памяти CUDA.

### Синхронизация и координация в CUDA

CUDA предоставляет примитивы синхронизации и координации для обеспечения сотрудничества и связи между потоками:

- Барьерная синхронизация: __syncthreads()Вот перевод на русский язык:

Функция s() действует как барьер, который гарантирует, что все потоки в блоке достигли одной и той же точки, прежде чем продолжить.
- Атомарные операции: CUDA поддерживает атомарные операции (например, atomicAdd, atomicExch), которые позволяют потокам выполнять операции чтения-модификации-записи в разделяемой или глобальной памяти без вмешательства других потоков.
- Примитивы на уровне варпа: CUDA предоставляет встроенные функции на уровне варпа (например, __shfl, __ballot), которые позволяют эффективно осуществлять обмен данными и синхронизацию внутри варпа.

Правильное использование примитивов синхронизации и координации является важным для написания правильных и эффективных параллельных программ на CUDA.

Пример 3.1 показывает простое ядро CUDA, которое выполняет сложение векторов.

```c
__global__ void vectorAdd(int *a, int *b, int *c, int n) {
    // Вычисление индекса текущего потока
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    // Проверка, что индекс находится в пределах размера векторов
    if (i < n) {
        // Выполнение сложения элементов векторов
        c[i] = a[i] + b[i];
    }
}

int main() {
    int *a, *b, *c;
    int n = 1024;
    
    // Выделение памяти на хосте
    a = (int*)malloc(n * sizeof(int));
    b = (int*)malloc(n * sizeof(int));
    c = (int*)malloc(n * sizeof(int));
    
    // Инициализация входных векторов
    for (int i = 0; i < n; i++) {
        a[i] = i;
        b[i] = i * 2;
    }
    
    // Выделение памяти на устройстве
    int *d_a, *d_b, *d_c;
    cudaMalloc(&d_a, n * sizeof(int));
    cudaMalloc(&d_b, n * sizeof(int));
    cudaMalloc(&d_c, n * sizeof(int));
    
    // Копирование входных векторов с хоста на устройство
    cudaMemcpy(d_a, a, n * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, b, n * sizeof(int), cudaMemcpyHostToDevice);
    
    // Запуск ядра
    int blockSize = 256;
    int numBlocks = (n + blockSize - 1) / blockSize;
    vectorAdd<<<numBlocks,blockSize>>>(d_a, d_b, d_c, n);
    
    // Копирование результирующего вектора с устройства на хост
    cudaMemcpy(c, d_c, n * sizeof(int), cudaMemcpyDeviceToHost);
    
    // Освобождение памяти на устройстве
    cudaFree(d_a);
    cudaFree(d_b); 
    cudaFree(d_c);
    
    // Освобождение памяти на хосте
    free(a); 
    free(b);
    free(c);
    
    return 0;
}
```Вот перевод на русский язык с сохранением оригинального кода:

```c
// Этот CUDA-код запускает ядро `vectorAdd` с `numBlocks` блоками и `blockSize` потоками на блок. Ядро выполняет поэлементное сложение входных векторов `a` и `b` и сохраняет результат в векторе `c`. Синтаксис `<<<...>>>` используется для указания размерности сетки и блоков при запуске ядра.

### CUDA-потоки и события

CUDA-потоки и события предоставляют механизм для параллельного выполнения и синхронизации ядер и операций с памятью:

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

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

Пример 3.2 демонстрирует использование CUDA-потоков для совмещения выполнения ядер и передачи данных.

```c
// Создаем два потока
cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);

// Асинхронно копируем входные данные на устройство
cudaMemcpyAsync(d_a, a, size, cudaMemcpyHostToDevice, stream1);
cudaMemcpyAsync(d_b, b, size, cudaMemcpyHostToDevice, stream2);

// Запускаем ядра в разных потоках
kernelA<<<blocks, threads, 0, stream1>>>(d_a);
kernelB<<<blocks, threads, 0, stream2>>>(d_b);

// Асинхронно копируем результаты обратно на хост
cudaMemcpyAsync(a, d_a, size, cudaMemcpyDeviceToHost, stream1);
cudaMemcpyAsync(b, d_b, size, cudaMemcpyDeviceToHost, stream2);

// Синхронизируем потоки
cudaStreamSynchronize(stream1);
cudaStreamSynchronize(stream2);

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

Фреймворк OpenCL

OpenCL (Open Computing Language) - это открытый, свободный от роялти стандарт для параллельного программирования на гетерогенных платформах, включая CPU, GPU, FPGA и другие ускорители. OpenCL предоставляет единую модель программирования и набор API, которые позволяют разработчикам писать портируемый и эффективный параллельный код.

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

Модель программирования OpenCL похожа на CUDA, с несколькими ключевыми различиями в терминологии и абстракциях:

• Ядро (Kernel): Функция, выполняемая параллельно большим количеством рабочих элементов (потоков) на устройстве OpenCL.
• Рабочий элемент (Work-item): Базовая единица выполнения в OpenCL, аналогичная потоку в CUDA.
• Рабочая группа (Work-group): Коллекция рабочих элементов, которые могут синхронизироваться и обмениваться данными через локальную память. Рабочие группы аналогичны блокам потоков в CUDA.
• NDRange: Определяет пространство индексов и организацию рабочих элементов для выполнения ядра. Оно может быть одно-, двух- или трехмерным.

OpenCL также определяет иерархическую модель памяти, аналогичную CUDA:

• Глобальная память (Global memory): Доступна всем рабочим элементам во всех рабочих группах, аналогична глобальной памяти в CUDA.
• Локальная память (Local memory): Общая для всех рабочих элементов в рабочей группе, аналогична разделяемой памяти в CUDA.
• Частная память (Private memory): Частная для одного рабочего элемента, аналогична регистрам в CUDA.
• Константная память (Constant memory): Только для чтения, доступна всем рабочим элементам.

Ядра OpenCL компилируются во время выполнения программы OpenCL-runtime. Хост-программа может запрашивать доступные устройства OpenCL, выбирать подходящее устройство, создавать контекст и компилировать ядро для этого конкретного устройства. Это позволяет приложениям OpenCL быть высоко портируемыми на различные аппаратные платформы.

Пример 3.3 показывает ядро OpenCL, которое выполняет сложение векторов, аналогично примеру CUDA в Примере 3.1.

__kernel void vectorAdd(__global const int *a, __global const int *b, __global int *c) {
    // Получаем индекс текущего рабочего элемента
    int i = get_global_id(0);
    
    // Выполняем сложение векторов
    c[i] = a[i] + b[i];
}
```Вот перевод на русский язык:
 
```c
__kernel void bal(const int *a, const int *b, __global int *c, int n) {
    int i = get_global_id(0);
    if (i < n) {
        c[i] = a[i] + b[i];
    }
}

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

Сопоставление алгоритмов с архитектурой GPU

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

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

• Минимизация расхождения ветвлений: Расходящиеся потоки управления внутри варпа/волны могут привести к последовательному выполнению и снижению эффективности SIMD. Алгоритмы должны быть структурированы так, чтобы минимизировать расхождение ветвлений, где это возможно.

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

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

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

Несколько шаблонов параллельного проектирования алгоритмов часто используются при разработке ядер GPU:

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

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

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

• Шаблон: Каждый поток вычисляет значение на основе соседних элементов данных. Вычисления по шаблону распространены в научных моделированиях и приложениях обработки изображений.

• Сбор/Рассеивание: Потоки читают из (сбор) или записывают в (рассеивание) произвольные места в глобальной памяти. Для эффективности требуется тщательная компоновка данных и схемы доступа.

Заключение

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

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

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