【发布时间】：2020-08-10 21:52:27
【问题描述】：
我需要将图像从 bgr 转换为 yuv420p，我首先使用 OpenCV 来做。
Mat img = imread("1.bmp");
Mat yuvImg;
cvtColor(img,yuvImg,COLOR_BGR2YUV_I420);
结果正常。但是，我的图像太大了，它的像素几乎是 6400 * 2000。 我发现使用 opencv api cvtcolor 将 bgr 转换为 yuv420p 会花费太多时间。
然后我决定自己转换它并用 cuda 加速它。
这是 cpu 中的代码：
void bgr_to_yuv420p(unsigned char* yuv420p, unsigned char* bgr, int width, int height)
{
if (yuv420p == NULL || bgr== NULL)
return;
int frameSize = width*height;
int chromaSize = frameSize / 4;
int yIndex = 0;
int uIndex = frameSize;
int vIndex = frameSize + chromaSize;
int R, G, B, Y, U, V;
for (int i = 0; i < height; i++)
{
for (int j = 0; j < width; j++)
{
B = bgr[(i * width + j) * 3 + 0];
G = bgr[(i * width + j) * 3 + 1];
R = bgr[(i * width + j) * 3 + 2];
//BGR to YUV
Y = ((66 * R + 129 * G + 25 * B + 128) >> 8) + 16;
U = ((-38 * R - 74 * G + 112 * B + 128) >> 8) + 128;
V = ((112 * R - 94 * G - 18 * B + 128) >> 8) + 128;
yuv420p[yIndex++] = (unsigned char)((Y < 0) ? 0 : ((Y > 255) ? 255 : Y));
if (i % 2 == 0 && j % 2 == 0)
{
yuv420p[uIndex++] = (unsigned char)((U < 0) ? 0 : ((U > 255) ? 255 : U));
yuv420p[vIndex++] = (unsigned char)((V < 0) ? 0 : ((V > 255) ? 255 : V));
}
}
}
}
我测试了代码bgr_to_yuv420p(...)，结果也正常。
然后我用 cuda 加速它。
这是我所有的代码，包括内核函数和测试函数。
#include <iostream>
#include <time.h>
#include <vector_types.h>
#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include "opencv2/highgui.hpp"
#include "opencv2/opencv.hpp"
using namespace cv;
using namespace std;
//kernel function to convert bgr to yuv420p
__global__ void bgr2yuv420p(uchar3 * d_in, unsigned char * d_out,
uint imgheight, uint imgwidth)
{
int col_num = blockIdx.x*blockDim.x+threadIdx.x;
int row_num = blockIdx.y*blockDim.y+threadIdx.y;
if ((row_num < imgheight) && (col_num < imgwidth))
{
// uint32_t a = *((uint32_t *)&dinput[global_offset*3]);
int global_offset = row_num*imgwidth+col_num;
int r,g,b;
r = int(d_in[global_offset].z);
g = int (d_in[global_offset].y);
b = int (d_in[global_offset].x);
d_out[row_num * imgwidth + col_num] = ((66*r + 129*g + 25*b) >> 8) + 16;
if(((threadIdx.x & 1) == 0) && ((threadIdx.y & 1) == 0)){
int uv_offset = imgwidth*imgheight+((row_num*imgwidth))+col_num;
d_out[uv_offset] = ((112*r + -94*g + -18*b) >> 8) + 128;
d_out[uv_offset+1] = ((-38*r + -74*g + 112*b) >> 8) + 128;
}
}
}
int main(void)
{
Mat srcImage = imread("1.bmp");
imshow("srcImage", srcImage);
const uint imgheight = srcImage.rows;
const uint imgwidth = srcImage.cols;
Mat nv12Image(imgheight * 3 / 2, imgwidth, CV_8UC1, Scalar(255));
//input and output
uchar3 *d_in;
unsigned char *d_out;
// malloc memo in gpu
cudaMalloc((void**)&d_in, imgheight*imgwidth*sizeof(uchar3));
cudaMalloc((void**)&d_out, imgheight*imgwidth*sizeof(unsigned char) * 3 / 2);
//copy image from cpu to gpu
cudaMemcpy(d_in, srcImage.data, imgheight*imgwidth*sizeof(uchar3), cudaMemcpyHostToDevice);
dim3 threadsPerBlock(32, 32);
dim3 blocksPerGrid((imgwidth + threadsPerBlock.x - 1) / threadsPerBlock.x,
(imgheight + threadsPerBlock.y - 1) / threadsPerBlock.y);
//run kernel function
bgr2yuv420p<<<blocksPerGrid, threadsPerBlock>>>(d_in, d_out, imgheight, imgwidth);
cudaDeviceSynchronize();
//copy yuv420p from gpu to cpu
cudaMemcpy(nv12Image.data, d_out, imgheight*imgwidth*sizeof(unsigned char) * 3 / 2, cudaMemcpyDeviceToHost);
imshow("nv12",nv12Image);
imwrite("cuda.bmp",nv12Image);
cudaFree(d_in);
cudaFree(d_out);
return 0;
}
带有cuda的代码可以运行但结果不正常。 YUV420p的Y是正常的，但是U和V有问题，我想原因在__global__ void bgr2yuv420p(...)
if(((threadIdx.x & 1) == 0) && ((threadIdx.y & 1) == 0)){
int uv_offset = imgwidth*imgheight+((row_num*imgwidth))+col_num;
d_out[uv_offset] = ((112*r + -94*g + -18*b) >> 8) + 128;
d_out[uv_offset+1] = ((-38*r + -74*g + 112*b) >> 8) + 128;
}
我尝试了很多，但仍然无法解决。我发现很少有关于将 rgb 转换为 yuv420p 的代码，更多的代码是关于将 yuv420p 转换为 rgb 的。所以我想知道是否有人遇到同样的问题或给我一些建议？
感谢 Robert Crovella。这是我的 update-1。
我听从 Robert Crovella 的建议，像这样更改内核函数：
//kernel function to convert bgr to yuv420p
__global__ void bgr2yuv420p(uchar3 * d_in, unsigned char * d_out,
uint imgheight, uint imgwidth)
{
int col_num = blockIdx.x*blockDim.x+threadIdx.x;
int row_num = blockIdx.y*blockDim.y+threadIdx.y;
if ((row_num < imgheight) && (col_num < imgwidth))
{
// uint32_t a = *((uint32_t *)&dinput[global_offset*3]);
int global_offset = row_num*imgwidth+col_num;
int r,g,b;
r = int(d_in[global_offset].z);
g = int (d_in[global_offset].y);
b = int (d_in[global_offset].x);
d_out[row_num * imgwidth + col_num] = ((66*r + 129*g + 25*b) >> 8) + 16;
if(((threadIdx.x & 1) == 0) && ((threadIdx.y & 1) == 0)){
int uv_offset = imgwidth*imgheight+((row_num>>1)*imgwidth)+col_num;
d_out[uv_offset] = ((112*r + -94*g + -18*b) >> 8) + 128;
d_out[uv_offset+1] = ((-38*r + -74*g + 112*b) >> 8) + 128;
}
}
}
我兴奋地测试了新内核，但结果也不正常。 这是我的更新内核函数的结果图像。 yuv420p image converted by myself
那么opencv api转换的正常结果图就到这里了。 yuv420p image converted by opencv api
我们可以看到，两个图像的区别是U和V。我已经改变了核函数中U和V的索引，即
if(((threadIdx.x & 1) == 0) && ((threadIdx.y & 1) == 0)){
int uv_offset = imgwidth*imgheight+((row_num >>1)*imgwidth)+col_num;
d_out[uv_offset] = ((112*r + -94*g + -18*b) >> 8) + 128;
d_out[uv_offset+1] = ((-38*r + -74*g + 112*b) >> 8) + 128;
}
我认为它会起作用，但它不会。还有什么建议吗？罗伯特·克罗维拉
编辑：解决方案是 Robert Crovella 的最新答案。我仔细检查了它，它真的很完美。
【问题讨论】：
-
你的内核定义是
bgr2yuv420p，但你的内核启动是rgb2yuv420p，所以你的代码不会编译。我猜这实际上不是您正在运行的代码。在您的内核中，您对uv_offset的计算会在后续行中生成非法的越界访问。您可能希望研究该索引并使用here 描述的方法来帮助调试。 -
仔细考虑您的
uv_offset计算。imgwidth*imgheight偏移让您越过 Y 区域（正确），但从那时起，使用row_num*imgwidth按行索引到 UV 平面区域是否正确？ （提示，它不是。你在 UV 平面区域没有那么多行，你只有一半的行）。 -
我认为如果您进行此更改，您将更接近正确的东西：
int uv_offset = imgwidth*imgheight+(((row_num>>1)*imgwidth))+col_num; -
比较你的
bgr_to_yuv420pCPU 代码和你的bgr2yuv420pGPU 内核，我们还看到U 和V 存储的顺序是相反的，并且还有一些其他的计算差异。如果我把它们全部整理出来，我可以得到它们之间的匹配结果。请注意，您的bgr_to_yuv420pCPU 代码意味着 Y、U、V 的平面存储，而您的 GPU 代码为 Y 平面和 UV 交错平面提供半平面存储。 -
对于内核函数 rgb2yuv420p() 和 bgr2yuv420p() 的低级错误，我们深表歉意。我只是在运行内核函数时输入错误。事实上，我定义了 bgr2yuv420p() 并且还运行了 bgr2yuv420p()。我会仔细检查并编辑我的问题。