За годы работы я подробно изучил, как центральные процессоры (CPU) выполняют код и как они устроены внутри. Дело в том, что я участвовал в разработке ядра Linux и ScyllaDB, а этот код очень близок к металлу. Я даже немного баловался с Verilog, безрезультатно попытавшись собрать моё собственное ядро RISC-V.

Графические процессоры (GPU) в отличие от обычных в основном оставались для меня чёрным ящиком, несмотря на то, что поработать с ними всё-таки довелось. Помню, что экспериментировал с NVIDIA RIVA 128 или чем-то подобным, проверяя, как там работает DirectX. Тогда такие процессоры ещё не выделялись на фоне ускорителей 3D-графики. Я также пытался идти в ногу со временем и немного упражнялся в программировании элементарных шейдеров на современных GPU. Но я никогда глубоко не вдавался в работу с GPU, и мои взгляды можно назвать CPU-центричными.

Однако, поскольку сегодня наблюдается всплеск рабочих нагрузок, связанных с ИИ, и, в частности, приходится работать с большими языковыми моделями (БЯМ), графические процессоры становятся незаменимыми для современных вычислений. К задачам, решаемым с применением ИИ, относятся масштабные прикладные тензорные операции, в том числе — сложение и перемножение матриц. А это уже работа для GPU. Но как современный GPU выполняет их, и насколько при этом возрастает эффективность по сравнению с выполнением таких же рабочих нагрузок на CPU?

❯ Определяем ядра для работы на GPU

Отправляясь в мир GPU, для начала рассмотрим CUDA. Это разработанный компанией NVIDIA язык программирования, который является расширением C и предназначен для эффективного распараллеливания данных на GPU. Притом, что сегодня для программирования ИИ в основном используются высокоуровневые библиотеки, среди которых — PyTorch и TensorFlow, всё равно не повредит знать CUDA, чтобы лучше понимать, как именно функционируют GPU.

На CUDA мы пишем код для CPU (хост-код) или для GPU (код ядра). Хост-код обычно пишется на чистом C, но CUDA расширяет этот язык в двух отношениях: во-первых, позволяет определять функции для GPU (так называемые «ядра», kernels), во-вторых — запускать эти ядра на GPU. Интересно, почему функции для GPU называются «ядрами». Дело в том, что код, написанный на CUDA, выполняется параллельно — в отличие от обычных функций, которые выполняются последовательно.

Рассмотрим для примера сложение векторов. В данном случае мы поэлементно суммируем векторы A и B, собирая результирующий вектор C. Тогда у нас получится ядро (функция для выполнения на GPU), которая будет выглядеть примерно так:

global
void vecAddKernel(float A, float B, float C, int n) {
    int i = blockDim.x  blockIdx.x + threadIdx.x;
    if (i < n) C[i] = A[i] + B[i];
}

Ключевое слово global в CUDA означает, что данная функция C является ядром и должна выполняться на GPU. Входные векторы A и B и результирующий вектор C передаются ей как параметры. Все векторы имеют размер по n элементов.

Вот та часть ядра, где, в сущности, происходит сложение векторов:

if (i < n) C[i] = A[i] + B[i];

Но что представляют собой переменные blockDim, blockIdx и threadIdx, содержащиеся в ядре и предназначенные здесь для вычисления индекса элемента i и проверки границ?

❯ Пуск ядер на GPU из CPU

Чтобы ответить на этот вопрос, давайте сначала разберёмся с хост-кодом, используемым для запуска ядра. Он будет иметь примерно следующий вид:

void vecAdd(float A, float B, float C, int n) {
    float A_d, B_d, C_d;
    int size = n * sizeof(float);

    cudaMalloc(&A_d, size);
    cudaMalloc(&B_d, size);
    cudaMalloc(&C_d, size);

    cudaMemcpy(A_d, A, size, cudaMemcpyHostToDevice);
    cudaMemcpy(B_d, B, size, cudaMemcpyHostToDevice);

    vecAddKernel<<<ceil(n/256.0), 256>>>(A_d, B_d, C_d, n);

    cudaMemcpy(C, C_d, size, cudaMemcpyDeviceToHost);

    cudaFree(A_d);
    cudaFree(B_d);
    cudaFree(C_d);
}

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

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

Вернёмся к определению ядра. Здесь мы наблюдаем, как используются уникальные переменные blockDim, blockIdx и threadIdx. От них, в частности, зависит, какую часть грида будет выполнять ядро.

global
void vecAddKernel(float A, float B, float C, int n) {
    int i = blockDim.x  blockIdx.x + threadIdx.x;
    if (i < n) C[i] = A[i] + B[i];
}

В переменной blockDim указывается размерность блока, поскольку в CUDA разрешено иметь трёхмерные блоки. В нашем примере с вектором используется лишь одномерный блок, так что здесь всё просто. В свою очередь, переменная blockIdx сообщает, какой именно блок выполняет ядро.

