学习OpenCL与卷积神经网络 Day5——基于OpenCL的图像旋转以及高斯过滤的程序构建

最新推荐文章于 2024-06-23 14:25:49 发布

Ivan'yang

最新推荐文章于 2024-06-23 14:25:49 发布

阅读量881

点赞数 3

分类专栏： OpenCL 文章标签： opencl 计算机视觉

本文链接：https://blog.csdn.net/qq_31049419/article/details/111322732

版权

OpenCL 专栏收录该内容

4 篇文章 0 订阅

订阅专栏

前言

OpenCL其实算是一种“加速语言”，其由设备上运行的kernel函数语言和控制平台的API组成，它通过将某些简单而又重复的工作转交给GPU/FPGA外接设备，实现异构并行来加速原本的工作。比如在OpenCV中就引入了OCL module，其编译模块要比CUDA快很多，它就起到了一个加速图像处理过程的作用。当然，今天的工作只是简单的熟悉一下如何在OpenCL中对图像进行处理。

采样器

要做图像采集，我们肯定会需要一个采样器Sampler，但对于采样器对象的创建一般有两种方式：
1.直接在设备端创建采样器对象。
2.在主机端创建采样器对象，而后将采样器对象作为参数传给内核函数。
我将在图像旋转与高斯过滤中分别采取不同方式来创建采样器对象供大家参考。

设备端采样器创建

采样器在设备端的内核代码申明如下：

__const sampler_t sampler=sampler_properites;

其中sampler_properites为采样器属性值，它主要有三种状态（规格化坐标、寻址模式、过滤模式），我们都要对其进行设置。在设备端其属性如下：

规格化坐标：
CLK_NORMALIZED_COORDS_FALSE
CLK_NORMALIZED_COORDS_TRUE
寻址模式：
CLK_ADDRESS_MIRRORED_REPEAT
CLK_ADDRESS_REPEAT
CLK_ADDRESS_CLAMP_TO_EDGE
CLK_ADDRESS_CLAMP
CLK_ADDRESS_NONE
过滤模式：
CLK_FILTER_NEAREST
CLK_FILTER_LINEAR

主机端采样器创建

第一步(采样器声明)
采样器在主机端的内核代码声明有两种方法：
第一种：

cl_sampler clCreateSamplerWithProperties(
					cl_context context,
					const cl_sampler_properties *sampler_properties,
					cl_int *errcode_ret)

参数context为一个有效的OpenCL上下文。
sampler_properties为采样器属性名称与它们相应值。
errcode_ret为错误码返回值。
采样器属性名与相应值如下：

CLK_SAMPLER_NORMALIZED_COORDS 
//默认值为CL_TRUE(一个布尔值用来指定图像坐标是否归一化)
CLK_SAMPLER_ADDRESSING_MODE
//CLK_ADDRESS_MIRRORED_REPEAT
//CLK_ADDRESS_REPEAT
//CLK_ADDRESS_CLAMP_TO_EDGE
//CLK_ADDRESS_CLAMP
//CLK_ADDRESS_NONE
//(指定当读取图像坐落在范围外时的处理)
CLK_SAMPLER_FILTER_MODE
//CLK_FILTER_NEAREST
//CLK_FILTER_LINEAR
//(指定读取图片时的滤波模式)

在使用clCreateSamplerWithProperties函数时，一般要对其属性先行申明。
比如这样:

const cl_sampler_properties *sampler_properties[]=
{
	CLK_SAMPLER_NORMALIZED_COORDS,CL_FALSE,
	CLK_SAMPLER_ADDRESSING_MODE,CLK_ADDRESS_CLAMP_TO_EDGE,
	CLK_SAMPLER_FILTER_MODE,CLK_FILTER_LINEAR,
	0
	};
cl_sampler sampler=clCreatSamplerWithProperties(context,
												sampler_properties,
												&errNum);

第二种：

cl_sampler clCreateSampler(
			cl_context context, 
 			cl_bool normalized_coords, 
 			cl_addressing_mode addressing_mode, 
 			cl_filter_mode filter_mode, 
 			cl_int *errcode_ret)

参数与clCreateSamplerWithProperties几乎是相同的,但使用起来会相对方便。

下一步(传给内核)
在主机创建完采样器之后要传给内核函数

