Библиотека ElcoreCL

Общее описание библиотеки

Библиотека ElcoreCL для DSP реализует аналогичный OpenCL функционал. Именование функций аналогично OpenCL, но вместо префикса cl используется префикс ecl. Аргументы почти во всех функциях передаются в том же порядке, что и в OpenCL и имеют те же типы, но с префиксом ecl, например: ecl_program вместо cl_program. Коды ошибок и другие макросы также соответствуют OpenCL, но с префиксом ECL.

ElcoreCL имеет следующие особенности:

• реализована поддержка кэшируемых буферов на основе указателей пространства пользователя;

• реализована поддержка импорта кэшируемых и некэшируемых dma-buf;

• не реализован компилятор языка OpenCL. Вместо этого используются предкомпилированные при помощи C++ компилятора DSP бинарные elf-файлы;

• гарантируется работа только следующих функций, аналогичных OpenCL:

Функция ElcoreCL	Функция OpenCL
eclGetPlatformIDs()	clGetPlatformIDs
eclGetPlatformInfo()	clGetPlatformInfo
eclGetDeviceIDs()	clGetDeviceIDs
eclGetDeviceInfo()	clGetDeviceInfo
eclCreateContext()	clCreateContext
eclCreateContextFromType()	clCreateContextFromType
eclGetContextInfo()	clGetContextInfo
eclCreateCommandQueueWithProperties()	clCreateCommandQueueWithProperties
eclGetCommandQueueInfo()	clGetCommandQueueInfo
eclCreateBuffer()	clCreateBuffer
eclGetMemObjectInfo()	clGetMemObjectInfo
eclSetMemObjectDestructorCallback()	clSetMemObjectDestructorCallback
eclCreateProgramWithBinary()	clCreateProgramWithBinary
eclGetProgramInfo()	clGetProgramInfo
eclSetProgramReleaseCallback()	clSetProgramReleaseCallback
eclCreateKernel()	clCreateKernel
eclCreateKernelsInProgram()	clCreateKernelsInProgram
eclCloneKernel()	clCloneKernel
eclGetKernelInfo()	clGetKernelInfo
eclSetKernelArg()	clSetKernelArg
eclEnqueueNDRangeKernel()	clEnqueueNDRangeKernel
eclCreateUserEvent()	clCreateUserEvent
eclSetUserEventStatus()	clSetUserEventStatus
eclWaitForEvents()	clWaitForEvents
eclGetEventInfo()	clGetEventInfo
eclSetEventCallback()	clSetEventCallback
eclFinish()	clFinish
eclRelease*(), eclRetain*()	clRelease*, clRetain*

• Следующие функции не имеют аналогов в OpenCL:

  • eclSetKernelArgELcoreMem()

  • eclSetKernelArgELcoreMemNonCached()

  • eclSetKernelArgELcoreLocalMem()

  • eclSetKernelArgDMAMem()

  • eclCreateDMABuffer()

  • eclCreateBufferFromDmaBuf()

  • eclCachesFlush()

  • eclCachesInvalidate()

  • eclCachesFlushInvalidate()

• для функции eclCreateBuffer() действует ограничение:

  • параметр flags обязательно должен включать флаг ECL_MEM_USE_HOST_PTR.

• для функции eclEnqueueNDRangeKernel() действуют ограничения:

  • параметр work_dim должен быть равен 1;

  • параметр global_work_offset должен быть равен NULL;

  • параметр global_work_size должен указывать на массив, содержащий один элемент, равный 1;

  • параметр local_work_size должен быть равен NULL.

• функция eclCreateProgramWithBinary() принимает в качестве указателя на бинарный код указатель на загруженный в память ELF-файл, собранный для DSP.

• реализована поддержка файловых системных вызовов DSP.

• для функции eclSetKernelArgELcoreLocalMem() действует ограничение:

  • возвращается ошибка ECL_INVALID_KERNEL, если в ELF-файле программы описана XYRAM-секция.

Описание функций и аргументов можно найти в спецификации OpenCL.

Функции, не имеющие аналогов в OpenCL

Создание программы с общими секциями данных

ecl_program eclCreateProgramWithSharedSections()

