diff --git a/c++/contourgl/clcontext.cpp b/c++/contourgl/clcontext.cpp
index b7ee20d..3a478dc 100644
--- a/c++/contourgl/clcontext.cpp
+++ b/c++/contourgl/clcontext.cpp
@@ -17,7 +17,6 @@
 
 #include <cassert>
 
-#include <string>
 #include <iostream>
 #include <fstream>
 
@@ -27,7 +26,7 @@
 using namespace std;
 
 
-ClContext::ClContext(): err(), context() {
+ClContext::ClContext(): err(), context(), queue() {
 
 	// platform
 
@@ -59,12 +58,32 @@ ClContext::ClContext(): err(), context() {
 
     context = clCreateContext(0, 1, &devices.front(), NULL, NULL, &err);
     assert(context);
+
+	// command queue
+
+	cl_command_queue queue = clCreateCommandQueue(context, devices[0], 0, NULL);
+	assert(queue);
+
 }
 
 ClContext::~ClContext() {
+	clReleaseCommandQueue(queue);
 	clReleaseContext(context);
 }
 
+cl_program ClContext::load_program(const std::string &filename) {
+	ifstream f(("cl/" + filename).c_str());
+	string text((istreambuf_iterator<char>(f)), istreambuf_iterator<char>());
+	const char *text_pointer = text.c_str();
+	cl_program program = clCreateProgramWithSource(context, 1, &text_pointer, NULL, NULL);
+	assert(program);
+
+	err = clBuildProgram(program, devices.size(), &devices.front(), "", NULL, NULL);
+	assert(!err);
+
+	return program;
+}
+
 void ClContext::hello() {
 
 	// data
@@ -78,14 +97,7 @@ void ClContext::hello() {
 
 	// program
 
-	ifstream f("cl/hello.cl");
-	string text((istreambuf_iterator<char>(f)), istreambuf_iterator<char>());
-	const char *text_pointer = text.c_str();
-	cl_program program = clCreateProgramWithSource(context, 1, &text_pointer, NULL, NULL);
-	assert(program);
-
-	err = clBuildProgram(program, devices.size(), &devices.front(), "", NULL, NULL);
-	assert(!err);
+	cl_program program = load_program("hello.cl");
 
 	// kernel
 
@@ -94,11 +106,6 @@ void ClContext::hello() {
 	err = clSetKernelArg(kernel, 0, sizeof(buffer), &buffer);
 	assert(!err);
 
-	// command queue
-
-	cl_command_queue queue = clCreateCommandQueue(context, devices[0], 0, NULL);
-	assert(queue);
-
 	size_t work_group_size = sizeof(data);
 	cl_event event = NULL;
 	err = clEnqueueNDRangeKernel(
@@ -120,4 +127,10 @@ void ClContext::hello() {
 	clEnqueueReadBuffer(queue, buffer, CL_TRUE, 0, sizeof(data), data, 0, NULL, &event);
 	clWaitForEvents(1, &event);
 	cout << data << endl;
+
+	// deinitialize
+
+	clReleaseKernel(kernel);
+	clReleaseProgram(program);
+	clReleaseMemObject(buffer);
 }
diff --git a/c++/contourgl/clcontext.h b/c++/contourgl/clcontext.h
index 2114f9a..6ced007 100644
--- a/c++/contourgl/clcontext.h
+++ b/c++/contourgl/clcontext.h
@@ -19,6 +19,7 @@
 #define _CLCONTEXT_H_
 
 #include <vector>
+#include <string>
 
 #include <CL/opencl.h>
 
@@ -28,11 +29,13 @@ public:
 	cl_int err;
 	cl_context context;
 	std::vector<cl_device_id> devices;
+	cl_command_queue queue;
 
 	ClContext();
 	~ClContext();
 
 	void hello();
+	cl_program load_program(const std::string &filename);
 };
 
 #endif
diff --git a/c++/contourgl/clrender.cpp b/c++/contourgl/clrender.cpp
index 18d0dae..d769b08 100644
--- a/c++/contourgl/clrender.cpp
+++ b/c++/contourgl/clrender.cpp
@@ -15,25 +15,280 @@
     along with this program.  If not, see <http://www.gnu.org/licenses/>.
 */
 
+#include <cassert>
+
+#include <algorithm>
+
 #include "clrender.h"
+#include "measure.h"
 
 
 using namespace std;
 
 
-ClRender::ClRender(ClContext &cl): cl(cl) {
-	// TODO: init programs
+ClRender::ClRender(ClContext &cl):
+	cl(cl),
+	contour_program(),
+	contour_lines_kernel(),
+	contour_fill_kernel(),
+	surface(),
+	rows_buffer(),
+	mark_buffer(),
+	surface_buffer(),
+	rows_count(),
+	even_rows_count(),
+	odd_rows_count()
+{
+	contour_program = cl.load_program("contour.cl");
+	contour_lines_kernel = clCreateKernel(contour_program, "lines", NULL);
+	assert(contour_lines_kernel);
+	contour_fill_kernel = clCreateKernel(contour_program, "fill", NULL);
+	assert(contour_fill_kernel);
 }
 
 ClRender::~ClRender() {
-	// TODO: release programs
+	send_surface(NULL);
+	clReleaseKernel(contour_fill_kernel);
+	clReleaseKernel(contour_lines_kernel);
+	clReleaseProgram(contour_program);
+}
+
+void ClRender::send_surface(Surface *surface) {
+	if (this->surface == surface) return;
+
+	cl.err = clFinish(cl.queue);
+	assert(cl.err);
+
+	if (this->surface) {
+		rows.clear();
+		clReleaseMemObject(rows_buffer);
+		clReleaseMemObject(mark_buffer);
+		clReleaseMemObject(surface_buffer);
+	}
+
+	this->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);
+
+		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(vec2f), NULL,
+			NULL );
+		assert(mark_buffer);
+
+		surface_buffer = clCreateBuffer(
+			cl.context, CL_MEM_READ_WRITE,
+			surface->data_size(), NULL,
+			NULL );
+		assert(surface_buffer);
+
+		cl.err |= clEnqueueWriteBuffer(
+			cl.queue, surface_buffer, CL_TRUE,
+			0, surface->data_size(), surface->data,
+			0, NULL, NULL );
+
+		cl.err |= clFinish(cl.queue);
+		assert(cl.err);
+	}
+}
+
+Surface* ClRender::receive_surface() {
+	if (surface) {
+		Measure t("ClRender::receive_surface");
+
+		cl.err |= clFinish(cl.queue);
+		cl.err |= clEnqueueReadBuffer(
+			cl.queue, surface_buffer, CL_TRUE,
+			0, sizeof(surface->data_size()), surface->data,
+			0, NULL, NULL );
+		cl.err |= clFinish(cl.queue);
+		assert(cl.err);
+	}
+	return surface;
 }
 
-void ClRender::contour(const Contour &contour, Surface &target) {
-	// TODO:
-	// split contour
-	// resterize lines
-	// fill
+
+void ClRender::contour(const Contour &contour, const Rect &rect, const Color &color) {
+	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");
+		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 = 1; 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);
+					int row = (int)floorf(min(l.p0.y, l.p1.y));
+					row = row % 2 ? row/2 : even_rows_count + row/2;
+					line_rows.push_back(row);
+					lines.push_back(l);
+					++rows[row].second;
+				}
+			}
+			prev = i->p1;
+		}
+
+		// 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)
+			sorted_lines[ --rows[line_rows[i]].first ] = lines[i];
+	}
+
+	cl_mem lines_buffer = NULL;
+
+	{
+		Measure t("create lines buffer");
+
+		lines_buffer = clCreateBuffer(
+			cl.context, CL_MEM_READ_ONLY,
+			sorted_lines.size()*sizeof(sorted_lines.front()), NULL,
+			NULL );
+		assert(lines_buffer);
+	}
+
+	{
+		Measure t("enqueue commands");
+
+		clFinish(cl.queue);
+
+		// kernel args
+
+		cl.err |= clSetKernelArg(contour_lines_kernel, 0, sizeof(lines_buffer), &lines_buffer);
+		cl.err |= clSetKernelArg(contour_lines_kernel, 1, sizeof(rows_buffer), &rows_buffer);
+		cl.err |= clSetKernelArg(contour_lines_kernel, 2, sizeof(mark_buffer), &mark_buffer);
+		assert(!cl.err);
+
+		cl.err |= clSetKernelArg(contour_fill_kernel, 0, sizeof(mark_buffer), &mark_buffer);
+		cl.err |= clSetKernelArg(contour_fill_kernel, 1, sizeof(surface_buffer), &surface_buffer);
+		cl.err |= clSetKernelArg(contour_fill_kernel, 2, sizeof(Color::type), &color.r);
+		cl.err |= clSetKernelArg(contour_fill_kernel, 3, sizeof(Color::type), &color.g);
+		cl.err |= clSetKernelArg(contour_fill_kernel, 4, sizeof(Color::type), &color.b);
+		cl.err |= clSetKernelArg(contour_fill_kernel, 5, sizeof(Color::type), &color.a);
+		assert(!cl.err);
+
+		// prepare buffers
+
+		cl_event prepare_buffers_events[3] = { };
+
+		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.err |= clEnqueueWriteBuffer(
+			cl.queue, rows_buffer, CL_TRUE,
+			0, rows.size()*sizeof(rows.front()), &rows.front(),
+			0, NULL, &prepare_buffers_events[1] );
+
+		vec2f pattern;
+		cl.err |= clEnqueueFillBuffer(
+			cl.queue, mark_buffer,
+			&pattern, sizeof(pattern),
+			0, surface->count()*sizeof(vec2f),
+			0, NULL, &prepare_buffers_events[2] );
+
+		// run kernels
+
+		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);
+
+		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);
+
+		cl_event fill_event = NULL;
+		cl.err = clEnqueueNDRangeKernel(
+			cl.queue,
+			contour_fill_kernel,
+			1,
+			NULL,
+			&rows_count,
+			NULL,
+			1,
+			&lines_even_event,
+			&fill_event );
+		assert(!cl.err);
+
+		clWaitForEvents(1, &fill_event);
+	}
+
+	{
+		Measure t("release lines buffer");
+		clReleaseMemObject(lines_buffer);
+	}
 }
 
 
