# Copy cv::cuda::GpuMat in Cuda Kernel

**URL:** <https://forum.opencv.org/t/copy-cv-gpumat-in-cuda-kernel/10436>\
**Category:** C++\
**Tags:** cuda\
**Created:** [October 4, 2022, 10:59pm UTC](https://forum.opencv.org/t/copy-cv-gpumat-in-cuda-kernel/10436 "2022-10-04T22:59:04Z")\
**Posts on this page:** 5\
**Page:** 1

<div class="post-metadata">

**Author:** ![Aidan\_Vickars](https://sea2.discourse-cdn.com/flex020/user_avatar/forum.opencv.org/aidan_vickars/32/6205_2.png) [@Aidan\_Vickars](https://forum.opencv.org/u/Aidan_Vickars)\
**Post date:** [October 4, 2022, 10:59pm UTC](https://forum.opencv.org/t/copy-cv-gpumat-in-cuda-kernel/10436/1 "2022-10-04T22:59:04Z")

</div>

I am looking to copy a GpuMat into a 1D array using a custom Kernel. I do know that opencv has function to perform the copy, I would like to do it this way since it will be extended to do some custom padding and convert the image from HWC to CHW. The issue I am having is that I am simply getting garbage after making the copy. My code is shown below.

In short, what I can tell is that when I am making the copy in the kernal from the GpuMat to the 1D array, I am misunderstanding how data is structured in the GpuMat, as a result it is not copying properly.

If anyone has any suggestions it would be greatly appreciated!!!

Thanks,

Aidan

```auto
int main() {
   /* .... Some other code that is working just fine ....*/

   cv::cuda::GpuMat d_frame_resized; // This is for context. You can assume that d_frame_resized contains a valid image (I verified by writing the image to disc) of dimension 320 * 640 * 3 (HWC).
  int out_height = 320;
  int out_width - 640;
    unsigned char * h_out = copy(&d_frame_resized, out_height, out_width);
  cv::Mat A(out_height, out_width, CV_8UC3, h_out);
  cv::imwrite("test.jpg", A); // This image is a bunch of garbage.
}

```

```auto
unsigned char * copy(
        cv::cuda::GpuMat* d_in,
        int in_height,
        int in_width
        ) {
    // Computing the output dimensions
    int out_height = in_height;
    int out_width = in_width;

    // Making sure the height and width is valid
    if (out_width <= 0 || out_height <= 0 || in_height <= 0 || in_width <= 0 || left_pad < 0 || right_pad < 0 || top_pad < 0 || bottom_pad < 0) {
        throw std::runtime_error(std::string("INVALID: out_width <= 0 || out_height <= 0 || in_height <= 0 || in_width <= 0 || left_pad < 0 || right_pad < 0 || top_pad < 0 || bottom_pad < 0 in Cuda_Frame::crop_frame"));
    }
    // NOTE: Size of image is dynamic, hence the dynamic determination of the threads. If there is a better way to do this, please do let me know.

    // Defining initial grid dimensions
    dim3 gridDim(32, 32);

    // Defining initial block dimensions
    dim3 blockDim(iDivUp(out_width,gridDim.x), iDivUp(out_height,gridDim.y), 3);

    // Computing the number of threads per block
    int num_threads_per_block = blockDim.x * blockDim.y * 3;

    // Making the number of threads per block fit requirements
    while (num_threads_per_block > 512 || num_threads_per_block < 32) {

        // If there are too many threads in a block, we decrease the number of threads per block
        if (num_threads_per_block > 512) {
            gridDim.x *= 2;
            gridDim.y *= 2;

            blockDim.x = iDivUp(out_width,gridDim.x);
            blockDim.y = iDivUp(out_height,gridDim.y);
        }
        else { // We increase the number of threads per block
            if (blockDim.x > 2 | blockDim.y > 2) {
                gridDim.x /= 2;
                gridDim.y /= 2;

                blockDim.x = iDivUp(out_width,gridDim.x);
                blockDim.y = iDivUp(out_height,gridDim.y);
            } else {
                break;
            }
        }

        num_threads_per_block = blockDim.x * blockDim.y * 3;
    }

    // Pointer to the frame on device memory
    unsigned char* d_out;

    gpuErrchk( cudaMalloc((void**)&d_out, sizeof(unsigned char) * (out_width) * (out_height) * 3));

    // Launching kernel
    pad_and_convert_to_chw_kernel<<<gridDim,blockDim>>>(
            *d_in,
            d_out,
            in_width,
            in_height,
            out_width,
            out_height
    );

    gpuErrchk( cudaPeekAtLastError() );

    // Forcing thread to wait for crop
    gpuErrchk( cudaDeviceSynchronize() );

    // Creating space for the cropped frame on the host
    unsigned char * h_out = (unsigned char*) malloc(sizeof(unsigned char) * out_width * out_height * 3);

    // Copying the cropped frame from the device to host
    gpuErrchk( cudaMemcpy(h_out, d_out, sizeof(unsigned char) * out_width * out_height * 3, cudaMemcpyDeviceToHost));

    // Freeing the device output
    gpuErrchk( cudaFree(d_out));

    return h_out;
}

```

```auto
__global__ void pad_and_convert_to_chw_kernel(
        cv::cuda::PtrStepSzb input,
        unsigned char * output,
        int in_width,
        int in_height,
        int out_width,
        int out_height
) {

    // Computing the x position of the thread relative to the frame
    int out_x = (blockIdx.x * blockDim.x) + threadIdx.x;

    // Computing the y position of the thread relative to the frame
    int out_y = (blockIdx.y * blockDim.y) + threadIdx.y;

    // Computing the z position of the thread relative to the frame
    int out_in_z = threadIdx.z;

    // Performing check to make sure thread position is inside frame
    if (out_x >= out_width || out_y >= out_height || out_in_z >= 3) {
        return;
    }

    // Computing the x position of the thread relative to the original frame
    int in_x = out_x

    // Computing the y position of the thread relative to the original frame
    int in_y = out_y

    // Computing the index of the thread for the flattened original frame
    int in_pos = (in_y * input.step) + (in_x * 3) + out_in_z;

    // Computing the index of the thread for the flattened new frame
    int out_pos = (out_y * out_width) + (out_x * 3) + out_in_z;

    output[out_pos] = input[in_pos];
}

```

---

<div class="post-metadata">

**Author:** ![cudawarped](https://avatars.discourse-cdn.com/v4/letter/c/9dc877/32.png) [@cudawarped](https://forum.opencv.org/u/cudawarped)\
**Post date:** [October 5, 2022, 11:23am UTC](https://forum.opencv.org/t/copy-cv-gpumat-in-cuda-kernel/10436/2 "2022-10-05T11:23:48Z")

</div>

I am not sure this is an OpenCV question, have you tried the same approach on the host (using loops instead of grids of threads)?

I’ve only had a quick look at your code but it looks like you are launching one thread per BGR pixel and only copying the B but I may be wrong. Also won’t out\_in\_z always be 0?

---

<div class="post-metadata">

**Author:** ![Aidan\_Vickars](https://sea2.discourse-cdn.com/flex020/user_avatar/forum.opencv.org/aidan_vickars/32/6205_2.png) [@Aidan\_Vickars](https://forum.opencv.org/u/Aidan_Vickars)\
**Post date:** [October 5, 2022, 3:14pm UTC](https://forum.opencv.org/t/copy-cv-gpumat-in-cuda-kernel/10436/3 "2022-10-05T15:14:01Z")

</div>

thanks so much for the response!

So in regards to the threads, the block dim in the z dimension is set to three. This is done initially here:

> [@Aidan\_Vickars](#):
>
> `dim3 blockDim(iDivUp(out_width,gridDim.x), iDivUp(out_height,gridDim.y), 3);`

. Thats how “out\_in\_z” will be 0,1, or 2, and as a result there will be 3 threads per BGR pixel.

One thing to note, I don’t believe it really is my “code” per say, be cause when I do an identical copy in the kernel with the exception that I copy the contents of the GpuMat to **another** GpuMat instead of a raw 1D array, it works just fine.

So that’s why I’m guessing I am miss-understanding the structure of a GpuMat in memory, specifically I think the issue is likely right around here inside the kernel.

> [@Aidan\_Vickars](#):
>
> ```auto
> // Computing the index of the thread for the flattened original frame
> int in_pos = (in_y * input.step) + (in_x * 3) + out_in_z;
> 
> // Computing the index of the thread for the flattened new frame
> int out_pos = (out_y * out_width) + (out_x * 3) + out_in_z;
> 
> output[out_pos] = input[in_pos];
> 
> ```

This is why I’ve posted my question **here** instead of on say nvidia’s forum as the issue is likely a missunderstanding of how GpuMat is stored in memory and as a result I am not copying it properly.

Another interesting thing I’ve noticed as that when I force the source GpuMat to be continuous, the results are better. For context I’ve attached 2 pictures, 1 of the result I am getting when the source GpuMat is not continuous (the really bad jumbled of picture) and the other is the result I am getting when the source GpuMat is continuous (the not jumbled up picture but is incomplete).

Anyways I’ll try on the host side as you suggest and post an update.  
 ![test](https://us1.discourse-cdn.com/flex020/uploads/opencv/original/2X/7/7d615dcf8eed1f4a940a9fb19a5552e08a35be45.jpeg)

So to confirm that I am understanding correctly how the underlying data is stored in a GpuMat, my understanding is that it is stored as follows:

1 Row has the following structure:

[B G R B G R … B G R unusedSpace] where each B/G/R occupies is single uchar and the unused space is there so that the row is a contiguous block of memory and that the memory occupied by the row matches the step size (i.e. Mat.step). Further in a single row we have a B G and R for every pixel in that row. I.e. if the image has width 640, then the row will be of size: sizeof(uchar) \* 640 \* 3 + some unused space.

Putting the rows together we have the following structure:

[B G R B G R … B G R unusedSpace]  
[B G R B G R … B G R unusedSpace]  
… (for every row in the image)  
[B G R B G R … B G R unusedSpace]

such that a pointer to the 1st element in a row be achieved using the follow: Mat.data + (row \* Mat.step)

Is this all correct? If something doesn’t make sense please let me know.

---

<div class="post-metadata">

**Author:** ![Aidan\_Vickars](https://sea2.discourse-cdn.com/flex020/user_avatar/forum.opencv.org/aidan_vickars/32/6205_2.png) [@Aidan\_Vickars](https://forum.opencv.org/u/Aidan_Vickars)\
**Post date:** [October 5, 2022, 3:14pm UTC](https://forum.opencv.org/t/copy-cv-gpumat-in-cuda-kernel/10436/4 "2022-10-05T15:14:47Z")

</div>

the other picture  
 ![test](https://us1.discourse-cdn.com/flex020/uploads/opencv/original/2X/b/b43c31606491add3633026cbbcf46a346fd026e1.jpeg)

---

<div class="post-metadata">

**Author:** ![Aidan\_Vickars](https://sea2.discourse-cdn.com/flex020/user_avatar/forum.opencv.org/aidan_vickars/32/6205_2.png) [@Aidan\_Vickars](https://forum.opencv.org/u/Aidan_Vickars)\
**Post date:** [October 5, 2022, 4:09pm UTC](https://forum.opencv.org/t/copy-cv-gpumat-in-cuda-kernel/10436/5 "2022-10-05T16:09:58Z")

</div>

UPDATE: Figured it out. This:

> [@Aidan\_Vickars](#):
>
> `int out_pos = (out_y * out_width) + (out_x * 3) + out_in_z;`

needed to be:

```auto
int out_pos = (out_y * out_width * 3) + (out_x * 3) + out_in_z;

```