__kernel void imagec(__global sampler_t sampler,...)

下一步(在主机中参数传入)

clSetKernelArg(kernel,0,sizeof(cl_sampler),&sampler);

图像旋转

要将一张图片围绕中心点旋转45°，我们首先进行计算：
在一个平面直角坐标系中，我们要求一条过原点的直线围绕原点旋转一个角度θ，其旋转后的坐标。
假设原坐标为(x1,y1)，初始角度为δ，旋转角为θ，直线长R，旋转后的坐标为(x2,y2)
有
x1=Rcosδ
y1=Rsinδ
x2=Rcos(θ+δ)=cosθx1-sinθy1
y2=Rsin(θ+δ)=cosθy1+sinθx1

内核创建

这次我们在设备端创建采样器

__constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE |
CLK_FILTER_NEAREST |
CLK_ADDRESS_CLAMP;
__kernel void image_rotate(__read_only image2d_t srcImg,
	__write_only image2d_t dstImg,
	float angle)
{
	//根据图像对象获得图像长和宽
	int width = get_image_width(srcImg);
	int height = get_image_height(srcImg);
	const int x = get_global_id(0);
	const int y = get_global_id(1);
	float sinma = sin(angle);
	float cosma = cos(angle);
	//计算旋转中心点
	int hwidth = width / 2;
	int hheight = height / 2;
	int xt = x - hwidth;
	int yt = y - hheight;
	//计算旋转后坐标
	float2 readCoord;
	readCoord.x = (cosma*xt - sinma * yt) + hwidth;
	readCoord.y = (sinma*xt + cosma * yt) + hheight;
	//根据旋转后坐标读取原始图像元素值
	float4 value = read_imagef(srcImg, sampler, readCoord);
	write_imagef(dstImg, (int2)(x, y), value);
}

其中srcImg为输入图片，dstImg为输出图片，angle为旋转角度。

主机函数创建

这里我有用到FreeImage库来存储读取图片，大家可以自行在CSDN上下载，也可以到我这里下载 FreeImage.

首先我们要创建一个用来读取图像的函数：

cl_mem LoadImage(cl_context context, const char *filename, int &width, int &height)
{
	FREE_IMAGE_FORMAT format = FreeImage_GetFileType(filename, 0);
	FIBITMAP *image = FreeImage_Load(format, filename);
	FIBITMAP *temp = image;
	image = FreeImage_ConvertTo32Bits(image);
	FreeImage_Unload(temp);
	width = FreeImage_GetWidth(image);
	height = FreeImage_GetHeight(image);
	char *buffer = new char[width*height * 4];
	memcpy(buffer, FreeImage_GetBits(image), width*height * 4);
	FreeImage_Unload(image);
	cl_image_format clImageFormat;
	clImageFormat.image_channel_order = CL_RGBA;
	clImageFormat.image_channel_data_type = CL_UNORM_INT8;
	cl_image_desc clImageDesc;
	memset(&clImageDesc, 0, sizeof(cl_image_desc));
	clImageDesc.image_type = CL_MEM_OBJECT_IMAGE2D;
	clImageDesc.image_width = width;
	clImageDesc.image_height = height;
	cl_int errNum;
	cl_mem clImage;
	clImage = clCreateImage(context,
		CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
		&clImageFormat,
		&clImageDesc,
		buffer,
		&errNum);
	return clImage;
}

然后再创建一个存储图像的函数：

bool SaveImage(const char *fileName, char *buffer, int width, int height)
{
	FREE_IMAGE_FORMAT format = FreeImage_GetFIFFromFilename(fileName);
	FIBITMAP *image = FreeImage_ConvertFromRawBits((BYTE*)buffer, width,
		height, width * 4, 32,
		0xFF000000, 0x00FF0000, 0x0000FF00);
	return (FreeImage_Save(format, image, fileName) == TRUE) ? true : false;
}

创建上下文函数：

