Tearing when processing image with cuda kernel, not present when processing with cpu

xlucny · January 19, 2025, 3:04pm

Hello,

I’m doing a depacking operation on raw image data with my custom cuda kernel and I’m experiencing a sort of tearing.

I have some raw image, if I’m doing the operation with CPU

 for (int i = 0; i < width*height/16; i++)
    {

        int src_i = i*20;
        int dst_i = i*16*2;
        
       //sixteen pixels of two bytes each are in twenty bytes
        dst_buffer[dst_i] = (src_buffer[src_i] & 0b00111111) << 2 | src_buffer[src_i+16] & 0b00000011 ;
        dst_buffer[dst_i+1] = (src_buffer[src_i] &  0b11000000) >> 6;

        dst_buffer[dst_i+2] = (src_buffer[src_i+1] & 0b00111111) << 2 | (src_buffer[src_i+16] & 0b00001100) >> 2;
        dst_buffer[dst_i+3] = (src_buffer[src_i+1] &  0b11000000) >> 6;

        dst_buffer[dst_i+4] = (src_buffer[src_i+2] & 0b00111111) << 2 | (src_buffer[src_i+16] & 0b00110000) >> 4;
        dst_buffer[dst_i+5] = (src_buffer[src_i+2] &  0b11000000) >> 6;

        dst_buffer[dst_i+6] = (src_buffer[src_i+3] & 0b00111111) << 2 | (src_buffer[src_i+16] & 0b11000000) >> 6;
        dst_buffer[dst_i+7] = (src_buffer[src_i+3] &  0b11000000) >> 6;

        dst_buffer[dst_i+8] = (src_buffer[src_i+4] & 0b00111111) << 2 | src_buffer[src_i+17] & 0b00000011;
        dst_buffer[dst_i+9] = (src_buffer[src_i+4] &  0b11000000) >> 6;

        dst_buffer[dst_i+10] = (src_buffer[src_i+5] & 0b00111111) << 2 | (src_buffer[src_i+17] & 0b00001100) >> 2;
        dst_buffer[dst_i+11] = (src_buffer[src_i+5] &  0b11000000) >> 6;

        dst_buffer[dst_i+12] = (src_buffer[src_i+6] & 0b00111111) << 2 | (src_buffer[src_i+17] & 0b00110000) >> 4;
        dst_buffer[dst_i+13] = (src_buffer[src_i+6] &  0b11000000) >> 6;

        dst_buffer[dst_i+14] = (src_buffer[src_i+7] & 0b00111111) << 2 | (src_buffer[src_i+17] & 0b11000000) >> 6;
        dst_buffer[dst_i+15] = (src_buffer[src_i+7] &  0b11000000) >> 6;

        dst_buffer[dst_i+16] = (src_buffer[src_i+8] & 0b00111111) << 2 | src_buffer[src_i+18] & 0b00000011;
        dst_buffer[dst_i+17] = (src_buffer[src_i+8] &  0b11000000) >> 6;

        dst_buffer[dst_i+18] = (src_buffer[src_i+9] & 0b00111111) << 2 | (src_buffer[src_i+18] & 0b00001100) >> 2;
        dst_buffer[dst_i+19] = (src_buffer[src_i+9] &  0b11000000) >> 6;

        dst_buffer[dst_i+20] = (src_buffer[src_i+10] & 0b00111111) << 2 | (src_buffer[src_i+18] & 0b00110000) << 4;
        dst_buffer[dst_i+21] = (src_buffer[src_i+10] &  0b11000000) >> 6;

        dst_buffer[dst_i+22] = (src_buffer[src_i+11] & 0b00111111) << 2 | (src_buffer[src_i+18] & 0b11000000) >> 6;
        dst_buffer[dst_i+23] = (src_buffer[src_i+11] &  0b11000000) >> 6;

        dst_buffer[dst_i+24] = (src_buffer[src_i+12] & 0b00111111) << 2 | src_buffer[src_i+19] & 0b00000011;
        dst_buffer[dst_i+25] = (src_buffer[src_i+12] &  0b11000000) >> 6;

        dst_buffer[dst_i+26] = (src_buffer[src_i+13] & 0b00111111) << 2 | (src_buffer[src_i+19] & 0b00001100) >> 2;
        dst_buffer[dst_i+27] = (src_buffer[src_i+13] &  0b11000000) >> 6;

        dst_buffer[dst_i+28] = (src_buffer[src_i+14] & 0b00111111) << 2 | (src_buffer[src_i+19] & 0b00110000) >> 4;
        dst_buffer[dst_i+29] = (src_buffer[src_i+14] &  0b11000000) >> 6;

        dst_buffer[dst_i+30] = (src_buffer[src_i+15] & 0b00111111) << 2 | (src_buffer[src_i+19] & 0b11000000) >> 6;
        dst_buffer[dst_i+31] = (src_buffer[src_i+15] &  0b11000000) >> 6;
    }

I get the following, correct output:

