c++ - CUDA 图像不显示输出

Question

这是使用CUDA内核和opencv读取显示图像来翻转图像的代码，在main函数中，显示了输入的图片，但输出显示的是黑窗。顺便说一句，代码没有错误，它可以编译运行但输出看起来很奇怪。以下是我到目前为止所做的尝试。

#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)
{
    //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;

    int thread_x = blockDim.x * blockIdx.x + threadIdx.x;
    int thread_y = blockDim.y * blockIdx.y + threadIdx.y;
    int thread_x_new = numCols-thread_x;
    int thread_y_new = thread_y;
    int mId = thread_y * numCols + thread_x;
    int mId_new = thread_y_new * numCols + thread_x_new;
    output[mId_new] = input[mId]; 
}

 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 numPixels = numrows * numcols;
    unsigned char *d_input, *d_output;

    cudaMalloc<unsigned char>(&d_input, numPixels);
    cudaMalloc<unsigned char>(&d_output,numPixels);
    //Copy data from OpenCV input image to device memory
    cudaMemcpy(d_input,input.ptr(), numPixels,cudaMemcpyHostToDevice);
    //Call mirror kernel.
    mirror<<<gridSize, blockSize>>>(d_input,d_output, numrows, numcols);
    cudaDeviceSynchronize(); 
    //copy output from device to host
    cudaMemcpy(output.ptr(), d_output,numPixels, 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;
}

Answer 1

TLDR:问题在于为图像分配的设备内存量以及用于访问内核内像素值的索引方案。使用此答案最后一个代码部分中更正的实现。

以下是对所提供实现的问题方面的解释。

1。图片总字节数

输入图像为8位RGB图像，因此其理论占用字节数等于width x height x number_of_channels。在这种情况下，它应该是 numRows * numCols * 3。但实际上，OpenCV 分配 aligned memory for image data ，因此无论图像类型和 channel 数如何，图像字节总数都应计算为 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。内核中的像素索引计算

由于图像内存是对齐的，因此应该使用 Mat 对象的 step 参数计算像素的实际索引。计算 OpenCV Mat 中像素起始索引的通用公式如下:

index = row * step/bytes_per_pixel_component + (channels * column)

对于8位的RGB图像，一个RGB像素的单个分量占用的字节数为1字节。这意味着单个 R 或 G 或 B 占用 1 个字节，而整个 RGB 像素为 3 个字节。所以起始索引计算为

int index = row * step + 3 * column;

由于这是起始索引，因此可以通过将此索引递增至 channel 数来访问此特定像素的每个单独 channel ，如下所示:

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

在 Ubuntu 16.04 上使用 OpenCV 4.0 和 CUDA 9 测试