OpenCL optimizations and alterations.

Makefile_linux

@@ -32,7 +32,6 @@ OPENCV_INCLUDEDIR  = -I $(OPENCV_INCLUDEPATH)
 OPENCV_LIBDIR      = -L $(OPENCV_LIBPATH)
 OPENCV_LIBRARIES   = -lopencv_highgui -lopencv_core -lopencv_imgproc -lopencv_legacy
 
-
 WITH_CUDA = 0
 WITH_OPENCL = 0
 
@@ -132,6 +131,7 @@ LINUXAMD64_DEMO_BENCHMARK_CV = $(CC) $(COMPILER_FLAGS) \
 LINUXAMD64_DEMO_BENCHMARK_CVOCL = $(CC) $(COMPILER_FLAGS) \
                         -o $(OUTPUT_BINPATH)/demo_$(BENCHMARK_CVOCL_NAME) \
                          src/demo/$(BENCHMARK_CVOCL_NAME).cpp src/demo/timer.cpp \
+>>>>>>> upstream/master
                          -lvisiongl -lopencv_ocl \
                          $(LINUXAMD64_DIRS_LIBS) \
 

src/cl2cpp_shaders.cpp

@@ -424,3 +424,92 @@ void vglClThreshold(VglImage* src, VglImage* dst, float thresh){
   vglSetContext(dst, VGL_CL_CONTEXT);
 }
 
+/** Erosion of src image by mask. Result is stored in dst image.
+
+  */
+void vglClErosion(VglImage* img_input, VglImage* img_output, float* convolution_window, int window_size_x, int window_size_y){
+
+  vglCheckContext(img_input, VGL_CL_CONTEXT);
+  vglCheckContext(img_output, VGL_CL_CONTEXT);
+
+  cl_int err;
+
+  cl_mem mobj_convolution_window = NULL;
+  mobj_convolution_window = clCreateBuffer(cl.context, CL_MEM_READ_ONLY, (window_size_x*window_size_y)*sizeof(float), NULL, &err);
+  vglClCheckError( err, (char*) "clCreateBuffer convolution_window" );
+  err = clEnqueueWriteBuffer(cl.commandQueue, mobj_convolution_window, CL_TRUE, 0, (window_size_x*window_size_y)*sizeof(float), convolution_window, 0, NULL, NULL);
+  vglClCheckError( err, (char*) "clEnqueueWriteBuffer convolution_window" );
+
+  cl_mem mobj_window_size_x = NULL;
+  mobj_window_size_x = clCreateBuffer(cl.context, CL_MEM_READ_ONLY, (1)*sizeof(int), NULL, &err);
+  vglClCheckError( err, (char*) "clCreateBuffer window_size_x" );
+  err = clEnqueueWriteBuffer(cl.commandQueue, mobj_window_size_x, CL_TRUE, 0, (1)*sizeof(int), &window_size_x, 0, NULL, NULL);
+  vglClCheckError( err, (char*) "clEnqueueWriteBuffer window_size_x" );
+
+  cl_mem mobj_window_size_y = NULL;
+  mobj_window_size_y = clCreateBuffer(cl.context, CL_MEM_READ_ONLY, (1)*sizeof(int), NULL, &err);
+  vglClCheckError( err, (char*) "clCreateBuffer window_size_y" );
+  err = clEnqueueWriteBuffer(cl.commandQueue, mobj_window_size_y, CL_TRUE, 0, (1)*sizeof(int), &window_size_y, 0, NULL, NULL);
+  vglClCheckError( err, (char*) "clEnqueueWriteBuffer window_size_y" );
+
+  static cl_program program = NULL;
+  if (program == NULL)
+  {
+    char* file_path = (char*) "CL/vglClErosion.cl";
+    printf("Compiling %s\n", file_path);
+    std::ifstream file(file_path);
+    if(file.fail())
+    {
+      std::string str("File not found: ");
+      str.append(file_path);
+      vglClCheckError(-1, (char*)str.c_str());
+    }
+    std::string prog( std::istreambuf_iterator<char>( file ), ( std::istreambuf_iterator<char>() ) );
+    const char *source_str = prog.c_str();
+#ifdef __DEBUG__
+    printf("Kernel to be compiled:\n%s\n", source_str);
+#endif
+    program = clCreateProgramWithSource(cl.context, 1, (const char **) &source_str, 0, &err );
+    vglClCheckError(err, (char*) "clCreateProgramWithSource" );
+    err = clBuildProgram(program, 1, cl.deviceId, NULL, NULL, NULL );
+    vglClBuildDebug(err, program);
+  }
+
+  static cl_kernel kernel = NULL;
+  if (kernel == NULL)
+  {
+    kernel = clCreateKernel( program, "vglClErosion", &err ); 
+    vglClCheckError(err, (char*) "clCreateKernel" );
+  }
+
+
+  err = clSetKernelArg( kernel, 0, sizeof( cl_mem ), (void*) &img_input->oclPtr );
+  vglClCheckError( err, (char*) "clSetKernelArg 0" );
+
+  err = clSetKernelArg( kernel, 1, sizeof( cl_mem ), (void*) &img_output->oclPtr );
+  vglClCheckError( err, (char*) "clSetKernelArg 1" );
+
+  err = clSetKernelArg( kernel, 2, sizeof( cl_mem ), (float*) &mobj_convolution_window );
+  vglClCheckError( err, (char*) "clSetKernelArg 2" );
+
+  err = clSetKernelArg( kernel, 3, sizeof( cl_mem ), (int*) &mobj_window_size_x );
+  vglClCheckError( err, (char*) "clSetKernelArg 3" );
+
+  err = clSetKernelArg( kernel, 4, sizeof( cl_mem ), (int*) &mobj_window_size_y );
+  vglClCheckError( err, (char*) "clSetKernelArg 4" );
+
+  size_t worksize[] = { img_input->width, img_input->height, 1 };
+  clEnqueueNDRangeKernel( cl.commandQueue, kernel, 2, NULL, worksize, 0, 0, 0, 0 );
+  vglClCheckError( err, (char*) "clEnqueueNDRangeKernel" );
+
+  err = clReleaseMemObject( mobj_convolution_window );
+  vglClCheckError(err, (char*) "clReleaseMemObject mobj_convolution_window");
+
+  err = clReleaseMemObject( mobj_window_size_x );
+  vglClCheckError(err, (char*) "clReleaseMemObject mobj_window_size_x");
+
+  err = clReleaseMemObject( mobj_window_size_y );
+  vglClCheckError(err, (char*) "clReleaseMemObject mobj_window_size_y");
+
+  vglSetContext(img_output, VGL_CL_CONTEXT);
+}

