Глава 11: Направления исследований GPU по скаляризации и аффинному исполнению
Как описано в Главе 2, API вычислений на GPU, такие как CUDA и OpenCL, имеют программную модель, похожую на MIMD, которая позволяет программисту запускать большой массив скалярных потоков на GPU. Хотя каждый из этих скалярных потоков может следовать своему уникальному пути выполнения и может получать доступ к произвольным местам памяти, в обычном случае они все следуют небольшому набору путей выполнения и выполняют похожие операции.
Сходящийся управляющий поток среди потоков GPU эксплуатируется на большинстве, если не на всех, современных GPU с помощью модели исполнения SIMT, где скалярные потоки объединяются в варпы, которые работают на аппаратном обеспечении SIMD (см. Раздел 3.1.1). Эта глава обобщает серию исследований, которые дополнительно используют сходство этих скалярных потоков с помощью скаляризации и аффинного исполнения.
Ключевое понимание этого исследования заключается в наблюдении за структурой значений [Kim et al., 2013] среди потоков, выполняющих один и тот же вычислительный ядро. Два типа структуры значений, равномерная и аффинная, проиллюстрированы в вычислительном ядре в Примере 11.1.
Равномерная переменная
Переменная, которая имеет одно и то же постоянное значение для каждого потока в вычислительном ядре. В Алгоритме 11.1 переменная a, а также литералы THRESHOLD и Y_MAX_VALUE имеют равномерные значения во всех потоках вычислительного ядра. Равномерная переменная может храниться в одном скалярном регистре и использоваться всеми потоками в вычислительном ядре.
Аффинная переменная
Переменная со значениями, которые являются линейной функцией идентификатора потока для каждого потока в вычислительном ядре. В Алгоритме 11.1 адрес памяти переменной y[idx] может быть представлен как аффинное преобразование идентификатора потока threadIdx.x:
&(y[idx]) = &(y[0]) + sizeof(int) * threadIdx.x;Это аффинное представление может храниться как пара скалярных значений, базовое и шаговое, что гораздо более компактно, чем полностью развернутый вектор.
__global__ void vsadd( int y[], int a ) {
i
```Вот перевод на русский язык с сохранением оригинального кода:
int idx = threadIdx.x; y[idx] = y[idx] + a; if ( y[idx] > THRESHOLD ) y[idx] = Y_MAX_VALUE; }
Алгоритм 11.1: Пример скалярных и аффинных операций в вычислительном ядре (из [Kim et al., 2013]).
Существует множество исследовательских предложений о том, как обнаруживать и использовать равномерные или аффинные переменные в GPU. Остальная часть этой главы обобщает эти предложения в этих двух аспектах.
## Обнаружение равномерных или аффинных переменных
Существует два основных подхода к обнаружению существования равномерных или аффинных переменных в вычислительном ядре GPU: обнаружение, управляемое компилятором, и обнаружение с помощью аппаратного обеспечения.
### Обнаружение, управляемое компилятором
Один из способов обнаружить существование равномерных или аффинных переменных в вычислительном ядре GPU - сделать это с помощью специального анализа компилятора. Это возможно, потому что существующие модели программирования GPU, CUDA и OpenCL, уже предоставляют средства для программиста объявить переменную как константную на протяжении всего вычислительного ядра, а также предоставляют специальную переменную для идентификатора потока. Компилятор может выполнить анализ зависимости от управления, чтобы обнаружить переменные, которые зависят исключительно от констант и идентификаторов потоков, и пометить их как равномерные/аффинные. Операции, работающие исключительно с равномерными/аффинными переменными, затем являются кандидатами для скаляризации.
AMD GCN [AMD, 2012] полагается на компилятор для обнаружения равномерных переменных и скалярных операций, которые могут быть сохранены и обработаны специализированным скалярным процессором.
Asanovic et al. [2013] представляют комбинированный анализ сходящихся и вариантных, который позволяет компилятору определять операции в произвольном вычислительном ядре, которые подходят для скаляризации и/или аффинного преобразования. Инструкции внутри сходящихся областей вычислительного ядра могут быть преобразованы в скалярные/аффинные инструкции. При любом переходе от расходящихся к сходящимся областям вычислительного ядра компилятор вставляет инструкцию `syncwarp` для обработки зависимостей регистров, вызванных управлением потоком, между двумя областями. Asanovic et al. [2013] принялиВот перевод на русский язык:
Этот анализ предназначен для генерации скалярных операций для архитектуры Temporal-SIMT [Keckler et al., 2011, Krashinsky, 2011].
Разделенные аффинные вычисления (DAC) [Wang и Lin, 2017] основаны на аналогичном компиляторном анализе для извлечения скалярных и аффинных кандидатов, которые должны быть разделены в отдельный варп. Wang и Lin [2017] дополняют процесс анализом расходящихся аффинных функций с целью извлечения последовательностей инструкций
# Код (комментарии переведены на русский язык):
# Этот код выполняет анализ для извлечения скалярных и аффинных операций
# Результаты этого анализа используются для разделения вычислений на отдельные варпы