`cuda::calcHist` only works when source data is 4 byte aligned

[cudawarped]

System information (version)

• OpenCV => 4.7
• Operating System / Platform => Windows 11 64 Bit
• Compiler => Visual Studio 2022

Detailed description

Internally

cv::cuda::calcHist(InputArray _src, OutputArray _hist, Stream& stream)

reads 4 byte integers in

__global__ void histogram256Kernel(const uchar* src, int cols, int rows, size_t step, int* hist)

from src regardless of its alignment

If _src is a sub roi of a larger matrix where the x offset is not a multiple of 4 then you get a misaligned pointer resulting in a CUDA_ERROR_MISALIGNED_ADDRESS error.

This could be addressed by asserting on the alignment requirement or modifying the kernel.

Steps to reproduce

GpuMat src(128, 128, CV_8U), hist;
cuda::calcHist(src(Rect(1,0,10,10)), hist);

results in the following exception

OpenCV(4.7.0-dev) D:\repos\opencv\contrib\modules\cudaimgproc\src\cuda\hist.cu:106: error: (-217:Gpu API call) misaligned address in function 'hist::histogram256'

Issue submission checklist

• I report the issue, it's not a question
• I checked the problem with documentation, FAQ, open issues,
  forum.opencv.org, Stack Overflow, etc and have not found any solution
• I updated to the latest OpenCV version and the issue is still there
• There is reproducer code and related data files: videos, images, onnx, etc