diff --git a/c++/contourgl/clrender.h b/c++/contourgl/clrender.h
index 065e10b..3dd269d 100644
--- a/c++/contourgl/clrender.h
+++ b/c++/contourgl/clrender.h
@@ -27,13 +27,31 @@
 
 
 class ClRender {
-public:
+private:
 	ClContext &cl;
+	cl_program contour_program;
+	cl_kernel contour_lines_kernel;
+	cl_kernel contour_fill_kernel;
+
+	Surface *surface;
+	cl_mem rows_buffer;
+	cl_mem mark_buffer;
+	cl_mem surface_buffer;
 
+	size_t rows_count;
+	size_t even_rows_count;
+	size_t odd_rows_count;
+
+	typedef std::pair<int, int> Row;
+	std::vector<Row> rows;
+
+public:
 	ClRender(ClContext &cl);
 	~ClRender();
 
-	void contour(const Contour &contour, Surface &target);
+	void send_surface(Surface *surface);
+	Surface* receive_surface();
+	void contour(const Contour &contour, const Rect &rect, const Color &color);
 };
 
 
diff --git a/c++/contourgl/geometry.h b/c++/contourgl/geometry.h
index 825b1c5..88691cc 100644
--- a/c++/contourgl/geometry.h
+++ b/c++/contourgl/geometry.h
@@ -58,7 +58,7 @@ public:
 
 	vec2():
 		x(), y() { }
