首先,为了处理16位图像,像素数据必须被解释为16位宽的数据类型,其可能是
unsigned short
要么
short
。请记住,我们只需要
的
译
</强>
图像数据为
unsigned short
类型;不要输入它。为此,我们将把图像数据指针转换为所需的类型,如下例所示:
unsigned short ptr16 = reinterpret_cast<unsigned short>(im1.ptr());
</code>
作为上述步骤的结果,我们还必须为16位数据类型创建单独的内核。我们可以通过将内核定义为C ++模板来巧妙地做到这一点。
所以内核可能如下所示:
template
global void binCUDAKernel(T input, T output, int binDim, int outputWidth, int outputHeight, int inputWstep, int outputWstep, int nChannels)
{
int outXind = blockIdx.x blockDim.x + threadIdx.x;
int outYind = blockIdx.y blockDim.y + threadIdx.y;
if ((outXind < outputWidth) && (outXind > outputWidth/2) && (outYind < outputHeight)) // Only run threads in output image coordinate range
{
if (nChannels == 1) // Test only for greyscale images
{
// Calculate x & y index of input binned pixels corresponding to current output pixel
int inXstart = outXind * binDim;
int inYstart = outYind * binDim;
// Perform binning on identified input pixels
float sum = 0;
for (int binY = inYstart; binY < (inYstart + binDim); binY++) {
for (int binX = inXstart; binX < (inXstart + binDim); binX++) {
int input_tid = binY * inputWstep + binX;
sum += float(input[input_tid]);
}
}
// Establish output thread index in current output pixel index
int output_tid = outYind * outputWstep + outXind;
// Assign binned pixel value to output pixel
output[output_tid] = static_cast<T>(sum / (binDim*binDim));
}
}
}
</code>
使用自定义CUDA内核处理OpenCV Mat时的另一个重要问题是图像步骤必须除以数据类型的大小(以字节为单位)。对于16位图像,单个像素的大小为16位(2个字节),因此内核中使用的步骤必须除以2.请记住,不应修改原始步骤。只是作为内核参数传递的步长值应该被划分。
结合上述修复,最终的CPU代码可能如下所示:
void binFilter(const cv::Mat input, std::string output_file)
{
// 2X2 binning
int binDim = 2;
// Create blank output image & calculate size of input and output
cv::Size outsize(input.size().width / binDim, input.size().height / binDim);
cv::Mat output(outsize, input.type());
const int inputBytes = input.step * input.rows;
const int outputBytes = output.step * output.rows;
// Allocate memory in device
unsigned char *d_input, *d_output;
gpuErrchk(cudaMalloc<unsigned char>(&d_input, inputBytes));
gpuErrchk(cudaMalloc<unsigned char>(&d_output, outputBytes));
// Copy input image to device
gpuErrchk(cudaMemcpy(d_input, input.ptr(), inputBytes, cudaMemcpyHostToDevice));
// Configure size of block and grid
const dim3 block(16, 16);
const dim3 grid((output.cols + block.x - 1) / block.x, (output.rows + block.y - 1) / block.y); // Additional block for rounding up
int depth = input.depth();
// Execute kernel
if (input.depth() == CV_16U)
{
typedef unsigned short t16;
t16* input16 = reinterpret_cast<t16*>(d_input);
t16* output16 = reinterpret_cast<t16*>(d_output);
int inputStep16 = input.step / sizeof(t16);
int outputStep16 = output.step / sizeof(t16);
binCUDAKernel <t16> <<<grid, block>>> (input16, output16, binDim, output.cols, output.rows, inputStep16, outputStep16, input.channels());
}
else
{
binCUDAKernel <unsigned char> <<<grid, block>>> (d_input, d_output, binDim, output.cols, output.rows, input.step, output.step, input.channels());
}
gpuErrchk(cudaPeekAtLastError());
// Wait for all threads to finish
//gpuErrchk(cudaDeviceSynchronize());
// Copy output image from device back to host (cudaMemcpy is a blocking instruction)
gpuErrchk(cudaMemcpy(output.ptr(), d_output, outputBytes, cudaMemcpyDeviceToHost));
// Free device memory
gpuErrchk(cudaFree(d_input));
gpuErrchk(cudaFree(d_output));
// Write image to specified output_file path
cv::imwrite(output_file, output);
}
</code>
由于分箱算法的逻辑,输出图像中的噪声似乎是引入的混叠。例如,它与使用最近邻方法重新采样图像非常相似。
更新:
上面提到的计算像素的存储器地址的方法没有记录,只是直觉的结果,所以它看起来有点不同寻常。
OpenCV和其他库使用的另一种方法避免了图像步骤划分的混淆。给定像素的x和y索引,它如下进行:
将图像数据指针重新解释为字节表示(
unsigned char
)。
使用y索引和图像步骤计算图像行的起始地址。
将行开始地址重新解释为所需类型(
unsigned short
)。
访问行开始指针的x索引。
</醇>
使用此方法,我们可以计算灰度图像的像素存储器地址,如下所示:
template
T getPixelAddress(unsigned char data, int x, int y, int step)
{
T row = (T)((unsigned char)(data) + y step);
return row + x;
}
</code>
在上述方法中,步长值是原始值,没有任何除法。