cl_context CreateContext()
{
	cl_int errNum;
	cl_uint numPlatforms;
	cl_platform_id firstPlatformId;
	cl_context context = NULL;

	// 选择第一个平台
	errNum = clGetPlatformIDs(1, &firstPlatformId, &numPlatforms);
	if (errNum != CL_SUCCESS || numPlatforms <= 0)
	{
		std::cerr << "Failed to find any OpenCL platforms." << std::endl;
		return NULL;
	}

	// 接下来尝试通过GPU设备建立上下文
	cl_context_properties contextProperties[] =
	{
		CL_CONTEXT_PLATFORM,
		(cl_context_properties)firstPlatformId,
		0
	};
	context = clCreateContextFromType(contextProperties, CL_DEVICE_TYPE_GPU,
		NULL, NULL, &errNum);
	if (errNum != CL_SUCCESS)
	{
		std::cout << "Could not create GPU context, trying CPU..." << std::endl;
		context = clCreateContextFromType(contextProperties, CL_DEVICE_TYPE_CPU,
			NULL, NULL, &errNum);
		if (errNum != CL_SUCCESS)
		{
			std::cerr << "Failed to create an OpenCL GPU or CPU context." << std::endl;
			return NULL;
		}
	}

	return context;
}

创建命令队列：

cl_command_queue CreateCommandQueue(cl_context context, cl_device_id *device)
{
	cl_int errNum;
	cl_device_id *devices;
	cl_command_queue commandQueue = NULL;
	size_t deviceBufferSize = -1;

	// 首先获得设备的信息
	errNum = clGetContextInfo(context, CL_CONTEXT_DEVICES, 0, NULL, &deviceBufferSize);
	if (errNum != CL_SUCCESS)
	{
		std::cerr << "Failed call to clGetContextInfo(...,GL_CONTEXT_DEVICES,...)";
		return NULL;
	}

	if (deviceBufferSize <= 0)
	{
		std::cerr << "No devices available.";
		return NULL;
	}

	//为设备分配内存
	devices = new cl_device_id[deviceBufferSize / sizeof(cl_device_id)];
	errNum = clGetContextInfo(context, CL_CONTEXT_DEVICES, deviceBufferSize, devices, NULL);
	if (errNum != CL_SUCCESS)
	{
		std::cerr << "Failed to get device IDs";
		return NULL;
	}

	// 选择第一个设备并为其创建命令队列
	commandQueue = clCreateCommandQueue(context, devices[0], 0, NULL);
	if (commandQueue == NULL)
	{
		std::cerr << "Failed to create commandQueue for device 0";
		return NULL;
	}

	//释放信息
	*device = devices[0];
	delete[] devices;
	return commandQueue;
}

创建OpenCL程序对象：

cl_program CreateProgram(cl_context context, cl_device_id device, const char* fileName)
{
	cl_int errNum;
	cl_program program;

	std::ifstream kernelFile(fileName, std::ios::in);
	if (!kernelFile.is_open())
	{
		std::cerr << "Failed to open file for reading: " << fileName << std::endl;
		return NULL;
	}

	std::ostringstream oss;
	oss << kernelFile.rdbuf();

	std::string srcStdStr = oss.str();
	const char *srcStr = srcStdStr.c_str();
	program = clCreateProgramWithSource(context, 1,
		(const char**)&srcStr,
		NULL, NULL);
	if (program == NULL)
	{
		std::cerr << "Failed to create CL program from source." << std::endl;
		return NULL;
	}

	errNum = clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
	if (errNum != CL_SUCCESS)
	{
		// 输出错误信息
		char buildLog[16384];
		clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG,
			sizeof(buildLog), buildLog, NULL);

		std::cerr << "Error in kernel: " << std::endl;
		std::cerr << buildLog;
		clReleaseProgram(program);
		return NULL;
	}

	return program;
}

最后将这些函数在main函数中引用：

