Black Lines on CUeglFrame Luma Surface

[wilkinsw]

I am pulling frames from a cuEGLStreamConsumer and creating surfaces from its luminance and chrominance plane. I then pass these through a cudaKernel to convert the NV12 to RGB.

Here’s the code:

CUgraphicsResource cudaResource = 0;
        CUstream cudaStream = 0;
        CUresult cuResult;

        cuResult = cuEGLStreamConsumerAcquireFrame(&cudaConnection, &cudaResource, &cudaStream, -1);
        if(cuResult != CUDA_SUCCESS) break;

        // Get the CUDA EGL frame.
        CUeglFrame cudaEGLFrame;
        cuResult = cuGraphicsResourceGetMappedEglFrame(&cudaEGLFrame, cudaResource, 0, 0);
        if(cuResult != CUDA_SUCCESS)
        {
            ORIGINATE_ERROR("Unable to get the CUDA EGL frame (CUresult %s).",
                getCudaErrorString(cuResult));
        }y

        CUDA_RESOURCE_DESC cudaResourceDesc;
        CUsurfObject cudaSurfLuma = 0;
        CUsurfObject cudaSurfChroma = 0;

        memset(&cudaResourceDesc, 0, sizeof(cudaResourceDesc));
        cudaResourceDesc.resType = CU_RESOURCE_TYPE_ARRAY;

        // Create a surface from the luminance plane
        cudaResourceDesc.res.array.hArray = cudaEGLFrame.frame.pArray[0];
        cuResult = cuSurfObjectCreate(&cudaSurfLuma, &cudaResourceDesc);
        if(cuResult != CUDA_SUCCESS)
        {
            ORIGINATE_ERROR("Unable to create the luma surface object (CUresult %s)",
                getCudaErrorString(cuResult));
        }

        // Create a surface from the chrominance plane
        cudaResourceDesc.res.array.hArray = cudaEGLFrame.frame.pArray[1];
        cuResult = cuSurfObjectCreate(&cudaSurfChroma, &cudaResourceDesc);
        if(cuResult != CUDA_SUCCESS)
        {
            ORIGINATE_ERROR("Unable to create the chroma surface object (CUresult %s)",
                getCudaErrorString(cuResult));
        }

        pthread_mutex_lock(ctx->frame_mutex);

        cudaConvertNV12(
            cudaSurfLuma,
            cudaSurfChroma,
            ctx->rgb_cuda,
            ctx->rgba_cuda,
            ctx->width,
            ctx->height);

        *ctx->new_frame = true;
        pthread_mutex_unlock(ctx->frame_mutex);

        cuResult = cuSurfObjectDestroy(cudaSurfLuma);
        if(cuResult != CUDA_SUCCESS)
        {
            ORIGINATE_ERROR("Unable to destroy the luma surface object (CUresult %s)",
                getCudaErrorString(cuResult));
        }

        cuResult = cuSurfObjectDestroy(cudaSurfChroma);
        if(cuResult != CUDA_SUCCESS)
        {
            ORIGINATE_ERROR("Unable to destroy the chroma surface object (CUresult %s)",
                getCudaErrorString(cuResult));
        }

        cuResult = cuEGLStreamConsumerReleaseFrame(&cudaConnection, cudaResource, &cudaStream);
        if (cuResult != CUDA_SUCCESS)
        {
            ORIGINATE_ERROR("Unable to release the last frame (CUresult %s).",
                getCudaErrorString(cuResult));
        }

I have noticed I get black horizontal streaks in my image after conversion. At first I thought it was my kernel, but I simplified my kernel down to creating the greyscale image just from the values of the lumanance plane and I still see the streaks.

Here’s my kernel:

__global__ void ConvertNV12(
    cudaSurfaceObject_t surfLuma,
    cudaSurfaceObject_t surfChroma,
    uchar3* RGB,
    float4* RGBAf,
    uint32_t width,
    uint32_t height)
{
    const uint32_t x = blockIdx.x * blockDim.x + threadIdx.x;
    const uint32_t y = blockIdx.y * blockDim.y + threadIdx.y;

    if(x >= width || y >= height) return;

    uint8_t Y = surf2Dread<uint8_t>(surfLuma, x, y);

    uint8_t red = Y;
    uint8_t green = Y;
    uint8_t blue = Y;

    RGB[y * width + x] = make_uchar3(red, green, blue);
    RGBAf[y * width + x] = make_float4(red, green, blue, 1.0f);
}

cudaError_t cudaConvertNV12(
    cudaSurfaceObject_t surfLuma,
    cudaSurfaceObject_t surfChroma,
    uchar3* RGB,
    float4* RGBAf,
    uint32_t width,
    uint32_t height )
{
    if(!RGB || !RGBAf)
        return cudaErrorInvalidDevicePointer;

    if(width == 0 || height == 0)
        return cudaErrorInvalidValue;

    const dim3 block(8, 8, 1);
    const dim3 grid(width/block.x+1, height/block.y+1, 1);

    ConvertNV12<<<grid, block>>>(surfLuma, surfChroma, RGB, RGBAf, width, height );
	
    return CUDA(cudaGetLastError());
}

Has anyone seen this before? Is it something I’m doing wrong? I don’t think it’s my sensor, as streaming with nvcamerasrc does not produce the lines.

External Media

[AastaLLL]

Hi,

Guess the issue is pitch format related.

Following is some information for your reference:
Memory format:
[url]https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__MEMORY.html#group__CUDART__MEMORY_1g32bd7a39135594788a542ae72217775c[/url]

Sample for NV12->ARGB:
[url]https://github.com/dusty-nv/jetson-inference/blob/master/util/cuda/cudaYUV-NV12.cu[/url]

Thanks

[wilkinsw]

Yes, I was using dusty’s NV12->ARGB kernel modified for the surfaces. Colors and everything were fine, I just got those streaks, which is why I simplified the kernel. I wanted to see if it was the image data itself or some problem with the processing.

I was under the impression that using surfaces you did not have to worry about pitch, that surf2Dread takes care of that already.

On the output side I’m writing into memory allocated with cudaAllocMapped which should be contiguous too.

Something else strange is the lines seem to be more pronounced over dark surfaces and don’t show up at all on white surfaces, such as the reflection near the bottom of the image.

Here’s the result from printCUDAEGLFrame showing that the pitch is zero

CUeglFrame:
 width: 1280
 height: 580
 depth: 0
 pitch: 0
 planeCount: 2
 numChannels: 1
 frameType: array
 colorFormat: YUV420 semi-planar
 cuFormat: uint8

[wilkinsw]

Color bars don’t seem to have the same problem
External Media

[wilkinsw]

Ok this is weird. If I go to color bars and then back to real image the lines go away! Maybe it is something with the sensor. Any ideas?

Completely repeatable. There are lines like in the image of post 1, I enable/disable colorbars, and now it looks fine. Works with the solid black test pattern too.

External Media

[AastaLLL]

Hi,

Thanks for sharing your experiment.

We want to reproduce this issue internally and find the corresponding team for suggestion.
Could you share a complete sample and the reproduce steps with us?

Thanks.

[wilkinsw]

Unfortunately the issue does not appear with the devboard camera, only our dual LI-M021C-MIPI cameras. We’d be happy to mail in the hardware and documentation for your team to take a look, but I expect that might be out of the scope of your support.

I’m happy with my hack solution for now

[wuyujie840]

Can you tell me the kernal using the luminance and chrominance plane?
Thank you very much!

[wuyujie840]

Hi,Atrer
Can you tell me the kernal using the luminance and chrominance plane?
Thank you very much!

[wilkinsw]

The code is in the first post.