diff --git a/c++/contourgl/cl/contour.cl b/c++/contourgl/cl/contour.cl
index d03cd67..cbe5244 100644
--- a/c++/contourgl/cl/contour.cl
+++ b/c++/contourgl/cl/contour.cl
@@ -15,71 +15,69 @@
     along with this program.  If not, see <http://www.gnu.org/licenses/>.
 */
 
-__kernel void clear2f(
-	__global float2 *buffer )
+kernel void clear(
+	global int2 *mark_buffer )
 {
-	const float2 v = { 0.f, 0.f };
-	buffer[get_global_id(0)] = v;
+	const int2 v = { 0, 0 };
+	mark_buffer[get_global_id(0)] = v;
 }
 
-__kernel void lines(
-	int width, 
-	__global float4 *lines,
-	__global int2 *rows,
-	__global float2 *mark_buffer )
+kernel void path(
+	int width,
+	int height,
+	global int *mark_buffer,
+	global float2 *path )
 {
 	const float e = 1e-6f;
-	int2 row = rows[get_global_id(0)];
-	sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE
-					  | CLK_ADDRESS_NONE
-			          | CLK_FILTER_NEAREST;
-	int w = width;
-	for(__global float4 *i = lines + row.x, *end = i + row.y; i < end; ++i) {
-		float4 line = *i; 
-		float2 p0 = { line.x, line.y };
-		float2 p1 = { line.z, line.w };
-		
-		int iy0 = (int)floor(fmin(p0.y, p1.y) + e);
-		int iy1 = (int)floor(fmax(p0.y, p1.y) - e);
-		
-		float2 d = p1 - p0;
-		float kx = fabs(d.y) < e ? 0.f : d.x/d.y;
-		float ky = fabs(d.x) < e ? 0.f : d.y/d.x;
-		
-		for(int r = iy0; r <= iy1; ++r) {
-			float y = (float)r;
-			float2 pya = { p0.x + kx*(y       - p0.y), y };
-			float2 pyb = { p0.x + kx*(y + 1.0 - p0.y), y + 1.f };
-			float2 pp0 = p0.y - y < -e      ? pya
-			           : (p0.y - y > 1.f + e ? pyb : p0);
-			float2 pp1 = p1.y - y < -e      ? pya
-			           : (p1.y - y > 1.f + e ? pyb : p1);
+	size_t id = get_global_id(0);
+	
+	float2 s = { (float)width, (float)height }; 
+	int w1 = width - 1;
+	int h1 = height - 1;
 
-			int ix0 = (int)floor(fmin(pp0.x, pp1.x) + e);
-			int ix1 = (int)floor(fmax(pp0.x, pp1.x) - e);
-			for(int c = ix0; c <= ix1; ++c) {
-				float x = (float)c;
-				float2 pxa = { x,       p0.y + ky*(x       - p0.x) };
-				float2 pxb = { x + 1.0, p0.y + ky*(x + 1.0 - p0.x) };
-				float2 ppp0 = pp0.x - x < -e      ? pxa
-				            : (pp0.x - x > 1.f + e ? pxb : pp0);
-				float2 ppp1 = pp1.x - x < -e      ? pxa
-				            : (pp1.x - x > 1.f + e ? pxb : pp1);
+	float2 p0 = path[id];
+	float2 p1 = path[id + 1];
+	bool flipx = p1.x < p0.x;
+	bool flipy = p1.y < p0.y;
+	if (flipx) { p0.x = s.x - p0.x; p1.x = s.x - p1.x; }
+	if (flipy) { p0.y = s.y - p0.y; p1.y = s.y - p1.y; }
+	float2 d = p1 - p0;
+	float kx = fabs(d.y) < e ? 1e10 : d.x/d.y;
+	float ky = fabs(d.x) < e ? 1e10 : d.y/d.x;
+	
+	while(p0.x != p1.x || p0.y != p1.y) {
+		int ix = (int)floor(p0.x + e);
+		int iy = (int)floor(p0.y + e);
+		if (iy > h1) break;
 
-				float2 m;
-				m.y = ppp1.y - ppp0.y;
-				m.x = (x + 1.f - 0.5f*(ppp0.x + ppp1.x))*m.y;
-				mark_buffer[r*w + c] += m;  
-			}
+		float2 px, py;
+		px.x = (float)(ix + 1);
+		px.y = p0.y + ky*(px.x - p0.x);
+		py.y = max((float)(iy + 1), 0.f);
+		py.x = p0.x + kx*(py.y - p0.y);
+		float2 pp1 = p1;
+		if (pp1.x > px.x) pp1 = px;
+		if (pp1.y > py.y) pp1 = py;
+		
+		if (iy >= 0) {
+			float cover = pp1.y - p0.y;
+			float area = px.x - 0.5f*(p0.x + pp1.x);
+			if (flipx) { ix = w1 - ix; area = 1.f - area; }
+			if (flipy) { iy = h1 - iy; cover = -cover; }
+			global int *mark = mark_buffer + 2*(iy*width + clamp(ix, 0, w1));
+			atomic_add(mark, (int)round(area*cover*65536));
+			atomic_add(mark + 1, (int)round(cover*65536));
 		}
+		
+		p0 = pp1;
 	}
 }
 
-__kernel void fill(
+kernel void fill(
 	int width,
-	__global float2 *mark_buffer,
-	__read_only image2d_t surface_read_image,
-	__write_only image2d_t surface_write_image,
+	global int2 *mark_buffer,
+	read_only image2d_t surface_read_image,
+	write_only image2d_t surface_write_image,
 	float4 color,
 	int invert,
 	int evenodd )
@@ -91,14 +89,15 @@ __kernel void fill(
 			                | CLK_FILTER_NEAREST;
 	float4 cl = color;
 	float cover = 0.f;
-	__global float2 *mark = mark_buffer + id*w;
+	global int2 *mark = mark_buffer + id*w;
+	const int2 izero = { 0, 0 };
 	for(int2 coord = { 0, id }; coord.x < w; ++coord.x, ++mark) {
-		float2 m = *mark;
-
-		float alpha = fabs(m.x + cover);
-		cover += m.y;
+		int2 im = *mark;
+		//*mark = izero;
+		float alpha = fabs((float)im.x/65536.f + cover);
+		cover += (float)im.y/65536.f;
 		alpha = evenodd ? (1.f - fabs(1.f - alpha - 2.f*floor(0.5f*alpha)))
-				        : fmin(alpha, 1.f);
+						: fmin(alpha, 1.f);
 		alpha *= cl.w;
 		if (invert) alpha = 1.f - alpha;
 		float alpha_inv = 1.f - alpha;
diff --git a/c++/contourgl/clcontext.cpp b/c++/contourgl/clcontext.cpp
index 15c876d..4cb448b 100644
--- a/c++/contourgl/clcontext.cpp
+++ b/c++/contourgl/clcontext.cpp
@@ -69,12 +69,15 @@ ClContext::ClContext(): err(), context(), queue() {
 
     // context
 
-    context = clCreateContext(0, 1, &devices.front(), NULL, NULL, &err);
+    cl_context_properties context_props[] = {
+    	CL_CONTEXT_PLATFORM, (cl_context_properties)platform,
+		CL_NONE };
+    context = clCreateContext(context_props, 1, &devices.front(), callback, NULL, &err);
     assert(context);
 
 	// command queue
 
-	queue = clCreateCommandQueue(context, devices[0], 0, NULL);
+	queue = clCreateCommandQueue(context, devices.front(), CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, NULL);
 	assert(queue);
 
 }
@@ -84,6 +87,8 @@ ClContext::~ClContext() {
 	clReleaseContext(context);
 }
 
+void ClContext::callback(const char *, const void *, size_t, void *) { }
+
 cl_program ClContext::load_program(const std::string &filename) {
 	ifstream f(("cl/" + filename).c_str());
 	string text((istreambuf_iterator<char>(f)), istreambuf_iterator<char>());
diff --git a/c++/contourgl/clcontext.h b/c++/contourgl/clcontext.h
index 6ced007..6061d43 100644
--- a/c++/contourgl/clcontext.h
+++ b/c++/contourgl/clcontext.h
@@ -36,6 +36,7 @@ public:
 
 	void hello();
 	cl_program load_program(const std::string &filename);
+	static void callback(const char *, const void *, size_t, void *);
 };
 
 #endif
diff --git a/c++/contourgl/clrender.cpp b/c++/contourgl/clrender.cpp
index f6335c6..ae918a4 100644
--- a/c++/contourgl/clrender.cpp
+++ b/c++/contourgl/clrender.cpp
@@ -29,44 +29,40 @@ using namespace std;
 ClRender::ClRender(ClContext &cl):
 	cl(cl),
 	contour_program(),
-	contour_lines_kernel(),
+	contour_path_kernel(),
 	contour_fill_kernel(),
 	surface(),
-	rows_buffer(),
+	path_buffer(),
 	mark_buffer(),
 	surface_image(),
-	prev_event(),
-	rows_count(),
-	even_rows_count(),
-	odd_rows_count()
+	prev_event()
 {
 	contour_program = cl.load_program("contour.cl");
-	contour_clear2f_kernel = clCreateKernel(contour_program, "clear2f", NULL);
-	assert(contour_clear2f_kernel);
-	contour_lines_kernel = clCreateKernel(contour_program, "lines", NULL);
-	assert(contour_lines_kernel);
+	contour_clear_kernel = clCreateKernel(contour_program, "clear", NULL);
+	assert(contour_clear_kernel);
+	contour_path_kernel = clCreateKernel(contour_program, "path", NULL);
+	assert(contour_path_kernel);
 	contour_fill_kernel = clCreateKernel(contour_program, "fill", NULL);
 	assert(contour_fill_kernel);
 }
 
 ClRender::~ClRender() {
 	send_surface(NULL);
-	clReleaseKernel(contour_clear2f_kernel);
+	send_path(NULL, 0);
+	clReleaseKernel(contour_clear_kernel);
 	clReleaseKernel(contour_fill_kernel);
-	clReleaseKernel(contour_lines_kernel);
+	clReleaseKernel(contour_path_kernel);
 	clReleaseProgram(contour_program);
 }
 
 void ClRender::send_surface(Surface *surface) {
-	if (this->surface == surface) return;
+	if (!surface && !this->surface) return;
 
-	cl.err = clFinish(cl.queue);
-	prev_event = NULL;
+	cl.err |= clFinish(cl.queue);
 	assert(!cl.err);
+	prev_event = NULL;
 
 	if (this->surface) {
-		rows.clear();
-		clReleaseMemObject(rows_buffer);
 		clReleaseMemObject(mark_buffer);
 		clReleaseMemObject(surface_image);
 	}
@@ -76,24 +72,28 @@ void ClRender::send_surface(Surface *surface) {
 	if (this->surface) {
 		//Measure t("ClRender::send_surface");
 
-		rows_count = surface->height;
-		even_rows_count = (rows_count+1)/2;
-		odd_rows_count = rows_count - even_rows_count;
-		rows.resize(rows_count);
-		marks.resize(surface->count());
-
-		rows_buffer = clCreateBuffer(
-			cl.context, CL_MEM_READ_ONLY,
-			rows.size()*sizeof(rows.front()), NULL,
-			NULL );
-		assert(rows_buffer);
-
 		mark_buffer = clCreateBuffer(
 			cl.context, CL_MEM_READ_WRITE,
-			surface->count()*sizeof(cl_float2), NULL,
+			surface->count()*sizeof(cl_int2), NULL,
 			NULL );
 		assert(mark_buffer);
 
+		cl.err |= clSetKernelArg(contour_clear_kernel, 0, sizeof(mark_buffer), &mark_buffer);
+		assert(!cl.err);
+
+		size_t pixels_count = (size_t)surface->count();
+		cl.err |= clEnqueueNDRangeKernel(
+			cl.queue,
+			contour_clear_kernel,
+			1,
+			NULL,
+			&pixels_count,
+			NULL,
+			0,
+			NULL,
+			NULL );
+		assert(!cl.err);
+
 		cl_image_format surface_format = { };
 		surface_format.image_channel_order = CL_RGBA;
 		surface_format.image_channel_data_type = CL_FLOAT;
@@ -107,10 +107,25 @@ void ClRender::send_surface(Surface *surface) {
 		size_t origin[3] = { };
 		size_t region[3] = { (size_t)surface->width, (size_t)surface->height, 1 };
 		cl.err |= clEnqueueWriteImage(
-			cl.queue, surface_image, CL_TRUE,
+			cl.queue, surface_image, CL_FALSE,
 			origin, region, 0, 0, surface->data,
-			0, NULL, &prev_event );
+			0, NULL, NULL );
+		assert(!cl.err);
+
+		int width = surface->width;
+		int height = surface->height;
+		cl.err |= clSetKernelArg(contour_path_kernel, 0, sizeof(width), &width);
+		cl.err |= clSetKernelArg(contour_path_kernel, 1, sizeof(height), &height);
+		cl.err |= clSetKernelArg(contour_path_kernel, 2, sizeof(mark_buffer), &mark_buffer);
+		assert(!cl.err);
+
+		cl.err |= clSetKernelArg(contour_fill_kernel, 0, sizeof(width), &width);
+		cl.err |= clSetKernelArg(contour_fill_kernel, 1, sizeof(mark_buffer), &mark_buffer);
+		cl.err |= clSetKernelArg(contour_fill_kernel, 2, sizeof(surface_image), &surface_image);
+		cl.err |= clSetKernelArg(contour_fill_kernel, 3, sizeof(surface_image), &surface_image);
+		assert(!cl.err);
 
+		cl.err |= clFinish(cl.queue);
 		assert(!cl.err);
 	}
 }
@@ -122,213 +137,111 @@ Surface* ClRender::receive_surface() {
 		size_t origin[3] = { };
 		size_t region[3] = { (size_t)surface->width, (size_t)surface->height, 1 };
 		cl.err |= clEnqueueReadImage(
-			cl.queue, surface_image, CL_TRUE,
+			cl.queue, surface_image, CL_FALSE,
 			origin, region, 0, 0, surface->data,
 			prev_event ? 1 : 0, &prev_event, NULL );
 		assert(!cl.err);
-		clFinish(cl.queue);
+
+		cl.err |= clFinish(cl.queue);
+		assert(!cl.err);
 		prev_event = NULL;
 	}
 	return surface;
 }
 
+void ClRender::send_path(const vec2f *path, int count) {
+	if (!path_buffer && (!path || count <=0)) return;
 
-void ClRender::contour(const Contour &contour, const Rect &rect, const Color &color, bool invert, bool evenodd) {
-	//Measure t("ClRender::contour");
-
-	Contour transformed, splitted;
-	Rect to(1.0, 1.0, surface->width - 1.0, surface->height - 1.0);
-
-	{
-		//Measure t("clone");
-		transformed = contour;
-	}
-
-	{
-		//Measure t("transform");
-		transformed.transform(rect, to);
-	}
-
-	{
-		Measure t("split");
-		splitted.allow_split_lines = true;
-		transformed.split(splitted, to, Vector(0.5, 0.5));
-	}
-
-	vector<line2f> lines;
-	vector<line2f> sorted_lines;
-	vector<int> line_rows;
-
-	{
-		Measure t("sort lines");
-
-		// reset rows
-		for(int i = 0; i < (int)rows_count; ++i)
-			rows[i].second = 0;
-
-		// count lines
-		Vector prev;
-		lines.reserve(splitted.get_chunks().size());
-		line_rows.reserve(splitted.get_chunks().size());
-		float x0 = (float)to.p0.x;
-		float x1 = (float)to.p1.x;
-		for(Contour::ChunkList::const_iterator i = splitted.get_chunks().begin(); i != splitted.get_chunks().end(); ++i) {
-			if ( i->type == Contour::LINE
-			  || i->type == Contour::CLOSE )
-			{
-				if (i->p1.y > to.p0.y && i->p1.y < to.p1.y) {
-					line2f l(vec2f(prev), vec2f(i->p1));
-					l.p0.x = min(max(l.p0.x, x0), x1);
-					l.p1.x = min(max(l.p1.x, x0), x1);
-					assert( (int)floorf(l.p0.x) >= 0 && (int)floorf(l.p0.x) < surface->width
-						 && (int)floorf(l.p1.x) >= 0 && (int)floorf(l.p1.x) < surface->width
-						 && (int)floorf(l.p0.y) >= 0 && (int)floorf(l.p1.y) < surface->height
-						 && (int)floorf(l.p1.y) >= 0 && (int)floorf(l.p1.y) < surface->height
-						 && abs((int)floorf(l.p1.x) - (int)floorf(l.p0.x)) <= 1
-						 && abs((int)floorf(l.p1.y) - (int)floorf(l.p0.y)) <= 1 );
-					int row = (int)floorf(min(l.p0.y, l.p1.y));
-					row = row % 2 ? row/2 : even_rows_count + row/2;
-					assert(row >= 0 && row < (int)rows_count);
-					line_rows.push_back(row);
-					lines.push_back(l);
-					++rows[row].second;
-				}
-			}
-			prev = i->p1;
-		}
+	cl.err |= clFinish(cl.queue);
+	assert(!cl.err);
+	prev_event = NULL;
 
-		// calc rows offsets
-		int lines_count = (int)lines.size();
-		rows[0].first = rows[0].second;
-		for(int i = 1; i < (int)rows_count; ++i)
-			rows[i].first = rows[i-1].first + rows[i].second;
-
-		// make sorted list
-		sorted_lines.resize(lines_count);
-		for(int i = 0; i < lines_count; ++i) {
-			assert(rows[line_rows[i]].first > 0 && rows[line_rows[i]].first <= lines_count);
-			sorted_lines[ --rows[line_rows[i]].first ] = lines[i];
-		}
+	if (path_buffer) {
+		clReleaseMemObject(path_buffer);
+		path_buffer = NULL;
 	}
 
-	if (sorted_lines.empty()) return;
+	if (path && count > 0) {
+		//Measure t("ClRender::send_path");
 
-	cl_mem lines_buffer = NULL;
-
-	{
-		Measure t("create lines buffer");
-
-		lines_buffer = clCreateBuffer(
+		path_buffer = clCreateBuffer(
 			cl.context, CL_MEM_READ_ONLY,
-			sorted_lines.size()*sizeof(sorted_lines.front()), NULL,
+			count*sizeof(*path), NULL,
 			NULL );
-		assert(lines_buffer);
-	}
-
-	{
-		Measure t("enqueue commands");
-
-		// kernel args
-
-		int width = surface->width;
-
-		cl.err |= clSetKernelArg(contour_clear2f_kernel, 0, sizeof(mark_buffer), &mark_buffer);
-		assert(!cl.err);
-
-		cl.err |= clSetKernelArg(contour_lines_kernel, 0, sizeof(width), &width);
-		cl.err |= clSetKernelArg(contour_lines_kernel, 1, sizeof(lines_buffer), &lines_buffer);
-		cl.err |= clSetKernelArg(contour_lines_kernel, 2, sizeof(rows_buffer), &rows_buffer);
-		cl.err |= clSetKernelArg(contour_lines_kernel, 3, sizeof(mark_buffer), &mark_buffer);
-		assert(!cl.err);
-
-		int iinvert = invert, ievenodd = evenodd;
-		cl.err |= clSetKernelArg(contour_fill_kernel, 0, sizeof(width), &width);
-		cl.err |= clSetKernelArg(contour_fill_kernel, 1, sizeof(mark_buffer), &mark_buffer);
-		cl.err |= clSetKernelArg(contour_fill_kernel, 2, sizeof(surface_image), &surface_image);
-		cl.err |= clSetKernelArg(contour_fill_kernel, 3, sizeof(surface_image), &surface_image);
-		cl.err |= clSetKernelArg(contour_fill_kernel, 4, sizeof(color), &color);
-		cl.err |= clSetKernelArg(contour_fill_kernel, 5, sizeof(int), &iinvert);
-		cl.err |= clSetKernelArg(contour_fill_kernel, 6, sizeof(int), &ievenodd);
-		assert(!cl.err);
-
-		// init buffers
-
-		cl_event prepare_buffers_events[3] = { };
+		assert(path_buffer);
 
 		cl.err |= clEnqueueWriteBuffer(
-			cl.queue, lines_buffer, CL_TRUE,
-			0, sorted_lines.size()*sizeof(sorted_lines.front()), &sorted_lines.front(),
-			0, NULL, &prepare_buffers_events[0] );
+			cl.queue, path_buffer, CL_FALSE,
+			0, count*sizeof(*path), path,
+			0, NULL, NULL );
 		assert(!cl.err);
 
-		cl.err |= clEnqueueWriteBuffer(
-			cl.queue, rows_buffer, CL_TRUE,
-			0, rows.size()*sizeof(rows.front()), &rows.front(),
-			0, NULL, &prepare_buffers_events[1] );
+		int path_size = count;
+		cl.err |= clSetKernelArg(contour_path_kernel, 3, sizeof(path_buffer), &path_buffer);
 		assert(!cl.err);
 
-		size_t count = (size_t)surface->count();
-		cl.err |= clEnqueueNDRangeKernel(
-			cl.queue,
-			contour_clear2f_kernel,
-			1,
-			NULL,
-			&count,
-			NULL,
-			prev_event ? 1 : 0,
-			&prev_event,
-			&prepare_buffers_events[2] );
+		cl.err |= clFinish(cl.queue);
 		assert(!cl.err);
+	}
+}
 
-		// build marks
-
-		cl_event lines_odd_event = NULL;
-		cl.err |= clEnqueueNDRangeKernel(
-			cl.queue,
-			contour_lines_kernel,
-			1,
-			NULL,
-			&even_rows_count,
-			NULL,
-			3,
-			prepare_buffers_events,
-			&lines_odd_event );
-		assert(!cl.err);
+void ClRender::path(int start, int count, const Color &color, bool invert, bool evenodd) {
+	//Measure t("ClRender::contour");
 
-		cl_event lines_even_event = NULL;
-		cl.err |= clEnqueueNDRangeKernel(
-			cl.queue,
-			contour_lines_kernel,
-			1,
-			&even_rows_count,
-			&odd_rows_count,
-			NULL,
-			1,
-			&lines_odd_event,
-			&lines_even_event );
-		assert(!cl.err);
+	if (count <= 1) return;
 
-		// fill
+	// kernel args
+	int iinvert = invert, ievenodd = evenodd;
+	cl.err |= clSetKernelArg(contour_fill_kernel, 4, sizeof(color), &color);
+	cl.err |= clSetKernelArg(contour_fill_kernel, 5, sizeof(int), &iinvert);
+	cl.err |= clSetKernelArg(contour_fill_kernel, 6, sizeof(int), &ievenodd);
+	assert(!cl.err);
 
-		cl.err |= clEnqueueNDRangeKernel(
-			cl.queue,
-			contour_fill_kernel,
-			1,
-			NULL,
-			&rows_count,
-			NULL,
-			1,
-			&lines_even_event,
-			&prev_event );
-		assert(!cl.err);
+	// clear
+
+	size_t pixels_count = (size_t)surface->count();
+	cl.err |= clEnqueueNDRangeKernel(
+		cl.queue,
+		contour_clear_kernel,
+		1,
+		NULL,
+		&pixels_count,
+		NULL,
+		0,
+		NULL,
+		NULL );
+	assert(!cl.err);
 
-		clWaitForEvents(1, &lines_even_event);
-	}
+	// build marks
+
+	cl_event path_event = NULL;
+	size_t sstart = start;
+	size_t scount = count-1;
+	cl.err |= clEnqueueNDRangeKernel(
+		cl.queue,
+		contour_path_kernel,
+		1,
+		&sstart,
+		&scount,
+		NULL,
+		prev_event ? 1 : 0,
+		&prev_event,
+		&path_event );
+	assert(!cl.err);
 
-	{
-		Measure t("release lines buffer");
-		clReleaseMemObject(lines_buffer);
-	}
+	// fill
+	size_t sheight = surface->height;
+	cl.err |= clEnqueueNDRangeKernel(
+		cl.queue,
+		contour_fill_kernel,
+		1,
+		NULL,
+		&sheight,
+		NULL,
+		1,
+		&path_event,
+		&prev_event );
+	assert(!cl.err);
 }
 
 
diff --git a/c++/contourgl/clrender.h b/c++/contourgl/clrender.h
index 703195b..649c89b 100644
--- a/c++/contourgl/clrender.h
+++ b/c++/contourgl/clrender.h
@@ -30,31 +30,25 @@ class ClRender {
 private:
 	ClContext &cl;
 	cl_program contour_program;
-	cl_kernel contour_clear2f_kernel;
-	cl_kernel contour_lines_kernel;
+	cl_kernel contour_clear_kernel;
+	cl_kernel contour_path_kernel;
 	cl_kernel contour_fill_kernel;
 
 	Surface *surface;
-	cl_mem rows_buffer;
+	cl_mem path_buffer;
 	cl_mem mark_buffer;
 	cl_mem surface_image;
 	cl_event prev_event;
 
-	size_t rows_count;
-	size_t even_rows_count;
-	size_t odd_rows_count;
-
-	typedef std::pair<int, int> Row;
-	std::vector<Row> rows;
-	std::vector<vec2f> marks;
-
 public:
 	ClRender(ClContext &cl);
 	~ClRender();
 
 	void send_surface(Surface *surface);
 	Surface* receive_surface();
-	void contour(const Contour &contour, const Rect &rect, const Color &color, bool invert, bool evenodd);
+	void send_path(const vec2f *path, int count);
+	void path(int start, int count, const Color &color, bool invert, bool evenodd);
+	void wait();
 };
 
 
diff --git a/c++/contourgl/glcontext.cpp b/c++/contourgl/glcontext.cpp
index ea3623a..218f5a1 100644
--- a/c++/contourgl/glcontext.cpp
+++ b/c++/contourgl/glcontext.cpp
@@ -48,8 +48,8 @@ GlContext::GlContext():
 	int framebuffer_width = 512;
 	int framebuffer_height = 512;
 	int framebuffer_samples = 16;
-	bool antialising = true;
-	bool hdr = false;
+	bool antialising = false;
+	bool hdr = true;
 
 	// display
 
diff --git a/c++/contourgl/test.cpp b/c++/contourgl/test.cpp
index 5995b82..cbe4a57 100644
--- a/c++/contourgl/test.cpp
+++ b/c++/contourgl/test.cpp
@@ -17,6 +17,7 @@
 
 #include <fstream>
 #include <iostream>
+#include <iomanip>
 
 #include "test.h"
 #include "contourbuilder.h"
@@ -340,6 +341,8 @@ void Test::test4() {
 	bounds_sw.p0 = Vector();
 	bounds_sw.p1 = frame_size;
 
+	Rect bounds_cl = bounds_sw;
+
 	Rect bounds_file;
 	bounds_file.p0 = Vector(0.0, 450.0);
 	bounds_file.p1 = Vector(500.0, -50.0);
@@ -422,7 +425,7 @@ void Test::test4() {
 			}
 
 			{
-				Measure t("test_4_gl_stencil_send_data");
+				//Measure t("test_4_gl_stencil_send_data");
 				glBufferSubData( GL_ARRAY_BUFFER,
 								 0,
 								 vertices.size()*sizeof(vertices.front()),
@@ -435,7 +438,7 @@ void Test::test4() {
 			}
 
 			{
-				Measure t("test_4_gl_stencil.tga");
+				Measure t("test_4_gl_stencil_render.tga");
 				for(int i = 0; i < (int)contours_gl.size(); ++i) {
 					draw_contour(
 						starts[i],
@@ -448,16 +451,16 @@ void Test::test4() {
 		}
 
 		// gl_triangles
-
 		/*
+
 		{
-			Measure t("test_4_gl_triangles.tga", true);
+			Measure t("test_4_gl_triangles", false);
 
 			GLuint index_buffer_id = 0;
 			vector<int> triangle_starts(contours_gl.size());
 			vector<int> triangle_counts(contours_gl.size());
 			vector<int> triangles;
-			vertices.clear();
+			vector<vec2f> vertices;
 			vertices.reserve(commands_count);
 
 			{
@@ -484,7 +487,7 @@ void Test::test4() {
 				}
 			}
 
-			//cout << triangles.size() << " triangles" << endl;
+			cout << triangles.size() << " triangles" << endl;
 
 			{
 				//Measure t("test_4_gl_triangles_prepare_vertices");
@@ -508,12 +511,13 @@ void Test::test4() {
 			}
 
 			{
-				Measure t("test_4_gl_triangles_render");
+				Measure t("test_4_gl_triangles.tga");
 				for(int i = 0; i < (int)contours_gl.size(); ++i) {
 					e.shaders.color(contours_gl[i].color);
 					glDrawElements(GL_TRIANGLES, triangle_counts[i], GL_UNSIGNED_INT, (char*)NULL + triangle_starts[i]*sizeof(int));
 				}
 			}
+
 		}
 		*/
 	}
@@ -522,22 +526,26 @@ void Test::test4() {
 		// software
 
 		Surface surface(width, height);
-		Measure t("test_4_sw.tga", surface, true);
+		Measure t("test_4_sw.tga", surface, false);
 
 		vector<ContourInfo> contours_sw = contours;
 		for(vector<ContourInfo>::iterator i = contours_sw.begin(); i != contours_sw.end(); ++i)
 			i->contour.transform(bounds_file, bounds_sw);
 
+		int count = 0;
 		vector<Polyspan> polyspans(contours_sw.size());
 		{
 			Measure t("test_4_sw_build_polyspans");
 			for(int i = 0; i < (int)contours_sw.size(); ++i) {
 				polyspans[i].init(0, 0, width, height);
 				contours_sw[i].contour.to_polyspan(polyspans[i]);
+				count += polyspans[i].get_covers().size();
 				polyspans[i].sort_marks();
 			}
 		}
 
+		cout << setbase(10) << count << endl;
+
 		{
 			Measure t("test_4_sw_render_polyspans");
 			for(int i = 0; i < (int)contours_sw.size(); ++i)
@@ -549,14 +557,40 @@ void Test::test4() {
 		// cl
 
 		Surface surface(width, height);
+
+		Measure t("test_4_cl.tga", surface, true);
+
 		vector<ContourInfo> contours_cl = contours;
+		vector<vec2f> paths;
+		vector<int> starts(contours_cl.size());
+		vector<int> counts(contours_cl.size());
+		for(int i = 0; i < (int)contours_cl.size(); ++i) {
+			contours_cl[i].contour.transform(bounds_file, bounds_cl);
+			starts[i] = paths.size();
+			for(Contour::ChunkList::const_iterator j = contours_cl[i].contour.get_chunks().begin(); j != contours_cl[i].contour.get_chunks().end(); ++j)
+				paths.push_back(vec2f(j->p1));
+			paths.push_back(paths[starts[i]]);
+			counts[i] = paths.size() - starts[i];
+		}
+
 		ClRender clr(e.cl);
+		clr.send_surface(&surface);
+		clr.send_path(&paths.front(), paths.size());
+
 		{
-			Measure t("test_4_cl.tga", surface, true);
-			clr.send_surface(&surface);
-			for(vector<ContourInfo>::const_iterator i = contours_cl.begin(); i != contours_cl.end(); ++i)
-				clr.contour(i->contour, bounds_file, i->color, i->invert, i->evenodd);
-			clr.receive_surface();
+			Measure t("test_4_cl_render");
+			for(int i = 0; i < (int)contours_cl.size(); ++i)
+				clr.path(starts[i], counts[i], contours_cl[i].color, contours_cl[i].invert, contours_cl[i].evenodd);
+			clr.wait();
 		}
+
+		clr.receive_surface();
+	}
+}
+
+void ClRender::wait() {
+	if (prev_event) {
+		clWaitForEvents(1, &prev_event);
+		prev_event = NULL;
 	}
 }