# How to address a cv::cuda::GpuMat variable, passed to kernel as cv::cudaPtrStepSz?

**URL:** <https://forum.opencv.org/t/how-to-address-a-cv-gpumat-variable-passed-to-kernel-as-cv-cudaptrstepsz/17792>\
**Category:** C++\
**Tags:** cuda\
**Created:** [May 30, 2024, 9:32am UTC](https://forum.opencv.org/t/how-to-address-a-cv-gpumat-variable-passed-to-kernel-as-cv-cudaptrstepsz/17792 "2024-05-30T09:32:56Z")\
**Posts on this page:** 2\
**Page:** 1

<div class="post-metadata">

**Author:** ![Abinas\_Chopdar](https://sea2.discourse-cdn.com/flex020/user_avatar/forum.opencv.org/abinas_chopdar/32/10019_2.png) [@Abinas\_Chopdar](https://forum.opencv.org/u/Abinas_Chopdar)\
**Post date:** [May 30, 2024, 9:32am UTC](https://forum.opencv.org/t/how-to-address-a-cv-gpumat-variable-passed-to-kernel-as-cv-cudaptrstepsz/17792/1 "2024-05-30T09:32:56Z")

</div>

Hi!  
I have an BGR image in device(cv::cuda::GpuMat) and I am passing this variable to my custom kernel function as cv::cuda::PtrStepSz. Below is the code of kernel.

```auto
__global__ void hist_mod(
                            int* BMap, int* GMap, int* RMap,
                            cv::cuda::PtrStepSz<uchar> input,
                            unsigned char* output, 
                            int color_width_step,
                            int rows, int cols
                        ){
    
    int xidx = threadIdx.x + blockIdx.x*blockDim.x;
    int yidx = threadIdx.y + blockIdx.y*blockDim.y;

    if(xidx < cols && yidx < rows){
        const int color_tid = yidx * color_width_step + xidx * 3; // location of the pixel in the image 3d
        // const int gray_tid = yidx * cols + xidx; // location of the pixel in the image 2d
        
        const unsigned char blue = input(yidx, xidx)[0];
        const unsigned char green = input(yidx, xidx)[1];
        const unsigned char red = input(yidx, xidx)[2];

        output[color_tid] = (uchar)BMap[(int)blue];
        output[color_tid + 1] = (uchar)GMap[(int)green];
        output[color_tid + 2] = (uchar)RMap[(int)red];

    }
}

```

My query is, as a 3 channel BGR Image, how to address inside the kernel?

```auto
        const unsigned char blue = input(yidx, xidx)[0];
        const unsigned char green = input(yidx, xidx)[1];
        const unsigned char red = input(yidx, xidx)[2];

```

---

<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:** [May 30, 2024, 10:24am UTC](https://forum.opencv.org/t/how-to-address-a-cv-gpumat-variable-passed-to-kernel-as-cv-cudaptrstepsz/17792/2 "2024-05-30T10:24:35Z")

</div>

> [@Abinas\_Chopdar](#):
>
> My query is, as a 3 channel BGR Image, how to address inside the kernel?

Unless I’m missing something it should be

```auto
const unsigned char blue = input(yidx, xidx);
const unsigned char green = input(yidx, xidx+1);
const unsigned char red = input(yidx, xidx+2);

```