dlib.image и OpenCL

Представляем вашему вниманию очередную гостевую статью. На этот раз у нас в гостях Тимур Гафаров — создатель и главный редактор электронного журнала FPS.

Пакет dlib.image, как известно, предоставляет удобный набор средств для обработки изображений – в том числе, коллекцию простейших фильтров (размытие, выделение границ, свертка и т.д.). Как автор dlib, могу честно сказать, что реализация фильтров там далека от оптимальной – для маленьких изображений она еще годится, но, например, уже на фотографиях из мыльниц даже boxBlur выполняется неприлично долго. Можно, конечно, попытаться распараллеливать фильтры на несколько потоков или задействовать SIMD, но, как мне кажется, при наличии современных видеокарт все это – мертвому припарки: куда интереснее было бы перенести вычисления на GPU!

В этой статье мы рассмотрим реализацию boxBlur (прямоугольного размытия), одного из самых простых оконных фильтров, средствами OpenCL – специализированного API для параллельных вычислений на графических процессорах. На наше счастье, для OpenCL существует полноценный биндинг для D в лице проекта DerelictCL. Правда, он не включает поддержку OpenCL 2.0, но для наших задач хватит и версии 1.1.

Итак, подключаем необходимые модули и инициализируем биндинг:

import std.stdio;
import std.string;
import std.file;
import dlib.image;
import derelict.opencl.cl;

void main()
{
    DerelictCL.load();
    DerelictCL.reload(CLVersion.CL11);
    
    // ...
}

Дальнейший код размещается в функции main. Для краткости я опущу проверку OpenCL-вызовов на ошибки, но в реальном приложении, конечно, ей пренебрегать не стоит.

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

auto inpImage = loadPNG("input.png");
if (inpImage.pixelFormat != PixelFormat.RGBA8)
{
    writeln("Only RGBA8 images are supported!");
    return;
}

Запрашиваем количество доступных платформ и их id’ы:

cl_platform_id[] platforms;
cl_uint numPlatforms = 0;
clGetPlatformIDs(5, null, &numPlatforms);
platforms = new cl_platform_id[numPlatforms];
clGetPlatformIDs(numPlatforms, platforms.ptr, null);

Платформа – это backend реализации OpenCL. Например, для GPU от NVIDIA платформой будет CUDA. В качестве рабочей берем первую полученную платформу:

cl_platform_id platform = platforms[0];

Запрашиваем количество доступных устройств и их id’ы:

cl_uint numDevices = 0;
clGetDeviceIDs(platform, CL_DEVICE_TYPE_DEFAULT, 1, null, &numDevices);
cl_device_id[] devices = new cl_device_id[numDevices];
clGetDeviceIDs(platform, CL_DEVICE_TYPE_DEFAULT, numDevices, devices.ptr, null);

В качестве рабочего берем первое устройство:

cl_device_id device = devices[0];

Создаем контекст OpenCL. Контекст – это непосредственно среда, в которой происходят вычисления:

cl_int ret;
cl_context context = clCreateContext(null, 1, &device, null, null, &ret);

Создаем очередь команд:

cl_command_queue command_queue = clCreateCommandQueue(context, device, 0, &ret);

Создаем входное изображение (OpenCL предоставляет встроенные средства для работы с изображениями, что очень удобно):

cl_image_format format = {
    CL_RGBA,
    CL_UNORM_INT8
};

cl_mem inImgMem = clCreateImage2D(context,
    CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
    &format, inpImage.width, inpImage.height, 0, inpImage.data.ptr, &ret);

Параметр CL_MEM_COPY_HOST_PTR означает, что входные данные должны быть скопированы в память устройства.

Аналогично создаем выходное изображение:

сl_mem outImgMem = clCreateImage2D(context,
    CL_MEM_WRITE_ONLY,
    &format, inpImage.width, inpImage.height, 0, null, &ret);

Теперь самое интересное – непосредственно фильтр. В OpenCL программы для GPU пишутся на специальном C-подобном языке, а сами программы называются ядрами (kernels). За основу для нашего блюра я взял встроенный boxBlur в dlib (dlib.image.filters.boxblur) с радиусом (полуразмером окна), равным 10. Ядро выглядит следующим образом:

__constant sampler_t sampler =
    CLK_NORMALIZED_COORDS_FALSE |
    CLK_ADDRESS_CLAMP_TO_EDGE |
    CLK_FILTER_NEAREST;

__kernel void blur(
    __read_only image2d_t inImage,  
    __write_only image2d_t outImage)
{
    const int2 pos = {get_global_id(0), get_global_id(1)};

    float4 sum = (float4)(0.0f, 0.0f, 0.0f, 0.0f);

    const int radius = 10;

    for(int a = -radius; a < radius+1; a++)
    {
        for(int b = -radius; b < radius+1; b++)
        {
            sum += read_imagef(inImage, sampler, pos + (int2)(b, a));
        }
    }
    
    sum /= radius * radius * 4.0f;

    write_imagef(outImage, pos, sum);
}

Радиус, в принципе, можно (и нужно) передавать в ядро в виде параметра – при необходимости это легко добавить.

Загружаем программу из файла «boxblur.cl» и компилируем ее:

string progSrc = readText("boxblur.cl");
auto progSrcPtr = progSrc.toStringz;

size_t* lengths = cast(size_t*)[progSrc.length];
char** ptrs = cast(char**)[progSrcPtr];

cl_program program = clCreateProgramWithSource(context, 1, ptrs, lengths, &ret);
clBuildProgram(program, 0, null, null, null, null);
cl_kernel kernel = clCreateKernel(program, "blur", &ret);

Перед тем, как выполнять ядро, нужно передать ему параметры (входное и выходное изображения):

clSetKernelArg(kernel, 0, cl_mem.sizeof, cast(void*)&inImgMem);
clSetKernelArg(kernel, 1, cl_mem.sizeof, cast(void*)&outImgMem);

И последнее – копируем данные из оперативной памяти в память устройства:

size_t[3] origin = [0, 0, 0];
size_t[3] region = [inpImage.width, inpImage.height, 1];
clEnqueueWriteImage(command_queue, inImgMem, CL_TRUE,
    origin.ptr, region.ptr, 0, 0, inpImage.data.ptr, 0, null, null);

Выполняем:

size_t[2] global_work_size = [inpImage.width, inpImage.height];
ret = clEnqueueNDRangeKernel(command_queue, kernel, 2, null,
    global_work_size.ptr, null, 0, null, null);

Теперь нам нужно получить результаты вычисления обратно в оперативную память. Для этого создаем новое изображение того же размера и используем clEnqueueReadImage:

auto outImage = inpImage.createSameFormat(inpImage.width, inpImage.height);
clEnqueueReadImage(command_queue, outImgMem, CL_TRUE,
    origin.ptr, region.ptr, 0, 0, outImage.data.ptr, 0, null, null);

Осталось сохранить результат в файл:

outImage.savePNG("output.png");

На GeForce GT 740 для изображения размером 3648х2736 вся эта процедура (за вычетом операций файлового ввода/вывода) занимает около 2 секунд.

Gecko