Example of accelerating image scaling using Cuda

First, the preface

This paper mainly explains a small example of cuda parallel acceleration, and accelerates the nearest neighbor interpolation algorithm for image scaling.

second, code implementation

Because each new pixel is calculated in the same way as it is scaled, parallel computations can be used, as do the resize in OpenCV.

main.cu////#include "cuda_runtime.h" #include <windows.h> #include <iostream> #include <opencv2/cor e/core.hpp> #include <opencv2/highgui/highgui.hpp> #include <opencv2/imgproc/imgproc.hpp> using
namespace CV;

using namespace Std;
	void Resizeimage (const Mat &AMP;_SRC, Mat &_dst, const Size &s) {_dst = Mat::zeros (s, CV_8UC3);
	Double frows = s.height/(float) _src.rows;
	Double fcols = s.width/(float) _src.cols;
	int PX = 0;
	int PY = 0;
			for (int i = 0; I!= _dst.rows, ++i) {for (int j = 0; J!= _dst.cols; ++j) {PX = Cvround (i/(double) frows);
			PY = Cvround (j/(double) fcols); if (PX < _src.rows && PX >= 0 && PY < _src.cols && PY >= 0) {_dst.at<vec3b> (
				I, J) [0] = _src.at<vec3b> (PX, PY) [0];
				_dst.at<vec3b> (i, J) [1] = _src.at<vec3b> (PX, PY) [1];
			_dst.at<vec3b> (i, J) [2] = _src.at<vec3b> (PX, PY) [2];
	BOOL Initcuda () {int count; CudagetdEvicecount (&count);
		if (count = = 0) {fprintf (stderr, "There is no device.\n");
	return false;
	} int i;
		for (i = 0; i < count; i++) {Cudadeviceprop prop;
			if (Cudagetdeviceproperties (&prop, i) = = cudasuccess) {if (prop.major >= 1) {break;
		}} if (i = = count) {fprintf (stderr, "There is no device supporting CUDA 1.x.\n");
	return false;
	} cudasetdevice (i);
return true; } __global__ void Kernel (uchar* _src_dev, Uchar * _dst_dev, int _src_step, int _dst_step, int _src_rows, int _src_cols
	, int _dst_rows, int _dst_cols) {int i = blockidx.x;

	int j = blockidx.y;
	Double frows = _dst_rows/(float) _src_rows;

	Double fcols = _dst_cols/(float) _src_cols;
	int PX = 0;

	int PY = 0;
	PX = (int) (i/frows);
	PY = (int) (j/fcols); if (PX < _src_rows && PX >= 0 && PY < _src_cols && PY >= 0) {* (_dst_dev + i*_dst_ste
		P + 3 * j + 0) = * (_src_dev + px*_src_step + 3 * pY); * (_dst_dev + i*_dst_step + 3 * j + 1) = * (_src_dev + px*_src_step + 3 * PY + 1);
	
	* (_dst_dev + i*_dst_step + 3 * j + 2) = * (_src_dev + px*_src_step + 3 * PY + 2);
	} void Resizeimagegpu (const Mat &AMP;_SRC, Mat &_dst, const Size &s) {_dst = Mat (s, CV_8UC3);
	Uchar *src_data = _src.data;
	int width = _src.cols;
	int height = _src.rows;

	Uchar *src_dev, *dst_dev;
	Cudamalloc ((void**) &src_dev, 3 * width*height * sizeof (UCHAR));
	Cudamalloc ((void**) &dst_dev, 3 * s.width * s.height * sizeof (UCHAR));

	cudamemcpy (Src_dev, Src_data, 3 * width*height * sizeof (UCHAR), cudamemcpyhosttodevice);
	Double frows = s.height/(float) _src.rows;
	Double fcols = s.width/(float) _src.cols;
	int src_step = _src.step;

	int dst_step = _dst.step;
	Dim3 Grid (S.height, s.width);

	Kernel << < grid, 1 >> > (Src_dev, Dst_dev, Src_step, dst_step, height, width, s.height, s.width);


cudamemcpy (_dst.data, Dst_dev, 3 * s.width * s.height * sizeof (UCHAR), cudamemcpydevicetohost); int main () {Mat SRc = Cv::imread ("e:\\ learning data \ Test Standard Chart \\lena.bmp", 1);

	Mat dst_cpu;
	Double start = GetTickCount ();
	Resizeimage (SRC, dst_cpu, Size (Src.cols * 2, Src.rows * 2));
	
	Double end = GetTickCount ();

	cout << "CPU Scaling Time:" << end-start << "\ n";

	Initcuda ();

	Mat Dst_gpu;
	Start = GetTickCount ();
	Resizeimagegpu (SRC, Dst_gpu, Size (Src.cols * 2, Src.rows * 2));
	End = GetTickCount ();

	cout << "GPU Scaling Time:" << end-start << "\ n";
	Cv::imshow ("Demo", DST_CPU);

	Waitkey (0);
return 0; }

iii. Results of the experiment

In this paper, the experimental environment is vs2013+cuda7.0+opencv2.4.9, the results can be obtained as follows, when the 512*512 Lena image is magnified to 1024*1024, the method of using GPU parallel computation accelerates by more than one time, but for narrowing operations, Using GPU acceleration is not necessarily fast because uploading data takes time.