src/cl2cpp_shaders.h

@@ -37,3 +37,5 @@ void vglClInvert(VglImage* img_input, VglImage* img_output);
   */
 void vglClThreshold(VglImage* src, VglImage* dst, float thresh);
 
+void vglClErosion(VglImage* img_input, VglImage* img_output, float* convolution_window, int window_size_x, int window_size_y);
+

src/demo/OpenCVCLBenchmark.cpp

@@ -0,0 +1,193 @@
+#include <opencv2/imgproc/imgproc.hpp>
+#include <opencv2/highgui/highgui.hpp>
+#include <opencv2/ocl/ocl.hpp>
+
+#include "demo/timer.h"
+#include <fstream>
+
+using namespace std;
+using namespace cv;
+
+
+/*
+	argv[1] = input_image_path
+	argv[2] = number of operations to execute
+*/
+int main(int argc, char* argv[])
+{
+	if (argc != 3)
+	{
+		printf("Usage: OpenCVCLBenchmark lena_1024.tif 1000\n");
+		printf("for this example, will run the program for lena_1024.tif using 1000 operations each\n");
+		printf("bad arguments, read usage again\n");
+		exit(1);
+	}
+
+	int limite = atoi(argv[2]);
+	char* image_path = argv[1];
+	printf("OpenCVCLBenchmark on %s for %d operations\n\n",image_path,limite);
+	cv::Mat imgxm = imread(image_path,1);
+
+	cvtColor(imgxm,imgxm,CV_BGR2RGBA);
+	ocl::oclMat img(imgxm), out(imgxm);
+
+	vector<int> saveparams;
+	saveparams.push_back(CV_IMWRITE_PNG_COMPRESSION);
+	saveparams.push_back(1);
+
+	if (img.data == NULL)
+	{
+		std::string str("cvLoadImage/File not found: ");
+		str.append(image_path);
+		printf("%s",str.c_str());
+	}
+
+	//Primeira chamada a blurSq3
+	Size ksize(3,3);
+	TimerStart();
+	ocl::GaussianBlur(img,out,Size(3,3),0);
+	//ocl::blur(img,out,ksize);
+	printf("Primeira chamada de Blur: %s \n", getTimeElapsedInSeconds());
+	//Mede o tempo para limite blur 3x3 sem a cria��o da opera��o
+	int p = 0;
+	TimerStart();
+	while (p < limite)
+	{
+		p++;
+		//ocl::blur(img,out,ksize);
+		ocl::GaussianBlur(img,out,Size(3,3),0);
+	}
+	printf("Tempo gasto para fazer %d Blur 3x3: %s\n", limite, getTimeElapsedInSeconds());
+
+		cv::Mat saveout(out);
+		cvtColor(saveout,saveout,CV_RGBA2BGR);
+        cv::imwrite("../images/lenaout_blur33.png", saveout,saveparams);
+
+	//Primeira chamada a Convolution
+
+    Mat cvkernel33 = (Mat_<float>(3,3) << 1/9.0, 1/9.0, 1/9.0, 
+                                          1/9.0, 1/9.0, 1/9.0, 
+	                                      1/9.0, 1/9.0, 1/9.0  );
+
+    Mat cvkernel55 = (Mat_<float>(5,5) << 1/25.0, 1/25.0, 1/25.0, 1/25.0, 1/25.0,
+                                          1/25.0, 1/25.0, 1/25.0, 1/25.0, 1/25.0,
+                                          1/25.0, 1/25.0, 1/25.0, 1/25.0, 1/25.0,
+                                          1/25.0, 1/25.0, 1/25.0, 1/25.0, 1/25.0,
+                                          1/25.0, 1/25.0, 1/25.0, 1/25.0, 1/25.0  );
+
+
+	TimerStart();
+	cv::Ptr<ocl::FilterEngine_GPU> filter33 = ocl::createLinearFilter_GPU(CV_8UC4,CV_8UC4, cvkernel33);
+	cv::Ptr<ocl::FilterEngine_GPU> filter55 = ocl::createLinearFilter_GPU(CV_8UC4,CV_8UC4, cvkernel55);
+	filter33->apply(img,out);
+	printf("Primeira chamada da Convolucao com kernel 3x3: %s\n", getTimeElapsedInSeconds());
+	//Mede o tempo para limite convolu��es 3x3 sem a cria��o da opera��o
+	p = 0;
+	TimerStart();
+	while (p < limite)
+	{
+		p++;
+		filter33->apply(img,out);
+	}
+	printf("Tempo gasto para fazer %d convolucoes 3x3: %s \n", limite, getTimeElapsedInSeconds());
+
+		saveout = *new cv::Mat(out);
+		cvtColor(saveout,saveout,CV_RGBA2BGR);
+		cv::imwrite("../images/lenaout_conv33.png", saveout,saveparams);
+
+	//Mede o tempo para limite convolu��es 5x5 sem a cria��o da opera��o
+	p = 0;
+	filter55->apply(img,out);
+	TimerStart();
+	while (p < limite)
+	{
+		p++;
+		filter55->apply(out,out);
+	}
+	printf("Tempo gasto para fazer %d convolucoes 5x5: %s\n", limite, getTimeElapsedInSeconds());
+
+		saveout = *new cv::Mat(out);
+		cvtColor(saveout,saveout,CV_RGBA2BGR);
+		cv::imwrite("../images/lenaout_conv55.png", saveout,saveparams);
+
+	//Primeira chamada a threshold
+	TimerStart();
+	Mat cverode33 = getStructuringElement(MORPH_CROSS, Size( 3, 3 ));
+	ocl::erode(img,out,cverode33);
+
+	printf("Primeira chamada da Erode: %s \n", getTimeElapsedInSeconds());
+	//Mede o tempo para limite thresholds sem a cria��o da opera��o
+	p = 0;
+	TimerStart();
+	while (p < limite)
+	{
+		p++;
+		ocl::erode(img,out,cverode33);
+	}
+	printf("Tempo gasto para fazer %d Erosions: %s\n", limite, getTimeElapsedInSeconds());
+
+		saveout = *new cv::Mat(out);
+		cvtColor(saveout,saveout,CV_RGBA2BGR);
+		cv::imwrite("../images/lenaout_erosion.png", saveout,saveparams);
+
+	//Primeira chamada a vglClInvert
+	TimerStart();
+	ocl::bitwise_not(img,out);
+	printf("Primeira chamada da Invert: %s \n", getTimeElapsedInSeconds());
+	//Mede o tempo para limite invert sem a cria��o da opera��o
+	p = 0;
+	TimerStart();
+	while (p < limite)
+	{
+		p++;
+		ocl::bitwise_not(img,out);
+	}
+	printf("Tempo gasto para fazer %d Invert: %s\n",limite, getTimeElapsedInSeconds());
+
+		saveout = *new cv::Mat(out);
+		cvtColor(saveout,saveout,CV_RGBA2BGR);
+		cv::imwrite("../images/lenaout_invert.png", saveout,saveparams);
+
+	TimerStart();
+
+	img.copyTo(out);
+	printf("Primeira chamada da Copy: %s \n", getTimeElapsedInSeconds());
+	//Mede o tempo para limite copia GPU->GPU
+	p = 0;
+	TimerStart();
+	ocl::oclMat x1(img.size(), CV_8UC4);
+	while (p < limite)
+	{
+		p++;
+		img.copyTo(x1);
+	}
+	printf("Tempo gasto para fazer %d copia GPU->GPU: %s\n", limite, getTimeElapsedInSeconds());
+
+        saveout = *new cv::Mat(out);
+		cvtColor(saveout,saveout,CV_RGBA2BGR);
+		cv::imwrite("../images/lenaout_copy.png", saveout,saveparams);
+
+	//Mede o tempo para limite copia CPU->GPU
+
+	ocl::oclMat x;
+	p = 0;
+	TimerStart();
+	while (p < limite)
+	{
+		p++;
+		img.upload(imgxm);
+	}
+	printf("Tempo gasto para fazer %d copia CPU->GPU(Upload) %s\n", limite, getTimeElapsedInSeconds());
+
+	//Mede o tempo para limite copia GPU->CPU
+	cv::Mat x2;
+	p = 0;
+	TimerStart();
+	while (p < limite)
+	{
+		p++;
+		img.download(x2);
+	}
+	printf("Tempo gasto para fazer %d copia GPU->CPU(Download) %s\n", limite, getTimeElapsedInSeconds());
+	return 0;
+}