当前位置：首页 > news >正文

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

news 来源：原创 2024/7/5 4:05:04

在使用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 boundsif (x < width && y < height) {// Calculate Y, U, and V indicesint 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 valuesuchar yValue = nv12[yIndex];uchar uValue = nv12[uvIndex];uchar vValue = nv12[uvIndex + 1];// Convert YUV to RGBuchar 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 valuesuchar bgrValue = (bValue << 2) | (gValue >> 4) | (rValue << 6);// Store BGR valuebgr[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 NEONuint8_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 pixelint 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]; // Blueuint8x16_t g = bgr_vector.val[1]; // Greenuint8x16_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);}
}