При запуске ядра была следующая картина, означающая, что у нас n/256 блоков с округлением до ближайшего целого числа, а также 256 потоков на блок.

void vecAdd(float A, float B, float *C, int n) {
    // ...
    vecAddKernel<<<ceil(n/256.0), 256>>>(A_d, B_d, C_d, n);
    // ...
}

Например, если размер вектора n равен 1000, у нас четыре блока и 256 потоков на блок, всего 1024 аппаратных потока. Именно здесь нам пригодится проверка границ i < n, указанная в определении ядра. Первые 1000 потоков выполняют сложение, а оставшиеся 24 потока, в сущности, ничего не делают.

❯ Итак, чем же выполнение на GPU отличается от выполнения на CPU?

Если бы мы попытались реализовать такое же сложение векторов на CPU, то начали бы примерно так: написали бы явный цикл, который перебирает все n элементов. При этом на каждой итерации цикла выполняется сложение:

void vecAdd(float A, float B, float *C, int n) {
    for (int i = 0; i < n; i++) {
        C[i] = A[i] + B[i];
    }
}

Затем можно попробовать воспользоваться инструкциями SIMD, чтобы ускорить функцию. Для этого можно, например, вручную разматывать цикл и использовать инструкции SIMD через внутренние функции компилятора или ассемблер. Другой способ — включить в компиляторе оптимизационные флаги, активирующие автовекторизацию; так мы переложим эту задачу на компилятор. В любом случае, цикл в функции всегда будет последовательно обрабатывать все элементы, объединённые в пакет.

Но, работая с GPU, мы описываем параллелизм в рамках запуска ядра и определяем только ту последовательную часть работы по сложению векторов в CUDA, которая выполняется на аппаратных потоках GPU. Это значит, что на GPU нет цикла для сложения векторов. Напротив, каждая операция в рамках поэлементного сложения векторов выполняется в собственном потоке. Если же вы привыкли программировать для CPU, вам интересно, как такая работа вообще может идти быстро.

Дело в том, что цена аппаратного потока на GPU по сравнению со стоимостью потоков операционной системы, используемых при программировании на CPU. Внутри GPU используется большой регистровый файл, он гораздо больше, чем в CPU. Таким образом, переключение между потоками почти ничего не стоит, поскольку не требуется сохранять и восстанавливать регистры. Более того, планирование потоков происходит на аппаратном уровне — таким образом, чтобы приступить к планированию, не требуется переключаться в привилегированный код ядра. В результате мы предполагаем, что на создание потока GPU и планирование его работы уйдёт несколько циклов GPU, тогда как на аналогичные операции с потоком POSIX может потребоваться всего несколько микросекунд.

Более того, GPU дополнительно оптимизированы для распараллеливания работы благодаря использованию варпов. Варп – это набор из 32 аппаратных потоков, и все эти потоки одновременно работают над одной и той же инструкцией. Только представьте, как 32 АЛУ параллельно выполняют поэлементное сложение, и в это время совместно используют одну и ту же управляющую логику. Если вы привыкли программировать для CPU, то можете сравнить эту совокупность с одной огромной SIMD-инструкцией. (Потоки управления в варпах могут и расходиться, но планировщик потоков в GPU вполне справляется и с этим, в любой момент времени выполняя лишь часть варпа в зависимости от того, как именно складываются потоки управления.)

Но ключевое отличие от SIMD заключается в следующем: если на GPU произойдёт кэш-промах, то процессор может быстро переключиться на другие варпы и параллельно выполнять их, а вся загрузка данных в память тем временем происходит на лету. Иными словами, GPU оптимизированы для повышения пропускной способности, где огромное количество потоков бросается на решение такой задачи, которая легко распараллеливается. В свою очередь, CPU оптимизированы для выполнения отдельных инструкций с минимальными задержками.

tl;dr;

Если вы привыкли программировать для CPU, то можете представить себе модель выполнения на GPU так: мы не перебираем набор данных в цикле, а разбиваем задачу на очень и очень много потоков, каждый из которых обрабатывает подмножество данных. Разумеется, существует ряд ключевых отличий между GPU и CPU. Прежде всего, на GPU приходится определять не функцию, а ядро. Это нужно для того, чтобы после запуска ядра явно задать параллелизм. В то же время, переключение контекста между потоками GPU обходится в разы дешевле, чем такое же переключение между потоками операционной системы — вот почему любую задачу можно «завалить числом». Наконец, GPU ещё лучше оптимизированы для работы с потоками, поскольку потоки группируются в варпы. Они выполняются на параллельных АЛУ и совместно используют общую управляющую логику. Если произойдёт кэш-промах, то GPU может быстро переключиться на другой варп, не отрываясь от загрузки данных в память, происходящей в фоновом режиме.

Новости, обзоры продуктов и конкурсы от команды Timeweb.Cloud — в нашем Telegram-канале ↩

Перед оплатой в разделе «Бонусы и промокоды» в панели управления активируйте промокод и получите кэшбэк на баланс.