TLDR: Проблема связана с количеством памяти устройства, выделенной для изображения, и схемой индексации, используемой для доступа к значениям пикселей внутри ядра.Используйте исправленную реализацию из последнего раздела кода этого ответа.
Ниже приведено объяснение проблемных аспектов предоставленной реализации.
1.Общее количество байтов изображения
Входное изображение является 8-битным RGB-изображением, поэтому теоретическое число байтов, занимаемых им, равно width x height x number_of_channels.В этом случае оно должно быть numRows * numCols * 3.Но практически OpenCV выделяет выровненную память для данных изображения , поэтому общее количество байтов изображения должно быть рассчитано как image.step * numrows независимо от типа изображения и количества каналов.При этом вызовы cudaMalloc и cudaMemcpy ожидают общее количество байтов, которые мы хотим выделить или скопировать соответственно.Исправьте вызовы следующим образом (адаптируя код из ответа @ micehlson):
const size_t numBytes = input.step * numrows;
cudaMalloc<unsigned char>(&d_input, numBytes);
^
cudaMalloc<unsigned char>(&d_output, numBytes);
^
//Copy data from OpenCV input image to device memory
cudaMemcpy(d_input, input.ptr(), numBytes, cudaMemcpyHostToDevice);
^
//copy output from device to host
cudaMemcpy(output.ptr(), d_output, numBytes, cudaMemcpyDeviceToHost);
^
2.Вычисление индекса пикселя в ядре
Поскольку память изображения выровнена, фактический индекс пикселя должен быть рассчитан с использованием параметра step объекта Mat.Общая формула для расчета начального индекса пикселя в OpenCV Mat выглядит следующим образом:
index = строка * step / bytes_per_pixel_component + (каналы * столбец)
Для 8-битного изображения RGB число байтов, занимаемых одним компонентом пикселя RGB, составляет 1 байт.Это означает, что отдельный R или G или B занимает 8 байтов, в то время как весь пиксель RGB составляет 24 байта.Таким образом, начальный индекс рассчитывается как
int index = row * step + 3 * column;
Поскольку это начальный индекс, доступ к каждому отдельному каналу этого конкретного пикселя можно получить путем увеличения этого индекса до количества каналов следующим образом:
int R = index;
int G = index + 1;
int B = index + 2;
Впоследствии индекс пикселя в перевернутом изображении можно рассчитать следующим образом (при условии переворота вокруг оси y):
int flipped_index = row * step + 3 * (numCols - column - 1);
Конечно, нам потребуется шаг изображения какаргумент ядру.
Конечное ядро может выглядеть так:
__global__ void mirror( unsigned char* input, unsigned char* output, int numRows, int numCols, int channels, int step)
{
//2D Index of current thread
const int col = blockIdx.x * blockDim.x + threadIdx.x;
const int row = blockIdx.y * blockDim.y + threadIdx.y;
if ( col >= numCols || row >= numRows ) return;
const int tid = row * step + (channels * col);
const int tid_flipped = row * step + (channels * (numCols - col - 1)); //Flip about y axis
//Copy each component of the current pixel
for(int i=0; i<channels; i++)
output[tid_flipped + i] = input[tid + i];
}
Со всеми исправлениями финальный код может выглядеть так:
#include<iostream>
#include<cstdio>
#include<opencv2/core.hpp>
#include<opencv2/imgcodecs.hpp>
#include<opencv2/highgui.hpp>
#include<cuda_runtime.h>
using std::cout;
using std::endl;
__global__ void mirror( unsigned char* input, unsigned char* output, int numRows, int numCols, int channels, int step)
{
//2D index of current thread
const int col = blockIdx.x * blockDim.x + threadIdx.x;
const int row = blockIdx.y * blockDim.y + threadIdx.y;
if ( col >= numCols || row >= numRows ) return;
const int tid = row * step + (3 * col);
const int tid_new = row * step + (3 * (numCols - col - 1)); //Flip about y axis
//Copy each component of the current pixel
for(int i=0; i<channels; i++)
output[tid_new + i] = input[tid + i];
}
void convert_to_mirror(const cv::Mat& input, cv::Mat& output,int numrows,int numcols)
{
const dim3 blockSize(1024,1,1);
int a=numcols/blockSize.x, b=numrows/blockSize.y;
const dim3 gridSize(a+1,b+1,1);
const size_t numBytes = input.step * input.rows;
unsigned char *d_input, *d_output;
cudaMalloc<unsigned char>(&d_input, numBytes);
cudaMalloc<unsigned char>(&d_output,numBytes);
//Copy data from OpenCV input image to device memory
cudaMemcpy(d_input,input.ptr(), numBytes, cudaMemcpyHostToDevice);
//Call mirror kernel.
mirror<<<gridSize, blockSize>>>(d_input,d_output, numrows, numcols, input.channels(), input.step);
assert(cudaSuccess == cudaDeviceSynchronize());
//copy output from device to host
cudaMemcpy(output.ptr(), d_output,numBytes, cudaMemcpyDeviceToHost);
cudaFree(d_input);
cudaFree(d_output);
}
int main()
{
//Read input image from the disk
cv::Mat input = cv::imread("C:/a.jpg", cv::IMREAD_COLOR);
const int rows = input.rows;
const int cols = input.cols;
if(input.empty())
{
std::cout<<"Image Not Found!"<<std::endl;
std::cin.get();
return -1;
}
//Create output image
cv::Mat output(rows,cols,CV_8UC3);
//Call the wrapper function
convert_to_mirror(input,output,rows,cols);
//Show the input and output
cv::imshow("Input",input);
cv::imshow("Output",output);
//Wait for key press
cv::waitKey();
return 0;
}
Скомпилировано с помощью следующей команды:
nvcc -o mirror -std = c ++ 11 mirror.cu -I / usr / local / include / opencv4 -L / usr/ local / lib -lopencv_core -lopencv_imgcodecs -lopencv_highgui
Протестировано с OpenCV 4.0 и CUDA 9 в Ubuntu 16.04