int main()
{
	cl_context context = 0;
	cl_command_queue commandQueue = 0;
	cl_program program = 0;
	cl_device_id device = 0;
	cl_kernel kernel = 0;
	cl_mem imageObjects[2] = { 0, 0 };
	cl_int errNum;
	context = CreateContext();
	commandQueue = CreateCommandQueue(context, &device);
	cl_bool imageSupport = CL_FALSE;
	clGetDeviceInfo(device, CL_DEVICE_IMAGE_SUPPORT, sizeof(cl_bool),
		&imageSupport, NULL);
	int width, height;
	imageObjects[0] = LoadImage(context, "D:\\pic\\picture.jpg", width, height);
	cl_image_format clImageFormat;
	clImageFormat.image_channel_order = CL_RGBA;
	clImageFormat.image_channel_data_type = CL_UNORM_INT8;
	imageObjects[1] = clCreateImage2D(context,
		CL_MEM_WRITE_ONLY,
		&clImageFormat,
		width,
		height,
		0,
		NULL,
		&errNum);
	program = CreateProgram(context, device, "image_rotate.cl");
	kernel = clCreateKernel(program, "image_rotate", NULL);
	float angle = 45 * PI / 180.0f;
	errNum = clSetKernelArg(kernel, 0, sizeof(cl_mem), &imageObjects[0]);
	errNum |= clSetKernelArg(kernel, 1, sizeof(cl_mem), &imageObjects[1]);	
	errNum |= clSetKernelArg(kernel, 2, sizeof(cl_float), &angle);
	size_t globalWorkSize[2] = {  width,height };
	errNum = clEnqueueNDRangeKernel(commandQueue, kernel, 2, NULL,
		globalWorkSize, NULL,
		0, NULL, NULL);
	char *buffer = new char[width * height * 4];
	size_t origin[3] = { 0, 0, 0 };
	size_t region[3] = { size_t(width), size_t(height), 1 };
	errNum = clEnqueueReadImage(commandQueue, imageObjects[1], CL_TRUE,
		origin, region, 0, 0, buffer,
		0, NULL, NULL);
	if (errNum != CL_SUCCESS)
	{
		std::cerr << "Error reading result buffer." << std::endl;
		return 1;
	}
	std::cout << std::endl;
	std::cout << "Executed program succesfully." << std::endl;
	if (!SaveImage("GPU_Rotate.png", buffer, width, height))
	{
		std::cerr << "Error writing output image: " << "GPU_Rotate.png" << std::endl;
		delete[] buffer;
		return 1;
	}

	delete[] buffer;

	return 0;
}

完整程序

#ifdef __APPLE__
#include <OpenCL/cl.h>
#else
#include <CL/cl.h>
#endif
#include<stdio.h>
#include<stdlib.h>
#define MAX_SOURCE_SIZE (0x10000)
#include<string.h>
#include <iostream>  
#include<math.h>
#include "FreeImage.h"
#include <opencv2/opencv.hpp>
using namespace cv;
using namespace std;
#define PI 3.1415926
#pragma warning(disable : 4996)

cl_mem LoadImage(cl_context context, const char *filename, int &width, int &height)
{
	FREE_IMAGE_FORMAT format = FreeImage_GetFileType(filename, 0);
	FIBITMAP *image = FreeImage_Load(format, filename);
	FIBITMAP *temp = image;
	image = FreeImage_ConvertTo32Bits(image);
	FreeImage_Unload(temp);
	width = FreeImage_GetWidth(image);
	height = FreeImage_GetHeight(image);
	char *buffer = new char[width*height * 4];
	memcpy(buffer, FreeImage_GetBits(image), width*height * 4);
	FreeImage_Unload(image);
	cl_image_format clImageFormat;
	clImageFormat.image_channel_order = CL_RGBA;
	clImageFormat.image_channel_data_type = CL_UNORM_INT8;
	cl_image_desc clImageDesc;
	memset(&clImageDesc, 0, sizeof(cl_image_desc));
	clImageDesc.image_type = CL_MEM_OBJECT_IMAGE2D;
	clImageDesc.image_width = width;
	clImageDesc.image_height = height;
	cl_int errNum;
	cl_mem clImage;
	clImage = clCreateImage(context,
		CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
		&clImageFormat,
		&clImageDesc,
		buffer,
		&errNum);
	return clImage;
}

bool SaveImage(const char *fileName, char *buffer, int width, int height)
{
	FREE_IMAGE_FORMAT format = FreeImage_GetFIFFromFilename(fileName);
	FIBITMAP *image = FreeImage_ConvertFromRawBits((BYTE*)buffer, width,
		height, width * 4, 32,
		0xFF000000, 0x00FF0000, 0x0000FF00);
	return (FreeImage_Save(format, image, fileName) == TRUE) ? true : false;
}