Функция eclCreateProgramWithSharedSections() аналогична функции eclCreateProgramWithBinary(), дополнительно позволяя передать список объектов памяти, которые не являются частью ELF-файла, но которые необходимо отобразить в адресное пространство соответствующих DSP. Список объектов памяти передается через структуру ecl_shared_section, которая дополнительно содержит виртуальный адрес, по которому должна располагаться секция в адресном пространстве DSP. Функция eclCreateProgramWithSharedSections() позволяет снизить накладные расходы при выполнении eclEnqueueNDRangeKernel().

Функции установки аргументов ядра

ElcoreCL при создании объекта ecl_program принимает в качестве аргумента ELF-файл, в котором отсутствует информация о типах принимаемых аргументов. Поэтому пользователю необходимо явно вызывать различные функции для различных типов аргументов:

Функция ElcoreCL	Аргумент CPU-части	Тип аргумента в DSP
eclSetKernelArgELcoreMem()	ecl_mem	Указатель на любой тип. Например, void\*
eclSetKernelArgELcoreMemNonCached()	ecl_mem	Указатель на любой тип. Например, void\*
eclSetKernelArgELcoreLocalMem()	размер требуемой XYRAM	Указатель на любой тип. Например, void\*
eclSetKernelArgDMAMem()	ecl_mem	64-битное целое. Например, uint64_t
eclSetKernelArg()	Указатель и размер аргумента	Произвольный тип, передаваемый по значению

Функция eclSetKernelArgELcoreMemNonCached() передаёт область памяти, привязанную к объекту ecl_mem, как некэшируемую с точки зрения DSP.

Функция eclSetKernelArgELcoreLocalMem() резервирует память в XYRAM, никак не инициализируя её. Соответствует вызову функции clSetKernelArg() для аргумента с модификатором __local.

Функции создания очередей задач

ecl_multi_command_queue eclCreateMultiCommandQueueWithProperties(ecl_context context, ecl_device_id device, const ecl_queue_properties *properties, ecl_uint num_spawn_devices, const ecl_device_id *spawn_device_list, ecl_int *errcode_ret)

Функция eclCreateMultiCommandQueueWithProperties() создает объект очереди задач типа ecl_multi_command_queue для последующего запуска ядер в режиме spawn/sync, используя eclEnqueueKernelWithSpawn(). В рамках модели spawn/sync device является основным устройством, а spawn_device_list — массив из num_spawn_devices вспомогательных устройств.

В рамках eclCreateMultiCommandQueueWithProperties() создаются служебные объекты памяти, размещаемые в некэшируемой памяти и используемые для взаимодействия между основным и вспомогательными устройствами.

Функции запуска ядер

ecl_int eclEnqueueKernelWithSpawn(ecl_multi_command_queue multi_command_queue, ecl_kernel kernel, ecl_uint num_events_in_wait_list, const ecl_event *event_wait_list, ecl_event *event)

Функция eclEnqueueKernelWithSpawn() ставит в очередь multi_command_queue, созданную функцией eclCreateMultiCommandQueueWithProperties() на выполнение kernel в режиме spawn/sync.

Функция eclEnqueueKernelWithSpawn() выполняет:

1. Создает специальное ядро _spawner_loop для запуска на каждом из вспомогательном устройстве очереди multi_command_queue.

2. Если в ядро kernel была передана в качестве одного из аргументов локальная память, передает в ядро _spawner_loop в качестве аргумента локальную память такого же размера.

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

4. Помещает в очереди вспомогательных устройств, созданные в рамках multi_command_queue, ядро _spawner_loop.

5. Помещает в очередь multi_command_queue на выполнение ядро kernel.

Примечание

При вызове функции eclEnqueueKernelWithSpawn() запрещено использовать аргументы, передаваемые функцией eclSetKernelArg() и имеющий размер больше 8 байт.

ecl_int eclEnqueueSharedKernel(ecl_command_queue command_queue, ecl_kernel kernel, ecl_uint work_dim, const size_t *global_work_offset, const size_t *global_work_size, const size_t *local_work_size, ecl_uint num_events_in_wait_list, const ecl_event *event_wait_list, ecl_event *event)