but If I’m doing it with the GPU:


void Depack10bitGrouping160(uint8_t* src_buffer, uint8_t* dst_buffer, int width, int height)
{
    int cycle_length = (width)*(height)/16;   
    int num_threads = 1024; 
    int num_blocks = (cycle_length/num_threads)+1; //+1 in case cycle length isn't perfectly divisible
                                                   // by num_threads,

    Depack10bitGrouping160Kernel<<<num_blocks,num_threads,0>>>(src_buffer,dst_buffer,cycle_length-1);  //index_max = cycle_length-1

    cudaDeviceSynchronize();

}


__global__ void Depack10bitGrouping160Kernel (uint8_t* src_buffer, uint8_t* dst_buffer, int index_max)
{
    int index = blockIdx.x * blockDim.x + threadIdx.x;
    if (index > index_max)
        return;

        int src_i = index*20;
        int dst_i = index*16*2;
        
        //sixteen pixels of two bytes each are in twenty bytes
        dst_buffer[dst_i] = (src_buffer[src_i] & 0b00111111) << 2 | src_buffer[src_i+16] & 0b00000011 ;
        dst_buffer[dst_i+1] = (src_buffer[src_i] &  0b11000000) >> 6;

        dst_buffer[dst_i+2] = (src_buffer[src_i+1] & 0b00111111) << 2 | (src_buffer[src_i+16] & 0b00001100) >> 2;
        dst_buffer[dst_i+3] = (src_buffer[src_i+1] &  0b11000000) >> 6;

        dst_buffer[dst_i+4] = (src_buffer[src_i+2] & 0b00111111) << 2 | (src_buffer[src_i+16] & 0b00110000) >> 4;
        dst_buffer[dst_i+5] = (src_buffer[src_i+2] &  0b11000000) >> 6;

        dst_buffer[dst_i+6] = (src_buffer[src_i+3] & 0b00111111) << 2 | (src_buffer[src_i+16] & 0b11000000) >> 6;
        dst_buffer[dst_i+7] = (src_buffer[src_i+3] &  0b11000000) >> 6;

        dst_buffer[dst_i+8] = (src_buffer[src_i+4] & 0b00111111) << 2 | src_buffer[src_i+17] & 0b00000011;
        dst_buffer[dst_i+9] = (src_buffer[src_i+4] &  0b11000000) >> 6;

        dst_buffer[dst_i+10] = (src_buffer[src_i+5] & 0b00111111) << 2 | (src_buffer[src_i+17] & 0b00001100) >> 2;
        dst_buffer[dst_i+11] = (src_buffer[src_i+5] &  0b11000000) >> 6;

        dst_buffer[dst_i+12] = (src_buffer[src_i+6] & 0b00111111) << 2 | (src_buffer[src_i+17] & 0b00110000) >> 4;
        dst_buffer[dst_i+13] = (src_buffer[src_i+6] &  0b11000000) >> 6;

        dst_buffer[dst_i+14] = (src_buffer[src_i+7] & 0b00111111) << 2 | (src_buffer[src_i+17] & 0b11000000) >> 6;
        dst_buffer[dst_i+15] = (src_buffer[src_i+7] &  0b11000000) >> 6;

        dst_buffer[dst_i+16] = (src_buffer[src_i+8] & 0b00111111) << 2 | src_buffer[src_i+18] & 0b00000011;
        dst_buffer[dst_i+17] = (src_buffer[src_i+8] &  0b11000000) >> 6;

        dst_buffer[dst_i+18] = (src_buffer[src_i+9] & 0b00111111) << 2 | (src_buffer[src_i+18] & 0b00001100) >> 2;
        dst_buffer[dst_i+19] = (src_buffer[src_i+9] &  0b11000000) >> 6;

        dst_buffer[dst_i+20] = (src_buffer[src_i+10] & 0b00111111) << 2 | (src_buffer[src_i+18] & 0b00110000) << 4;
        dst_buffer[dst_i+21] = (src_buffer[src_i+10] &  0b11000000) >> 6;

        dst_buffer[dst_i+22] = (src_buffer[src_i+11] & 0b00111111) << 2 | (src_buffer[src_i+18] & 0b11000000) >> 6;
        dst_buffer[dst_i+23] = (src_buffer[src_i+11] &  0b11000000) >> 6;

        dst_buffer[dst_i+24] = (src_buffer[src_i+12] & 0b00111111) << 2 | src_buffer[src_i+19] & 0b00000011;
        dst_buffer[dst_i+25] = (src_buffer[src_i+12] &  0b11000000) >> 6;

        dst_buffer[dst_i+26] = (src_buffer[src_i+13] & 0b00111111) << 2 | (src_buffer[src_i+19] & 0b00001100) >> 2;
        dst_buffer[dst_i+27] = (src_buffer[src_i+13] &  0b11000000) >> 6;

        dst_buffer[dst_i+28] = (src_buffer[src_i+14] & 0b00111111) << 2 | (src_buffer[src_i+19] & 0b00110000) >> 4;
        dst_buffer[dst_i+29] = (src_buffer[src_i+14] &  0b11000000) >> 6;

        dst_buffer[dst_i+30] = (src_buffer[src_i+15] & 0b00111111) << 2 | (src_buffer[src_i+19] & 0b11000000) >> 6;
        dst_buffer[dst_i+31] = (src_buffer[src_i+15] &  0b11000000) >> 6;

}

