opencl+opencv实现sobel算法

浅浅的花香味﹌ 2022-08-01 13:51 152阅读 0赞

这几天在看opencl编程指南，照着书中的例子实现了sobel算法：

1.结合opencv读取图像，保存到缓冲区中；

2.编写和编译内核，并保存显示处理后的结果。

内核：

const sampler_t sampler = CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;
    kernel void sobel_rgb(read_only image2d_t src,write_only image2d_t dst)
    {
    	int x = (int)get_global_id(0);
    	int y = (int)get_global_id(1);
    	if (x >= get_image_width(src) || y >= get_image_height(src))
    		return;
    	float4 p00 = read_imagef(src, sampler, (int2)(x - 1, y - 1));
    	float4 p10 = read_imagef(src, sampler, (int2)(x, y - 1));
    	float4 p20 = read_imagef(src, sampler, (int2)(x + 1, y - 1));
    	float4 p01 = read_imagef(src, sampler, (int2)(x - 1, y));
    	float4 p21 = read_imagef(src, sampler, (int2)(x + 1, y));
    	float4 p02 = read_imagef(src, sampler, (int2)(x - 1, y + 1));
    	float4 p12 = read_imagef(src, sampler, (int2)(x, y + 1));
    	float4 p22 = read_imagef(src, sampler, (int2)(x + 1, y + 1));
    	float3 gx = -p00.xyz + p20.xyz + 2.0*(p21.xyz - p01.xyz) - p02.xyz + p22.xyz;
    	float3 gy = -p00.xyz + p02.xyz + 2.0*(p21.xyz - p10.xyz) - p20.xyz + p22.xyz;
    	float3 g = native_sqrt(gx*gx + gy*gy);
    	write_imagef(dst,(int2)(x,y),(float4)(g.x,g.y,g.z,1.0f));
    }
    // TODO: Add OpenCL kernel code here.

c++源码：

