diff --git a/c++/contourgl/Makefile b/c++/contourgl/Makefile
index b550e82..9a6d776 100644
--- a/c++/contourgl/Makefile
+++ b/c++/contourgl/Makefile
@@ -1,9 +1,9 @@
-DEPLIBS = gl x11
+DEPLIBS = gl x11 OpenCL
-CXXFLAGS = -O3 -Wall -fmessage-length=0 `pkg-config --libs $(DEPLIBS)` -DGL_GLEXT_PROTOTYPES
-DEPS = contour.h contourbuilder.h geometry.h polyspan.h rendersw.h shaders.h test.h \
- contour.cpp contourbuilder.cpp polyspan.cpp rendersw.cpp shaders.cpp test.cpp
-OBJS = contour.o contourgl.o contourbuilder.o polyspan.o rendersw.o shaders.o test.o
+CXXFLAGS = -O0 -g -Wall -fmessage-length=0 `pkg-config --cflags $(DEPLIBS)` -DGL_GLEXT_PROTOTYPES
+DEPS = clcontext.h contour.h contourgl.h contourbuilder.h geometry.h polyspan.h rendersw.h shaders.h test.h triangulator.h \
+ clcontext.cpp contour.cpp contourgl.cpp contourbuilder.cpp polyspan.cpp rendersw.cpp shaders.cpp test.cpp triangulator.cpp
+OBJS = clcontext.o contour.o contourgl.o contourbuilder.o polyspan.o rendersw.o shaders.o test.o triangulator.o
LIBS = `pkg-config --libs $(DEPLIBS)`
TARGET = contourgl
diff --git a/c++/contourgl/cl/hello.cl b/c++/contourgl/cl/hello.cl
new file mode 100644
index 0000000..7d453ad
--- /dev/null
+++ b/c++/contourgl/cl/hello.cl
@@ -0,0 +1,6 @@
+//#pragma OPENCL EXTENSION cl_khr_byte_addressable_store : enable
+__constant char s[] = "Hello!";
+__kernel void hello(__global char *out) {
+ size_t tid = get_global_id(0);
+ out[tid] = s[tid];
+}
\ No newline at end of file
diff --git a/c++/contourgl/clcontext.cpp b/c++/contourgl/clcontext.cpp
new file mode 100644
index 0000000..b7ee20d
--- /dev/null
+++ b/c++/contourgl/clcontext.cpp
@@ -0,0 +1,123 @@
+/*
+ ......... 2015 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 .
+*/
+
+#include
+
+#include
+#include
+#include
+
+#include "clcontext.h"
+
+
+using namespace std;
+
+
+ClContext::ClContext(): err(), context() {
+
+ // platform
+
+ cl_uint platform_count = 0;
+ clGetPlatformIDs(0, NULL, &platform_count);
+ assert(platform_count);
+ cout << platform_count << " platforms" << endl;
+ vector platforms(platform_count);
+ clGetPlatformIDs(platforms.size(), &platforms.front(), NULL);
+ cl_platform_id platform = platforms[0];
+
+ char vendor[256] = { };
+ err = clGetPlatformInfo(platform, CL_PLATFORM_VENDOR, sizeof(vendor), vendor, NULL);
+ assert(!err);
+ cout << "Use CL platform 0 by " << vendor << endl;
+
+ // devices
+
+ cl_uint device_count = 0;
+ err = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 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);
+ assert(!err);
+
+ // context
+
+ context = clCreateContext(0, 1, &devices.front(), NULL, NULL, &err);
+ assert(context);
+}
+
+ClContext::~ClContext() {
+ clReleaseContext(context);
+}
+
+void ClContext::hello() {
+
+ // data
+
+ char data[] = "......";
+
+ // buffer
+
+ cl_mem buffer = clCreateBuffer(context, CL_MEM_WRITE_ONLY | CL_MEM_USE_HOST_PTR, sizeof(data), data, NULL);
+ assert(buffer);
+
+ // program
+
+ ifstream f("cl/hello.cl");
+ string text((istreambuf_iterator(f)), istreambuf_iterator());
+ 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);
+
+ // kernel
+
+ cl_kernel kernel = clCreateKernel(program, "hello", NULL);
+ assert(kernel);
+ 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(
+ queue,
+ kernel,
+ 1,
+ NULL,
+ &work_group_size,
+ NULL,
+ 0,
+ NULL,
+ &event );
+ assert(!err);
+
+ clWaitForEvents(1, &event);
+
+ // read
+
+ clEnqueueReadBuffer(queue, buffer, CL_TRUE, 0, sizeof(data), data, 0, NULL, &event);
+ clWaitForEvents(1, &event);
+ cout << data << endl;
+}
diff --git a/c++/contourgl/clcontext.h b/c++/contourgl/clcontext.h
new file mode 100644
index 0000000..57fe157
--- /dev/null
+++ b/c++/contourgl/clcontext.h
@@ -0,0 +1,33 @@
+/*
+ ......... 2015 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 .
+*/
+
+#include
+
+#include
+
+
+class ClContext {
+public:
+ cl_int err;
+ cl_context context;
+ std::vector devices;
+
+ ClContext();
+ ~ClContext();
+
+ void hello();
+};
diff --git a/c++/contourgl/contourgl.cpp b/c++/contourgl/contourgl.cpp
index a6f6514..3e8655f 100644
--- a/c++/contourgl/contourgl.cpp
+++ b/c++/contourgl/contourgl.cpp
@@ -25,6 +25,7 @@
#include
#include
+#include "clcontext.h"
#include "test.h"
#include "shaders.h"
@@ -93,12 +94,16 @@ int main() {
int framebuffer_width = 512;
int framebuffer_height = 512;
int framebuffer_samples = 16;
- bool antialising = true;
+ bool antialising = false;
+ bool hdr = false;
+
+ GLenum internal_format = hdr ? GL_RGBA16F : GL_RGBA;
+ GLenum color_type = hdr ? GL_FLOAT : GL_UNSIGNED_BYTE;
GLuint multisample_texture_id = 0;
glGenTextures(1, &multisample_texture_id);
glBindTexture(GL_TEXTURE_2D_MULTISAMPLE, multisample_texture_id);
- glTexImage2DMultisample(GL_TEXTURE_2D_MULTISAMPLE, framebuffer_samples, GL_RGBA16F, framebuffer_width, framebuffer_height, GL_TRUE);
+ glTexImage2DMultisample(GL_TEXTURE_2D_MULTISAMPLE, framebuffer_samples, internal_format, framebuffer_width, framebuffer_height, GL_TRUE);
GLuint multisample_renderbuffer_id = 0;
glGenRenderbuffers(1, &multisample_renderbuffer_id);
@@ -114,7 +119,7 @@ int main() {
GLuint texture_id = 0;
glGenTextures(1, &texture_id);
glBindTexture(GL_TEXTURE_2D, texture_id);
- glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA32F, framebuffer_width, framebuffer_height, 0, GL_RGBA, GL_FLOAT, NULL);
+ glTexImage2D(GL_TEXTURE_2D, 0, internal_format, framebuffer_width, framebuffer_height, 0, GL_RGBA, color_type, NULL);
GLuint renderbuffer_id = 0;
glGenRenderbuffers(1, &renderbuffer_id);
@@ -146,7 +151,9 @@ int main() {
//Test::test1();
//Test::test2();
//Test::test3();
- Test::test4();
+ //Test::test4();
+
+ ClContext().hello();
Shaders::deinitialize();