CUDA C/C++: Async kernel error: an illegal memory access was encountered

I'm beginner in parallel programming for GPU. I'm trying to implement a convolutional operation between a 512x512 RGB image size and a 9x9 filter.

I got an error: illegal memory access was encountered.

Here is my kernel function:

__global__ void blurImgKernel(uchar3 * inPixels, int width, int height, 
        float * filter, int filterWidth, 
        uchar3 * outPixels)
{
    // TODO
    int c = threadIdx.x + blockIdx.x * blockDim.x;
    int r = threadIdx.y + blockIdx.y * blockDim.y;
    int padding = filterWidth / 2;
    if (r < height && c < width)
    {
        int idx = r * width + c;        
        float red = 0;
        float green = 0;
        float blue = 0;
        
        // Do some calculation here....
        // ............................

        
        outPixels[idx].x = (uint8_t)red;
        outPixels[idx].y = (uint8_t)green;
        outPixels[idx].z = (uint8_t)blue;
    }

}

How I call this function:

dim3 blockSize(32, 32); 
uchar3 *d_inPixels, *d_outPixels;
CHECK(cudaMalloc(&d_inPixels, width*height*sizeof(uchar3)));
CHECK(cudaMalloc(&d_outPixels, width*height*sizeof(uchar3)));


// Copy data to device memories
CHECK(cudaMemcpy(d_inPixels, inPixels, 
    width*height*sizeof(uchar3), cudaMemcpyHostToDevice));

// Set grid size and call kernel (remember to check kernel error)
dim3 gridSize((height - 1) / blockSize.x + 1, (width - 1) / blockSize.y + 1);
blurImgKernel<<<gridSize, blockSize>>>(d_inPixels, width, height, filter, filterWidth, d_outPixels);

cudaError_t errSync  = cudaGetLastError();
cudaError_t errAsync = cudaDeviceSynchronize();
if (errSync != cudaSuccess) 
    printf("Sync kernel error: %s\n", cudaGetErrorString(errSync));
if (errAsync != cudaSuccess)
    printf("Async kernel error: %s\n", cudaGetErrorString(errAsync));

// Copy result from device memories
CHECK(cudaMemcpy(outPixels, d_outPixels, 
    width*height*sizeof(uchar3), cudaMemcpyDeviceToHost));

When I use cuda-memcheck to getting more error details, I got a lot of errors like this:

=========     Invalid __global__ read of size 4
=========     at 0x00000490 in blurImgKernel(uchar3*, int, int, float*, int, uchar3*)
=========     by thread (0,10,0) in block (0,0,0)
=========     Address 0x562bcf426000 is out of bounds
=========     Device Frame:blurImgKernel(uchar3*, int, int, float*, int, uchar3*) (blurImgKernel(uchar3*, int, int, float*, int, uchar3*) : 0x490)
=========     Saved host backtrace up to driver entry point at kernel launch time

I see my code look fine but I think there is something wrong in dividing blocksize and gridsize step. Could anybody help me to figure out?

Actually, I forgot to allocate filter on the device. After I allocate it, everything works fine.