【CUDA GPU 实现 Sobel 边缘

使用 CUDA GPU 实现 Sobel 边缘 轮廓提取

// 使用 GPU CUDA 实现 Sobel 算法

#include <iostream>
#include <cmath>
#include <cuda_runtime.h>
#include <cudnn.h>
#include <cuda.h>
#include <device_functions.h>
#include <device_launch_parameters.h>

#include <opencv.hpp>

// GUP 实现 Sobel 边缘检测
//                 x0  x1 x2
//                 x3  x4 x5
//                 x6  x7 x8

__global__ void Sobel_gpu(unsigned char* in, unsigned char* out, int imgHeigeht, int imgWidth)
{
	// 找到线程索引
	int IdxX = threadIdx.x + blockDim.x * blockIdx.x;
	int IdxY = threadIdx.y + blockDim.y * blockIdx.y;

	// 把二维图像拉成一维向量
	int Idx = IdxY * imgWidth + IdxX;		// 全局一维的索引
	int Gx = 0;
	int Gy = 0;

	unsigned char x0, x1, x2, x3, x4, x5, x6, x7, x8;	// 定义卷积数
	if (Idx > 0 && Idx < imgHeigeht && IdxY > 0 && IdxY < imgHeigeht) // 设置越界
	{
		x0 = in[(IdxY - 1) * imgWidth + IdxX - 1];	// 左上角
		x1 = in[(IdxY - 1) * imgWidth + IdxX];		// 上中
		x2 = in[(IdxY - 1) * imgWidth + IdxX + 1];	// 右上角
		x3 = in[IdxY * imgWidth + IdxX - 1];		// 左中
		x4 = in[IdxY * imgWidth + IdxX];			// 中心
		x5 = in[IdxY * imgWidth + IdxX + 1];		// 右中
		x6 = in[(IdxY + 1) * imgWidth + IdxX - 1];	// 左下角
		x7 = in[(IdxY + 1) * imgWidth + IdxX - 1];	// 下中
		x8 = in[(IdxY + 1) * imgWidth + IdxX - 1];	// 右下角
		Gx = (x0 + x3 * 2 + x6) - (x2 + x5 * 2 + x8);
		Gy = (x0 - x1 * 2 + x2) - (x6 + x7 * 2 + x8);
		out[Idx] = (abs(Gx) + abs(Gy)) / 2;
	}
}


__global__ void sobelFilter(const unsigned char* input, unsigned char* output, int width, int height) {
	int x = blockIdx.x * blockDim.x + threadIdx.x;
	int y = blockIdx.y * blockDim.y + threadIdx.y;

	if (x < 1 || x >= width - 1 || y < 1 || y >= height - 1) {
		return;
	}

	int Gx = -1 * input[(y - 1) * width + (x - 1)] - 2 * input[y * width + (x - 1)] - input[(y + 1) * width + (x - 1)]
		+ input[(y - 1) * width + (x + 1)] + 2 * input[y * width + (x + 1)] + input[(y + 1) * width + (x + 1)];

	int Gy = -1 * input[(y - 1) * width + (x - 1)] - 2 * input[(y - 1) * width + x] - input[(y - 1) * width + (x + 1)]
		+ input[(y + 1) * width + (x - 1)] + 2 * input[(y + 1) * width + x] + input[(y + 1) * width + (x + 1)];

	//output[y * width + x] = sqrt(Gx * Gx + Gy * Gy);
	output[y * width + x] = (abs(Gx) + abs(Gy)) / 2;
}


int main()
{
	// 利用 OpenCV 的接口读取图片
	cv::Mat img = cv::imread("./lena.jpg", 0);
	int imgHeight = img.rows;
	int imgWidth = img.cols;
	
	cv::Mat dst_gpu = cv::Mat(imgHeight, imgWidth, CV_8UC1, cv::Scalar(0));
	// 申请指针并指向 GPU 的空间
	size_t N = imgHeight * imgWidth * sizeof(unsigned char);
	unsigned char* in_gpu;
	unsigned char* out_gpu;
	cudaMalloc((void**)&in_gpu, N);
	cudaMalloc((void**)&out_gpu, N);
	// 定义 grid 和 block 的维度和形状
	// 
	dim3 threadPerBlock(32, 32, 1);
	dim3 threadPerGrid((imgWidth + threadPerBlock.x - 1) / threadPerBlock.x, 
		(imgHeight + threadPerBlock.y - 1) / threadPerBlock.y);
	// 将数据从 CPU 传输到 GPU
	cudaMemcpy(in_gpu, img.data, N, cudaMemcpyHostToDevice);
	// 调用GPU 上的函数
	sobelFilter <<<threadPerBlock, threadPerGrid>>>(in_gpu, out_gpu, imgHeight, imgWidth);

	// 将计算结果回传 CPU
	cudaMemcpy(dst_gpu.data, out_gpu, N, cudaMemcpyDeviceToHost);
	// 显示计算结果

	cv::imshow("dst_gpu", dst_gpu);
	cv::waitKey(0);



	// 释放内存
	cudaFree(in_gpu);
	cudaFree(out_gpu);
	return 0;
}