圖像二值化

main.cu

#include <iostream>
#include <stdio.h>
#include <opencv2/opencv.hpp>
#include <cuda_runtime.h>
#include "utils.cuh"

using namespace cv;

// 圖像二值化
__global__ void binary_kernel(const uchar *img_in, uchar *img_out, uchar threashold, int w, int h)
{
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    if (x < w && y < h)
    {
        img_out[y * w + x] = (img_in[y * w + x] > threashold) ? 255 : 0;
    }
}

int main()
{
    Mat img_src = imread("/mnt/data3/lena.jpg", 0);
    if (img_src.empty())
    {
        std::cout << "Failed to load image !" << std::endl;
        return -1;
    }

    uchar *d_in;
    uchar *d_out;
    int M = img_src.rows * img_src.cols * sizeof(uchar);
    CHECK(cudaMalloc((void **)&d_in, M));
    CHECK(cudaMalloc((void **)&d_out, M));
    CHECK(cudaMemcpy(d_in, img_src.data, M, cudaMemcpyHostToDevice));

    const int nx = 32;
    const int ny = 32;
    dim3 block_size(nx, ny);
    dim3 grid_size((img_src.cols + nx - 1) / nx, (img_src.rows + ny - 1) / ny);
    binary_kernel<<<grid_size, block_size>>>(d_in, d_out, 128, img_src.cols, img_src.rows);
    CUDA_CHECK_KERNEL

    Mat img_dst(img_src.size(), img_src.type());
    CHECK(cudaMemcpy(img_dst.data, d_out, M, cudaMemcpyDeviceToHost));

    imshow("src", img_src);
    imshow("dst", img_dst);
    waitKey(0);
}

utils.cuh

#pragma once
#include <stdio.h>
#include <cuda_runtime.h>

#define CHECK(call)                                     \
    do                                                  \
    {                                                   \
        const cudaError_t error_code = call;            \
        if (error_code != cudaSuccess)                  \
        {                                               \
            printf("CUDA Error:\n");                    \
            printf("    File:       %s\n", __FILE__);   \
            printf("    Line:       %d\n", __LINE__);   \
            printf("    Error code: %d\n", error_code); \
            printf("    Error text: %s\n",              \
                   cudaGetErrorString(error_code));     \
            exit(1);                                    \
        }                                               \
    } while (0)


#ifdef STRONG_DEBUG
#define CUDA_CHECK_KERNEL               \
    {                                   \
        CHECK(cudaGetLastError());      \
        CHECK(cudaDeviceSynchronize()); \
    }
#else
#define CUDA_CHECK_KERNEL          \
    {                              \
        CHECK(cudaGetLastError()); \
    }
#endif // CUDA_CHECK_KERNEL

class GPUTimer
{
public:
    GPUTimer()
    {
        cudaEventCreate(&m_start);
        cudaEventCreate(&m_end);
    }
    ~GPUTimer()
    {
        cudaEventDestroy(m_start);
        cudaEventDestroy(m_end);
    }

    void start()
    {
        cudaEventRecord(m_start);
    }

    void stop()
    {
        cudaEventRecord(m_end);
        cudaEventSynchronize(m_end);
    }

    float elpased_ms()
    {
        float ms = 0.0f;
        cudaEventElapsedTime(&ms, m_start, m_end);
        return ms;
    }

private:
    cudaEvent_t m_start;
    cudaEvent_t m_end;
};

CMakeLists.txt

cmake_minimum_required(VERSION 3.10)
project(TEST_BINARY)

enable_language(CXX CUDA)

# OpenCV
set(CMAKE_PREFIX_PATH "/home/wei/ubuntu/Libs/opencv-4.5.1/INSTALL")
find_package(OpenCV REQUIRED)
if(OpenCV_FOUND)
    message(STATUS "Found OpenCV")
endif()

# Debug
add_definitions(-DSTRONG_DEBUG)

add_executable(main "main.cu")
target_include_directories(main PRIVATE ${OpenCV_INCLUDE_DIRS})
target_link_libraries(main PRIVATE ${OpenCV_LIBS})

結(jié)果

image.png