From 0281542d406b4457bef37d0d79463ad46b678942 Mon Sep 17 00:00:00 2001
From: Ivan Mahonin <bh@icystar.com>
Date: Aug 29 2015 11:50:07 +0000
Subject: Fix CL queue synchronization


---

diff --git a/c++/contourgl/clrender.cpp b/c++/contourgl/clrender.cpp
index 81e3de4..6a68143 100644
--- a/c++/contourgl/clrender.cpp
+++ b/c++/contourgl/clrender.cpp
@@ -35,6 +35,7 @@ ClRender::ClRender(ClContext &cl):
 	rows_buffer(),
 	mark_buffer(),
 	surface_buffer(),
+	prev_event(),
 	rows_count(),
 	even_rows_count(),
 	odd_rows_count()
@@ -57,6 +58,7 @@ void ClRender::send_surface(Surface *surface) {
 	if (this->surface == surface) return;
 
 	cl.err = clFinish(cl.queue);
+	prev_event = NULL;
 	assert(!cl.err);
 
 	if (this->surface) {
@@ -95,14 +97,11 @@ void ClRender::send_surface(Surface *surface) {
 			NULL );
 		assert(surface_buffer);
 
-		cl_event event = NULL;
 		cl.err |= clEnqueueWriteBuffer(
 			cl.queue, surface_buffer, CL_TRUE,
 			0, surface->data_size(), surface->data,
-			0, NULL, &event );
-		clWaitForEvents(1, &event);
+			0, NULL, &prev_event );
 
-		cl.err |= clFinish(cl.queue);
 		assert(!cl.err);
 	}
 }
@@ -111,13 +110,13 @@ Surface* ClRender::receive_surface() {
 	if (surface) {
 		//Measure t("ClRender::receive_surface");
 
-		cl_event event = NULL;
 		cl.err |= clEnqueueReadBuffer(
 			cl.queue, surface_buffer, CL_TRUE,
 			0, surface->data_size(), surface->data,
-			0, NULL, &event );
+			prev_event ? 1 : 0, &prev_event, NULL );
 		assert(!cl.err);
-		clWaitForEvents(1, &event);
+		clFinish(cl.queue);
+		prev_event = NULL;
 	}
 	return surface;
 }
@@ -218,8 +217,6 @@ void ClRender::contour(const Contour &contour, const Rect &rect, const Color &co
 	{
 		//Measure t("enqueue commands");
 
-		clFinish(cl.queue);
-
 		// kernel args
 		int width = surface->width;
 
@@ -260,7 +257,7 @@ void ClRender::contour(const Contour &contour, const Rect &rect, const Color &co
 		cl.err |= clEnqueueWriteBuffer(
 			cl.queue, mark_buffer, CL_TRUE,
 			0, marks.size()*sizeof(marks.front()), &marks.front(),
-			0, NULL, &prepare_buffers_events[2] );
+			prev_event ? 1 : 0, &prev_event, &prepare_buffers_events[2] );
 		assert(!cl.err);
 
 		// run kernels
@@ -291,7 +288,6 @@ void ClRender::contour(const Contour &contour, const Rect &rect, const Color &co
 			&lines_even_event );
 		assert(!cl.err);
 
-		cl_event fill_event = NULL;
 		cl.err |= clEnqueueNDRangeKernel(
 			cl.queue,
 			contour_fill_kernel,
@@ -301,10 +297,10 @@ void ClRender::contour(const Contour &contour, const Rect &rect, const Color &co
 			NULL,
 			1,
 			&lines_even_event,
-			&fill_event );
+			&prev_event );
 		assert(!cl.err);
 
-		clWaitForEvents(1, &fill_event);
+		clWaitForEvents(1, &lines_even_event);
 	}
 
 	{
diff --git a/c++/contourgl/clrender.h b/c++/contourgl/clrender.h
index 097c9f0..4e7a1d9 100644
--- a/c++/contourgl/clrender.h
+++ b/c++/contourgl/clrender.h
@@ -37,6 +37,7 @@ private:
 	cl_mem rows_buffer;
 	cl_mem mark_buffer;
 	cl_mem surface_buffer;
+	cl_event prev_event;
 
 	size_t rows_count;
 	size_t even_rows_count;
diff --git a/c++/contourgl/test.cpp b/c++/contourgl/test.cpp
index a9129f1..344c7ff 100644
--- a/c++/contourgl/test.cpp
+++ b/c++/contourgl/test.cpp
@@ -447,6 +447,7 @@ void Test::test4() {
 
 		// gl_triangles
 
+		/*
 		GLuint index_buffer_id = 0;
 		vector<int> triangle_starts(contours_gl.size());
 		vector<int> triangle_counts(contours_gl.size());
@@ -508,6 +509,7 @@ void Test::test4() {
 				glDrawElements(GL_TRIANGLES, triangle_counts[i], GL_UNSIGNED_INT, (char*)NULL + triangle_starts[i]*sizeof(int));
 			}
 		}
+		*/
 	}
 
 	{