Интеграция CUDA-буферов с QImage/QPixmap для GPU-обработки изображений в Qt

Введение

При построении Qt-приложений с GPU-ускоренной обработкой изображений (фильтры, свёртки, цветокоррекция) возникает задача эффективной передачи данных между CUDA-буферами на GPU и Qt-классами для отображения, такими как QImage. В статье рассмотрим типичный пайплайн: загрузка изображения через Qt, передача пикселей на GPU для обработки CUDA-ядром, и возврат результата обратно в QImage для отображения в интерфейсе.

Концепция

QImage хранит пиксельные данные в линейном буфере, доступном через bits()/constBits(), с форматом, заданным QImage::Format (например, Format_RGBA8888). Этот буфер можно напрямую скопировать в память GPU через cudaMemcpy, обработать CUDA-ядром и скопировать результат обратно. Для интерактивных сценариев (живой предпросмотр, видео) повторные выделения и освобождения памяти на каждый кадр создают накладные расходы — лучше выделять буферы GPU один раз и переиспользовать их, либо использовать pinned host-память для ускорения передачи между Qt и CUDA.

Пример кода

#include <cuda_runtime.h>
#include <QImage>
#include <QDebug>

// CUDA-ядро: простая инверсия цветов для RGBA8888-изображения
__global__ void invertColorsKernel(unsigned char *pixels, int width, int height)
{
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    if (x >= width || y >= height) return;

    int idx = (y * width + x) * 4; // 4 байта на пиксель: R, G, B, A
    pixels[idx + 0] = 255 - pixels[idx + 0]; // R
    pixels[idx + 1] = 255 - pixels[idx + 1]; // G
    pixels[idx + 2] = 255 - pixels[idx + 2]; // B
    // Альфа-канал (idx + 3) оставляем без изменений
}

QImage invertImageOnGpu(const QImage &input)
{
    QImage image = input.convertToFormat(QImage::Format_RGBA8888);
    const int width = image.width();
    const int height = image.height();
    const size_t bytes = static_cast<size_t>(width) * height * 4;

    unsigned char *deviceBuffer;
    cudaMalloc(&deviceBuffer, bytes);
    cudaMemcpy(deviceBuffer, image.constBits(), bytes, cudaMemcpyHostToDevice);

    dim3 blockSize(16, 16);
    dim3 gridSize((width + blockSize.x - 1) / blockSize.x,
                  (height + blockSize.y - 1) / blockSize.y);

    invertColorsKernel<<<gridSize, blockSize>>>(deviceBuffer, width, height);
    cudaDeviceSynchronize();

    QImage result(width, height, QImage::Format_RGBA8888);
    cudaMemcpy(result.bits(), deviceBuffer, bytes, cudaMemcpyDeviceToHost);

    cudaFree(deviceBuffer);
    return result;
}

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

invertImageOnGpu сначала приводит входное изображение к фиксированному формату Format_RGBA8888 через convertToFormat() — это гарантирует предсказуемый порядок и количество байт на пиксель, необходимое для корректной индексации внутри CUDA-ядра. Буфер GPU выделяется один раз под весь объём данных изображения, заполняется через cudaMemcpy из image.constBits(), обрабатывается двумерной сеткой блоков (по аналогии с любой другой двумерной задачей CUDA — каждый поток обрабатывает один пиксель), а результат копируется в новый объект QImage через result.bits().

Использование двумерной конфигурации сетки и блока (dim3 blockSize(16, 16)) — естественный выбор для изображений, поскольку напрямую отображает координаты пикселя (x, y) на индексы потока без дополнительных преобразований.

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

QImage::bits() вызывает глубокое копирование (detach), если изображение разделяет данные с другой копией (механизм copy-on-write Qt) — при многократном вызове bits() на одном и том же объекте в цикле обработки кадров стоит убедиться, что это не приводит к лишним копированиям.
Постоянное выделение/освобождение GPU-памяти на каждый кадр в реальном времени (видео) создаёт заметные накладные расходы — для потоковой обработки буферы GPU стоит выделять один раз при инициализации и переиспользовать между кадрами одинакового размера.
Несоответствие порядка байт каналов. Разные форматы QImage (Format_RGB888, Format_ARGB32, Format_RGBA8888) имеют разный порядок байт в памяти — попытка применить CUDA-ядро, рассчитанное на один порядок каналов, к изображению другого формата даст визуально неверный, но не сразу очевидный результат (например, перепутанные синий и красный каналы).
Отсутствие интеграции с GPU-рендерингом через OpenGL/Vulkan interop. Описанный подход всегда копирует данные через системную память (host), что добавляет латентность — для по-настоящему высокопроизводительных сценариев (реального видео в реальном времени) рассматривают CUDA-OpenGL interop, позволяющий CUDA напрямую писать в текстуру, используемую для рендеринга, без промежуточного копирования через CPU.