cl_context CreateContext()
{
	cl_int errNum;
	cl_uint numPlatforms;
	cl_platform_id firstPlatformId;
	cl_context context = NULL;

	// 选择第一个平台
	errNum = clGetPlatformIDs(1, &firstPlatformId, &numPlatforms);
	if (errNum != CL_SUCCESS || numPlatforms <= 0)
	{
		std::cerr << "Failed to find any OpenCL platforms." << std::endl;
		return NULL;
	}

	// 接下来尝试通过GPU设备建立上下文
	cl_context_properties contextProperties[] =
	{
		CL_CONTEXT_PLATFORM,
		(cl_context_properties)firstPlatformId,
		0
	};
	context = clCreateContextFromType(contextProperties, CL_DEVICE_TYPE_GPU,
		NULL, NULL, &errNum);
	if (errNum != CL_SUCCESS)
	{
		std::cout << "Could not create GPU context, trying CPU..." << std::endl;
		context = clCreateContextFromType(contextProperties, CL_DEVICE_TYPE_CPU,
			NULL, NULL, &errNum);
		if (errNum != CL_SUCCESS)
		{
			std::cerr << "Failed to create an OpenCL GPU or CPU context." << std::endl;
			return NULL;
		}
	}

	return context;
}


cl_command_queue CreateCommandQueue(cl_context context, cl_device_id *device)
{
	cl_int errNum;
	cl_device_id *devices;
	cl_command_queue commandQueue = NULL;
	size_t deviceBufferSize = -1;

	// 首先获得设备的信息
	errNum = clGetContextInfo(context, CL_CONTEXT_DEVICES, 0, NULL, &deviceBufferSize);
	if (errNum != CL_SUCCESS)
	{
		std::cerr << "Failed call to clGetContextInfo(...,GL_CONTEXT_DEVICES,...)";
		return NULL;
	}

	if (deviceBufferSize <= 0)
	{
		std::cerr << "No devices available.";
		return NULL;
	}

	//为设备分配内存
	devices = new cl_device_id[deviceBufferSize / sizeof(cl_device_id)];
	errNum = clGetContextInfo(context, CL_CONTEXT_DEVICES, deviceBufferSize, devices, NULL);
	if (errNum != CL_SUCCESS)
	{
		std::cerr << "Failed to get device IDs";
		return NULL;
	}

	// 选择第一个设备并为其创建命令队列
	commandQueue = clCreateCommandQueue(context, devices[0], 0, NULL);
	if (commandQueue == NULL)
	{
		std::cerr << "Failed to create commandQueue for device 0";
		return NULL;
	}

	//释放信息
	*device = devices[0];
	delete[] devices;
	return commandQueue;
}


cl_program CreateProgram(cl_context context, cl_device_id device, const char* fileName)
{
	cl_int errNum;
	cl_program program;

	std::ifstream kernelFile(fileName, std::ios::in);
	if (!kernelFile.is_open())
	{
		std::cerr << "Failed to open file for reading: " << fileName << std::endl;
		return NULL;
	}

	std::ostringstream oss;
	oss << kernelFile.rdbuf();

	std::string srcStdStr = oss.str();
	const char *srcStr = srcStdStr.c_str();
	program = clCreateProgramWithSource(context, 1,
		(const char**)&srcStr,
		NULL, NULL);
	if (program == NULL)
	{
		std::cerr << "Failed to create CL program from source." << std::endl;
		return NULL;
	}

	errNum = clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
	if (errNum != CL_SUCCESS)
	{
		// 输出错误信息
		char buildLog[16384];
		clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG,
			sizeof(buildLog), buildLog, NULL);

		std::cerr << "Error in kernel: " << std::endl;
		std::cerr << buildLog;
		clReleaseProgram(program);
		return NULL;
	}

	return program;
}
int main()
{
	cl_context context = 0;
	cl_command_queue commandQueue = 0;
	cl_program program = 0;
	cl_device_id device = 0;
	cl_kernel kernel = 0;
	cl_mem imageObjects[2] = { 0, 0 };
	cl_int errNum;
	context = CreateContext();
	commandQueue = CreateCommandQueue(context, &device);
	cl_bool imageSupport = CL_FALSE;
	clGetDeviceInfo(device, CL_DEVICE_IMAGE_SUPPORT, sizeof(cl_bool),
		&imageSupport, NULL);
	int width, height;
	imageObjects[0] = LoadImage(context, "D:\\pic\\picture.jpg", width, height);
	cl_image_format clImageFormat;
	clImageFormat.image_channel_order = CL_RGBA;
	clImageFormat.image_channel_data_type = CL_UNORM_INT8;
	imageObjects[1] = clCreateImage2D(context,
		CL_MEM_WRITE_ONLY,
		&clImageFormat,
		width,
		height,
		0,
		NULL,
		&errNum);
	program = CreateProgram(context, device, "image_rotate.cl");
	kernel = clCreateKernel(program, "image_rotate", NULL);
	float angle = 45 * PI / 180.0f;
	errNum = clSetKernelArg(kernel, 0, sizeof(cl_mem), &imageObjects[0]);
	errNum |= clSetKernelArg(kernel, 1, sizeof(cl_mem), &imageObjects[1]);	
	errNum |= clSetKernelArg(kernel, 2, sizeof(cl_float), &angle);
	size_t globalWorkSize[2] = {  width,height };
	errNum = clEnqueueNDRangeKernel(commandQueue, kernel, 2, NULL,
		globalWorkSize, NULL,
		0, NULL, NULL);
	char *buffer = new char[width * height * 4];
	size_t origin[3] = { 0, 0, 0 };
	size_t region[3] = { size_t(width), size_t(height), 1 };
	errNum = clEnqueueReadImage(commandQueue, imageObjects[1], CL_TRUE,
		origin, region, 0, 0, buffer,
		0, NULL, NULL);
	if (errNum != CL_SUCCESS)
	{
		std::cerr << "Error reading result buffer." << std::endl;
		return 1;
	}
	std::cout << std::endl;
	std::cout << "Executed program succesfully." << std::endl;
	if (!SaveImage("GPU_Rotate.png", buffer, width, height))
	{
		std::cerr << "Error writing output image: " << "GPU_Rotate.png" << std::endl;
		delete[] buffer;
		return 1;
	}
	delete[] buffer;
	return 0;
}

