Модель параллельного программирования spawn/sync
spawn/sync использует модель параллельных вычислений Fork-Join, при которой выполнение программы разветвляется в определенных точках и объединяется в последующих точках для возобновления последовательного выполнения программы.
Порождение и слияние параллельных ветвей программы выполняется с помощью функций
elcore_spawn() и elcore_sync(). Данные функции
доступны пользователю после подключения заголовочного файла elcore-runtime.h в коде программы
DSP.
В рамках spawn/sync модели существуют основной и вспомогательные DSP:
Основной DSP — ядро, которое выполняет программу (кернел), заданную пользователем. Вспомогательный DSP — ядро, которое выполняет специальный код и время от времени переключается в адресное пространство основного DSP для параллельного выполнения некоторой функции.
При вызове основным DSP функции elcore_spawn() выполняется разветвление исходной последовательной программы и делегирование выполнения участка кода одному из свободных вспомогательных DSP (при их наличии).
При использовании модели программирования spawn/sync для запуска кернелов в ELcoreCL необходимо
использовать функцию eclEnqueueKernelWithSpawn() вместо функции
eclEnqueueNDRangeKernel().
Ограничения
Не допускается использование функций из стандартных библиотек, в которых используются глобальные переменные (например,
printf()иmalloc()), в распараллеливаемой функции.Передача через
elcore_spawn()структур и прочих сложных объектов, размер которых превышает 8 байт, должна осуществлятся через указатель.Не реализована возможность рекурсивного вызова
elcore_spawn().Очередь заданий — циклическая. Размер очереди — 256 заданий.
Аргументы-указатели, передаваемые в функцию
elcore_spawn()должны быть кратны 64 байтам.Примечание
Выравнивание указателей, передаваемых только для чтения, необязательно.
При динамическом выделении памяти, выравнивания можно добиться, например, используя следующие стандартные функции:
posix_memalign()memalign()
При статическом выделении памяти выравнивание адресов обеспечивается за счет использования атрибута компилятора
aligned. Пример выделения выравненного статического массива размером 1024 целочисленных элементов:int intarray[1024] __attribute__((aligned(64)));
Примечание
Выравненность указателей, передаваемых с CPU через ELcoreCL функциями
eclSetKernelArgELcore*(), гарантируется драйвером DSP.
Примечание
Из-за отсутствия аппаратной когерентности между несколькими DSP часть времени уходит на
программное согласование данных во время вызовов elcore_spawn() и elcore_sync().
При распараллеливании небольших функций производительность может быть ограничена накладными
расходами, связанными с обеспечением согласованности данных.
Пример параллельного умножения двух векторов
В примере выполняется параллельное умножение двух векторов на всех DSP, доступных в системе с использованием spawn/sync модели запуска.
Исходные вектора и их размер, а также количество используемых ядер передаются из CPU в качестве входных аргументов, результирующий вектор — в качестве выходного аргумента. После завершения обработки проверяется корректность результата.
CPU-часть программы:
// Copyright 2022 RnD Center "ELVEES", JSC
/*
Пример параллельного умножения двух векторов на DSP
*/
#include <fstream>
#include <vector>
#include <err.h>
#include <error.h>
#include <stdio.h>
#include <unistd.h>
#include <elcorecl/elcorecl.h>
/* Путь к заранее скомпилированному ELF-файлу для DSP */
#define DSP_ELF_PATH "/usr/share/elcore50/dsp-parallel-vecmul.elf"
const size_t VECSIZE = (1024 * 1024 * 32); // 32 MB
int main(int argc, char **argv) {
ecl_int ret;
/* Получаем id платформы */
ecl_platform_id platform_id;
ret = eclGetPlatformIDs(1, &platform_id, nullptr);
if (ret != ECL_SUCCESS) errx(1, "Failed to get platform id. Error code: %d", ret);
/* Получаем список достуных DSP-устройств */
ecl_uint ndevs;
ret = eclGetDeviceIDs(platform_id, ECL_DEVICE_TYPE_CUSTOM, 0, nullptr, &ndevs);
if (ret != ECL_SUCCESS) errx(1, "Failed to get device count. Error code: %d", ret);
if (ndevs == 0) errx(1, "No available ECL_DEVICE_TYPE_CUSTOM devices");
/* Получаем массив устройств типа ecl_device_id */
std::vector<ecl_device_id> devices(ndevs);
ret = eclGetDeviceIDs(platform_id, ECL_DEVICE_TYPE_CUSTOM, ndevs, &devices[0], nullptr);
if (ret != ECL_SUCCESS) errx(1, "Failed to get device id. Error code: %d", ret);
/* Создаем контекст */
ecl_context context;
context = eclCreateContext(nullptr, ndevs, &devices[0], nullptr, nullptr, &ret);
if (context == nullptr || ret != ECL_SUCCESS)
errx(1, "Failed to create context. Error code: %d", ret);
/* Считываем ELF-файл в буфер */
std::vector<char> elf_buffer;
size_t elf_sizes[ndevs];
{
std::ifstream file(DSP_ELF_PATH, std::ios::binary | std::ios::ate);
elf_sizes[0] = file.tellg();
file.seekg(0, std::ios::beg);
elf_buffer.resize(elf_sizes[0]);
file.read(elf_buffer.data(), elf_sizes[0]);
}
const unsigned char *elfs[ndevs];
for (int i = 0; i < ndevs; ++i) {
elf_sizes[i] = elf_sizes[0];
elfs[i] = reinterpret_cast<unsigned char *>(elf_buffer.data());
}
/* Создаем программу */
ecl_program program;
program = eclCreateProgramWithBinary(context, ndevs, &devices[0], elf_sizes, elfs,
nullptr, &ret);
if (program == nullptr || ret != ECL_SUCCESS)
errx(1, "Failed to create program. Error code: %d", ret);
/* Создаём кернел ElcoreCL. Код функции vecmul_parallel, содержащий вызовы функций
elcore_spawn() и elcore_sync(), содержится в файле dsp-kernel/vecmul.c */
ecl_kernel kernel;
kernel = eclCreateKernel(program, "_vecmul_parallel", &ret);
if (kernel == nullptr || ret != ECL_SUCCESS)
errx(1, "Failed to create kernel. Error code: %d", ret);
/* Создаем очередь задач типа ecl_multi_command_queue.
DSP0 - основное ядро, остальные - вспомогательные */
ecl_multi_command_queue multi_command_queue;
multi_command_queue =
eclCreateMultiCommandQueueWithProperties(context, devices[0], nullptr,
ndevs - 1, &devices[1], &ret);
if (multi_command_queue == nullptr || ret != ECL_SUCCESS)
errx(1, "Failed to create multi_command_queue. Error code: %d", ret);
// Подготавливаем аргументы
uint8_t *input_a, *input_b, *output;
// Пользовательские указатели, передаваемые в DSP, должны быть кратны 64.
// Выравниваем по размеру страницы с запасом
size_t page_size = getpagesize();
if (posix_memalign((void **)&input_a, page_size, VECSIZE) != 0)
errx(1, "Failed to allocate memory. Error code: %d", ret);
if (posix_memalign((void **)&input_b, page_size, VECSIZE) != 0)
errx(1, "Failed to allocate memory. Error code: %d", ret);
if (posix_memalign((void **)&output, page_size, VECSIZE) != 0)
errx(1, "Failed to allocate memory. Error code: %d", ret);
// Заполним входные буфера случайными значениями из файла /dev/urandom
FILE *urandom_file = fopen("/dev/urandom", "r");
if (urandom_file == nullptr) errx(1, "Failed to open /dev/urandom");
ret = fread(input_a, 1, VECSIZE, urandom_file);
if (ret != VECSIZE) errx(1, "Failed to read %u bytes from /dev/urandom", VECSIZE);
ret = fread(input_b, 1, VECSIZE, urandom_file);
if (ret != VECSIZE) errx(1, "Failed to read %u bytes from /dev/urandom", VECSIZE);
// Создаем объекты памяти ELcoreCL
ecl_mem input_a_mem = eclCreateBuffer(context, ECL_MEM_USE_HOST_PTR, VECSIZE, input_a, &ret);
if (input_a_mem == nullptr || ret != ECL_SUCCESS)
errx(1, "Failed to create memory object. Error code: %d", ret);
ecl_mem input_b_mem = eclCreateBuffer(context, ECL_MEM_USE_HOST_PTR, VECSIZE, input_b, &ret);
if (input_b_mem == nullptr || ret != ECL_SUCCESS)
errx(1, "Failed to create memory object. Error code: %d", ret);
ecl_mem output_mem = eclCreateBuffer(context, ECL_MEM_USE_HOST_PTR, VECSIZE, output, &ret);
if (output_mem == nullptr || ret != ECL_SUCCESS)
errx(1, "Failed to create memory object. Error code: %d", ret);
// Передаем необходимые аргументы в ELcoreCL-ядро
ret = eclSetKernelArgELcoreMem(kernel, 0, input_a_mem);
if (ret != ECL_SUCCESS) errx(1, "Failed to set arg0. Error code: %d", ret);
ret = eclSetKernelArgELcoreMem(kernel, 1, input_b_mem);
if (ret != ECL_SUCCESS) errx(1, "Failed to set arg1. Error code: %d", ret);
ret = eclSetKernelArg(kernel, 2, sizeof(size_t), &VECSIZE);
if (ret != ECL_SUCCESS) errx(1, "Failed to set arg2. Error code: %d", ret);
ret = eclSetKernelArg(kernel, 3, sizeof(size_t), &ndevs);
if (ret != ECL_SUCCESS) errx(1, "Failed to set arg3. Error code: %d", ret);
ret = eclSetKernelArgELcoreMem(kernel, 4, output_mem);
if (ret != ECL_SUCCESS) errx(1, "Failed to set arg4. Error code: %d", ret);
ecl_event event;
/* Запускаем ELcoreCL-ядро на выполнение */
ret = eclEnqueueKernelWithSpawn(multi_command_queue, kernel, 0, nullptr, &event);
if (ret != ECL_SUCCESS) errx(1, "Failed to enqueue kernel. Error code: %d", ret);
/* Ожидаем завершения задания */
ret = eclWaitForEvents(1, &event);
if (ret != ECL_SUCCESS) errx(1, "Failed to wait kernel. Error code: %d", ret);
/* Инвалидируем CPU-кэши выходного буфера */
eclEnqueueMapBuffer((ecl_command_queue) multi_command_queue, output_mem, ECL_TRUE,
ECL_MAP_READ, 0, VECSIZE, 0, NULL, NULL, &ret);
if (ret != ECL_SUCCESS) errx(1, "Failed to map output_mem. Error code: %d", ret);
/* Проверяем результат */
int errors = 0;
for (int i = 0; i < VECSIZE; ++i) {
uint8_t expect_value = input_a[i] * input_b[i];
if (expect_value != output[i]) {
errors = 1;
printf("failed at %d: %x * %x != %x\n", i, input_a[i], input_b[i], output[i]);
break;
}
}
/* Освобождаем ресурсы */
eclReleaseMultiCommandQueue(multi_command_queue);
eclReleaseKernel(kernel);
eclReleaseProgram(program);
eclReleaseContext(context);
for (int i = 0; i < ndevs; ++i) eclReleaseDevice(devices[i]);
printf("Results validation %s\n", errors ? "failed" : "succeeded");
return errors;
}
DSP-часть программы:
// Copyright 2022 RnD Center "ELVEES", JSC
#include <stdio.h>
#include <stdint.h>
#include <malloc.h>
#include <elcore-runtime.h> // elcore_spawn(), elcore_sync()
void vecmul(uint8_t *a, uint8_t *b, uint8_t *c, size_t size) {
for (int i = 0; i < size; ++i) c[i] = a[i] * b[i];
}
void vecmul_parallel(uint8_t *input_a, uint8_t *input_b, const size_t vecsize, const size_t cores,
uint8_t *output) {
int task_ids[cores - 1];
/* Предполагаем, что vecsize делится без остатка на cores,
а результат деления кратен 64 (размеру кэш-линии) */
size_t size_per_task = vecsize / cores;
/* elcore_spawn() позволяет порождать до 256 параллельных ветвей (заданий), независимо от того
сколько вспомогательных DSP было запущено.
Если число параллельных ветвей превышает количество запущенных DSP, они будут выполняться
по мере освобождения или в момент вызова elcore_sync() текущим DSP.
В данном примере будем считать, что всего доступно 16 DSP. */
for (int i = 1; i < cores; ++i) {
task_ids[i - 1] = elcore_spawn((void *)vecmul, 4, input_a + i * size_per_task,
input_b + i * size_per_task, output + i * size_per_task,
size_per_task);
}
// На текущем DSP в это время обработаем свою часть данных
vecmul(input_a, input_b, output, size_per_task);
/* Ожидаем завершения параллельных ветвей.
Если одна из параллельных ветвей все еще не начала выполняться (нет свободных DSP),
текущий DSP в момент вызова elcore_sync() выполниит ее. */
for (int i = 1; i < cores; ++i)
elcore_sync(task_ids[i - 1], NULL);
}
Пример собирается в рамках пакета elcorecl-examples в Buildroot. Утилита доступна на отладочном модуле по имени ecl-parallel-vecmul.