diff --git a/c++/contourgl/Makefile b/c++/contourgl/Makefile
index ba6608e..98eca6f 100644
--- a/c++/contourgl/Makefile
+++ b/c++/contourgl/Makefile
@@ -1,6 +1,6 @@
 DEPLIBS = gl x11 OpenCL
 
-CXXFLAGS = -O0 -g -Wall -fmessage-length=0 `pkg-config --cflags $(DEPLIBS)` -DGL_GLEXT_PROTOTYPES
+CXXFLAGS = -O3 -Wall -fmessage-length=0 `pkg-config --cflags $(DEPLIBS)` -DGL_GLEXT_PROTOTYPES
 
 HEADERS = \
 	clcontext.h \
diff --git a/c++/contourgl/cl/contour.cl b/c++/contourgl/cl/contour.cl
index 6c9f926..d03cd67 100644
--- a/c++/contourgl/cl/contour.cl
+++ b/c++/contourgl/cl/contour.cl
@@ -24,21 +24,20 @@ __kernel void clear2f(
 
 __kernel void lines(
 	int width, 
-	__global float *lines,
-	__global int *rows,
+	__global float4 *lines,
+	__global int2 *rows,
 	__global float2 *mark_buffer )
 {
 	const float e = 1e-6f;
-	__global int *row = rows + 2*get_global_id(0);
+	int2 row = rows[get_global_id(0)];
 	sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE
 					  | CLK_ADDRESS_NONE
 			          | CLK_FILTER_NEAREST;
 	int w = width;
-	__global float *i = lines + 4*row[0];
-	__global float *end = i + 4*row[1];
-	for(; i < end; i += 4) {
-		float2 p0 = { i[0], i[1] };
-		float2 p1 = { i[2], i[3] };
+	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);
@@ -81,10 +80,7 @@ __kernel void fill(
 	__global float2 *mark_buffer,
 	__read_only image2d_t surface_read_image,
 	__write_only image2d_t surface_write_image,
-	float color_r,
-	float color_g,
-	float color_b,
-	float color_a,
+	float4 color,
 	int invert,
 	int evenodd )
 {
@@ -93,7 +89,7 @@ __kernel void fill(
 	const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE
 							| CLK_ADDRESS_NONE
 			                | CLK_FILTER_NEAREST;
-	float4 color = { color_r, color_g, color_b, color_a };
+	float4 cl = color;
 	float cover = 0.f;
 	__global float2 *mark = mark_buffer + id*w;
 	for(int2 coord = { 0, id }; coord.x < w; ++coord.x, ++mark) {
@@ -103,14 +99,14 @@ __kernel void fill(
 		cover += m.y;
 		alpha = evenodd ? (1.f - fabs(1.f - alpha - 2.f*floor(0.5f*alpha)))
 				        : fmin(alpha, 1.f);
-		alpha *= color.w;
+		alpha *= cl.w;
 		if (invert) alpha = 1.f - alpha;
 		float alpha_inv = 1.f - alpha;
-
+		
 		float4 c = read_imagef(surface_read_image, sampler, coord);
-		c.x = c.x*alpha_inv + color.x*alpha;
-		c.y = c.y*alpha_inv + color.y*alpha;
-		c.z = c.z*alpha_inv + color.z*alpha;
+		c.x = c.x*alpha_inv + cl.x*alpha;
+		c.y = c.y*alpha_inv + cl.y*alpha;
+		c.z = c.z*alpha_inv + cl.z*alpha;
 		c.w = min(c.w + alpha, 1.f);
 		write_imagef(surface_write_image, coord, c);
 	}
diff --git a/c++/contourgl/clrender.cpp b/c++/contourgl/clrender.cpp
index 1694f38..f6335c6 100644
--- a/c++/contourgl/clrender.cpp
+++ b/c++/contourgl/clrender.cpp
@@ -150,7 +150,7 @@ void ClRender::contour(const Contour &contour, const Rect &rect, const Color &co
 	}
 
 	{
-		//Measure t("split");
+		Measure t("split");
 		splitted.allow_split_lines = true;
 		transformed.split(splitted, to, Vector(0.5, 0.5));
 	}
@@ -160,7 +160,7 @@ void ClRender::contour(const Contour &contour, const Rect &rect, const Color &co
 	vector<int> line_rows;
 
 	{
-		//Measure t("sort lines");
+		Measure t("sort lines");
 
 		// reset rows
 		for(int i = 0; i < (int)rows_count; ++i)
@@ -216,7 +216,7 @@ void ClRender::contour(const Contour &contour, const Rect &rect, const Color &co
 	cl_mem lines_buffer = NULL;
 
 	{
-		//Measure t("create lines buffer");
+		Measure t("create lines buffer");
 
 		lines_buffer = clCreateBuffer(
 			cl.context, CL_MEM_READ_ONLY,
@@ -226,7 +226,7 @@ void ClRender::contour(const Contour &contour, const Rect &rect, const Color &co
 	}
 
 	{
-		//Measure t("enqueue commands");
+		Measure t("enqueue commands");
 
 		// kernel args
 
@@ -246,12 +246,9 @@ void ClRender::contour(const Contour &contour, const Rect &rect, const Color &co
 		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::type), &color.r);
-		cl.err |= clSetKernelArg(contour_fill_kernel, 5, sizeof(Color::type), &color.g);
-		cl.err |= clSetKernelArg(contour_fill_kernel, 6, sizeof(Color::type), &color.b);
-		cl.err |= clSetKernelArg(contour_fill_kernel, 7, sizeof(Color::type), &color.a);
-		cl.err |= clSetKernelArg(contour_fill_kernel, 8, sizeof(int), &iinvert);
-		cl.err |= clSetKernelArg(contour_fill_kernel, 9, sizeof(int), &ievenodd);
+		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
@@ -329,7 +326,7 @@ void ClRender::contour(const Contour &contour, const Rect &rect, const Color &co
 	}
 
 	{
-		//Measure t("release lines buffer");
+		Measure t("release lines buffer");
 		clReleaseMemObject(lines_buffer);
 	}
 }
diff --git a/c++/contourgl/contourgl.cpp b/c++/contourgl/contourgl.cpp
index 207bfa4..8f87517 100644
--- a/c++/contourgl/contourgl.cpp
+++ b/c++/contourgl/contourgl.cpp
@@ -27,7 +27,7 @@ int main() {
 	Environment e;
 	Test test(e);
 
-	e.cl.hello();
+	//e.cl.hello();
 	//test.test2();
 	//test.test3();
 	test.test4();
diff --git a/c++/contourgl/glcontext.cpp b/c++/contourgl/glcontext.cpp
index 06cfcca..ea3623a 100644
--- a/c++/contourgl/glcontext.cpp
+++ b/c++/contourgl/glcontext.cpp
@@ -48,7 +48,7 @@ GlContext::GlContext():
 	int framebuffer_width = 512;
 	int framebuffer_height = 512;
 	int framebuffer_samples = 16;
-	bool antialising = false;
+	bool antialising = true;
 	bool hdr = false;
 
 	// display
@@ -80,11 +80,9 @@ GlContext::GlContext():
 
 	// pbuffer
 
-	int pbuffer_width = 256;
-	int pbuffer_height = 256;
 	int pbuffer_attribs[] = {
-		GLX_PBUFFER_WIDTH, pbuffer_width,
-		GLX_PBUFFER_HEIGHT, pbuffer_height,
+		GLX_PBUFFER_WIDTH, 256,
+		GLX_PBUFFER_HEIGHT, 256,
 		None };
 	pbuffer = glXCreatePbuffer(display, config, pbuffer_attribs);
 	assert(pbuffer);
@@ -145,22 +143,23 @@ GlContext::GlContext():
 	// view port
 
 	glViewport(0, 0, framebuffer_width, framebuffer_height);
+
+	check();
 }
 
 GlContext::~GlContext() {
 	glBindFramebuffer(GL_READ_FRAMEBUFFER, 0);
-	glDeleteFramebuffers(1, &framebuffer_id);
-
+	glBindFramebuffer(GL_DRAW_FRAMEBUFFER, 0);
 	glBindTexture(GL_TEXTURE_2D, 0);
+	glBindTexture(GL_TEXTURE_2D_MULTISAMPLE, 0);
+	glBindRenderbuffer(GL_RENDERBUFFER, 0);
+
+	glDeleteFramebuffers(1, &framebuffer_id);
+	glDeleteRenderbuffers(1, &renderbuffer_id);
 	glDeleteTextures(1, &texture_id);
 
-	glBindFramebuffer(GL_DRAW_FRAMEBUFFER, 0);
 	glDeleteFramebuffers(1, &multisample_framebuffer_id);
-
-	glBindRenderbuffer(GL_RENDERBUFFER, 0);
 	glDeleteRenderbuffers(1, &multisample_renderbuffer_id);
-
-	glBindTexture(GL_TEXTURE_2D_MULTISAMPLE, 0);
 	glDeleteTextures(1, &multisample_texture_id);
 
 	glXMakeContextCurrent(display, None, None, NULL);
diff --git a/c++/contourgl/measure.cpp b/c++/contourgl/measure.cpp
index f5ff2d1..6009ddd 100644
--- a/c++/contourgl/measure.cpp
+++ b/c++/contourgl/measure.cpp
@@ -29,26 +29,15 @@ using namespace std;
 std::vector<Measure*> Measure::stack;
 
 
-Measure::Measure(const std::string &filename):
-	filename(filename),
-	surface(),
-	tga(filename.size() > 4 && filename.substr(filename.size()-4, 4) == ".tga"),
-	sub_tasks(),
-	t()
-{
-	cout << string(stack.size()*2, ' ') << "begin             " << filename << endl << flush;
-	stack.push_back(this);
-	t = clock();
-}
-
-Measure::Measure(const std::string &filename, Surface &surface):
-	filename(filename),
-	surface(&surface),
-	tga(filename.size() > 4 && filename.substr(filename.size()-4, 4) == ".tga"),
-	sub_tasks(),
-	t()
-{
-	cout << string(stack.size()*2, ' ') << "begin             " << filename << endl << flush;
+void Measure::init() {
+	hide = !stack.empty() && stack.back()->hide_subs;
+	hide_subs |= hide;
+	tga = filename.size() > 4 && filename.substr(filename.size()-4, 4) == ".tga";
+	if (!hide)
+		cout << string(stack.size()*2, ' ')
+		     << "begin             "
+			 << filename
+			 << endl << flush;
 	stack.push_back(this);
 	t = clock();
 }
@@ -56,14 +45,15 @@ Measure::Measure(const std::string &filename, Surface &surface):
 Measure::~Measure() {
 	if (!surface && tga) glFinish();
 
-	clock_t dt = sub_tasks ? sub_tasks : clock() - t;
+	clock_t dt = subs ? subs : clock() - t;
 	Real ms = 1000.0*(Real)dt/(Real)(CLOCKS_PER_SEC);
 
-	cout << string((stack.size()-1)*2, ' ') << "end "
-		 << setw(8) << fixed << setprecision(3)
-		 << ms << " ms - "
-		 << filename
-		 << endl << flush;
+	if (!hide)
+		cout << string((stack.size()-1)*2, ' ') << "end "
+			 << setw(8) << fixed << setprecision(3)
+			 << ms << " ms - "
+			 << filename
+			 << endl << flush;
 
 	if (tga) {
 		if (surface)
@@ -80,6 +70,6 @@ Measure::~Measure() {
 	}
 
 	stack.pop_back();
-	if (!stack.empty()) stack.back()->sub_tasks += dt;
+	if (!stack.empty()) stack.back()->subs += dt;
 }
 
diff --git a/c++/contourgl/measure.h b/c++/contourgl/measure.h
index 1903d74..ec5156a 100644
--- a/c++/contourgl/measure.h
+++ b/c++/contourgl/measure.h
@@ -32,14 +32,23 @@ private:
 	std::string filename;
 	Surface *surface;
 	bool tga;
-	clock_t sub_tasks;
+	bool hide;
+	bool hide_subs;
+	clock_t subs;
 	clock_t t;
 
-	Measure(const Measure&): surface(), tga(), sub_tasks(), t() { }
+	Measure(const Measure&): surface(), tga(), hide(), hide_subs(), subs(), t() { }
 	Measure& operator= (const Measure&) { return *this; }
+	void init();
 public:
-	Measure(const std::string &filename);
-	Measure(const std::string &filename, Surface &surface);
+	Measure(const std::string &filename, bool hide_subs = false):
+		filename(filename), surface(), tga(), hide(), hide_subs(hide_subs), subs(), t()
+	{ init(); }
+
+	Measure(const std::string &filename, Surface &surface, bool hide_subs = false):
+		filename(filename), surface(&surface), tga(), hide(), hide_subs(hide_subs), subs(), t()
+	{ init(); }
+
 	~Measure();
 };
 
diff --git a/c++/contourgl/test.cpp b/c++/contourgl/test.cpp
index 8185fe4..5995b82 100644
--- a/c++/contourgl/test.cpp
+++ b/c++/contourgl/test.cpp
@@ -33,9 +33,7 @@ void Test::draw_contour(int start, int count, bool even_odd, bool invert, const 
 	glEnable(GL_STENCIL_TEST);
 
 	// render mask
-	GLint draw_buffer;
-	glGetIntegerv(GL_DRAW_BUFFER, &draw_buffer);
-	glDrawBuffer(GL_NONE);
+	glColorMask(GL_FALSE, GL_FALSE, GL_FALSE, GL_FALSE);
 	glClear(GL_STENCIL_BUFFER_BIT);
 	glStencilFunc(GL_ALWAYS, 0, 0);
 	if (even_odd) {
@@ -46,7 +44,7 @@ void Test::draw_contour(int start, int count, bool even_odd, bool invert, const 
 	}
 	e.shaders.simple();
 	glDrawArrays(GL_TRIANGLE_STRIP, start, count);
-	glDrawBuffer((GLenum)draw_buffer);
+	glColorMask(GL_TRUE, GL_TRUE, GL_TRUE, GL_TRUE);
 
 	// fill mask
 	glStencilOp(GL_KEEP, GL_KEEP, GL_KEEP);
@@ -164,7 +162,7 @@ void Test::test2() {
 	GLuint buffer_id = 0;
 	GLuint array_id = 0;
 	int count = 0;
-	vector< vec2<float> > vertices;
+	vector<vec2f> vertices;
 
 	{
 		Measure t("test_2_init_buffer");
@@ -172,11 +170,9 @@ void Test::test2() {
 		glGenBuffers(1, &buffer_id);
 		glBindBuffer(GL_ARRAY_BUFFER, buffer_id);
 		glBufferData( GL_ARRAY_BUFFER,
-				      vertices.size()*sizeof(vec2<float>),
+				      vertices.size()*sizeof(vec2f),
 					  &vertices.front(),
 					  GL_DYNAMIC_DRAW );
-		vertices.clear();
-		vertices.reserve(4+4*chunks.size());
 
 		glGenVertexArrays(1, &array_id);
 		glBindVertexArray(array_id);
@@ -189,26 +185,29 @@ void Test::test2() {
 		glFinish();
 		glClear(GL_COLOR_BUFFER_BIT);
 		glFinish();
+
+		vertices.clear();
+		vertices.reserve(4+4*chunks.size());
 	}
 
 	{
 		Measure t("test_2_prepare_data");
-		vertices.push_back(vec2<float>(bounds.p0.x, bounds.p0.y));
-		vertices.push_back(vec2<float>(bounds.p0.x, bounds.p1.y));
-		vertices.push_back(vec2<float>(bounds.p1.x, bounds.p0.y));
-		vertices.push_back(vec2<float>(bounds.p1.x, bounds.p1.y));
-		vertices.push_back(vec2<float>());
-		vertices.push_back(vec2<float>());
+		vertices.push_back(vec2f(bounds.p0.x, bounds.p0.y));
+		vertices.push_back(vec2f(bounds.p0.x, bounds.p1.y));
+		vertices.push_back(vec2f(bounds.p1.x, bounds.p0.y));
+		vertices.push_back(vec2f(bounds.p1.x, bounds.p1.y));
+		vertices.push_back(vec2f());
+		vertices.push_back(vec2f());
 		for(Contour::ChunkList::const_iterator i = chunks.begin(); i != chunks.end(); ++i) {
 			if ( i->type == Contour::LINE
 			  || i->type == Contour::CLOSE)
 			{
-				vertices.push_back(vec2<float>(i->p1));
-				vertices.push_back(vec2<float>(-1.f, (float)i->p1.y));
+				vertices.push_back(vec2f(i->p1));
+				vertices.push_back(vec2f(-1.f, (float)i->p1.y));
 			} else {
 				vertices.push_back(vertices.back());
 				vertices.push_back(vertices.back());
-				vertices.push_back(vec2<float>(i->p1));
+				vertices.push_back(vec2f(i->p1));
 				vertices.push_back(vertices.back());
 			}
 		}
@@ -360,153 +359,160 @@ void Test::test4() {
 
 		// gl_stencil
 
-		GLuint buffer_id = 0;
-		GLuint array_id = 0;
-		vector< vec2<float> > vertices;
-		vector<int> starts(contours_gl.size());
-		vector<int> counts(contours_gl.size());
-
 		{
-			Measure t("test_4_gl_init_buffer");
-			vertices.resize(4 + 4*commands_count + 2*contours_gl.size());
-			glGenBuffers(1, &buffer_id);
-			glBindBuffer(GL_ARRAY_BUFFER, buffer_id);
-			glBufferData( GL_ARRAY_BUFFER,
-						  vertices.size()*sizeof(vec2<float>),
-						  &vertices.front(),
-						  GL_DYNAMIC_DRAW );
-			vertices.clear();
-			vertices.reserve(4 + 4*commands_count);
+			Measure t("test_4_gl_stencil", true);
 
-			glGenVertexArrays(1, &array_id);
-			glBindVertexArray(array_id);
+			GLuint buffer_id = 0;
+			GLuint array_id = 0;
+			vector<vec2f> vertices;
+			vector<int> starts(contours_gl.size());
+			vector<int> counts(contours_gl.size());
 
-			glEnableVertexAttribArray(0);
-			glVertexAttribPointer(0, 2, GL_FLOAT, GL_TRUE, 0, NULL);
-
-			e.shaders.color(Color(0.f, 0.f, 1.f, 1.f));
-			glDrawArrays(GL_TRIANGLE_STRIP, 0, vertices.size());
-			glFinish();
-			glClear(GL_COLOR_BUFFER_BIT);
-			glFinish();
-		}
+			{
+				//Measure t("test_4_gl_init_buffer");
+				vertices.resize(4 + 4*commands_count + 2*contours_gl.size());
+				glGenBuffers(1, &buffer_id);
+				glBindBuffer(GL_ARRAY_BUFFER, buffer_id);
+				glBufferData( GL_ARRAY_BUFFER,
+							  vertices.size()*sizeof(vec2f),
+							  &vertices.front(),
+							  GL_DYNAMIC_DRAW );
+				vertices.clear();
+				vertices.reserve(4 + 4*commands_count);
+
+				glGenVertexArrays(1, &array_id);
+				glBindVertexArray(array_id);
+
+				glEnableVertexAttribArray(0);
+				glVertexAttribPointer(0, 2, GL_FLOAT, GL_TRUE, 0, NULL);
+
+				e.shaders.color(Color(0.f, 0.f, 1.f, 1.f));
+				glDrawArrays(GL_TRIANGLE_STRIP, 0, vertices.size());
+				glFinish();
+				glClear(GL_COLOR_BUFFER_BIT);
+				glFinish();
+			}
 
-		{
-			Measure t("test_4_gl_stencil_prepare_data");
-			vertices.push_back(vec2<float>(bounds_gl.p0.x, bounds_gl.p0.y));
-			vertices.push_back(vec2<float>(bounds_gl.p0.x, bounds_gl.p1.y));
-			vertices.push_back(vec2<float>(bounds_gl.p1.x, bounds_gl.p0.y));
-			vertices.push_back(vec2<float>(bounds_gl.p1.x, bounds_gl.p1.y));
-			for(int i = 0; i < (int)contours_gl.size(); ++i) {
-				starts[i] = (int)vertices.size();
-				const Contour::ChunkList &chunks = contours_gl[i].contour.get_chunks();
-				for(Contour::ChunkList::const_iterator j = chunks.begin(); j != chunks.end(); ++j) {
-					if (j->type == Contour::LINE) {
-						vertices.push_back(vec2<float>(j->p1));
-						vertices.push_back(vec2<float>(-1.f, (float)j->p1.y));
-					} else
-					if (j->type == Contour::CLOSE) {
-						vertices.push_back(vec2<float>(j->p1));
-						vertices.push_back(vec2<float>(-1.f, (float)j->p1.y));
-					} else {
-						vertices.push_back(vertices.back());
-						vertices.push_back(vec2<float>(j->p1));
-						vertices.push_back(vertices.back());
-						vertices.push_back(vec2<float>(-1.f, (float)j->p1.y));
+			{
+				//Measure t("test_4_gl_stencil_prepare_data");
+				vertices.push_back(vec2f(bounds_gl.p0.x, bounds_gl.p0.y));
+				vertices.push_back(vec2f(bounds_gl.p0.x, bounds_gl.p1.y));
+				vertices.push_back(vec2f(bounds_gl.p1.x, bounds_gl.p0.y));
+				vertices.push_back(vec2f(bounds_gl.p1.x, bounds_gl.p1.y));
+				for(int i = 0; i < (int)contours_gl.size(); ++i) {
+					starts[i] = (int)vertices.size();
+					const Contour::ChunkList &chunks = contours_gl[i].contour.get_chunks();
+					for(Contour::ChunkList::const_iterator j = chunks.begin(); j != chunks.end(); ++j) {
+						if (j->type == Contour::LINE) {
+							vertices.push_back(vec2f(j->p1));
+							vertices.push_back(vec2f(-1.f, (float)j->p1.y));
+						} else
+						if (j->type == Contour::CLOSE) {
+							vertices.push_back(vec2f(j->p1));
+							vertices.push_back(vec2f(-1.f, (float)j->p1.y));
+						} else {
+							vertices.push_back(vertices.back());
+							vertices.push_back(vec2f(j->p1));
+							vertices.push_back(vertices.back());
+							vertices.push_back(vec2f(-1.f, (float)j->p1.y));
+						}
 					}
+					counts[i] = (int)vertices.size() - starts[i];
 				}
-				counts[i] = (int)vertices.size() - starts[i];
 			}
-		}
 
-		{
-			Measure t("test_4_gl_stencil_send_data");
-			glBufferSubData( GL_ARRAY_BUFFER,
-							 0,
-						     vertices.size()*sizeof(vertices.front()),
-						     &vertices.front() );
-		}
+			{
+				Measure t("test_4_gl_stencil_send_data");
+				glBufferSubData( GL_ARRAY_BUFFER,
+								 0,
+								 vertices.size()*sizeof(vertices.front()),
+								 &vertices.front() );
+			}
 
-		{
-			Measure t("test_4_gl_stencil_points.tga");
-			glDrawArrays(GL_POINTS, 0, vertices.size());
-		}
+			{
+				//Measure t("test_4_gl_stencil_points.tga");
+				//glDrawArrays(GL_POINTS, 0, vertices.size());
+			}
 
-		{
-			Measure t("test_4_gl_stencil_render.tga");
-			for(int i = 0; i < (int)contours_gl.size(); ++i) {
-				draw_contour(
-					starts[i],
-					counts[i],
-					contours_gl[i].invert,
-					contours_gl[i].evenodd,
-					contours_gl[i].color );
+			{
+				Measure t("test_4_gl_stencil.tga");
+				for(int i = 0; i < (int)contours_gl.size(); ++i) {
+					draw_contour(
+						starts[i],
+						counts[i],
+						contours_gl[i].invert,
+						contours_gl[i].evenodd,
+						contours_gl[i].color );
+				}
 			}
-			// glDrawArrays(GL_POINTS, 0, vertices.size());
 		}
 
 		// gl_triangles
 
 		/*
-		GLuint index_buffer_id = 0;
-		vector<int> triangle_starts(contours_gl.size());
-		vector<int> triangle_counts(contours_gl.size());
-		vector<int> triangles;
-		vertices.clear();
-		vertices.reserve(commands_count);
-
 		{
-			Measure t("test_4_gl_init_index_buffer");
-			triangles.resize(3*commands_count);
-			glGenBuffers(1, &index_buffer_id);
-			glBindBuffer(GL_ELEMENT_ARRAY_BUFFER, index_buffer_id);
-			glBufferData( GL_ELEMENT_ARRAY_BUFFER,
-						  triangles.size()*sizeof(triangles.front()),
-						  &triangles.front(),
-						  GL_DYNAMIC_DRAW );
-			triangles.clear();
-			triangles.reserve(3*commands_count);
-		}
+			Measure t("test_4_gl_triangles.tga", true);
 
-		{
-			Measure t("test_4_gl_triangulate");
-			int index_offset = 4;
-			for(int i = 0; i < (int)contours_gl.size(); ++i) {
-				triangle_starts[i] = (int)triangles.size();
-				Triangulator::triangulate(contours_gl[i].contour, triangles, index_offset);
-				triangle_counts[i] = (int)triangles.size() - triangle_starts[i];
-				index_offset += (int)contours_gl[i].contour.get_chunks().size();
+			GLuint index_buffer_id = 0;
+			vector<int> triangle_starts(contours_gl.size());
+			vector<int> triangle_counts(contours_gl.size());
+			vector<int> triangles;
+			vertices.clear();
+			vertices.reserve(commands_count);
+
+			{
+				//Measure t("test_4_gl_init_index_buffer");
+				triangles.resize(3*commands_count);
+				glGenBuffers(1, &index_buffer_id);
+				glBindBuffer(GL_ELEMENT_ARRAY_BUFFER, index_buffer_id);
+				glBufferData( GL_ELEMENT_ARRAY_BUFFER,
+							  triangles.size()*sizeof(triangles.front()),
+							  &triangles.front(),
+							  GL_DYNAMIC_DRAW );
+				triangles.clear();
+				triangles.reserve(3*commands_count);
+			}
+
+			{
+				Measure t("test_4_gl_triangulate");
+				int index_offset = 4;
+				for(int i = 0; i < (int)contours_gl.size(); ++i) {
+					triangle_starts[i] = (int)triangles.size();
+					Triangulator::triangulate(contours_gl[i].contour, triangles, index_offset);
+					triangle_counts[i] = (int)triangles.size() - triangle_starts[i];
+					index_offset += (int)contours_gl[i].contour.get_chunks().size();
+				}
 			}
-		}
 
-		cout << triangles.size() << " triangles" << endl;
+			//cout << triangles.size() << " triangles" << endl;
 
-		{
-			Measure t("test_4_gl_triangles_prepare_vertices");
-			for(int i = 0; i < (int)contours_gl.size(); ++i) {
-				const Contour::ChunkList &chunks = contours_gl[i].contour.get_chunks();
-				for(Contour::ChunkList::const_iterator j = chunks.begin(); j != chunks.end(); ++j)
-					vertices.push_back(vec2<float>(j->p1));
+			{
+				//Measure t("test_4_gl_triangles_prepare_vertices");
+				for(int i = 0; i < (int)contours_gl.size(); ++i) {
+					const Contour::ChunkList &chunks = contours_gl[i].contour.get_chunks();
+					for(Contour::ChunkList::const_iterator j = chunks.begin(); j != chunks.end(); ++j)
+						vertices.push_back(vec2f(j->p1));
+				}
 			}
-		}
 
-		{
-			Measure t("test_4_gl_triangles_send_data");
-			glBufferSubData( GL_ARRAY_BUFFER,
-							 4*sizeof(vertices.front()),
-						     vertices.size()*sizeof(vertices.front()),
-						     &vertices.front() );
-			glBufferSubData( GL_ELEMENT_ARRAY_BUFFER,
-							 0,
-						     triangles.size()*sizeof(triangles.front()),
-						     &triangles.front() );
-		}
+			{
+				Measure t("test_4_gl_triangles_send_data");
+				glBufferSubData( GL_ARRAY_BUFFER,
+								 4*sizeof(vertices.front()),
+								 vertices.size()*sizeof(vertices.front()),
+								 &vertices.front() );
+				glBufferSubData( GL_ELEMENT_ARRAY_BUFFER,
+								 0,
+								 triangles.size()*sizeof(triangles.front()),
+								 &triangles.front() );
+			}
 
-		{
-			Measure t("test_4_gl_triangles_render.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));
+			{
+				Measure t("test_4_gl_triangles_render");
+				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));
+				}
 			}
 		}
 		*/
@@ -515,6 +521,9 @@ void Test::test4() {
 	{
 		// software
 
+		Surface surface(width, height);
+		Measure t("test_4_sw.tga", surface, true);
+
 		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);
@@ -529,24 +538,21 @@ void Test::test4() {
 			}
 		}
 
-		Surface surface(width, height);
-
 		{
-			Measure t("test_4_sw_render_polyspans.tga", surface);
+			Measure t("test_4_sw_render_polyspans");
 			for(int i = 0; i < (int)contours_sw.size(); ++i)
 				SwRender::polyspan(surface, polyspans[i], contours_sw[i].color, contours_sw[i].evenodd, contours_sw[i].invert);
 		}
 	}
 
-	Surface surface(width, height);
-
 	{
 		// cl
 
+		Surface surface(width, height);
 		vector<ContourInfo> contours_cl = contours;
 		ClRender clr(e.cl);
 		{
-			Measure t("test_4_cl.tga", surface);
+			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);