旋转输出结果

在这里插入图片描述

高斯过滤

高斯滤波矩阵为：在这里插入图片描述
对于高斯过滤（模糊处理），我将采样器在主机端创建。

内核创建

__kernel void gs(__read_only image2d_t srcImg,
	__write_only image2d_t dstImg,
	sampler_t sampler,
	int width, int height)
{
	float kernelWeights[9] = { 1.0f, 2.0f, 1.0f,
							   2.0f, 4.0f, 2.0f,
							   1.0f, 2.0f, 1.0f };

	int2 startImageCoord = (int2) (get_global_id(0) - 1, get_global_id(1) - 1);
	int2 endImageCoord = (int2) (get_global_id(0) + 1, get_global_id(1) + 1);
	int2 outImageCoord = (int2) (get_global_id(0), get_global_id(1));

	if (outImageCoord.x < width && outImageCoord.y < height)
	{
		int weight = 0;
		float4 outColor = (float4)(0.0f, 0.0f, 0.0f, 0.0f);
		for (int y = startImageCoord.y; y <= endImageCoord.y; y++)
		{
			for (int x = startImageCoord.x; x <= endImageCoord.x; x++)
			{
				outColor += (read_imagef(srcImg, sampler, (int2)(x, y)) *(kernelWeights[weight] / 16.0f));
				weight += 1;
			}
		}
		write_imagef(dstImg, outImageCoord, outColor);
	}
}

主机函数创建

步骤与之前的图像旋转基本一致，但这次的采样器在主机端创建。

sampler = clCreateSampler(context,
		CL_FALSE, 
		CL_ADDRESS_CLAMP_TO_EDGE,
		CL_FILTER_NEAREST,
		&errNum);

完整程序

#ifdef __APPLE__
#include <OpenCL/cl.h>
#else
#include <CL/cl.h>
#endif
#include<stdio.h>
#include<stdlib.h>
#define MAX_SOURCE_SIZE (0x10000)
#include<string.h>
#include <iostream>  
#include<math.h>
#include "FreeImage.h"
#include <opencv2/opencv.hpp>
using namespace cv;
using namespace std;
#define PI 3.1415926
#pragma warning(disable : 4996)