//
    // Book:      OpenCL(R) Programming Guide
    // Authors:   Aaftab Munshi, Benedict Gaster, Timothy Mattson, James Fung, Dan Ginsburg
    // ISBN-10:   0-321-74964-2
    // ISBN-13:   978-0-321-74964-2
    // Publisher: Addison-Wesley Professional
    // URLs:      http://safari.informit.com/9780132488006/
    //            http://www.openclprogrammingguide.com
    //
    
    // ImageFilter2D.cpp
    //
    //    This example demonstrates performing gaussian filtering on a 2D image using
    //    OpenCL
    //
    //    Requires FreeImage library for image I/O:
    //      http://freeimage.sourceforge.net/
    
    #include <iostream>
    #include <fstream>
    #include <sstream>
    #include <string.h>
    #include <opencv.hpp>
    
    #ifdef __APPLE__
    #include <OpenCL/cl.h>
    #else
    #include <CL/cl.h>
    #endif
    
    #include "FreeImage.h"
    
    ///
    //  Create an OpenCL context on the first available platform using
    //  either a GPU or CPU depending on what is available.
    //
    cl_context CreateContext()
    {
    	cl_int errNum;
    	cl_uint numPlatforms;
    	cl_platform_id firstPlatformId;
    	cl_context context = NULL;
    
    	// First, select an OpenCL platform to run on.  For this example, we
    	// simply choose the first available platform.  Normally, you would
    	// query for all available platforms and select the most appropriate one.
    	errNum = clGetPlatformIDs(1, &firstPlatformId, &numPlatforms);
    	if (errNum != CL_SUCCESS || numPlatforms <= 0)
    	{
    		std::cerr << "Failed to find any OpenCL platforms." << std::endl;
    		return NULL;
    	}
    
    	// Next, create an OpenCL context on the platform.  Attempt to
    	// create a GPU-based context, and if that fails, try to create
    	// a CPU-based context.
    	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;
    }
    
    ///
    //  Create a command queue on the first device available on the
    //  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;
    
    	// First get the size of the devices buffer
    	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;
    	}
    
    	// Allocate memory for the devices buffer
    	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;
    	}
    
    	// In this example, we just choose the first available device.  In a
    	// real program, you would likely use all available devices or choose
    	// the highest performance device based on OpenCL device queries
    	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;
    }
    
    ///
    //  Create an OpenCL program from the kernel source file
    //
    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)
    	{
    		// Determine the reason for the error
    		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;
    }
    
    
    ///
    //  Cleanup any created OpenCL resources
    //
    void Cleanup(cl_context context, cl_command_queue commandQueue,
    	cl_program program, cl_kernel kernel, cl_mem imageObjects[2],
    	cl_sampler sampler)
    {
    	for (int i = 0; i < 2; i++)
    	{
    		if (imageObjects[i] != 0)
    			clReleaseMemObject(imageObjects[i]);
    	}
    	if (commandQueue != 0)
    		clReleaseCommandQueue(commandQueue);
    
    	if (kernel != 0)
    		clReleaseKernel(kernel);
    
    	if (program != 0)
    		clReleaseProgram(program);
    
    	if (sampler != 0)
    		clReleaseSampler(sampler);
    
    	if (context != 0)
    		clReleaseContext(context);
    
    }
    
    ///
    //  Load an image using the FreeImage library and create an OpenCL
    //  image out of it
    //
    cl_mem LoadImage(cl_context context, char *fileName, int &width, int &height)
    {
    	//FREE_IMAGE_FORMAT format = FreeImage_GetFileType(fileName, 0);
    	//FIBITMAP* image = FreeImage_Load(format, fileName);
    
    	 Convert to 32-bit image
    	//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);*/
    	cv::Mat image1 = cv::imread(fileName);
    	width = image1.cols;
    	height = image1.rows;
    	char *buffer = new char[width * height * 4];
    	int w = 0;
    	for (int v = height - 1; v >= 0; v--)
    	{
    		for (int u = 0; u <width; u++)
    		{
    			buffer[w++] = image1.at<cv::Vec3b>(v, u)[0];
    			buffer[w++] = image1.at<cv::Vec3b>(v, u)[1];
    			buffer[w++] = image1.at<cv::Vec3b>(v, u)[2];
    			w++;
    		}
    	}
    	
    
    	// Create OpenCL image
    	cl_image_format clImageFormat;
    	clImageFormat.image_channel_order = CL_RGBA;
    	clImageFormat.image_channel_data_type = CL_UNORM_INT8;
    
    	cl_int errNum;
    	cl_mem clImage;
    	clImage = clCreateImage2D(context,
    		CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
    		&clImageFormat,
    		width,
    		height,
    		0,
    		buffer,
    		&errNum);
    
    	if (errNum != CL_SUCCESS)
    	{
    		std::cerr << "Error creating CL image object" << std::endl;
    		return 0;
    	}
    
    	return clImage;
    }
    
    ///
    //  Save an image using the FreeImage library
    //
    bool SaveImage(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);
    }
    
    ///
    //  Round up to the nearest multiple of the group size
    //
    size_t RoundUp(int groupSize, int globalSize)
    {
    	int r = globalSize % groupSize;
    	if (r == 0)
    	{
    		return globalSize;
    	}
    	else
    	{
    		return globalSize + groupSize - r;
    	}
    }
    
    ///
    //	main() for HelloBinaryWorld example
    //
    int main(int argc, char** argv)
    {
    	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;
    
    
    	/*if (argc != 3)
    	{
    		std::cerr << "USAGE: " << argv[0] << " <inputImageFile> <outputImageFiles>" << std::endl;
    		return 1;
    	}*/
    
    	// Create an OpenCL context on first available platform
    	context = CreateContext();
    	if (context == NULL)
    	{
    		std::cerr << "Failed to create OpenCL context." << std::endl;
    		return 1;
    	}
    
    	// Create a command-queue on the first device available
    	// on the created context
    	commandQueue = CreateCommandQueue(context, &device);
    	if (commandQueue == NULL)
    	{
    		Cleanup(context, commandQueue, program, kernel, imageObjects, sampler);
    		return 1;
    	}
    
    	// Make sure the device supports images, otherwise exit
    	cl_bool imageSupport = CL_FALSE;
    	clGetDeviceInfo(device, CL_DEVICE_IMAGE_SUPPORT, sizeof(cl_bool),
    		&imageSupport, NULL);
    	if (imageSupport != CL_TRUE)
    	{
    		std::cerr << "OpenCL device does not support images." << std::endl;
    		Cleanup(context, commandQueue, program, kernel, imageObjects, sampler);
    		return 1;
    	}
    
    	// Load input image from file and load it into
    	// an OpenCL image object
    	int width, height;
    	char *src0 = "C:/Users/jiang/Desktop/image/tu1.jpg";
    	imageObjects[0] = LoadImage(context, src0, width, height);
    	if (imageObjects[0] == 0)
    	{
    		std::cerr << "Error loading: " << std::string(src0) << std::endl;
    		Cleanup(context, commandQueue, program, kernel, imageObjects, sampler);
    		return 1;
    	}
    
    	// Create ouput image object
    	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);
    
    	if (errNum != CL_SUCCESS)
    	{
    		std::cerr << "Error creating CL output image object." << std::endl;
    		Cleanup(context, commandQueue, program, kernel, imageObjects, sampler);
    		return 1;
    	}
    
    
    	// Create sampler for sampling image object
    	sampler = clCreateSampler(context,
    		CL_FALSE, // Non-normalized coordinates
    		CL_ADDRESS_CLAMP_TO_EDGE,
    		CL_FILTER_NEAREST,
    		&errNum);
    
    	if (errNum != CL_SUCCESS)
    	{
    		std::cerr << "Error creating CL sampler object." << std::endl;
    		Cleanup(context, commandQueue, program, kernel, imageObjects, sampler);
    		return 1;
    	}
    
    	// Create OpenCL program
    	//program = CreateProgram(context, device, "ImageFilter2D.cl");
    	program = CreateProgram(context, device, "Sobel.cl");
    	if (program == NULL)
    	{
    		Cleanup(context, commandQueue, program, kernel, imageObjects, sampler);
    		return 1;
    	}
    	// Create OpenCL kernel
    	kernel = clCreateKernel(program, "sobel_rgb", NULL);
    	if (kernel == NULL)
    	{
    		std::cerr << "Failed to create kernel" << std::endl;
    		Cleanup(context, commandQueue, program, kernel, imageObjects, sampler);
    		return 1;
    	}
    
    	// Set the kernel arguments
    	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);*/
    	if (errNum != CL_SUCCESS)
    	{
    		std::cerr << "Error setting kernel arguments." << std::endl;
    		Cleanup(context, commandQueue, program, kernel, imageObjects, sampler);
    		return 1;
    	}
    
    	size_t localWorkSize[2] = { 16, 16 };
    	size_t globalWorkSize[2] = { RoundUp(localWorkSize[0], width),
    		RoundUp(localWorkSize[1], height) };
    
    	// Queue the kernel up for execution
    	errNum = clEnqueueNDRangeKernel(commandQueue, kernel, 2, NULL,
    		globalWorkSize, localWorkSize,
    		0, NULL, NULL);
    	if (errNum != CL_SUCCESS)
    	{
    		std::cerr << "Error queuing kernel for execution." << std::endl;
    		Cleanup(context, commandQueue, program, kernel, imageObjects, sampler);
    		return 1;
    	}
    
    	// Read the output buffer back to the Host
    	char *buffer = new char[width * height * 4];
    	size_t origin[3] = { 0, 0, 0 };
    	size_t region[3] = { width, 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;
    		Cleanup(context, commandQueue, program, kernel, imageObjects, sampler);
    		return 1;
    	}
    
    	std::cout << std::endl;
    	std::cout << "Executed program succesfully." << std::endl;
    
    	//memset(buffer, 0xff, width * height * 4);
    	// Save the image out to disk
    	char *saveImage = "C:/Users/jiang/Desktop/image/tu2.jpg";
    	//std::cout << buffer << std::endl;
    	cv::Mat imageColor = cv::imread(src0);
    	cv::Mat imageColor2;
    	imageColor2.create(imageColor.rows, imageColor.cols, imageColor.type());
    	int w = 0;
    	for (int v = imageColor2.rows-1; v >=0; v--)
    	{
    		for (int u =0 ; u <imageColor2.cols; u++)
    		{
    			imageColor2.at<cv::Vec3b>(v, u)[0] = buffer[w++];
    			imageColor2.at<cv::Vec3b>(v, u)[1] = buffer[w++];
    			imageColor2.at<cv::Vec3b>(v, u)[2] = buffer[w++];
    			w++;
    		}
    	}
    	cv::imshow("image", imageColor2);
    	cv::imwrite(saveImage, imageColor2);
    	cv::waitKey(0);
    	/*if (!SaveImage(saveImage, buffer, width, height))
    	{
    		std::cerr << "Error writing output image: " << saveImage<< std::endl;
    		Cleanup(context, commandQueue, program, kernel, imageObjects, sampler);
    		delete[] buffer;
    		return 1;
    	}*/
    
    	delete[] buffer;
    	Cleanup(context, commandQueue, program, kernel, imageObjects, sampler);
    	return 0;
    }