如何在 CUDA 中实现最近邻图像大小调整算法？

Question

我的主要目的是使用 OpenCV 从视频中加载帧，然后将其复制到 Nvidia Gpu 内存，使用基于 Cuda 的最近邻算法调整大小，然后将其复制回主机端并使用 [=12 进行可视化=]

不幸的是，我总是遇到分段错误。定义要复制的字节数或数据转换可能存在问题。下面，您可以找到源代码的主要部分，但这里是完整项目的回购协议： https://github.com/foxakarmi/imageResize

主要功能：

#include <iostream>
#include "cuda_utils.h"
#include "yololayer.h"
#include <opencv2/highgui/highgui.hpp>

void *buffers[3];

int main() {

    cv::VideoCapture capture;
    cv::Mat frame;

    capture.open("/p.mp4");

    if (!capture.isOpened()) {
        std::cout << "can not open" << std::endl;
        return -1;
    }
    capture.read(frame);

    CUDA_CHECK(cudaMalloc(&buffers[0], frame.cols * frame.step[0]));
    CUDA_CHECK(cudaMalloc(&buffers[1], 3 * 640 * 640));
    buffers[2] = malloc(3 * 640 * 640);

    while (capture.read(frame)) {
        CUDA_CHECK(cudaMemcpy(buffers[0], frame.ptr(), frame.step[0] * frame.rows, cudaMemcpyHostToDevice))

        cudaNearestResize((uchar *) buffers[0], (uchar *) buffers[1], frame.cols, frame.rows, 640, 640);

        CUDA_CHECK(cudaMemcpy(buffers[2], buffers[1], 640 * 640 * 3, cudaMemcpyDeviceToHost))

        cv::Mat foo;
        foo.data = static_cast<uchar *>(buffers[2]);
        cv::imshow("img", foo);
        cv::waitKey(1);
    }

    capture.release();
    return 0;
}

包含内核和包装函数的 .cu 文件：

#include <opencv2/core/hal/interface.h>
#include "yololayer.h"
#include "cuda_utils.h"

__global__ void kernelNearestNeighbourResize(uchar *src_img, uchar *dst_img, int src_w, int src_h, int dst_w, int dst_h) {
    int i = blockDim.y * blockIdx.y + threadIdx.y;
    int j = blockDim.x * blockIdx.x + threadIdx.x;

    int channel = 3;

    if (i < dst_h && j < dst_w) {
        int iIn = i * src_h / dst_h;
        int jIn = j * src_w / dst_h;

        dst_img[(i * dst_w + j) * channel + 0] = src_img[(iIn * src_w + jIn) * channel + 0];
        dst_img[(i * dst_w + j) * channel + 1] = src_img[(iIn * src_w + jIn) * channel + 1];
        dst_img[(i * dst_w + j) * channel + 2] = src_img[(iIn * src_w + jIn) * channel + 2];
    }
}

cudaError_t cudaNearestResize(uchar *src_img, uchar *dst_img, int src_w, int src_h, int dst_w, int dst_h) {
    if (!src_img || !dst_img)
        return cudaErrorInvalidDevicePointer;

    if (src_w == 0 || src_h == 0 || dst_w == 0 || dst_h == 0)
        return cudaErrorInvalidValue;

    kernelNearestNeighbourResize <<< 3600, 256>>>(
            src_img, dst_img, src_w,
            src_h, dst_w, dst_h);

    return cudaGetLastError();
}

Answer 1

您可以在下面看到一个完整的工作解决方案。

您的代码中存在 3 个主要问题：

CUDA 网格的设置不正确。请参阅下面我的代码中如何设置它的示例（只是您可以进一步改进的初始工作版本）。在此处查看一些一般信息：The CUDA Programming Model.
注意：网格设置可以对整体性能产生有意义的影响，并且优化起来并非易事。在此处查看更多信息：How do I choose grid and block dimensions for CUDA kernels?.
将数据复制到设备时，您使用了 frame.ptr() 而不是 frame.data。
您只设置了输出 cv::Mat foo 的数据指针，而没有正确初始化它。因此 cv::Mat 元数据（行、列等）未设置，cv::imshow 无法正确显示。在我的代码中它不是必需的 - 见下文。

请注意，您的代码跳过了第一帧。我保持了这种行为。您可以通过检查 dst_img 是否已经初始化来包括第一帧，如果没有（因为它是第一帧） - 初始化它和 CUDA 缓冲区。

关于下面代码的更多注释：