-	vec2(type x, type y):
+	vec2(const type &x, const type &y):
 		x(x), y(y) { }
 
 	template<typename TT>
@@ -81,9 +81,9 @@ public:
 	vec2 operator/(const vec2 &a) const
 		{ return vec2(x/a.x, y/a.y); }
 
-	vec2 operator*(type a) const
+	vec2 operator*(const type &a) const
 		{ return vec2(x*a, y*a); }
-	vec2 operator/(type a) const
+	vec2 operator/(const type &a) const
 		{ return vec2(x/a, y/a); }
 
 	type dot(const vec2 &a) const
@@ -92,33 +92,52 @@ public:
 	static vec2 zero() { return vec2(); }
 };
 
-typedef vec2<Real> Vector;
+template<typename T>
+class line2 {
+public:
+	typedef T type;
+	vec2<type> p0, p1;
+	line2() { }
+	line2(const vec2<type> &p0, const vec2<type> &p1): p0(p0), p1(p1) { }
+	line2(const type &x0, const type &y0, const type &x1, const type &y1): p0(x0, y0), p1(x1, y1) { }
+};
 
-class Rect {
+template<typename T>
+class rect {
 public:
-	Vector p0, p1;
+	typedef T type;
+	vec2<type> p0, p1;
 
-    bool intersects(const Rect &other) const
+    bool intersects(const rect &other) const
         { return ::intersects(p0.x, p1.x, other.p0.x, other.p1.x)
               && ::intersects(p0.y, p1.y, other.p0.y, other.p1.y); }
 
-    Rect expand(const Vector &p) const {
-    	Rect r;
+    rect expand(const vec2<type> &p) const {
+    	rect r;
     	r.p0.x = std::min(std::min(p0.x, p1.x), p.x);
     	r.p0.y = std::min(std::min(p0.y, p1.y), p.y);
     	r.p1.x = std::max(std::max(p0.x, p1.x), p.x);
     	r.p1.y = std::max(std::max(p0.y, p1.y), p.y);
     	return r;
     }
+
+    rect() { }
+    rect(const vec2<type> &p0, const vec2<type> &p1): p0(p0), p1(p1) { }
+    rect(const type &x0, const type &y0, const type &x1, const type &y1): p0(x0, y0), p1(x1, y1) { }
 };
 
+typedef vec2<Real> Vector;
+typedef line2<Real> Line;
+typedef rect<Real> Rect;
+
+typedef vec2<float> vec2f;
+typedef line2<float> line2f;
+typedef rect<float> rectf;
+
 class ContextRect {
 public:
 	int minx, miny, maxx, maxy;
 	ContextRect(): minx(), miny(), maxx(), maxy() { }
 };
 
-inline bool intersects(const Rect &a, const Rect &b)
-    { return a.intersects(b); }
-
 #endif
diff --git a/c++/contourgl/measure.cpp b/c++/contourgl/measure.cpp
index b57a5c8..06db589 100644
--- a/c++/contourgl/measure.cpp
+++ b/c++/contourgl/measure.cpp
@@ -56,8 +56,12 @@ Measure::~Measure() {
 
 	clock_t dt = sub_tasks ? sub_tasks : clock() - t;
 	Real ms = 1000.0*(Real)(clock() - t)/(Real)(CLOCKS_PER_SEC);
+
 	cout << setw(8) << fixed << setprecision(3)
-	     << ms << " ms - " << filename << flush << endl;
+	     << string(stack.size()*2, ' ') << "|"
+		 << ms << " ms - "
+		 << filename
+		 << endl << flush;
 
 	if (tga) {
 		if (surface)
diff --git a/c++/contourgl/swrender.h b/c++/contourgl/swrender.h
index 679ae7e..3a4530f 100644
--- a/c++/contourgl/swrender.h
+++ b/c++/contourgl/swrender.h
@@ -49,6 +49,7 @@ public:
 
 	void clear() { memset(data, 0, count()*sizeof(Color)); }
 	int count() const { return width*height; }
+	int data_size() const { return count()*sizeof(Color); }
 	Color* operator[] (int row) { return data + row*width; }
 	const Color* operator[] (int row) const { return data + row*width; }
 };