Функция eclEnqueueSharedKernel() аналогична функции eclEnqueueNDRangeKernel() за исключением:

• ELcoreCL-ядро kernel на уровне драйвера elcore50 ставится в глобальную очередь заданий, а не локальную очередь DSP-ядра как в случае функции eclEnqueueNDRangeKernel().

• Устройство, соответствующее очереди команд command_queue, интерпретируется как виртуальное. Выбор конкретного DSP-ядра осуществляется планировщиком внутри драйвера elcore50.

Подробнее про функционирование очередей в драйвере elcore50 см. Локальная и глобальная очереди задач.

Функции для поддержки 64-битного VDMA

DSP имеет следующие аппаратные особенности доступа к памяти:

• процессор работает в 32-битном адресном пространстве и, соответственно, не имеет возможности обращаться к памяти с адресами более 4ГБ;

• VDMA работает в 64-битном адресном пространстве и, соответственно, иммеет возможность обращаться к адресам более 4ГБ, а также размером более 4ГБ.

Таким образом, для поддержки буферов, размером более 4ГБ, реализованы следующие функции:

eclCreateDMABuffer()

ecl_mem eclCreateDMABuffer()

Функция eclCreateDMABuffer() аналогична функции eclCreateBuffer(). Единственное отличие — поддержка размера памяти более 4ГБ.

Возвращённый объект ecl_mem рекомендуется устанавливать в качестве аргумента ядру при помощи функции eclSetKernelArgDMAMem().

eclSetKernelArgDMAMem()

cl_int eclSetKernelArgDMAMem()

Функция eclSetKernelArgDMAMem() устанавливает аргумент ядра типа ecl_mem. В виртуальном адресном пространстве DSP адрес маппированной области памяти будет > 4ГБ. Поэтому соответствующий аргумент ядра должен быть не void\*, а uint64_t. Доступ к этой памяти возможен только через VDMA.

Функции для поддержки импортируемых dma-buf

dma-buf — подсистема ядра Linux, которая реализует интерфейс для совместного использования буферов между различными драйверами и подсистемами без копирования, а также для синхронизации асинхронного доступа.

Для поддержки импортируемых dma-buf реализованы функции:

Импорт dma-buf

ecl_mem eclCreateBufferFromDmaBuf()

Функция eclCreateBufferFromDmaBuf() создает объект памяти ecl_mem на основе файлового дескриптора импортируемого dma-buf.

Возвращённый объект ecl_mem рекомендуется устанавливать в качестве аргумента ядру при помощи функций eclSetKernelArgELcoreMem() и eclSetKernelArgELcoreMemNonCached().

Функции управления кэш-память CPU

ecl_int eclCachesFlush(ecl_mem memobj, size_t offset, size_t cb)

Функция eclCachesFlush() выполняет сброс кэш-памяти CPU для cb байт объекта памяти memobj со смещением, равным offset байт.

ecl_int eclCachesInvalidate(ecl_mem memobj, size_t offset, size_t cb)

Функция eclCachesInvalidate() выполняет инвалидацию кэш-памяти CPU для cb байт объекта памяти memobj со смещением, равным offset байт.

ecl_int eclCachesFlushInvalidate(ecl_mem memobj, size_t offset, size_t cb)

Функция eclCachesInvalidate() выполняет сброс и инвалидацию кэш-памяти CPU для cb байт объекта памяти memobj со смещением, равным offset байт.

Запуск ELcoreCL-ядер в режиме отладки

Библиотека elcorecl позволяет запускать ELcoreCL-ядра, помещаемые в очередь с помощью функций eclEnqueueNDRangeKernel() и eclEnqueueKernelWithSpawn(), в режиме отладки, при котором их выполнение будет прервано на первой инструкции для дальнейшей отладки в соответствии с Отладка заданий DSP. Управление режимом отладки осуществляется с помощью переменной окружения ELCORE_DEBUG_ENABLE.

Допустимые значения переменной окружения ELCORE_DEBUG_ENABLE:

• 0 — запуск в обычном режиме;

• 1 — запуск в режиме отладки.