cl_mem LoadImage(cl_context context, const char *filename, int &width, int &height)
{
	FREE_IMAGE_FORMAT format = FreeImage_GetFileType(filename, 0);
	FIBITMAP *image = FreeImage_Load(format, filename);
	FIBITMAP *temp = image;
	image = FreeImage_ConvertTo32Bits(image);
	FreeImage_Unload(temp);
	width = FreeImage_GetWidth(image);
	height = FreeImage_GetHeight(image);
	char *buffer = new char[width*height * 4];
	memcpy(buffer, FreeImage_GetBits(image), width*height * 4);
	FreeImage_Unload(image);
	cl_image_format clImageFormat;
	clImageFormat.image_channel_order = CL_RGBA;
	clImageFormat.image_channel_data_type = CL_UNORM_INT8;
	cl_image_desc clImageDesc;
	memset(&clImageDesc, 0, sizeof(cl_image_desc));
	clImageDesc.image_type = CL_MEM_OBJECT_IMAGE2D;
	clImageDesc.image_width = width;
	clImageDesc.image_height = height;
	cl_int errNum;
	cl_mem clImage;
	clImage = clCreateImage(context,
		CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
		&clImageFormat,
		&clImageDesc,
		buffer,
		&errNum);
	return clImage;
}

bool SaveImage(const char *fileName, char *buffer, int width, int height)
{
	FREE_IMAGE_FORMAT format = FreeImage_GetFIFFromFilename(fileName);
	FIBITMAP *image = FreeImage_ConvertFromRawBits((BYTE*)buffer, width,
		height, width * 4, 32,
		0xFF000000, 0x00FF0000, 0x0000FF00);
	return (FreeImage_Save(format, image, fileName) == TRUE) ? true : false;
}
size_t RoundUp(int groupSize, int globalSize)
{
	int r = globalSize % groupSize;
	if (r == 0)
	{
		return globalSize;
	}
	else
	{
		return globalSize + groupSize - r;
	}
}

cl_context CreateContext()
{
	cl_int errNum;
	cl_uint numPlatforms;
	cl_platform_id firstPlatformId;
	cl_context context = NULL;

	// 选择第一个平台
	errNum = clGetPlatformIDs(1, &firstPlatformId, &numPlatforms);
	if (errNum != CL_SUCCESS || numPlatforms <= 0)
	{
		std::cerr << "Failed to find any OpenCL platforms." << std::endl;
		return NULL;
	}

	// 接下来尝试通过GPU设备建立上下文
	cl_context_properties contextProperties[] =
	{
		CL_CONTEXT_PLATFORM,
		(cl_context_properties)firstPlatformId,
		0
	};
	context = clCreateContextFromType(contextProperties, CL_DEVICE_TYPE_GPU,
		NULL, NULL, &errNum);
	if (errNum != CL_SUCCESS)
	{
		std::cout << "Could not create GPU context, trying CPU..." << std::endl;
		context = clCreateContextFromType(contextProperties, CL_DEVICE_TYPE_CPU,
			NULL, NULL, &errNum);
		if (errNum != CL_SUCCESS)
		{
			std::cerr << "Failed to create an OpenCL GPU or CPU context." << std::endl;
			return NULL;
		}
	}

	return context;
}


cl_command_queue CreateCommandQueue(cl_context context, cl_device_id *device)
{
	cl_int errNum;
	cl_device_id *devices;
	cl_command_queue commandQueue = NULL;
	size_t deviceBufferSize = -1;

	// 首先获得设备的信息
	errNum = clGetContextInfo(context, CL_CONTEXT_DEVICES, 0, NULL, &deviceBufferSize);
	if (errNum != CL_SUCCESS)
	{
		std::cerr << "Failed call to clGetContextInfo(...,GL_CONTEXT_DEVICES,...)";
		return NULL;
	}

	if (deviceBufferSize <= 0)
	{
		std::cerr << "No devices available.";
		return NULL;
	}

	//为设备分配内存
	devices = new cl_device_id[deviceBufferSize / sizeof(cl_device_id)];
	errNum = clGetContextInfo(context, CL_CONTEXT_DEVICES, deviceBufferSize, devices, NULL);
	if (errNum != CL_SUCCESS)
	{
		std::cerr << "Failed to get device IDs";
		return NULL;
	}

	// 选择第一个设备并为其创建命令队列
	commandQueue = clCreateCommandQueue(context, devices[0], 0, NULL);
	if (commandQueue == NULL)
	{
		std::cerr << "Failed to create commandQueue for device 0";
		return NULL;
	}

	//释放信息
	*device = devices[0];
	delete[] devices;
	return commandQueue;
}