无需为主机输出图像分配 buffer[2]。相反，我用适当的大小初始化了 cv::Mat 并使用它分配的缓冲区。
我重命名了设备缓冲区，并为它们添加了 cudaFree。
将通道数传递给内核比让它假定为 3 更安全。
我将图像的步长（AKA 步幅）传递给了内核。这将支持图像有填充的情况（参见此处：）。

main的代码：

#include <iostream>
#include <opencv2/highgui/highgui.hpp>
#include "cuda_runtime.h"
#include <assert.h>

#define CUDA_CHECK(x) { cudaError_t cudaStatus = x; assert(cudaStatus == cudaSuccess); }

cudaError_t cudaNearestResize(unsigned char *src_img, unsigned char *dst_img, int channel,
    int src_w, int src_h, int src_step, int dst_w, int dst_h, int dst_step);

int main()
{
    cv::VideoCapture capture;
    cv::Mat frame;
    capture.open("/p.mp4");
    if (!capture.isOpened()) 
    {
        std::cout << "can not open" << std::endl;
        return -1;
    }
    capture.read(frame);

    int src_w = frame.cols;
    int src_h = frame.rows;
    int src_step = (int)frame.step[0];
    int channels = frame.channels();
    int data_type = frame.type();
    assert((data_type & CV_MAT_DEPTH_MASK) == CV_8U);   // assert that it is a uchar image

    // Parameters you can change:
    int dst_w = 640;
    int dst_h = 640;

    cv::Mat dst_img(dst_h, dst_w, data_type);
    int dst_step = (int)dst_img.step[0];

    void * src_dev_buffer;
    void * dst_dev_buffer;
    CUDA_CHECK(cudaMalloc(&src_dev_buffer, src_h * src_step));
    CUDA_CHECK(cudaMalloc(&dst_dev_buffer, dst_h * dst_step));

    while (capture.read(frame))
    {
        // assert that the current frame has the same type and dimensions as the first one (should be guaranteed by the video decoder):
        assert(frame.cols == src_w);
        assert(frame.rows == src_h);
        assert((int)frame.step[0] == src_step);
        assert(frame.type() == data_type);

        CUDA_CHECK(cudaMemcpy(src_dev_buffer, frame.data, src_h * src_step, cudaMemcpyHostToDevice));
        CUDA_CHECK(cudaNearestResize((unsigned char *)src_dev_buffer, (unsigned char *)dst_dev_buffer, channels, src_w, src_h, src_step, dst_w, dst_h, dst_step));
        CUDA_CHECK(cudaMemcpy(dst_img.data, dst_dev_buffer, dst_h * dst_step, cudaMemcpyDeviceToHost));
        cv::imshow("dst_img", dst_img);
        cv::waitKey(1);
    }

    CUDA_CHECK(cudaFree(src_dev_buffer));
    CUDA_CHECK(cudaFree(dst_dev_buffer));

    capture.release();
    return 0;
}

CUDA内核代码和包装函数：

#include "cuda_runtime.h"

__global__ void kernelNearestNeighbourResize(unsigned char *src_img, unsigned char *dst_img, int channels,
    int src_w, int src_h, int src_step, int dst_w, int dst_h, int dst_step)
{
    int i = blockDim.y * blockIdx.y + threadIdx.y;
    int j = blockDim.x * blockIdx.x + threadIdx.x;

    if (i < dst_h && j < dst_w) 
    {
        int iIn = i * src_h / dst_h;
        int jIn = j * src_w / dst_w;

        int src_offset = i * dst_step + j * channels;
        int dst_offset = iIn * src_step + jIn * channels;
        for (int c = 0; c < channels; ++c) 
        {
            dst_img[src_offset + c] = src_img[dst_offset + c];
        }
    }
}

cudaError_t cudaNearestResize(unsigned char *src_img, unsigned char *dst_img, int channels,
    int src_w, int src_h, int src_step, int dst_w, int dst_h, int dst_step)
{
    if (!src_img || !dst_img)
        return cudaErrorInvalidDevicePointer;

    if (src_w == 0 || src_h == 0 || dst_w == 0 || dst_h == 0)
        return cudaErrorInvalidValue;

    // The grid dimensions
    dim3 dimBlock(32, 32);
    dim3 dimGrid(dst_w / 32 + 1, dst_h / 32 + 1);

    kernelNearestNeighbourResize << < dimGrid, dimBlock >> >(
        src_img, dst_img, channels,
        src_w, src_h, src_step, dst_w, dst_h, dst_step);

    return cudaGetLastError();
}

如何在 CUDA 中实现最近邻图像大小调整算法？

How to implement nearest neighbours image resizing algorithm in CUDA?

opencv

cuda

image-resizing