opencv_cuda版本+cuda学习 （二）图像反向

先放代码

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include<cuda.h>
#include<cuda_device_runtime_api.h>
#include <opencv2/core/cuda.hpp>
#include<opencv2/opencv.hpp>
#include<opencv.hpp>
#include <stdio.h>
#include<iostream>
#include "error.cuh"

using namespace std;
using namespace cv;
using namespace cuda;

template<int nthreads>
__global__ void compute_kernel(int height, int width, const PtrStepb img, PtrStepb dst) {

	const int x = blockIdx.x * blockDim.x + threadIdx.x;
	const int y = blockIdx.y * blockDim.y + threadIdx.y;

	const uchar* src_y = (const uchar*)(img + y * img.step); //原图像
	uchar* dst_y = (uchar*)(dst + y * dst.step);

	if (x < width && y < height) {
		dst_y[3 * x] = 255 - src_y[3 * x];
		dst_y[3 * x + 1] = 255 - src_y[3 * x + 1];
		dst_y[3 * x + 2] = 255 - src_y[3 * x + 2];   //三通道
	}


}


int main() {

	Mat a = imread("C:/Users/1/Pictures/Saved Pictures/1.jpg");
	GpuMat d_a(a);
	GpuMat d_dst(d_a.size(), CV_8UC3);


	int width = a.size().width;
	int height = a.size().height;

	const int nthreads = 256;
	dim3 bdim(nthreads, 1);
	dim3 gdim(divUp(width, bdim.x), divUp(height, bdim.y));

	cudaEvent_t start, stop;
	CHECK(cudaEventCreate(&start));
	CHECK(cudaEventCreate(&stop));
	CHECK(cudaEventRecord(start));
	while (cudaEventQuery(start) != cudaSuccess) {}


	compute_kernel<nthreads> << <gdim, bdim >> > (height, width, d_a, d_dst);


	CHECK(cudaEventRecord(stop));
	CHECK(cudaEventSynchronize(stop));
	float elapsed_time;
	CHECK(cudaEventElapsedTime(&elapsed_time, start, stop));
	std::cout << "Time = " << elapsed_time << " ms." << std::endl;

	CHECK(cudaEventDestroy(start));
	CHECK(cudaEventDestroy(stop));


	namedWindow("原始图像",WINDOW_NORMAL);
	imshow("原始图像", a);
	resizeWindow("原始图像",1000,700);

	Mat dst(d_dst);
	namedWindow("反向图像", WINDOW_NORMAL);
	imshow("反向图像", dst);
	resizeWindow("反向图像", 1000, 700);

	waitKey();

	return 0;
}

其中CHECK来自定义的error.cuh文件，用来锁定那块代码出错；

#pragma once
#include <stdio.h>
#include <stdlib.h>  //exit()函数在这个库


#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)

//#endif

//CHECK函数首先是用了一个do while
//定义cudaError_t类型变量error_code，并初始化函数call的返回值
//这里判断是否返回成功，如果不成功，就输出错误信息。
//cudaGetErrorString 作用是将错误代号转化为错误的文字描述。

现在来讲讲代码。

GpuMat 是OpenCV的CUDA模块中的一个类，用于在GPU上存储图像和图像数据。它类似于Mat，但数据存储在GPU内存中，可以在GPU上执行各种图像处理操作，以提高性能。GpuMat 类提供了一种在GPU上进行图像处理的方法，而不需要将数据从主机内存传输到GPU内存。

不能在主机端使用imshow输出GpuMat类的参数。

（1）创建对象

cv::cuda::GpuMat d_image; // 创建一个空的GpuMat对象

（2）从主机Mat创建对象

cv::Mat h_image = cv::imread("image.jpg");
cv::cuda::GpuMat d_image(h_image); // 从主机Mat对象创建GpuMat对象

（3）从GpuMat对象复制到主机

cv::cuda::GpuMat d_image; // 创建GpuMat对象
cv::Mat h_image;
d_image.download(h_image); // 从GPU复制到主机

（4）在GPU上执行图像处理操作：

cv::cuda::GpuMat d_image; // 创建GpuMat对象
// 在GPU上执行图像处理操作，如高斯模糊
cv::cuda::GaussianBlur(d_image, d_image, cv::Size(5, 5), 0);

（5）在主机和GPU之间传输数据：

cv::Mat h_image = cv::imread("image.jpg");
cv::cuda::GpuMat d_image;
d_image.upload(h_image); // 从主机传输到GPU
d_image.download(h_image); // 从GPU传输到主机

然后其实有些内容上一节已经讲过了，这里再来看看模板核函数。

template<int nthreads>
__global__ void compute_kernel(int height, int width, const PtrStepb img, PtrStepb dst) {

	const int x = blockIdx.x * blockDim.x + threadIdx.x;
	const int y = blockIdx.y * blockDim.y + threadIdx.y;

	const uchar* src_y = (const uchar*)(img + y * img.step); //原图像
	uchar* dst_y = (uchar*)(dst + y * dst.step);

	if (x < width && y < height) {
		dst_y[3 * x] = 255 - src_y[3 * x];
		dst_y[3 * x + 1] = 255 - src_y[3 * x + 1];
		dst_y[3 * x + 2] = 255 - src_y[3 * x + 2];   //三通道
	}


}

首先核函数的第一个参数前两个参数height，width，是图像的高和宽。img是待处理图像在设备上的副本，d_dst是预先申请处理后的图像。前两个参数都是int类型，而后面两个是参数是PtrStepb类型，这在上节也有提到，它是OpenCV中的一个封装类，允许你以指针方式访问图像数据，同时提供了图像的宽度和步幅信息。这里采用的是一个指针的形式标注img与dst图像在Gpu中的位置。

我们想象一下，我们将一个3通道的图像存放在GPU中，我们对图像的操作本质上还是在一个3维的立体空间中，只不过实际情况下电脑是将图像数据展开成一维。每个像素点由一个线程标号。

再来看看核函数中对图像像素点的操作：

	if (x < width && y < height) {
		dst_y[3 * x] = 255 - src_y[3 * x];
		dst_y[3 * x + 1] = 255 - src_y[3 * x + 1];
		dst_y[3 * x + 2] = 255 - src_y[3 * x + 2];   //三通道
	}

我们所用的指引还是图像的行和列，x，y。

而里面的操作索引却采用 3*x，3*x+1，3*x+2，基于这样的索引可以明白，图像在GPU的存储方式应该是这样的。

前3个红绿蓝表示第一个像素点的值，依次类推。

至于反向操作，就是用255-原来图片的值。

来看看效果。