Media Pipeline Inter-operation and Memory Sharing
Media engine capabilities exposed in low-level OS-specific interfaces, such as
VA-API (Video Acceleration API) for Linux OS
Microsoft DirectX® Video Acceleration for Windows OS
as well as various high-level media frameworks built on top of low-level interfaces, such as
oneVPL
FFmpeg and libav
GStreamer
Each media framework defines own interfaces for device and context creation, memory allocation and task submission. Most frameworks also expose export/import interfaces to convert memory objects to/from other memory handles
high-level media frameworks (FFMpeg, GStreamer) support conversion to/from low-level media handles (VA-API and DirectX surfaces)
low-level media interfaces (VA-API, DirectX) support conversion to/ OS-specific general-purpose GPU memory handles such as DMA buffers on Linux and NT handles on Windows
Level-zero support conversion between DMA buffers / NT handles and USM device pointers
Together these interfaces allow zero-copy memory sharing between media operations submitted via media frameworks and SYCL compute kernels submitted into SYCL queue, assuming the SYCL queue created on same GPU device as media framework and SYCL device uses Level-zero backend (not OpenCL backend).
Despite multiple stages of memory handles conversion (FFmpeg/GStreamer, VA-API/DirectX, DMA/NT, Level-Zero, SYCL), all converted memory handles refer to same physical memory block. Thus writing data into one memory handle makes the data available in all other memory handles, assuming proper synchronization between write and read operations.
Below is reference to interfaces used for zero-copy buffer sharing between media frameworks and SYCL
(Linux) VA-API to DMA-BUF
(Windows) DirectX to NT-Handle
The memory pointer created by Level-zero from DMA-BUF or NT-Handle (#3 above) is USM device pointer only accessible by SYCL kernels running on same GPU device as used for media memory allocation and media operations. This USM pointer is not accessible from host and not accessible from SYCL kernels running on CPU or other XPU devices.
Example in next section demonstrates zero-copy buffer sharing between VA-API and SYCL using interfaces 1 and 3 from list above and synthetic video data (moving rectangle). For more advanced examples with FFmpeg/GStreamer video decode/encode on GPU media engine and SYCL kernels on GPU compute engines please refer to Intel® DL Streamer memory interoperability API (preview) and Intel® DL Streamer samples
VA-API and SYCL memory sharing example
The example
allocates shared VA-API surfaces and USM device pointers for NUM_FRAMES frames
submits VA-API calls to draw moving rectangle on frames
submits SYCL kernels to draw sub-rectangle inside rectangle created by VA-API on step 2
synchronize all frames and write RGB data into file
Output frames generated by this example look like picture below
The example supports Linux OS and requires installation of the following additional packages besides oneAPI packages (installation example via apt package manager on Ubuntu OS)
sudo apt install intel-level-zero-gpu level-zero-dev sudo apt install intel-media-va-driver-non-free libva-dev libva-drm2
and requires linkage with Level-zero and VA-API libraries
icpx -fsycl memory-sharing-with-media.cpp -lze_loader -lva -lva-drm
Example execution generates file output.bgra which could be directly played by some media players (ex, ffplay) or transcoded to compressed video format, for example using the following ffmpeg command:
ffmpeg -f rawvideo -pix_fmt bgra -s 320x240 -i output.bgra output.mp4
and then played by any media player, for example
ffplay output.mp4
//==============================================================
// Copyright © 2022 Intel Corporation
//
// SPDX-License-Identifier: MIT
// =============================================================
// SYCL
#include <CL/sycl.hpp>
// SYCL oneAPI extension
#include <sycl/ext/oneapi/backend/level_zero.hpp>
// Level-zero
#include <level_zero/ze_api.h>
// VA-API
#include <va/va_drm.h>
#include <va/va_drmcommon.h>
#include <cstdio>
#include <fcntl.h>
#include <unistd.h>
#include <vector>
#define OUTPUT_FILE "output.bgra"
#define VAAPI_DEVICE "/dev/dri/renderD128"
#define FRAME_WIDTH 320
#define FRAME_HEIGHT 240
#define RECT_WIDTH 160
#define RECT_HEIGHT 160
#define RECT_Y (FRAME_HEIGHT - RECT_HEIGHT) / 2
#define NUM_FRAMES (FRAME_WIDTH - RECT_WIDTH)
#define VA_FORMAT VA_FOURCC_BGRA
#define RED 0xffff0000
#define GREEN 0xff00ff00
#define BLUE 0xff0000ff
#define CHECK_STS(_FUNC) \
{ \
auto _sts = _FUNC; \
if (_sts != 0) { \
printf("Error %d calling " #_FUNC, (int)_sts); \
return -1; \
} \
}
VASurfaceID alloc_va_surface(VADisplay va_display, int width, int height) {
VASurfaceID va_surface;
VASurfaceAttrib surface_attrib{};
surface_attrib.type = VASurfaceAttribPixelFormat;
surface_attrib.flags = VA_SURFACE_ATTRIB_SETTABLE;
surface_attrib.value.type = VAGenericValueTypeInteger;
surface_attrib.value.value.i = VA_FORMAT;
vaCreateSurfaces(va_display, VA_RT_FORMAT_RGB32, width, height, &va_surface,
1, &surface_attrib, 1);
return va_surface;
}
int main() {
// Create SYCL queue on GPU device and Level-zero backend, and query
// Level-zero context and device
sycl::queue sycl_queue{sycl::ext::oneapi::filter_selector(
"level_zero")}; // { sycl::gpu_selector() }
auto ext_level_zero = sycl::backend::ext_oneapi_level_zero;
auto ze_context = sycl::get_native<ext_level_zero>(sycl_queue.get_context());
auto ze_device = sycl::get_native<ext_level_zero>(sycl_queue.get_device());
// Create VA-API context (VADisplay)
VADisplay va_display = vaGetDisplayDRM(open(VAAPI_DEVICE, O_RDWR));
if (!va_display) {
printf("Error creating VADisplay on device %s\n", VAAPI_DEVICE);
return -1;
}
int major = 0, minor = 0;
CHECK_STS(vaInitialize(va_display, &major, &minor));
// Create VA-API surfaces
VASurfaceID surfaces[NUM_FRAMES];
for (int i = 0; i < NUM_FRAMES; i++) {
surfaces[i] = alloc_va_surface(va_display, FRAME_WIDTH, FRAME_HEIGHT);
}
// Convert each VA-API surface into USM device pointer (zero-copy buffer
// sharing between VA-API and Level-zero)
void *device_ptr[NUM_FRAMES];
size_t stride;
for (int i = 0; i < NUM_FRAMES; i++) {
// Export DMA-FD from VASurface
VADRMPRIMESurfaceDescriptor prime_desc{};
CHECK_STS(vaExportSurfaceHandle(va_display, surfaces[i],
VA_SURFACE_ATTRIB_MEM_TYPE_DRM_PRIME_2,
VA_EXPORT_SURFACE_READ_WRITE, &prime_desc));
auto dma_fd = prime_desc.objects->fd;
auto dma_size = prime_desc.objects->size;
stride = prime_desc.layers[0].pitch[0] / sizeof(uint32_t);
// Import DMA-FD into Level-zero device pointer
ze_external_memory_import_fd_t import_fd = {
ZE_STRUCTURE_TYPE_EXTERNAL_MEMORY_IMPORT_FD,
nullptr, // pNext
ZE_EXTERNAL_MEMORY_TYPE_FLAG_DMA_BUF, dma_fd};
ze_device_mem_alloc_desc_t alloc_desc = {};
alloc_desc.stype = ZE_STRUCTURE_TYPE_DEVICE_MEM_ALLOC_DESC;
alloc_desc.pNext = &import_fd;
CHECK_STS(zeMemAllocDevice(ze_context, &alloc_desc, dma_size, 1, ze_device,
&device_ptr[i]));
// Close DMA-FD
close(dma_fd);
}
// Create VA-API surface with size 1x1 and write GREEN pixel
VASurfaceID surface1x1 = alloc_va_surface(va_display, 1, 1);
VAImage va_image;
void *data = nullptr;
CHECK_STS(vaDeriveImage(va_display, surface1x1, &va_image));
CHECK_STS(vaMapBuffer(va_display, va_image.buf, &data));
*(uint32_t *)data = GREEN;
CHECK_STS(vaUnmapBuffer(va_display, va_image.buf));
CHECK_STS(vaDestroyImage(va_display, va_image.image_id));
// VA-API call to fill background with BLUE color and upscale 1x1 surface into
// moving GREEN rectangle
VAConfigID va_config_id;
VAContextID va_context_id;
CHECK_STS(vaCreateConfig(va_display, VAProfileNone, VAEntrypointVideoProc,
nullptr, 0, &va_config_id));
CHECK_STS(vaCreateContext(va_display, va_config_id, 0, 0, VA_PROGRESSIVE,
nullptr, 0, &va_context_id));
for (int i = 0; i < NUM_FRAMES; i++) {
VAProcPipelineParameterBuffer param{};
param.output_background_color = BLUE;
param.surface = surface1x1;
VARectangle output_region = {int16_t(i), RECT_Y, RECT_WIDTH, RECT_HEIGHT};
param.output_region = &output_region;
VABufferID param_buf;
CHECK_STS(vaCreateBuffer(va_display, va_context_id,
VAProcPipelineParameterBufferType, sizeof(param),
1, ¶m, ¶m_buf));
CHECK_STS(vaBeginPicture(va_display, va_context_id, surfaces[i]));
CHECK_STS(vaRenderPicture(va_display, va_context_id, ¶m_buf, 1));
CHECK_STS(vaEndPicture(va_display, va_context_id));
CHECK_STS(vaDestroyBuffer(va_display, param_buf));
}
#if 0
// Synchronization is optional on Linux OS as i915 KMD driver synchronizes
// write/read commands submitted from Intel media and compute drivers
for (int i = 0; i < NUM_FRAMES; i++) {
CHECK_STS(vaSyncSurface(va_display, surfaces[i]));
}
#endif
// Submit SYCL kernels to write RED sub-rectangle inside GREEN rectangle
std::vector<sycl::event> sycl_events(NUM_FRAMES);
for (int i = 0; i < NUM_FRAMES; i++) {
uint32_t *ptr = (uint32_t *)device_ptr[i] +
(RECT_Y + RECT_HEIGHT / 4) * stride + (i + RECT_WIDTH / 4);
sycl_events[i] = sycl_queue.parallel_for(
sycl::range<2>(RECT_HEIGHT / 2, RECT_WIDTH / 2), [=](sycl::id<2> idx) {
auto y = idx.get(0);
auto x = idx.get(1);
ptr[y * stride + x] = RED;
});
}
// Synchronize all SYCL kernels
sycl::event::wait(sycl_events);
// Map VA-API surface to system memory and write to file
FILE *file = fopen(OUTPUT_FILE, "wb");
if (!file) {
printf("Error creating file %s\n", OUTPUT_FILE);
return -1;
}
for (int i = 0; i < NUM_FRAMES; i++) {
CHECK_STS(vaDeriveImage(va_display, surfaces[i], &va_image));
CHECK_STS(vaMapBuffer(va_display, va_image.buf, &data));
fwrite(data, 1, FRAME_HEIGHT * FRAME_WIDTH * 4, file);
CHECK_STS(vaUnmapBuffer(va_display, va_image.buf));
CHECK_STS(vaDestroyImage(va_display, va_image.image_id));
}
fclose(file);
printf("Created file %s\n", OUTPUT_FILE);
// Free device pointers and VA-API surfaces
for (int i = 0; i < NUM_FRAMES; i++)
zeMemFree(ze_context, device_ptr[i]);
vaDestroySurfaces(va_display, surfaces, NUM_FRAMES);
return 0;
}