Cuda Kernel CU_EGL_COLOR_FORMAT_Y10V10U10_420_SEMIPLANAR

I need some help with Cuda Kernels when eglFrame.eglColorFormat == CU_EGL_COLOR_FORMAT_Y10V10U10_420_SEMIPLANAR
I am getting unexpected results.

To help explain, I am attaching screenshots comparing results when I use P010_10LE frame vs. NV12 frame.

I am using the following functions to get yElement[0], uvElement[0], and uvElement[1]:

device char *
get_pixel_uv(char *pDevPtr, const int pitch, const int row, const int col, const int width, const int height) {
if (col >= 0 && row >= 0 && col < width && row < height) {
return pDevPtr + row * pitch + col * 2;
} else {
return NULL;
}
}

device char *
get_pixel_y(char *pDevPtr, const int pitch, const int row, const int col, const int width, const int height) {
if (col >= 0 && row >= 0 && col < width && row < height) {
return pDevPtr + row * pitch + col;
} else {
return NULL;
}
}

global void addLabelsKernel(char* pDevYPtr, char* pDevUVPtr, const int pitch, const int width, const int height)
{
int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;

char * yElement = get_pixel_y(pDevYPtr, pitch, row, col, width, height);
if (NULL == yElement) {
return;
}
char * uvElement = get_pixel_uv(pDevUVPtr, pitch, row/2, col/2, width, height);
if (NULL == uvElement) {
return;
}
/* THESE ARE THE THREE TESTS */
yElement[0] -= 800.0f;
uvElement[0] = 0.0f;
uvElement[1] = 0.0f;

return;
}

extern ā€œCā€ int addLabels(CUdeviceptr pDevYPtr, CUdeviceptr pDevUVPtr, int pitch, int width, int height)
{
dim3 threadsPerBlock(BOX_W, BOX_H);
dim3 blocks(((width + BOX_W-1) / BOX_W), (height + BOX_H-1) / BOX_H);
addLabelsKernel<<<blocks,threadsPerBlock>>>((char*)pDevYPtr, (char*)pDevUVPtr, pitch, width, height);
return 0;
}

The three tests I run simply change values of each yElement[0], uvElement[0], and uvElement[1] and compare using formats NV12 and P010_10LE


NV12, yElement[0] -=800


P010_10LE,yElement[0] -=800


NV12,uvElement[0]=0


P010_10LE,uvElement[0]=0


NV12,uvElement[1]=0


P010_10LE,uvElement[1]=0

and lastly, NV12 uvElement[0]=0 and uvElement[1]=0

Can somebody help me understand what is going on?

Hi,

We need to check with our internal team and update more info with you.
Thanks.

1 Like

Hi,
The data layout is like

NV12: two planes. One Y plane and one UV interleaved plane. Each pixel is 1 byte, 8-bit valid data.
For one 1920x1080 frame, Y plane is 1920x1080 bytes, and UV plane is 1920x1080x0.5 bytes

P010_10LE: two planes. One Y plane and one UV interleaved plane. Each pixel is 2 bytes, 10-bit valid data and 6-bit padding.
For one 1920x1080 frame, Y plane is 1920x1080x2 bytes, and UV plane is 1920x1080x2x0.5 bytes

1 Like

This topic was automatically closed 14 days after the last reply. New replies are no longer allowed.