opencl色域变换，处理传递显存数据

在使用ffmpeg解码后的多路解码数据非常慢，还要给AI做行的加速方式是在显存处理数据，在视频拼接融合产品的产品与架构设计中，提出了比较可靠的方式是使用cuda，那么没有cuda的显卡如何处理呢
，比较好的方式是使用opencl来提高数据传输效率

核函数

在OpenCL中，将NV12格式转换为BGR格式通常涉及到对UV分量的处理，nv12 是使用ffmpeg等解码后的直接数据，注意linesize对齐

#define GROUP_SIZE 16
 
// OpenCL kernel to convert NV12 to BGR
__kernel void nv12_to_bgr(__global const uchar *nv12,
                          __global uchar *bgr,
                          int width, int height) {
    int x = get_global_id(0);
    int y = get_global_id(1);
 
    // Make sure we are not out of bounds
    if (x < width && y < height) {
        // Calculate Y, U, and V indices
        int yIndex = y * width + x;
        int uvIndex = width * height + (y / 2) * (width) + (x & ~1); // Use '& ~1' to get even X indices for U/V
 
        // Load Y, U, and V values
        uchar yValue = nv12[yIndex];
        uchar uValue = nv12[uvIndex];
        uchar vValue = nv12[uvIndex + 1];
 
        // Convert YUV to RGB
        uchar bValue = (uchar)((yValue                  + 1.732446 * (uValue - 128));
        uchar gValue = (uchar)((yValue - 0.344134 * (vValue - 128) - 0.714136 * (uValue - 128));
        uchar rValue = (uchar)((yValue + 1.402225 * (vValue - 128));
 
        // Pack BGR values
        uchar bgrValue = (bValue << 2) | (gValue >> 4) | (rValue << 6);
 
        // Store BGR value
        bgr[yIndex] = bgrValue;
    }
}

cpu上继续

注意错误处理

// 设置OpenCL内核参数
size_t global_work_size[2] = {width, height};
cl_kernel nv12_to_bgr_kernel = ...; // 获取你编译的内核
 
// 设置内核参数
clSetKernelArg(nv12_to_bgr_kernel, 0, sizeof(cl_mem), &nv12_buffer);
clSetKernelArg(nv12_to_bgr_kernel, 1, sizeof(cl_mem), &bgr_buffer);
clSetKernelArg(nv12_to_bgr_kernel, 2, sizeof(int), &width);
clSetKernelArg(nv12_to_bgr_kernel, 3, sizeof(int), &height);
 
// 执行内核
cl_event event;
clEnqueueNDRangeKernel(command_queue, nv12_to_bgr_kernel, 2, NULL, global_work_size, NULL, 0, NULL, &event);
 
// 等待命令执行完毕
clWaitForEvents(1, &event);

针对arm，非显存

用128位的寄存器进行处理。
vld1_u8 从内存中读取88位数据到寄存器
vld1q_u8 从内存中读取168位数据到寄存器
vld3q_u8 从内存中读取3个168位数据到寄存器中
vst3q_u8 将三个128位寄存器的数据写到内存中
vld4_u8 从内存中读取4个88位数据到寄存器中
vmull_u8 执行两个8*8位无符号整数的乘法操作
vshrn_n_u16 16位无符号整数右移指定的位数
vst1_u8 将128位寄存器中的8位无符号整数元素存储到内存中
vshrq_n_s16 16位整数右移指定的位数
举例


void bgr_to_rgb(uint8_t *bgr, uint8_t *rgb, int width, int height)
{
    // Ensure BGR and BGR buffers are 16-byte aligned for NEON
    uint8_t *bgr_aligned = (uint8_t *)(((uintptr_t)bgr + 15) & ~15);
    uint8_t *rgb_aligned = (uint8_t *)(((uintptr_t)rgb + 15) & ~15);

    for (int q = 0; q < height * width / 16; q++)
    {
        // Calculate the index for the current pixel
        int index = q * 16 * 3;

        // Load 16 BGR pixels into three vectors.
        uint8x16x3_t bgr_vector = vld3q_u8(bgr_aligned + index);

        // Shuffle the bytes to convert from BGR to BGR.
        uint8x16_t b = bgr_vector.val[2]; // Blue
        uint8x16_t g = bgr_vector.val[1]; // Green
        uint8x16_t r = bgr_vector.val[0]; // Red

        // Combine the shuffled bytes into a single vector.
        uint8x16x3_t rgb_vector = {b, g, r};

        // Store the result.
        vst3q_u8(rgb_aligned + index, rgb_vector);
    }
}