Работа с OpenCL

cl-context.cpp

/*
 * Вариант с использованием Kronos C++ bindings.
 * Если файл CL/cl.hpp не входит в OpenCL SDK, его можно взять тут:
 * http://www.khronos.org/registry/cl/api/1.1/cl.hpp
 *
 * Преимущество перед C-API:
 * не нужно вручную отслеживать время жизни объектов OpenCL.
 */
#include <iostream>
#include <vector>
#include <string>
#include <stdexcept>

// Чтобы не нужно было проверять коды ошибок после каждой операции
// (в случае ошибки будет выброшено исключение типа cl::Error):
#define __CL_ENABLE_EXCEPTIONS

#include <CL/cl.hpp>


//---------------------------------------------------------------------------
// Проверка критического условия:
//---------------------------------------------------------------------------
void precondition(bool cond, const std::string &msg) {
    if (!cond) throw std::runtime_error(msg);
}

//---------------------------------------------------------------------------
// Класс, ответственный за работу с OpenCL.
// Инициализирует контекст, компилирует ядра, предоставляет функции вызова
// ядер.
//---------------------------------------------------------------------------
class CL_Context {
    public:
        CL_Context() {
            // Получаем подходящее устройство:
            std::vector<cl::Device> dev(1, get_device());

            // Создаем контекст и очередь заданий:
            ctx   = cl::Context(dev);
            queue = cl::CommandQueue(ctx, dev[0]);

            // Готовим ядра:
            create_add_krn(dev);
        }

        // Вызов ядра add
        void add(size_t n, cl::Buffer a, cl::Buffer b, cl::Buffer c) {
            size_t nt = 256;
            size_t nb = (n + nt - 1) / nt;

            unsigned arg_pos = 0;
            add_krn.setArg(arg_pos++, static_cast<cl_ulong>(n));
            add_krn.setArg(arg_pos++, a);
            add_krn.setArg(arg_pos++, b);
            add_krn.setArg(arg_pos++, c);

            queue.enqueueNDRangeKernel(add_krn, cl::NullRange, nt * nb, nt);
        }

        // Создание буфера на устройстве
        template <class T>
        cl::Buffer buffer(size_t n, const T *host = 0) {
            if (host) {
                return cl::Buffer(ctx, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR,
                        n * sizeof(T), const_cast<T*>(host));
            } else {
                return cl::Buffer(ctx, CL_MEM_READ_WRITE, n * sizeof(T));
            }
        }

        // Копирование данных с устройства на хост:
        template <class T>
        void D2H(cl::Buffer src, T *dst, size_t n, bool blocking = false) {
            queue.enqueueReadBuffer(src, blocking, 0, n * sizeof(T), dst);
        }

        // Копирование данных с хоста на устройство:
        template <class T>
        void H2D(const T *src, cl::Buffer dst, size_t n, bool blocking = false) {
            queue.enqueueWriteBuffer(dst, blocking, 0, n * sizeof(T), src);
        }

    private:
        cl::Context      ctx;
        cl::CommandQueue queue;
        cl::Kernel       add_krn;

        // Получить одно устройство, поддерживающее двойную точность.
        cl::Device get_device() {
            // Получаем список  OpenCL-платформ:
            std::vector<cl::Platform> platforms;
            cl::Platform::get(&platforms);

            precondition(!platforms.empty(), "No OpenCL platforms.");

            for(auto p = platforms.begin(); p != platforms.end(); ++p) {
                // Список устройств, поддерживаемых платформой:
                std::vector<cl::Device> devices;
                p->getDevices(CL_DEVICE_TYPE_ALL, &devices);

                // Находим первое подходящее устройство.
                for(auto d = devices.begin(); d != devices.end(); ++d) {
                    if (!d->getInfo<CL_DEVICE_AVAILABLE>()) continue;

                    // Поддерживаемые расширения:
                    std::string ext = d->getInfo<CL_DEVICE_EXTENSIONS>();

                    // Нас интересует двойная точность:
                    if (ext.find("cl_khr_fp64") == std::string::npos) continue;

                    return *d;
                }
            }

            precondition(false, "No compute devices found.");
        }

        // Компиляция программы
        cl::Program build_program(
                const std::vector<cl::Device> &dev,
                const std::string             &src
                )
        {
            cl::Program program(ctx,
                    cl::Program::Sources(1, std::make_pair(src.c_str(), src.size())));

            // Пытаемся скомпилировать программу:
            try {
                program.build(dev);
            } catch(const cl::Error &err) {
                // Выводим ошибки компиляции.
                std::cerr
                    << "OpenCL compilation error" << std::endl
                    << program.getBuildInfo<CL_PROGRAM_BUILD_LOG>(dev[0])
                    << std::endl;
                throw err;
            }

            return program;
        }

        // Комипиляция ядра суммирования.
        void create_add_krn(const std::vector<cl::Device> &dev) {
            add_krn = cl::Kernel(
                    build_program(dev,
                        "#pragma OPENCL EXTENSION cl_khr_fp64: enable\n"
                        "kernel void add(\n"
                        "       ulong n,\n"
                        "       global const double *a,\n"
                        "       global const double *b,\n"
                        "       global double *c\n"
                        "       )\n"
                        "{\n"
                        "    for(size_t i = get_global_id(0); i < n; i += get_global_size(0)) {\n"
                        "        c[i] = a[i] + b[i];\n"
                        "    }\n"
                        "}\n"
                        ),
                    "add"
                    );
        }
};

//---------------------------------------------------------------------------
int main() {
    try {
        // Готовим OpenCL контекст.
        CL_Context ctx;

        // Создаем входные данные, передаем их на устройство.
        const size_t n = 1 << 20;

        std::vector<double> a(n, 0);
        std::vector<double> b(n, 1);

        cl::Buffer A = ctx.buffer(n, a.data());
        cl::Buffer B = ctx.buffer(n, b.data());

        // Вызываем ядро в цикле (A = A + B):
        for(int i = 0; i < 100; ++i)
            ctx.add(n, A, B, A);

        // Забираем результат:
        ctx.D2H(A, a.data(), n, /*blocking=*/true);

        precondition(a[42] == 100, "Wrong result!");

        // Все объекты будут уничтожены автоматически при выходе из области
        // видимости.
    } catch(const cl::Error &err) {
        std::cerr << "OpenCL error: " << err.what() << "(" << err.err() << ")" << std::endl;
        return 1;
    }
}