OpenCL ? CUDA ? Используем GPU на полную мощь с VexCL
Оба SDK между собой не совместимы, но их можно использовать на заднем фоне, но иметь один высокоуровневый интерфейс, написаный на языке высокого уровня и подменять реализацию по мере необходимости. Это и есть библиотека VexCL
VexCL
VexCL может использовать в качестве бекэнда как CUDA SDK, так и OpenCL C++ обертку, так и Boost.Compute - другую обертку над OpenCL C API.VexCL берет на себя всю работу по инициализации SDK.
Модель памяти, которую предоставляет VexCL очень похожа на модель памяти OpenCL. Чтобы сделать любые вычисления необходимо скопировать данные в память GPU с помощью функции vex::copy. Далее сделать что-либо с данными - сложить/умнодить вектора сделать быстрое преобразование Фурье, или вызвать кастомное ядро и возвратить результаты на хост с помощью vex::copy.
 |
| Вычисления с помощью VexCL |
Настройка Xcode
Начнем. Для начала скачиваем VexCL
git clone https://github.com/ddemidov/vexcl.git
И устанавливаем буст (если он еще не установлен).
brew install boost
Далее в настройках проекта Target -> Build Phases добавляем необходимые библиотеки (boost.thread, boost.filesystem, boost.system, boost.chrono и boost.date-time) и OpenCL.framework
 |
| Линкуем библиотеки Boost и OpenCL фреймворк |
Начинаем писать код.
Подключим заголовки и определяем какой бэкэнд будет использовать VexCL:
#define VEXCL_BACKEND_OPENCL
//#define VEXCL_BACKEND_COMPUTE
//#define VEXCL_BACKEND_CUDA
#include <vexcl/vexcl.hpp>
Проверим, что в системе вообще есть GPU, которые мы можем использовать. Фильтры определяют какие вычислители мы будем использовать:
vex::Context ctx(vex::Filter::GPU);
if (!ctx)
throw std::runtime_error("No devices available.");
Далее создаем и заполняем вектора с данными:
const size_t N = 100000;
const size_t M = 1000;
std::vector<int> nodes(N);
for (size_t i = 0; i < N; ++i) {
nodes[i] = i % 2 == 0;
}
std::vector<int> edges(N * M);
for (size_t i = 0; i < N * M; ++i) {
edges[i] = rand() % 42;
}
std::vector<int> res(N);
Ну, и собственно сами подсчеты. Для каждой вершины подсчитываем сумму всех ребер, умножаем да значение узла графа и записываем результат по соответствующему индексу:
for (size_t i = 0; i < N; i++) {
int sum = 0;
for (size_t j = i * M; j < i * M + M; ++j) {
sum += edges[j];
}
res[i] = nodes[i] * sum;
}
Тут все просто. Теперь будем использовать GPU. Создадим инстансы класса vex::vector, который выделит память на GPU и будет ей оперировать, и скопируем туда данные из наших векторов. Вектора создаются в контекстах, :
vex::vector<int> device_nodes(ctx, N);
vex::copy(nodes, device_nodes);
vex::vector<int> device_edges(ctx, N * M);
vex::copy(edges, device_edges);
vex::vector<int> device_res(ctx, N);
vex::backend::kernel my_kernel(ctx.queue(0),
VEX_STRINGIZE_SOURCE(
kernel void my_kernel(ulong N, ulong M,
global int *nodes,
global int *edges,
global int *res)
{
for (size_t i = get_global_id(0);
i < N;
i += get_global_size(0))
{
int sum = 0;
for (size_t j = i * M; j < i * M + M; ++j) {
sum += edges[j];
}
res[i] = nodes[i] * sum;
}
}),
"my_kernel");
Цикл в ядре немного усложнился. Поскольку функция будет выполнятся многими ядрами, то нам нужно определить кусок работы, который будет выполнять каждое ядро. get_global_id - это специальная функция, которая возвращает индекс ядра для размерности (поскольку у нас одномерный массив, то параметр равен 0). get_global_size возвращает размер участка.
Дальше все что нам нужно сделать, это вызвать ядро как обычную функцию, и передать ей все параметры. vex::vector<T>.operator() возвращает указатель на буфер на GPU, который был выделен при инициализации:
my_kernel(ctx.queue(0), N, M, device_nodes(), device_edges(), device_res());
Это запустит процесс выполнения ядра на графической карте, после чего нужно будет получить результаты. Просто скопируем данные из device_res в вектор res:
vex::copy(device_res, res);
Результаты
Скорость выполнения каждого из вариантов:
CPU : 301.921 ms
OpenMP: 84.7895 ms
OpenCL: 0.0063035 ms
Где OpenMP - это первый вариант, но цикл окнужен #pragma omp parallel for, как описано в предыдущей статье для автоматического распараллеливания. Выполняется на 4 ядерном i7 в 8 потоков.
Ложка дегтя
Вроде бы все оочень быстро, но не все так радужно. Эти замеры сделаны без учета времени на копирование данных на GPU и обратно. Для примера:
copy to device nodes: 0.160803 ms
copy to device edges: 64.4936 ms
OpenCL : 0.0061901 ms
copy from device res: 33.5865 ms
copy to device edges: 64.4936 ms
OpenCL : 0.0061901 ms
copy from device res: 33.5865 ms
Комментарии
Отправить комментарий