Прежде всего, для обработки 16-битных изображений данные пикселей должны интерпретироваться как тип данных 16-битной ширины, который может быть unsigned short или short. Помните, что нам нужно только интерпретировать данные изображения как unsigned short тип; не типа брось. Для этого мы просто преобразуем указатель данных изображения в требуемый тип, как показано в следующем примере:
unsigned short* ptr16 = reinterpret_cast<unsigned short*>(im1.ptr());
Как следствие вышеприведенного шага, мы также должны создать отдельное ядро для 16-битного типа данных. Мы можем сделать это разумно, определив ядро как шаблон C ++.
Таким образом, ядро может выглядеть следующим образом:
template<typename T>
__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));
}
}
}
Еще одна важная особенность при обработке OpenCV Mat с использованием собственного ядра CUDA заключается в том, что шаг изображения должен быть разделен на размер типа данных в байтах. Для 16-битного изображения размер одного пикселя составляет 16 бит (2 байта), поэтому шаг, используемый внутри ядра, нужно разделить на 2. Помните, что исходный шаг не следует изменять. Нужно разделить только значение шага, переданное в качестве аргумента ядра.
С учетом вышеуказанных исправлений окончательный код ЦП может выглядеть следующим образом:
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);
}
Похоже, что шум в выходном изображении вносит искажения из-за логики алгоритма бинирования Например, это очень похоже на повторную выборку изображения с использованием подхода ближайшего соседа.
Обновление:
Вышеупомянутый подход к вычислению адреса памяти пикселя недокументирован и является просто результатом интуиции, поэтому он может показаться немного нестандартным.
Другой метод, используемый OpenCV и другими библиотеками, позволяет избежать путаницы при шаговом разделении изображения. Это происходит следующим образом с учетом индекса x и y пикселя:
- Переинтерпретировать указатель данных изображения как представление байтов (
unsigned char*).
- Рассчитать начальный адрес строки изображения, используя индекс y и шаг изображения.
- Повторно интерпретировать начальный адрес строки как требуемый тип (
unsigned short*).
- Доступ к индексу x указателя начала строки.
Используя этот метод, мы можем вычислить адрес памяти пикселя для изображения в градациях серого следующим образом:
template<typename T>
T* getPixelAddress(unsigned char* data, int x, int y, int step)
{
T* row = (T*)((unsigned char*)(data) + y * step);
return row + x;
}
В приведенном выше методе значение шага является исходным без деления.