cuda视频拼接

什么叫YUV

YUV的讲解网上一堆，给个链接：https://mp.weixin.qq.com/s/KKfkS5QpwPAdYcEwFAN9VA

这里着重CUDA解码后的数据是NV12（是YUV420SP的一种）

问题1：YUV420在内存中占几个字节？

1个像素点由RGB组成，R、G、B各占用一个字节，故1个像素点是3个字节.

YUV420对RGB压缩率一半，故是1.5个字节

问题2：如何在内存中表示YUV420数据呢?

这里要说到内存对接的概念，为了提升计算效率，需要作内存对齐，由于计算机的位数是8的整数倍，所以内存对接以8的整数倍为基准。这里对齐(align)取值512。

以1920x1080为例：由于内存要对齐，所以要重新计算宽度(width + align - 1) / align * align = (1920 + 512 - 1) /512 * 512 = 2048，2048这个值叫linesize(ffmpeg的叫法)或者pitch(GPU的叫法)。

然后计算下存储1920x1080需要的内存大小是 linesize * height *3 / 2 = 2048 * 1080 * 3 / 2 = 3317760

在内存中NV12存放顺序是先Y后UV或者VU

由于Y未被压缩，故Y的内存空间是大小是 linesize * height = 2048 * 1080 = 2211840，起始地址是内存首地址

UV被压缩了一半，故UV的内存空间大小是 linesize * height / 2 = 2048 * 1080 / 2 = 1105920，起始地址是内存首地址偏移Y的内存空间

问题3：宽和间距啥关系？内存怎么存储YUV？

用图表示：所以实际内存中可能有部分空间是不存储数据的

 

CUDA核函数

废话不多说，讲解链接：https://cloud.tencent.com/developer/article/1656068

以1920x1080 4分屏拼接为例子：

• 分屏坐标关系图

 

• 由于我们使用CUDA一个block最大线程数1024个，故设定block(32,32)，这里为啥使用二维不使用一维，因为在6分屏(1大5小)的这个1大分屏是1280x720 需设置kernel_video_mosaic<<<720, 1280>>>明显1280大于1024，CUDA核函数无法执行.

• 在核函数中，x方向对应宽度 y方向对应高度。

• 4分屏左上角视频的宽高是960x540，间距是1024。因为是将960x540这个YUV数据内存拷到1920x1080这个YUV内存空间里，gridx = (((width - (block.x-1))/block.x) + 1) + 1 = (((960 - (32-1))/32) + 1) + 1 = 31;gridy = (((height - (block.y-1))/block.y) + 1) + 1 = (((540 - (32-1))/32) + 1) + 1 = 17;grid(gridx, gridy);故设置kernel_video_mosaic<<< grid, block >>>

• 计算x方向的线程编号公式tidx = blockIdx.x * blockDim.x + threadIdx.x，blockDim.x=gridx，最大值 31 * 31 + 31 = 992

• 计算y方向的线程编号公式tidy = blockIdx.y * blockDim.y + threadIdx.y，blockDim.y=gridy，最大值 31 * 17 + 31 = 558

• 根据内存存储的YUV关系，4分屏左上角，1920x1080中(0, 0)=0 (1, 0)=2048 ... (960,540)=960x2048+540 和 960x540(0,0)=0 (1,0)=1024...(960,540)=960x1024+540 坐标点一一对应。四分屏左上角坐标x=0 y =0 左下角坐标x=0 y=540 右上角坐标x=960 y=0 右下角坐标x=960 y=540

• 换算关系见如下代码：

//CUDA核函数
__global__ void kernel_video_mosaic(unsigned char* dst_y, unsigned char* dst_uv, int dst_linesize
    , unsigned char* src_y, unsigned char* src_uv, int src_linesize, int src_width, int src_height, int x, int y)
{
    //线程编号转换成数据对应编号
    const int tidx = blockIdx.x * blockDim.x + threadIdx.x;
    const int tidy = blockIdx.y * blockDim.y + threadIdx.y;

    //printf("thread_id (%d,%d) block_id (%d,%d) coordinate (%d, %d)\n", idx, threadIdx.y, idy, blockIdx.y, idx, idy);

    if(tidy < src_height && tidx < src_width){//需要tidy的最大值大于等于高度, tidx的最大值大于等于宽度
        int dst_index_y = (tidy + y) * dst_linesize + tidx + x;
        int src_index_y = tidy * src_linesize + tidx;
        dst_y[dst_index_y] = src_y[src_index_y];
        if(tidy < (src_height / 2)){
            int dst_index_uv = (tidy + y / 2) * dst_linesize + tidx + x;
            int src_index_uv = tidy * src_linesize + tidx;
            dst_uv[dst_index_uv] = src_uv[src_index_uv];
        }
    }
}

int sc_video_mosaic(int deviceid, ScVMosicDev *ptdev_dst, ScVMosicDev *ptdev_src, int x, int y)
{
    if(NULL == ptdev_dst || NULL == ptdev_src){
        return -1;
    }

    cudaSetDevice(deviceid);

    dim3 block(32,32);//一个block用1024个线程
    int gridx = (((ptdev_src->m_width - (block.x-1))/block.x) + 1) + 1;
    int gridy = (((ptdev_src->m_height - (block.y-1))/block.y) + 1) + 1;
    dim3 grid(gridx, gridy);

    unsigned char* dst_y = ptdev_dst->m_dev;
    unsigned char* dst_uv = ptdev_dst->m_dev + (ptdev_dst->m_linesize * ptdev_dst->m_height);
    unsigned char* src_y = ptdev_src->m_dev;
    unsigned char* src_uv = ptdev_src->m_dev + (ptdev_src->m_linesize * ptdev_src->m_height);

        //, ptdev_src->m_width, ptdev_src->m_height, ptdev_src->m_linesize, ptdev_dst->m_width, ptdev_dst->m_height, ptdev_dst->m_linesize);
    
    kernel_video_mosaic<<<grid, block>>>(dst_y, dst_uv, ptdev_dst->m_linesize
        , src_y, src_uv, ptdev_src->m_linesize, ptdev_src->m_width, ptdev_src->m_height, x, y);
    
    return 0;
}