Случай, когда переменная окружения ELCORE_DEBUG_ENABLE не определена, аналогичен случаю, когда переменная окружения равна 0.

Примеры работы с ElcoreCL при программировании DSP

Сценарий использования кэшируемых буферов на основе пользовательских указателей

/* Получаем доступные платформы */
ret = eclGetPlatformIDs(1, &platform_id, &ret_num_platforms);

/* Получаем доступные устройства */
/* В результате в массиве device_ids получим список идентификаторов,
   соответствующих каждому отдельному DSP-ядру */
ret = eclGetDeviceIDs(platform_id, CL_DEVICE_TYPE_DEFAULT, 16, &device_ids,
  &ret_num_devices);

/* Создаём контекст */
/* Оставляем в массиве device_ids те DSP-ядра, которые мы хотим использовать в
   программе. Поскольку идентификаторы устройств выбираются из списка, полученного при помощи
   eclGetDeviceIDs, а device_num соответствует конечному количеству выбранных
   устройств, то num_devices <= ret_num_devices */
context = eclCreateContext(NULL, num_devices, &device_ids, NULL, NULL,
  &ret_errcode);

/* Загружаем ELF (или несколько ELF'ов) */
program = eclCreateProgramWithBinary(context, num_devices, &device_ids,
  &lengths, &binaries, &binary_status, &ret_errcode);

/* Получаем список доступных ядер (если не знаем названий без этого) */
eclGetProgramInfo(program, ECL_PROGRAM_KERNEL_NAMES, param_value_size,
  &kernel_names, &ret_kernel_names_size);

/* Создаём буфер, который можно будет использовать в ядре */
void host_ptr = aligned_alloc(alignment, mem_size);

/* Записываем данные в буфер host_ptr */
memset(host_ptr, 0xFE, mem_size);

/* Создаём ElcoreCL memory object */
memory = eclCreateBuffer(context, CL_MEM_USE_HOST_PTR, mem_size, host_ptr,
  &ret_errcode);

/* Создаем очередь команд */
command_queue = eclCreateCommandQueueWithProperties(context, device_ids[0], NULL,
  &ret_errcode);

/* Создаём ядро ElcoreCL */
kernel = eclCreateKernel(program, kernel_name, &ret_errcode);

/* Устанавливаем параметры ядра */
/* Если аргументом ядра является локальная память, то в качестве указателя
   передаётся NULL. Эта память никак не инициализируется. За её инициализацию
   отвечает ядро ElcoreCL. Для ядра EcloreCL - это просто непрерывная область в
   XYRAM */
/* На данном этапе библиотека ElcoreCL проверяет, уложится ли вся память в область
   адресов DSP без конфликтов. Если не уложится, то будет возвращена ошибка
   */
eclSetKernelArg(kernel, 0, mem_size, memory);

/* Запускаем ядро на выполнение */
/* Именно здесь формируются окончательные данные для драйвера на запуск задания:
   все параметры ядра ElcoreCL вносятся в структуру для ioctl и передаются в
   драйвер вместе с адресом ELF-файла и entry point для ядра ElcoreCL.
   Драйвер в свою очередь для всех memory object'ов вызывает get_user_pages (что
   заблокирует соответствующую память от swap'инга), формирует образ стека и
   будущие значения регистров в соответствии с соглашением о вызовах */
ret = eclEnqueueNDRangeKernel(command_queue, kernel, 1, NULL /*work_offset*/,
  &[1] /*global_work_size*/, NULL /*local_work_size*/, num_events_in_wait_list,
  &event_wait_list, &ret_event);
/* Ожидаем завершения задания */
eclWaitForEvents(1, &ret_event);

/* Инвалидируем кэш, чтобы host получил доступ к актуальным данным */
host_ptr = eclEnqueueMapBuffer(command_queue, memory, blocking_map, map_flags,
  offset, cb, num_events_in_wait_list, &event_wait_list, &ret_event, &ret);
/* Ожидаем завершения инвалидации */
eclWaitForEvents(1, &ret_event);

Сценарий использования импортируемого dma-buf