I get this tearing:

To me, this seems like some sort of synchronization, caching issue that I don’t understand. Could you guys please tell me what’s causing it?

Thank you,
xlucny

Full code:
SAMPLE_INV.zip (2.9 MB)
System:
NVIDIA Jetson Orin NX (16GB ram)
Jetpack 6.2, L4T 36.4.3
CUDA 12.6.68

AastaLLL · January 20, 2025, 4:44am

Hi,

Could you check if the launched threads/blocks have covered all the data points?

Thanks.

xlucny · January 22, 2025, 5:05pm

It is sufficient.

if I add some other commands to the kernel, the result is suddenly OK.

        dst_buffer[dst_i+30] = (src_buffer[src_i+15] & 0b00111111) << 2 | (src_buffer[src_i+19] & 0b11000000) >> 6;
        dst_buffer[dst_i+31] = (src_buffer[src_i+15] &  0b11000000) >> 6;

 +       if (((dst_i/2) /2048)>1058)
 +       {
 +          int dummy = 0;
 +           printf("A\n");
 +       }

Do you know what would cause such a behaviour?

AastaLLL · January 23, 2025, 7:28am

Hi,

We will try to reproduce this issue internally.
Will provide more info to you later.

Thanks.

AastaLLL · January 24, 2025, 5:31am

Hi,

Confirmed that we can reproduce the same issue in our environment.
We need to check this with our internal team and provide more info to you later.

Thanks.

AastaLLL · January 24, 2025, 6:27am

Hi,

Please try the below modification:

diff --git a/main.cpp b/main.cpp
index 8ca4dea..371bd5a 100644
--- a/main.cpp
+++ b/main.cpp
@@ -21,10 +21,10 @@ int main()
 
     void* depacked_buffer, *raw_buffer;
 
-    ret = cudaMallocManaged (&raw_buffer, raw_buffer_length, cudaMemAttachHost);
+    ret = cudaMallocManaged (&raw_buffer, raw_buffer_length);//, cudaMemAttachHost);
     CHECK_CUDA_RET("Allocating unified memory to raw_buffer variable")
     
-    ret = cudaMallocManaged (&depacked_buffer, width*height*2, cudaMemAttachHost);
+    ret = cudaMallocManaged (&depacked_buffer, width*height*2);// cudaMemAttachHost);
     CHECK_CUDA_RET("Allocating unified memory to depacked_buffer variable")
 
     FILE *f;

Change the cudaMallocManaged flag from cudaMemAttachHost to cudaMemAttachGlobal (the default value) can fix the issue.
You can find the corresponding doc for that flag below:
https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__MEMORY.html#group__CUDART__MEMORY_1gd228014f19cc0975ebe3e0dd2af6dd1b

Thanks.

xlucny · January 27, 2025, 9:35am

I confirm that this solves the issue.

Is this an issue with the cudaMemAttachHost flag or an expected behaviour?

AastaLLL · January 28, 2025, 3:50pm

Hi,

As you need to access the buffer with CUDA, you will need to use the cudaMemAttachGlobal to make it accessible for GPU.

Thanks.

system · February 24, 2025, 7:35am

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

Topic		Replies	Views
Processing image with a CUDA kernel gives me different result than a seemingly equivalent CPU function CUDA Programming and Performance opencv , cuda	27	1462	May 8, 2024
Fifo Reset Issue in GPU CUDA Programming and Performance	5	1446	August 3, 2015
Memory Fragmentation: 3d memory not completely freed CUDA Programming and Performance	1	978	November 29, 2014
How to share NvBufSurface with Cuda efficiently, without overhead of cuGraphicsEGLRegisterImage/cuGraphicsUnregisterResource? Jetson AGX Orin nvbugs	24	1093	May 13, 2024
Two memory questions, and a CPU question CUDA Programming and Performance	5	2529	March 24, 2008
Converting a kernel from floats and ints to halfs is 6x slower CUDA Programming and Performance cuda	14	1013	October 16, 2023
Concurrent Kernel executions & Data Transfers CUDA Programming and Performance cuda	3	608	March 8, 2023
CUDA runtime on Jetson Orin AGX Jetson AGX Orin cuda	46	6185	September 1, 2023
Unexplained stalls in CUDA API calls - reproducer attached Jetson TK1	27	2936	October 18, 2021
Gpu error on boot Jetson Nano boot	11	3299	October 15, 2021

Tearing when processing image with cuda kernel, not present when processing with cpu

Related topics