diff --git a/c++/contourgl/.gitignore b/c++/contourgl/.gitignore
index d87bd8d..38804c5 100644
--- a/c++/contourgl/.gitignore
+++ b/c++/contourgl/.gitignore
@@ -1,2 +1,2 @@
 /contourgl
-
+/*.d
diff --git a/c++/contourgl/Makefile b/c++/contourgl/Makefile
index 98eca6f..09ef1f2 100644
--- a/c++/contourgl/Makefile
+++ b/c++/contourgl/Makefile
@@ -1,22 +1,29 @@
+# config
+
+# just comment following line to disable CUDA
+CUDA = cuda-9.2
+CUDA_PATH := /opt/$(CUDA)
+CUDA_BIN := $(CUDA_PATH)/bin
+CUDA_PKGCONFIG := $(CUDA_PATH)/pkgconfig
+
 DEPLIBS = gl x11 OpenCL
 
-CXXFLAGS = -O3 -Wall -fmessage-length=0 `pkg-config --cflags $(DEPLIBS)` -DGL_GLEXT_PROTOTYPES
-
-HEADERS = \
-	clcontext.h \
-	clrender.h \
-	contour.h \
-	contourbuilder.h \
-	environment.h \
-	geometry.h \
-	glcontext.h \
-	measure.h \
-	polyspan.h \
-	shaders.h \
-	swrender.h \
-	test.h \
-	triangulator.h \
-	utils.h
+
+# compute build options
+
+CXXFLAGS := $(CXXFLAGS) -O3 -Wall -fmessage-length=0 -DGL_GLEXT_PROTOTYPES
+CXXFLAGS := $(CXXFLAGS) $(shell pkg-config --cflags $(DEPLIBS))
+LIBS := $(LIBS) $(shell pkg-config --libs $(DEPLIBS))
+
+ifdef CUDA
+	CXXFLAGS := $(CXXFLAGS) -DCUDA $(shell PKG_CONFIG_PATH=$(CUDA_PKGCONFIG) pkg-config --cflags $(CUDA))
+	LIBS := $(LIBS) $(shell PKG_CONFIG_PATH=$(CUDA_PKGCONFIG) pkg-config --libs $(CUDA))
+endif
+
+
+# files lists
+
+TARGET = contourgl
 
 SOURCES = \
 	contourgl.cpp \
@@ -35,31 +42,44 @@ SOURCES = \
 	triangulator.cpp \
 	utils.cpp
 
-OBJS = \
-	contourgl.o \
-	clcontext.o \
-	clrender.o \
-	contour.o \
-	contourbuilder.o \
-	environment.o \
-	geometry.o \
-	glcontext.o \
-	measure.o \
-	polyspan.o \
-	shaders.o \
-	swrender.o \
-	test.o \
-	triangulator.o \
-	utils.o
-
-DEPS = $(HEADERS) $(SOURCES) 
-LIBS = `pkg-config --libs $(DEPLIBS)`
-TARGET = contourgl
+ifdef CUDA
+	SOURCES += cudacontext.cpp
+	CUDA_SOURCES = \
+		cuda/contour.cu \
+		cuda/hello.cu
+endif
 
-$(TARGET):	$(OBJS)
+
+# files lists postprocessing 
+
+OBJS = $(SOURCES:.cpp=.o)
+DEPS = $(SOURCES:.cpp=.d)
+PTXS = $(CUDA_SOURCES:.cu=.ptx)
+
+# internal targets
+
+%.ptx: %.cu
+	$(CUDA_PATH)/bin/nvcc -ptx $< -o $@
+
+# rule for make *.d files with include (.h) dependencies
+%.d: %.cpp
+	@set -e; rm -f $@; \
+	 $(CXX) -MM $(CXXFLAGS) $< > $@.$$$$; \
+	 sed 's,\($*\)\.o[ :]*,\1.o $@ : ,g' < $@.$$$$ > $@; \
+	 rm -f $@.$$$$
+
+$(TARGET): $(PTXS) $(DEPS) $(OBJS)
 	$(CXX) -o $(TARGET) $(OBJS) $(LIBS)
 
-all:	$(TARGET)
+# include rules with include (.h) dependencies
+-include $(DEPS)
+
+# actual targets
+
+# declare that 'all', 'cuda' and 'clean' are cannot be a filenames 
+.PHONY: all clean
+
+all: $(TARGET)
 
 clean:
-	rm -f $(OBJS) $(TARGET)
+	rm -f $(PTXS) $(OBJS) $(DEPS) $(TARGET)
diff --git a/c++/contourgl/cuda/.gitignore b/c++/contourgl/cuda/.gitignore
new file mode 100644
index 0000000..a224e7f
--- /dev/null
+++ b/c++/contourgl/cuda/.gitignore
@@ -0,0 +1 @@
+/*.ptx
diff --git a/c++/contourgl/cuda/contour.cu b/c++/contourgl/cuda/contour.cu
new file mode 100644
index 0000000..e69de29
--- /dev/null
+++ b/c++/contourgl/cuda/contour.cu
diff --git a/c++/contourgl/cuda/hello.cu b/c++/contourgl/cuda/hello.cu
new file mode 100644
index 0000000..e494eeb
--- /dev/null
+++ b/c++/contourgl/cuda/hello.cu
@@ -0,0 +1,11 @@
+extern "C" {
+
+__constant__ char s[] = "Hello!";
+
+
+__global__ void hello(char *out) {
+	int i = threadIdx.x;
+	out[i] = s[i];
+}
+
+}
\ No newline at end of file
diff --git a/c++/contourgl/cudacontext.cpp b/c++/contourgl/cudacontext.cpp
new file mode 100644
index 0000000..3901098
--- /dev/null
+++ b/c++/contourgl/cudacontext.cpp
@@ -0,0 +1,134 @@
+/*
+    ......... 2018 Ivan Mahonin
+
+    This program is free software: you can redistribute it and/or modify
+    it under the terms of the GNU General Public License as published by
+    the Free Software Foundation, either version 3 of the License, or
+    (at your option) any later version.
+
+    This program is distributed in the hope that it will be useful,
+    but WITHOUT ANY WARRANTY; without even the implied warranty of
+    MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+    GNU General Public License for more details.
+
+    You should have received a copy of the GNU General Public License
+    along with this program.  If not, see <http://www.gnu.org/licenses/>.
+*/
+
+#include <cassert>
+#include <cstring>
+
+#include <iostream>
+
+#include "cudacontext.h"
+
+
+using namespace std;
+
+
+CudaParams::CudaParams(): params_buffer_size() {
+	params_extra.push_back(CU_LAUNCH_PARAM_BUFFER_POINTER);
+	params_extra.push_back(NULL);
+	params_extra.push_back(CU_LAUNCH_PARAM_BUFFER_SIZE);
+	params_extra.push_back(&params_buffer_size);
+	params_extra.push_back(CU_LAUNCH_PARAM_END);
+}
+
+void CudaParams::reset() {
+	params_buffer.clear();
+	params_offsets.clear();
+	params_pointers.clear();
+	params_extra.clear();
+	params_buffer_size = 0;
+}
+
+CudaParams& CudaParams::add(const void* data, int size, int align) {
+	assert(align > 0);
+
+	int index = params_buffer.empty() ? 0 : ((params_buffer.size() - 1)/align + 1)*align;
+	params_buffer.resize(index + size);
+	memcpy(&params_buffer[index], data, size);
+	params_buffer_size = params_buffer.size();
+
+	params_offsets.push_back(index);
+
+	char *root = &params_buffer.front();
+	params_pointers.push_back(root + index);
+	if (params_pointers.front() != root) {
+		params_pointers.clear();
+		for(std::vector<int>::iterator i = params_offsets.begin(); i != params_offsets.end(); ++i)
+			params_pointers.push_back(root + *i);
+	}
+	params_extra[1] = root;
+
+	return *this;
+}
+
+
+CudaContext::CudaContext():
+	device(),
+	context(),
+	err()
+{
+	const int device_index = 0;
+
+	err = cuInit(0);
+    assert(!err);
+
+    err = cuDeviceGet(&device, device_index);
+    assert(!err);
+
+    char device_name[1024] = {};
+    err = cuDeviceGetName(device_name, sizeof(device_name), device);
+    assert(!err);
+    //cout << "CUDA device " << device_index << ": " << device_name << endl;
+
+    err = cuCtxCreate(&context, CU_CTX_SCHED_AUTO, device);
+    assert(!err);
+
+    hello();
+}
+
+CudaContext::~CudaContext() {
+	cuCtxDestroy(context);
+}
+
+void CudaContext::hello() {
+	CUmodule module;
+	err = cuModuleLoad(&module, "cuda/hello.ptx");
+	assert(!err);
+
+	CUfunction kernel;
+	err = cuModuleGetFunction(&kernel, module, "hello");
+	assert(!err);
+
+	char data[] = "......";
+
+	CUdeviceptr buffer;
+	err = cuMemAlloc(&buffer, sizeof(data));
+
+	CudaParams params;
+	params.add(buffer);
+
+	err = cuLaunchKernel(
+		kernel,
+		1, 1, 1,
+		sizeof(data), 1, 1,
+		0, 0, 0,
+		params.get_extra() );
+	assert(!err);
+
+	err = cuStreamSynchronize(0);
+	assert(!err);
+
+	err = cuMemcpyDtoH(data, buffer, sizeof(data));
+	assert(!err);
+
+	err = cuMemFree(buffer);
+	assert(!err);
+
+	err = cuModuleUnload(module);
+	assert(!err);
+
+	cout << data << endl;
+}
diff --git a/c++/contourgl/cudacontext.h b/c++/contourgl/cudacontext.h
new file mode 100644
index 0000000..f30f7a2
--- /dev/null
+++ b/c++/contourgl/cudacontext.h
@@ -0,0 +1,68 @@
+/*
+    ......... 2018 Ivan Mahonin
+
+    This program is free software: you can redistribute it and/or modify
+    it under the terms of the GNU General Public License as published by
+    the Free Software Foundation, either version 3 of the License, or
+    (at your option) any later version.
+
+    This program is distributed in the hope that it will be useful,
+    but WITHOUT ANY WARRANTY; without even the implied warranty of
+    MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+    GNU General Public License for more details.
+
+    You should have received a copy of the GNU General Public License
+    along with this program.  If not, see <http://www.gnu.org/licenses/>.
+*/
+
+#ifndef _CUDACONTEXT_H_
+#define _CUDACONTEXT_H_
+
+#include <vector>
+
+#include <cuda.h>
+
+
+class CudaParams {
+private:
+	std::vector<char> params_buffer;
+	std::vector<int> params_offsets;
+	std::vector<void*> params_pointers;
+	std::vector<void*> params_extra;
+	size_t params_buffer_size;
+
+public:
+	CudaParams();
+
+	void reset();
+
+	CudaParams& add(const void* data, int size, int align);
+
+	template<typename T>
+	CudaParams& add(const T &data, int align)
+		{ return add(&data, sizeof(data), align); }
+
+	template<typename T>
+	CudaParams& add(const T &data)
+		{ return add(data, __alignof(T)); }
+
+	void** get_params() const
+		{ return params_pointers.empty() ? NULL : &(const_cast<CudaParams*>(this)->params_pointers.front()); }
+	void** get_extra() const
+		{ return params_extra.empty() ? NULL : &(const_cast<CudaParams*>(this)->params_extra.front()); }
+};
+
+
+class CudaContext {
+public:
+	CUdevice device;
+	CUcontext context;
+	CUresult err;
+
+	CudaContext();
+	~CudaContext();
+
+	void hello();
+};
+
+#endif
diff --git a/c++/contourgl/environment.h b/c++/contourgl/environment.h
index dd8b0e0..a9b3783 100644
--- a/c++/contourgl/environment.h
+++ b/c++/contourgl/environment.h
@@ -22,12 +22,20 @@
 #include "clcontext.h"
 #include "shaders.h"
 
+#ifdef CUDA
+#include "cudacontext.h"
+#endif
+
 class Environment {
 public:
 	GlContext gl;
 	ClContext cl;
 	Shaders shaders;
 
+	#ifdef CUDA
+	CudaContext cu;
+	#endif
+
 	Environment(int width, int height, bool hdr, bool multisample, int samples):
 		gl(width, height, hdr, multisample, samples) { }
 	void use() { gl.use(); }