ret = eclGetPlatformIDs(1, &platform_id, &ret_num_platforms);
ret = eclGetDeviceIDs(platform_id, CL_DEVICE_TYPE_DEFAULT, 16, &device_ids,
  &ret_num_devices);
context = eclCreateContext(NULL, num_devices, &device_ids, NULL, NULL,
  &ret_errcode);
program = eclCreateProgramWithBinary(context, num_devices, &device_ids,
  &lengths, &binaries, &binary_status, &ret_errcode);

/* Создаём ElcoreCL memory object на основе импортируемого dmabuf */
memory = eclCreateBufferFromDmaBuf(context, flags, imported_dmabuf_fd, mem_size,
  &ret_errcode);

/* Создаем очередь команд */
command_queue = eclCreateCommandQueueWithProperties(context, device_ids[0], NULL,
  &ret_errcode);

/* Создаём ядро ElcoreCL */
kernel = eclCreateKernel(program, kernel_name, &ret_errcode);

/* Устанавливаем параметры ядра */
eclSetKernelArg(kernel, 0, mem_size, memory);

/* Запускаем ядро на выполнение */
ret = eclEnqueueNDRangeKernel(command_queue, kernel, 1, NULL /*work_offset*/,
  &[1] /*global_work_size*/, NULL /*local_work_size*/, num_events_in_wait_list,
  &event_wait_list, &ret_event);
/* Ожидаем завершения задания */
eclWaitForEvents(1, &ret_event);

/* Инвалидируем кэш, чтобы host получил доступ к актуальным данным */
host_ptr = eclEnqueueMapBuffer(command_queue, memory, blocking_map, map_flags,
  offset, cb, num_events_in_wait_list, &event_wait_list, &ret_event, &ret);
/* Ожидаем завершения инвалидации */
eclWaitForEvents(1, &ret_event);

Обработка задачи на нескольких DSP-ядрах

В данном разделе продемонстрировано, как можно запустить конвейерную обработку на нескольких DSP-ядрах.

ret = eclGetPlatformIDs(1, &platform_id, &ret_num_platforms);
ret = eclGetDeviceIDs(platform_id, CL_DEVICE_TYPE_DEFAULT, 16, &device_ids,
  &ret_num_devices);
context = eclCreateContext(NULL, num_devices, &device_ids, NULL, NULL,
  &ret_errcode);
program = eclCreateProgramWithBinary(context, num_devices, &device_ids,
  &lengths, &binaries, &binary_status, &ret_errcode);

/* Создаем очередь команд для первого DSP-ядра */
command_queue_1 = eclCreateCommandQueueWithProperties(context, device_ids[0],
  NULL, &ret_errcode);

/* Создаем очередь команд для второго DSP-ядра */
command_queue_2 = eclCreateCommandQueueWithProperties(context, device_ids[1],
  NULL, &ret_errcode);

/* Создаём ядро выполняемое на первом DSP-ядре */
kernel_1 = eclCreateKernel(program, kernel_name_1, &ret_errcode);

/* Создаём ядро выполняемое на втором DSP-ядре */
kernel_2 = eclCreateKernel(program, kernel_name_2, &ret_errcode);

/* Ставим в очередь на выполнение ядро для первого DSP-ядра. В ret_event будет
   возвращён ecl_event, сигнализирующий о том, что данный экземпляр ядра закончил
   свою работу */
ret = eclEnqueueNDRangeKernel(command_queue_1, kernel_1, 1, NULL /*work_offset*/,
  &[1] /*global_work_size*/, NULL /*local_work_size*/, num_events_in_wait_list,
  &event_wait_list, &ret_event_1);

/* Ставим в очередь на выполнение ядро для второго DSP-ядра. При этом в
   event_wait_list передаём ранее полученный ret_event_1: это будет означать,
   что выполнение kernel_2 начнётся только после того как закончится выполнение
   kernel_1 */
num_events_in_wait_list = 1; event_wait_list[0] = ret_event_1;
ret = eclEnqueueNDRangeKernel(command_queue_2, kernel_2, 1, NULL /*work_offset*/,
  &[1] /*global_work_size*/, NULL /*local_work_size*/, num_events_in_wait_list,
  &event_wait_list, &ret_event_2);