From 2517ebb19aa14c75f1436330f4dd65d64196f717 Mon Sep 17 00:00:00 2001 From: Ivan Mahonin Date: Jul 13 2018 04:23:37 +0000 Subject: contourgl: improve cl code --- diff --git a/c++/contourgl/cl/contour.cl b/c++/contourgl/cl/contour.cl index 5f6dffc..84d7061 100644 --- a/c++/contourgl/cl/contour.cl +++ b/c++/contourgl/cl/contour.cl @@ -31,60 +31,57 @@ kernel void draw( - global char *paths_buffer, + const int width, + const int height, global int *mark_buffer, - read_only image2d_t read_image, - write_only image2d_t write_image ) // assumed that read and write image is the same object + global float4 *image, + global const char *paths_buffer ) { const float e = 1e-6f; - int id = (int)get_global_id(0); - int count = (int)get_global_size(0); + int id = (int)get_local_id(0); + int count = (int)get_local_size(0); int paths_count = *(global int *)paths_buffer; - global char *paths = paths_buffer + sizeof(int); + global const char *paths = paths_buffer + sizeof(int); - int width = get_image_width(write_image); - int height = get_image_height(write_image); int pixels_count = width*height; float2 size = (float2)((float)width, (float)height); int w1 = width - 1; int h1 = height - 1; - global int *bound_minx = (global int *)(mark_buffer + 2*pixels_count); - global int *bound_miny = bound_minx + 1; - global int *bound_maxx = bound_minx + 2; - global int *bound_maxy = bound_minx + 3; - - // clear marks - for(int i = id; i < 2*pixels_count; i += count) - mark_buffer[i] = 0; - barrier(CLK_LOCAL_MEM_FENCE); + local int bound_minx; + local int bound_miny; + local int bound_maxx; + local int bound_maxy; // draw paths for(int p = 0; p < paths_count; ++p) { - int points_count = *(global int *)paths; paths += sizeof(int); - int flags = *(global int *)paths; paths += sizeof(int); + int points_count = *(global const int *)paths; paths += sizeof(int); + int flags = *(global const int *)paths; paths += sizeof(int); float4 color; - color.x = *(global float *)paths; paths += sizeof(float); - color.y = *(global float *)paths; paths += sizeof(float); - color.z = *(global float *)paths; paths += sizeof(float); - color.w = *(global float *)paths; paths += sizeof(float); + color.x = *(global const float *)paths; paths += sizeof(float); + color.y = *(global const float *)paths; paths += sizeof(float); + color.z = *(global const float *)paths; paths += sizeof(float); + color.w = *(global const float *)paths; paths += sizeof(float); - global float *points = (global float *)paths; + global const float *points = (global const float *)paths; paths += 2*points_count*sizeof(float); int segments_count = points_count - 1; if (segments_count <= 0) continue; - int invert = flags & 1; - int evenodd = flags & 2; + bool invert = flags & 1; + bool evenodd = flags & 2; - *bound_minx = invert ? 0 : (int)floor(clamp(points[0] + e, 0.f, size.x - 1.f + e)); - *bound_miny = invert ? 0 : (int)floor(clamp(points[1] + e, 0.f, size.y - 1.f + e)); - *bound_maxx = invert ? w1 : *bound_minx; - *bound_maxy = invert ? h1 : *bound_miny; + if (id == 0) { + bound_minx = invert ? 0 : (int)floor(points[0] + e); + bound_miny = invert ? 0 : (int)floor(points[1] + e); + bound_maxx = invert ? w1 : bound_minx; + bound_maxy = invert ? h1 : bound_miny; + } + barrier(CLK_GLOBAL_MEM_FENCE | CLK_LOCAL_MEM_FENCE); // trace path for(int i = id; i < segments_count; i += count) { @@ -92,12 +89,12 @@ kernel void draw( float2 p0 = { points[ii + 0], points[ii + 1] }; float2 p1 = { points[ii + 2], points[ii + 3] }; - int p1x = (int)floor(clamp(p1.x + e, 0.f, size.x - 1.f + e)); - int p1y = (int)floor(clamp(p1.y + e, 0.f, size.y - 1.f + e)); - atomic_min(bound_minx, p1x - 1); - atomic_min(bound_miny, p1y - 1); - atomic_max(bound_maxx, p1x + 1); - atomic_max(bound_maxy, p1y + 1); + int p1x = (int)floor(p1.x + e); + int p1y = (int)floor(p1.y + e); + atomic_min(&bound_minx, p1x); + atomic_min(&bound_miny, p1y); + atomic_max(&bound_maxx, p1x); + atomic_max(&bound_maxy, p1y); bool flipx = p1.x < p0.x; bool flipy = p1.y < p0.y; @@ -112,19 +109,16 @@ kernel void draw( int iy = (int)floor(p0.y + e); if (iy > h1) break; - 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); + float px = (float)(ix + 1); + float py = (float)(iy + 1); float2 pp1 = p1; - if (pp1.x > px.x) pp1 = px; - if (pp1.y > py.y) pp1 = py; + if (pp1.x > px) { pp1.x = px; pp1.y = p0.y + ky*(px - p0.x); } + if (pp1.y > py) { pp1.y = py; pp1.x = p0.x + kx*(py - p0.y); } if (iy >= 0) { // calc values float cover = pp1.y - p0.y; - float area = px.x - 0.5f*(p0.x + pp1.x); + float area = px - 0.5f*(p0.x + pp1.x); if (flipx) { ix = w1 - ix; area = 1.f - area; } if (flipy) { iy = h1 - iy; cover = -cover; } ix = clamp(ix, 0, w1); @@ -138,38 +132,38 @@ kernel void draw( p0 = pp1; } } - barrier(CLK_LOCAL_MEM_FENCE); - + barrier(CLK_GLOBAL_MEM_FENCE | CLK_LOCAL_MEM_FENCE); + + // read bounds + int minx = max(bound_minx, 0); + int miny = max(bound_miny, 0); + int maxx = min(bound_maxx, w1); + int maxy = min(bound_maxy, h1); + barrier(CLK_GLOBAL_MEM_FENCE | CLK_LOCAL_MEM_FENCE); + // fill - int2 coord; - int minx = max(*bound_minx, 0); - int miny = max(*bound_miny, 0); - int maxx = min(*bound_maxx, w1); - int maxy = min(*bound_maxy, h1); - for(coord.y = miny + id; coord.y <= maxy; coord.y += count) { - global int *mark = mark_buffer + (coord.y*width + minx)*2; - - float cover = 0.f; - for(coord.x = minx; coord.x <= maxx; ++coord.x) { + for(int row = miny + id; row <= maxy; row += count) { + global int *mark = mark_buffer + (row*width + minx)*2; + global float4 *pixel = image + row*width + minx; + global float4 *pixel_end = pixel - minx + maxx + 1; + int icover = 0; + + while(pixel < pixel_end) { // read mark (alpha, cover) - float alpha = fabs(cover + *mark/65536.f); *mark = 0; ++mark; - cover += *mark/65536.f; *mark = 0; ++mark; + int ialpha = abs(icover + *mark); *mark = 0; ++mark; + icover += *mark; *mark = 0; ++mark; + + if (evenodd) ialpha = 65536 - abs(ialpha%131072 - 65536); + if (invert) ialpha = 65536 - ialpha; - //if (evenodd) alpha = 1.f - fabs(fmod(alpha, 2.f) - 1.f); - //if (invert) alpha = 1.f - alpha; - alpha *= color.w; + //if (!ialpha) continue; // write color - float alpha_inv = 1.f - alpha; - float4 cl = read_imagef(read_image, coord); - cl.x = cl.x*alpha_inv + color.x*alpha; - cl.y = cl.y*alpha_inv + color.y*alpha; - cl.z = cl.z*alpha_inv + color.z*alpha; - cl.w = min(cl.w + alpha, 1.f); - write_imagef(write_image, coord, cl); + float alpha = (float)ialpha/65536.f*color.w; + *pixel = *pixel*(1.f - alpha) + color*alpha; + ++pixel; } } - barrier(CLK_LOCAL_MEM_FENCE); } } diff --git a/c++/contourgl/clcontext.cpp b/c++/contourgl/clcontext.cpp index 7f6f84d..1d35e13 100644 --- a/c++/contourgl/clcontext.cpp +++ b/c++/contourgl/clcontext.cpp @@ -28,68 +28,89 @@ using namespace std; ClContext::ClContext(): err(), + device(), context(), queue(), max_compute_units(), max_group_size() { + const int platform_index = 0; + const int device_index = 0; // platform - cl_uint platform_count = 0; - clGetPlatformIDs(0, NULL, &platform_count); - assert(platform_count); + err |= clGetPlatformIDs(0, NULL, &platform_count); + assert(!err); //cout << platform_count << " platforms" << endl; + vector platforms(platform_count); - clGetPlatformIDs(platforms.size(), &platforms.front(), NULL); - cl_platform_id platform = platforms[0]; + err |= clGetPlatformIDs(platforms.size(), &platforms.front(), NULL); + assert(!err); + + assert(platform_index < (int)platform_count); + cl_platform_id platform = platforms[platform_index]; char vendor[256] = { }; - err = clGetPlatformInfo(platform, CL_PLATFORM_VENDOR, sizeof(vendor), vendor, NULL); + err |= clGetPlatformInfo(platform, CL_PLATFORM_VENDOR, sizeof(vendor), vendor, NULL); assert(!err); - //cout << "Use CL platform 0 by " << vendor << endl; + //cout << "Use CL platform " << platform_index << " by " << vendor << endl; char platform_version[256]; - err = clGetPlatformInfo(platform, CL_PLATFORM_VERSION, sizeof(platform_version), platform_version, NULL); + err |= clGetPlatformInfo(platform, CL_PLATFORM_VERSION, sizeof(platform_version), platform_version, NULL); assert(!err); - //cout << "Platform 0 OpenCL version " << platform_version << endl; + //cout << "Platform " << platform_index << " OpenCL version " << platform_version << endl; // devices cl_uint device_count = 0; - err = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 0, NULL, &device_count); + err |= clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, 0, NULL, &device_count); assert(!err); //cout << device_count << " devices" << endl; - devices.resize(device_count); - err = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, devices.size(), &devices.front(), NULL); + vector devices(device_count); + err |= clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, devices.size(), &devices.front(), NULL); assert(!err); + assert(device_index < (int)device_count); + device = devices[device_index]; + char device_name[256]; - clGetDeviceInfo(devices.front(), CL_DEVICE_NAME, sizeof(device_name), device_name, NULL); - //cout << "Device 0 name " << device_name << endl; + err |= clGetDeviceInfo(device, CL_DEVICE_NAME, sizeof(device_name), device_name, NULL); + assert(!err); + //cout << "Device " << device_index << " name " << device_name << endl; char device_version[256]; - clGetDeviceInfo(devices.front(), CL_DEVICE_VERSION, sizeof(device_version), device_version, NULL); - //cout << "Device 0 OpenCL version " << device_version << endl; + err |= clGetDeviceInfo(device, CL_DEVICE_VERSION, sizeof(device_version), device_version, NULL); + assert(!err); + //cout << "Device " << device_index << " OpenCL version " << device_version << endl; + + err |= clGetDeviceInfo(device, CL_DEVICE_MAX_COMPUTE_UNITS, sizeof(max_compute_units), &max_compute_units, NULL); + assert(!err); + //cout << "Device " << device_index << " max compute units " << max_compute_units << endl; - clGetDeviceInfo(devices.front(), CL_DEVICE_MAX_COMPUTE_UNITS, sizeof(max_compute_units), &max_compute_units, NULL); - //cout << "Device 0 max compute units " << max_compute_units << endl; + unsigned int max_dimensions; + err |= clGetDeviceInfo(device, CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS, sizeof(max_dimensions), &max_dimensions, NULL); + assert(!err); + assert(max_dimensions); + //cout << "Device " << device_index << " max work dimensions " << max_dimensions << endl; - clGetDeviceInfo(devices.front(), CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(max_group_size), &max_group_size, NULL); - //cout << "Device 0 max group size " << max_group_size << endl; + vector max_group_sizes(max_dimensions); + err |= clGetDeviceInfo(device, CL_DEVICE_MAX_WORK_GROUP_SIZE, max_group_sizes.size()*sizeof(size_t), &max_group_sizes.front(), NULL); + assert(!err); + max_group_size = max_group_sizes.front(); + //cout << "Device " << device_index << " max group size " << max_group_size << endl; // context cl_context_properties context_props[] = { CL_CONTEXT_PLATFORM, (cl_context_properties)platform, CL_NONE }; - context = clCreateContext(context_props, 1, &devices.front(), callback, NULL, &err); + context = clCreateContext(context_props, 1, &device, callback, NULL, &err); assert(context); // command queue - queue = clCreateCommandQueue(context, devices.front(), CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, NULL); + queue = clCreateCommandQueue(context, device, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, NULL); assert(queue); //hello(); @@ -109,12 +130,14 @@ cl_program ClContext::load_program(const std::string &filename) { cl_program program = clCreateProgramWithSource(context, 1, &text_pointer, NULL, NULL); assert(program); - err = clBuildProgram(program, 1, &devices.front(), "", NULL, NULL); + const char options[] = " -cl-fast-relaxed-math -Werror "; + + err = clBuildProgram(program, 1, &device, options, NULL, NULL); if (err) { size_t size; - clGetProgramBuildInfo(program, devices.front(), CL_PROGRAM_BUILD_LOG, 0, NULL, &size); + clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, 0, NULL, &size); char *log = new char[size]; - clGetProgramBuildInfo(program, devices.front(), CL_PROGRAM_BUILD_LOG, size, log, NULL); + clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, size, log, NULL); cout << log << endl; delete[] log; } diff --git a/c++/contourgl/clcontext.h b/c++/contourgl/clcontext.h index f73beae..0a67407 100644 --- a/c++/contourgl/clcontext.h +++ b/c++/contourgl/clcontext.h @@ -27,8 +27,8 @@ class ClContext { public: cl_int err; + cl_device_id device; cl_context context; - std::vector devices; cl_command_queue queue; unsigned int max_compute_units; diff --git a/c++/contourgl/clrender.cpp b/c++/contourgl/clrender.cpp index fe1f159..d29863b 100644 --- a/c++/contourgl/clrender.cpp +++ b/c++/contourgl/clrender.cpp @@ -31,6 +31,7 @@ ClRender::ClRender(ClContext &cl): cl(cl), contour_program(), contour_draw_kernel(), + contour_draw_workgroup_size(), surface(), paths_buffer(), mark_buffer(), @@ -38,8 +39,19 @@ ClRender::ClRender(ClContext &cl): prev_event() { contour_program = cl.load_program("contour.cl"); + assert(contour_program); + contour_draw_kernel = clCreateKernel(contour_program, "draw", NULL); assert(contour_draw_kernel); + + cl.err |= clGetKernelWorkGroupInfo( + contour_draw_kernel, + cl.device, + CL_KERNEL_WORK_GROUP_SIZE, + sizeof(contour_draw_workgroup_size), + &contour_draw_workgroup_size, + NULL ); + assert(!cl.err); } ClRender::~ClRender() { @@ -68,35 +80,29 @@ void ClRender::send_surface(Surface *surface) { mark_buffer = clCreateBuffer( cl.context, CL_MEM_READ_WRITE, - (surface->count() + 2)*sizeof(cl_int2), NULL, // extra two values to store contour bounds - NULL ); + (surface->count() + 2)*sizeof(cl_int2), NULL, + &cl.err ); + assert(!cl.err); assert(mark_buffer); - cl_image_format surface_format = { }; - surface_format.image_channel_order = CL_RGBA; - surface_format.image_channel_data_type = CL_FLOAT; - - cl_image_desc surface_desc = { }; - surface_desc.image_type = CL_MEM_OBJECT_IMAGE2D; - surface_desc.image_width = surface->width; - surface_desc.image_height = surface->height; - - surface_image = clCreateImage( - cl.context, CL_MEM_READ_WRITE, - &surface_format, &surface_desc, - NULL, NULL ); - assert(surface_image); - - 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_FALSE, - origin, region, 0, 0, surface->data, + char zero = 0; + cl.err |= clEnqueueFillBuffer( + cl.queue, mark_buffer, + &zero, 1, + 0, surface->count()*sizeof(cl_int2), 0, NULL, NULL ); assert(!cl.err); - cl.err |= clSetKernelArg(contour_draw_kernel, 1, sizeof(mark_buffer), &mark_buffer); - cl.err |= clSetKernelArg(contour_draw_kernel, 2, sizeof(surface_image), &surface_image); + surface_image = clCreateBuffer( + cl.context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, + surface->count()*sizeof(Color), surface->data, + &cl.err ); + assert(!cl.err); + assert(surface_image); + + cl.err |= clSetKernelArg(contour_draw_kernel, 0, sizeof(surface->width), &surface->width); + cl.err |= clSetKernelArg(contour_draw_kernel, 1, sizeof(surface->width), &surface->height); + cl.err |= clSetKernelArg(contour_draw_kernel, 2, sizeof(mark_buffer), &mark_buffer); cl.err |= clSetKernelArg(contour_draw_kernel, 3, sizeof(surface_image), &surface_image); assert(!cl.err); @@ -109,11 +115,9 @@ Surface* ClRender::receive_surface() { if (surface) { //Measure t("ClRender::receive_surface"); - size_t origin[3] = { }; - size_t region[3] = { (size_t)surface->width, (size_t)surface->height, 1 }; - cl.err |= clEnqueueReadImage( + cl.err |= clEnqueueReadBuffer( cl.queue, surface_image, CL_FALSE, - origin, region, 0, 0, surface->data, + 0, surface->count()*sizeof(Color), surface->data, prev_event ? 1 : 0, prev_event ? &prev_event : NULL, NULL ); @@ -126,37 +130,33 @@ Surface* ClRender::receive_surface() { return surface; } -void ClRender::send_paths(const void *paths, int size) { - if (!paths_buffer && (!paths || size <= 0)) return; - - cl.err |= clFinish(cl.queue); - assert(!cl.err); - prev_event = NULL; - +void ClRender::remove_paths() { if (paths_buffer) { + cl.err |= clFinish(cl.queue); + assert(!cl.err); + prev_event = NULL; + clReleaseMemObject(paths_buffer); paths_buffer = NULL; } +} + +void ClRender::send_paths(const void *paths, int size) { + if (!paths_buffer && (!paths || size <= 0)) return; + + remove_paths(); if (paths && size > 0) { //Measure t("ClRender::send_path"); paths_buffer = clCreateBuffer( - cl.context, CL_MEM_READ_ONLY, - size, NULL, - NULL ); - assert(paths_buffer); - - cl.err |= clEnqueueWriteBuffer( - cl.queue, paths_buffer, CL_FALSE, - 0, size, paths, - 0, NULL, NULL ); - assert(!cl.err); - - cl.err |= clSetKernelArg(contour_draw_kernel, 0, sizeof(paths_buffer), &paths_buffer); + cl.context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, + size, const_cast(paths), + &cl.err ); assert(!cl.err); + assert(paths_buffer); - cl.err |= clFinish(cl.queue); + cl.err |= clSetKernelArg(contour_draw_kernel, 4, sizeof(paths_buffer), &paths_buffer); assert(!cl.err); } } @@ -165,15 +165,14 @@ void ClRender::draw() { //Measure t("ClRender::contour"); cl_event event = prev_event; - - size_t start = 0; + size_t count = contour_draw_workgroup_size; cl.err |= clEnqueueNDRangeKernel( cl.queue, contour_draw_kernel, 1, - &start, - &cl.max_group_size, - &cl.max_group_size, + NULL, + &count, + &count, event ? 1 : 0, event ? &event : NULL, &prev_event ); @@ -188,72 +187,3 @@ void ClRender::wait() { } } - -void SwRenderAlt::line(const Vector &p0, const Vector &p1) { - int iy0 = min(max((int)floor(p0.y), 0), height); - int iy1 = min(max((int)floor(p1.y), 0), height); - if (iy1 < iy0) swap(iy0, iy1); - - Vector d = p1 - p0; - Vector k( fabs(d.y) < 1e-6 ? 0.0 : d.x/d.y, - fabs(d.x) < 1e-6 ? 0.0 : d.y/d.x ); - - for(int r = iy0; r <= iy1; ++r) { - Real y = (Real)iy0; - - Vector pp0 = p0; - pp0.y -= y; - if (pp0.y < 0.0) { - pp0.y = 0.0; - pp0.x = p0.x - k.x*y; - } else - if (pp0.y > 1.0) { - pp0.y = 1.0; - pp0.x = p0.x - k.x*(y - 1.0); - } - - Vector pp1 = p1; - pp1.y -= y; - if (pp1.y < 0.0) { - pp1.y = 0.0; - pp1.x = p0.x - k.x*y; - } else - if (pp1.y > 1.0) { - pp1.y = 1.0; - pp1.x = p0.x - k.x*(y - 1.0); - } - - int ix0 = min(max((int)floor(pp0.x), 0), width); - int ix1 = min(max((int)floor(pp1.x), 0), width); - if (ix1 < ix0) swap(ix0, ix1); - for(int c = ix0; c <= ix1; ++c) { - Real x = (Real)ix0; - - Vector ppp0 = pp0; - ppp0.x -= x; - if (ppp0.x < 0.0) { - ppp0.x = 0.0; - ppp0.y = pp0.y - k.y*x; - } else - if (ppp0.x > 1.0) { - ppp0.x = 1.0; - ppp0.y = pp0.y - k.y*(x - 1.0); - } - - Vector ppp1 = pp1; - ppp1.x -= x; - if (ppp1.x < 0.0) { - ppp1.x = 0.0; - ppp1.y = pp0.y - k.y*x; - } else - if (ppp1.x > 1.0) { - ppp1.x = 1.0; - ppp1.y = pp0.y - k.y*(x - 1.0); - } - - Real cover = ppp0.y - ppp1.y; - Real area = (0.5*(ppp1.x + ppp1.x) - 1.0)*cover; - (*this)[r][c].add(area, cover); - } - } -} diff --git a/c++/contourgl/clrender.h b/c++/contourgl/clrender.h index db74532..2724a39 100644 --- a/c++/contourgl/clrender.h +++ b/c++/contourgl/clrender.h @@ -31,6 +31,7 @@ private: ClContext &cl; cl_program contour_program; cl_kernel contour_draw_kernel; + size_t contour_draw_workgroup_size; Surface *surface; cl_mem paths_buffer; @@ -45,33 +46,10 @@ public: void send_surface(Surface *surface); Surface* receive_surface(); void send_paths(const void *paths, int size); + void remove_paths(); void draw(); void wait(); }; -class SwRenderAlt { -public: - struct Pixel { - Real area; - Real cover; - Pixel(): area(), cover() { } - void add(Real area, Real cover) { this->area += area; this->cover += cover; } - }; - -private: - std::vector data; - -public: - const int width; - const int height; - - SwRenderAlt(int width, int height): data(width*height), width(width), height(height) { } - - Pixel* operator[] (int row) { return &data.front() + row*width; } - const Pixel* operator[] (int row) const { return &data.front() + row*width; } - - void line(const Vector &p0, const Vector &p1); -}; - #endif diff --git a/c++/contourgl/test.cpp b/c++/contourgl/test.cpp index 90766ba..f351efa 100644 --- a/c++/contourgl/test.cpp +++ b/c++/contourgl/test.cpp @@ -282,12 +282,19 @@ void Test::test_cl(Environment &e, Data &data, Surface &surface) { ClRender clr(e.cl); clr.send_surface(&surface); + // warm-up + //clr.send_paths(&paths.front(), paths.size()); + //for(int i = 0; i < 1000; ++i) + // clr.draw(); + clr.remove_paths(); + + // actual task + clr.send_surface(&surface); { Measure t("render"); clr.send_paths(&paths.front(), paths.size()); clr.draw(); clr.wait(); } - clr.receive_surface(); }