Custom Kernel GpuMat with float

Question

I'm trying to write a custom kernel using `GpuMat` data to find the arc cosine of an image's pixels. I can upload, download, and change values when I upload data when the GPU has `CV_8UC1` data but chars cannot be used to calculate arc cosines. However, when I try to convert my GPU to `CV_32FC1` type (floats) I get an illegal memory access error during the download part. Here is my code: ``` //.cu code #include #include #include #include __global__ void funcKernel(const float* srcptr, float* dstptr, size_t srcstep, const size_t dststep, int cols, int rows){ int rowInd = blockIdx.y*blockDim.y+threadIdx.y; int colInd = blockIdx.x*blockDim.x+threadIdx.x; if(rowInd >= rows || colInd >= cols) return; const float* rowsrcptr=srcptr+rowInd*srcstep; float* rowdstPtr= dstptr+rowInd*dststep; float val = rowsrcptr[colInd]; if((int) val % 90 == 0) rowdstPtr[colInd] = -1 ; else{ float acos_val = acos(val); rowdstPtr[colInd] = acos_val; } } int divUp(…

Accepted Answer

You are treating image `step` as if it is a `float` offset. It is a byte offset from one row to the next.

Try something like this instead:

```
const float* rowsrcptr= (const float *)(((char *)srcptr)+rowInd*srcstep);
float* rowdstPtr=  (float *) (((char *)dstptr)+rowInd*dststep);
```

from the documentation:

step – Number of bytes each matrix row occupies.

It's also a good idea to add proper cuda error checking to your code (e.g. to `func`). And you can run your code with `cuda-memcheck` to see the actual kernel failure generating the invalid reads/writes.

Custom Kernel GpuMat with float

Solution

Problem

Related problems