cl_program CreateProgram(cl_context context, cl_device_id device, const char* fileName)
{
	cl_int errNum;
	cl_program program;

	std::ifstream kernelFile(fileName, std::ios::in);
	if (!kernelFile.is_open())
	{
		std::cerr << "Failed to open file for reading: " << fileName << std::endl;
		return NULL;
	}

	std::ostringstream oss;
	oss << kernelFile.rdbuf();

	std::string srcStdStr = oss.str();
	const char *srcStr = srcStdStr.c_str();
	program = clCreateProgramWithSource(context, 1,
		(const char**)&srcStr,
		NULL, NULL);
	if (program == NULL)
	{
		std::cerr << "Failed to create CL program from source." << std::endl;
		return NULL;
	}

	errNum = clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
	if (errNum != CL_SUCCESS)
	{
		// 输出错误信息
		char buildLog[16384];
		clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG,
			sizeof(buildLog), buildLog, NULL);

		std::cerr << "Error in kernel: " << std::endl;
		std::cerr << buildLog;
		clReleaseProgram(program);
		return NULL;
	}

	return program;
}

	int main()
{
	cl_context context = 0;
	cl_command_queue commandQueue = 0;
	cl_program program = 0;
	cl_device_id device = 0;
	cl_kernel kernel = 0;
	cl_mem imageObjects[2] = { 0, 0 };
	cl_sampler sampler = 0;
	cl_int errNum;
	context = CreateContext();
	commandQueue = CreateCommandQueue(context, &device);
	cl_bool imageSupport = CL_FALSE;
	clGetDeviceInfo(device, CL_DEVICE_IMAGE_SUPPORT, sizeof(cl_bool),
		&imageSupport, NULL);
	int width, height;
	imageObjects[0] = LoadImage(context, "D:\\pic\\picture.jpg", width, height);
	cl_image_format clImageFormat;
	clImageFormat.image_channel_order = CL_RGBA;
	clImageFormat.image_channel_data_type = CL_UNORM_INT8;
	imageObjects[1] = clCreateImage2D(context,
		CL_MEM_WRITE_ONLY,
		&clImageFormat,
		width,
		height,
		0,
		NULL,
		&errNum);

	// 创建采样器对象
	sampler = clCreateSampler(context,
		CL_FALSE, 
		CL_ADDRESS_CLAMP_TO_EDGE,
		CL_FILTER_NEAREST,
		&errNum);

	program = CreateProgram(context, device, "gs.cl");
	kernel = clCreateKernel(program, "gs", NULL);
	errNum = clSetKernelArg(kernel, 0, sizeof(cl_mem), &imageObjects[0]);
	errNum |= clSetKernelArg(kernel, 1, sizeof(cl_mem), &imageObjects[1]);
	errNum |= clSetKernelArg(kernel, 2, sizeof(cl_sampler), &sampler);
	errNum |= clSetKernelArg(kernel, 3, sizeof(cl_int), &width);
	errNum |= clSetKernelArg(kernel, 4, sizeof(cl_int), &height);
	size_t localWorkSize[2] = { 16, 16 };
	size_t globalWorkSize[2] = { RoundUp(localWorkSize[0], width),
								  RoundUp(localWorkSize[1], height) };
	errNum = clEnqueueNDRangeKernel(commandQueue, kernel, 2, NULL,
		globalWorkSize, localWorkSize,
		0, NULL, NULL);
	char *buffer = new char[width * height * 4];
	size_t origin[3] = { 0, 0, 0 };
	size_t region[3] = { size_t(width), size_t(height), 1 };
	errNum = clEnqueueReadImage(commandQueue, imageObjects[1], CL_TRUE,
		origin, region, 0, 0, buffer,
		0, NULL, NULL);
	if (errNum != CL_SUCCESS)
	{
		std::cerr << "Error reading result buffer." << std::endl;		
		return 1;
	}
	std::cout << std::endl;
	std::cout << "Executed program succesfully." << std::endl;
	if (!SaveImage("gs_out.png", buffer, width, height))
	{
		std::cerr << "Error writing output image: " << "gs_out.png" << std::endl;
		delete[] buffer;
		return 1;
	}
	delete[] buffer;
	return 0;