CUDA + Thrust: высокоуровневые алгоритмы вместо ручных ядер

Введение

Не всегда необходимо писать собственные CUDA-ядра вручную — для многих стандартных операций (сортировка, поиск, свёртка, преобразование коллекций) есть библиотека Thrust, которая предоставляет интерфейс, похожий на STL, но с реализациями, оптимизированными для GPU. Thrust позволяет писать параллельный код на уровне абстракции алгоритмов, а не низкоуровневых ядер, существенно сокращая объём кода и снижая риск ошибок. В статье рассмотрим основные возможности Thrust на практических примерах.

Концепция

Thrust оперирует контейнерами thrust::host_vector (память CPU) и thrust::device_vector (память GPU), автоматически управляющими выделением и освобождением памяти (аналог std::vector, но для GPU). Алгоритмы Thrust (thrust::sort, thrust::reduce, thrust::transform, thrust::copy_if и многие другие) принимают итераторы этих контейнеров и автоматически параллелизуют выполнение на GPU, либо могут выполняться на CPU при передаче host-итераторов — один и тот же вызов алгоритма работает на обоих «бэкендах» в зависимости от типа переданных итераторов.

Пример кода

#include <thrust/host_vector.h>
#include <thrust/device_vector.h>
#include <thrust/sort.h>
#include <thrust/reduce.h>
#include <thrust/transform.h>
#include <thrust/functional.h>
#include <thrust/copy.h>
#include <iostream>

struct SquareFunctor
{
    __host__ __device__ float operator()(float x) const { return x * x; }
};

int main()
{
    const int n = 1'000'000;

    thrust::host_vector<float> hostData(n);
    for (int i = 0; i < n; ++i) {
        hostData[i] = static_cast<float>(n - i); // обратный порядок для демонстрации сортировки
    }

    // Копирование данных на GPU происходит автоматически при присваивании
    thrust::device_vector<float> deviceData = hostData;

    // Параллельная сортировка на GPU
    thrust::sort(deviceData.begin(), deviceData.end());

    // Параллельное преобразование: возведение каждого элемента в квадрат
    thrust::device_vector<float> squared(n);
    thrust::transform(deviceData.begin(), deviceData.end(), squared.begin(), SquareFunctor());

    // Параллельная свёртка (сумма всех элементов)
    float sum = thrust::reduce(squared.begin(), squared.end(), 0.0f, thrust::plus<float>());

    std::cout << "Минимальный элемент после сортировки: " << deviceData[0] << std::endl;
    std::cout << "Сумма квадратов: " << sum << std::endl;

    return 0;
}

Пояснения к коду

thrust::device_vector<float> deviceData = hostData автоматически выполняет копирование данных с host на device одной строкой — без явных вызовов cudaMalloc/cudaMemcpy, как это требовалось бы при работе с «сырым» CUDA API. thrust::sort параллельно сортирует данные прямо на GPU, используя оптимизированную реализацию (как правило, основанную на radix sort для примитивных типов).

SquareFunctor — функтор с квалификаторами __host__ __device__, позволяющими ему выполняться как на CPU, так и на GPU; thrust::transform применяет его к каждому элементу параллельно, аналогично std::transform, но на GPU. thrust::reduce выполняет параллельную свёртку (в данном случае — суммирование) — типичная задача, ручная реализация которой на CUDA требует кода для редукции внутри блока и между блоками, тогда как Thrust скрывает все эти детали.

Подводные камни

Накладные расходы на копирование между host_vector и device_vector при частых небольших операциях могут превышать выигрыш от параллелизации — Thrust эффективен при больших объёмах данных и минимизации числа пересечений границы CPU/GPU.
Меньший контроль над деталями реализации алгоритмов. В отличие от ручного ядра, разработчик не может точно настроить tile-размер, использование shared memory или паттерн доступа к памяти внутри алгоритмов Thrust — для специфичных задач, где нужна максимальная производительность под конкретную архитектуру GPU, ручное ядро может оказаться быстрее.
Время компиляции. Thrust интенсивно использует шаблоны C++, и интенсивное использование сложных комбинаций алгоритмов и функторов может заметно увеличить время компиляции CUDA-кода по сравнению с написанием ядер вручную.
Совместимость версий Thrust с версией CUDA Toolkit. Thrust поставляется как часть CUDA Toolkit, и использование отдельно установленной версии Thrust, не совпадающей с версией Toolkit, иногда приводит к трудно диагностируемым ошибкам компиляции или линковки.