Avancee sur le C++. Boring parts disparait presque, au profit d'une classe OpenCLMeshKit.

Il ne manque que execKernel() a code dans cette classe.
Dans gpudataviz.cpp, il faut changer le code OpenGL pour afficher le maillage et pas un isodecaheron de test.


git-svn-id: file:///var/svn/2013-gpudataviz/trunk@18 371a6b4a-a258-45f8-9dcc-bdd82ce0ac9d

7 files changed, 313 insertions, 157 deletions

diff --git a/src/boring_parts.cpp b/src/boring_parts.cpp
index 0f4ee1a..447ae60 100644
--- a/src/boring_parts.cpp
+++ b/src/boring_parts.cpp
@@ -3,26 +3,7 @@
 // TODO : only need OpenGL things, not GTK ones for now
 //#include "gtk_includes.hpp"
 
-#define RETURN_IF_FAIL(expr) do { \
-	int res=(expr); \
-	if ( res != 0 ) return res; \
-} while(0)
-
-// TODO : print streamsdk::getOpenCLErrorCodeStr(res)
-#define CL_RETURN_VAL_IF_FAIL(val, expr) do { \
-	cl_int res=(expr); \
-	if ( res != CL_SUCCESS ) {	\
-		std::cerr << "file " << __FILE__ << ": line " << __LINE__ << " (" << __PRETTY_FUNCTION__ \
-			<< "): '" << "expr" << "' failed (return code : " << res << ")" << std::endl; \
-		return val; \
-	} \
-} while(0)
-
-
 /* From http://stackoverflow.com/questions/4317062/opengl-how-to-check-if-the-user-supports-glgenbuffers
-#ifndef STRINGIFY
-  #define STRINGIFY(x) #x
-#endif
 #ifdef WIN32
   #include <windows.h>
   #define glGetProcAddress(a) wglGetProcAddress(a)
@@ -54,109 +35,6 @@
 #endif
 */
 
-#ifdef HAS_OPENCL
-int initOpenCL(intptr_t gl_display, intptr_t gl_context, intptr_t gl_vbo) {
-	cl_uint id, numPlatforms;
-	char pbuf[100];
-	std::string dTypeStr;
-	cl_platform_id *platforms, platform;
-	cl_device_id /* *devices, */device;
-	cl_context cl_ctx;
-	cl_command_queue cl_commandQueue;
-	bool usableDeviceFound=false;
-
-	// Get platform count
-	CL_RETURN_VAL_IF_FAIL(1,
-		clGetPlatformIDs(0, NULL, &numPlatforms)
-	);
-	
-	std::cout << "Detected " << numPlatforms << " platform(s)" << std::endl;
-	if ( ! ( numPlatforms > 0 ) ) return 2;
-
-	// Allocate room for all platform IDs
-	platforms = new cl_platform_id[numPlatforms];
-
-	// Get platform IDs
-	CL_RETURN_VAL_IF_FAIL(3,
-		clGetPlatformIDs(numPlatforms, platforms, &numPlatforms)
-	);
-
-	// Enumerate platforms and grab informations
-	for(id=0;id<numPlatforms;id++) {
-		cl_int res;
-		platform=platforms[id];
-
-		CL_RETURN_VAL_IF_FAIL(4,
-			clGetPlatformInfo(platform, CL_PLATFORM_VENDOR, sizeof(pbuf), pbuf, NULL)
-		);
-		std::cout << "Platform " << id << " : " << pbuf << std::endl;
-
-		// Dynamically get the function pointer for clGetGLConetextInfoKHR
-		clGetGLContextInfoKHR_fn clGetGLContextInfoKHR_proc = (clGetGLContextInfoKHR_fn) clGetExtensionFunctionAddressForPlatform(platform, "clGetGLContextInfoKHR");
-		if (!clGetGLContextInfoKHR_proc) {
-			std::cerr << "clGetExtensionFunctionAddressForPlatform(platform, clGetGLContextInfoKHR) failed" << std::endl;
-			continue;
-		}
-
-		// Try to get the device corresponding to the GL context/display on this platform
-		cl_context_properties cpsGL[] = {
-			CL_CONTEXT_PLATFORM, (cl_context_properties)platform,
-			CL_GLX_DISPLAY_KHR, gl_display,
-			CL_GL_CONTEXT_KHR, gl_context,
-			0
-		};
-
-		std::cout << "cl_context_properties cpsGL :" << std::endl;
-		std::cout << "\tCL_CONTEXT_PLATFORM :" << (void *)cpsGL[1] << std::endl;
-		std::cout << "\tCL_GLX_DISPLAY_KHR :"  << (void *)cpsGL[3] << std::endl;
-		std::cout << "\tCL_GL_CONTEXT_KHR :"   << (void *)cpsGL[5] << std::endl;
-
-    size_t deviceSize=0;
-    // get deviceSize (should be 1*sizeof(cl_device_id) with CL_CURRENT_DEVICE_FOR_GL_CONTEXT_KHR)
-    res=clGetGLContextInfoKHR_proc(cpsGL,CL_CURRENT_DEVICE_FOR_GL_CONTEXT_KHR,0,NULL,&deviceSize);
-    if ( res!=CL_SUCCESS || deviceSize!=1*sizeof(cl_device_id)) {
-      std::cerr << "clGetGLContextInfoKHR_proc(cpsGL,CL_CURRENT_DEVICE_FOR_GL_CONTEXT_KHR,0,...) failed" << std::endl;
-			std::cerr << " (return code : " << res << ")" << std::endl;
-      continue;
-    }
-
-		device=0;
-		res=clGetGLContextInfoKHR_proc(cpsGL,CL_CURRENT_DEVICE_FOR_GL_CONTEXT_KHR,deviceSize,&device,NULL);
-		if ( res!=CL_SUCCESS || device==0 ) {
-			std::cerr << "clGetGLContextInfoKHR_proc(cpsGL,CL_CURRENT_DEVICE_FOR_GL_CONTEXT_KHR," << deviceSize << ",...) failed" << std::endl;
-			std::cerr << " (return code : " << res << ")" << std::endl;
-			continue;
-		}
-
-		std::cout << "cl_device :" << (void *)device << std::endl;
-
-		cl_ctx = clCreateContext(cpsGL,1,&device,0,0,&res);
-		if ( res!=CL_SUCCESS ) {
-			std::cerr << "clCreateContext() failed" << std::endl; 
-			std::cerr << " (return code : " << res << ")" << std::endl;
-			continue;
-		}
-
-		cl_commandQueue = clCreateCommandQueue(cl_ctx,device,0,&res);
-		if ( res!=CL_SUCCESS ) {
-			std::cerr << "clCreateCommandQueue() failed" << std::endl; 
-			std::cerr << " (return code : " << res << ")" << std::endl;
-			continue;
-		}
-
-		usableDeviceFound=true;
-		break;
-	}
-
-	if (! usableDeviceFound) {
-		std::cerr << "No OpenCL device has been successfully initialized" << std::endl;
-		return 1;
-	}
-		std::cout << "OpenCL initialization done." << std::endl;
-		return 0;
-}
-#endif /*HAS_OPENCL*/
-
 bool updateGLProjectionMatrix(Glib::RefPtr<Gdk::GL::Context> glCtx, Glib::RefPtr<Gdk::GL::Window> glWin, int width, int height) {
 
 	GLdouble aspect = (GLdouble) width/height;
diff --git a/src/boring_parts.hpp b/src/boring_parts.hpp
index b4cd58d..09404d8 100644
--- a/src/boring_parts.hpp
+++ b/src/boring_parts.hpp
@@ -9,11 +9,8 @@
 #include <GL/glu.h>
 //#include <CL/cl_gl.h>
 
-int initLibs();
-
 #ifdef HAS_OPENCL
-#include <CL/opencl.h>
-int initOpenCL(intptr_t gl_display, intptr_t gl_context, intptr_t gl_vbo);
+// Blabla
 #else
 // Quick and dirty cl_float4 replacement
 typedef union {
diff --git a/src/compil.sh b/src/compil.sh
index f3ca2a0..1f5db4c 100755
--- a/src/compil.sh
+++ b/src/compil.sh
@@ -45,7 +45,7 @@ function link_cxx() {
 
 rm -v $BUILD_PATH/* || true
 
-build_cxx gl_core_1_5_vbo.o gl_core_1_5_vbo.cpp
+build_cxx opencl_mesh_kit.o opencl_mesh_kit.cpp
 build_cxx gpudataviz.o gpudataviz.cpp
 build_cxx boring_parts.o boring_parts.cpp
 build_cxx gtk_win_main.o gtk_win_main.cpp
diff --git a/src/gpudataviz.cpp b/src/gpudataviz.cpp
index a9fdb8e..72ee472 100644
--- a/src/gpudataviz.cpp
+++ b/src/gpudataviz.cpp
@@ -1,6 +1,7 @@
 #include <iostream>
 
 
+#include <GL/glew.h>
 //#include "gtk_includes.h"
 #include "gtk_win_main.hpp"
 //#include "my_gtk_gl_scene_widget.hpp"
@@ -45,85 +46,78 @@ int main(int argc, char* argv[]) {
 	// Could exit() the program if problem with OpenGL or OpenCL
 	GTKWinMain gtkwinmain(glScene);
 
-	// Initialize OpenCL (only after the MyGTKGLSceneWidget realize)
-	//EXIT_IF_FAIL(3, initLibs()==0 ); // See boring_parts.cc
-
+	// Run the app
 	gtkKit.run(gtkwinmain);
 
 	return 0;
 }
 
 /* MyGTKGLSceneWidget implementation
-	I want to keep interesting code part in this file
-		in natural reading order
+	I want to keep interesting code part in this file in natural reading order
 */
 MyGTKGLSceneWidget::MyGTKGLSceneWidget(Glib::RefPtr<Gdk::GL::Config> &glconfig) {
 	set_gl_capability(glconfig);
 	Gdk::EventMask mask = Gdk::POINTER_MOTION_MASK | Gdk::BUTTON_MOTION_MASK | Gdk::BUTTON_PRESS_MASK | Gdk::BUTTON_RELEASE_MASK;
 	set_events(mask); // The containing window should have those attributes too
-	this->camera.rx = this->camera.ry = 0.0f; this->camera.tz = -3.0f;
+	this->camera.rx = 0.0f; this->camera.ry = 0.0f; this->camera.tz = -3.0f;
 }
 
 MyGTKGLSceneWidget::~MyGTKGLSceneWidget() { }
 
 void MyGTKGLSceneWidget::on_size_request(Gtk::Requisition* requisition) {
 	CALL_TRACE; // Technical stuff : GTK call this to ask the widget minimal size 
-  	*requisition = Gtk::Requisition();
+ 	*requisition = Gtk::Requisition();
 	requisition->width = 320; requisition->height = 240;
 }
 
 void MyGTKGLSceneWidget::on_realize() {
 	CALL_TRACE; // This one runs once at window creation time
 	// It's time to setup GL things that don't change on each frame
-
-	GLenum gl_res;
-
 	Gtk::DrawingArea::on_realize();
 	Glib::RefPtr<Gdk::GL::Window> glwindow = get_gl_window();
 
-
 	// *** OpenGL BEGIN ***
+	GLenum gl_res;
 	if (!glwindow->gl_begin(get_gl_context())) {
 		std::cerr << "Oups : glwindow->gl_begin(get_gl_context())" << std::endl;
 		return;
 	}
 
 	EXIT_IF_FAIL(3, Gdk::GL::query_gl_extension("GL_ARB_vertex_buffer_object") );
+	EXIT_IF_FAIL(4, glewInit() == 0 );
 	
-	size_t mesh_width=64;
-	size_t mesh_height=64;
+	size_t mesh_width=512, mesh_height=512, group_size=256; // TODO : not here
+
+	GLuint gl_vbo=0;
 	GLsizeiptr gl_vbo_data_size = mesh_width * mesh_height * sizeof(cl_float4);
-	intptr_t gl_vbo=0;
+	std::cout << "gl_vbo_data_size==" << gl_vbo_data_size << std::endl;
 
 	glGenBuffers(1, &gl_vbo);
 	glBindBuffer(GL_ARRAY_BUFFER, gl_vbo);
-	/*
-		STREAM : The data store contents will be modified once and used at most a few times.
-		STATIC : The data store contents will be modified once and used many times.
-		DYNAMIC : The data store contents will be modified repeatedly and used many times. 
-	*/
 	glBufferData(GL_ARRAY_BUFFER, gl_vbo_data_size, NULL, GL_DYNAMIC_DRAW);
-	if ( gl_res=glGetError() ) {
-		std::cerr << "glBufferData(). Unable to allocate " << gl_vbo_data_size << "bytes in VRAM";
+	gl_res=glGetError();
+	if ( gl_res != GL_NO_ERROR ) {
+		std::cerr << "glBufferData(). Unable to allocate " << gl_vbo_data_size << "bytes in VRAM" << std::endl;
 		std::cerr << gluErrorString(gl_res);
+		EXIT_IF_FAIL(5, false);
 	}
 
-#ifdef HAS_OPENCL
-//	static bool isOpenCLInitialized=false;
-
-//	if (! isOpenCLInitialized) {
+//#ifdef HAS_OPENCL
 //	#ifdef X11
 		intptr_t gl_context = (intptr_t)glXGetCurrentContext();
 		intptr_t gl_display = (intptr_t)glXGetCurrentDisplay();
 
-//		std::cerr << "DEBUG : begin initOpenCL()" << std::endl;
-		initOpenCL(gl_display, gl_context, gl_vbo); /* See boring_parts.cc */
-//		isOpenCLInitialized=true;
+		int cl_res = clKit.initCL(gl_display, gl_context, gl_vbo, mesh_width, mesh_height, group_size);
+		EXIT_IF_FAIL(cl_res, cl_res==0);
+
+//	std::cerr << "DEBUG : begin initOpenCL()" << std::endl;
+//		int cl_res = initOpenCL(gl_display, gl_context, gl_vbo); /* See boring_parts.cpp */
+//		EXIT_IF_FAIL(cl_res, cl_res==0);
 //	#else
 //		#error initOpenCL works only for X11 systems for now
 //	#endif
 //	}
-#endif
+//#endif
 
 
 	// Programmatically create rendering lists : opengl will able to replay that efficiently
@@ -155,7 +149,7 @@ void MyGTKGLSceneWidget::on_realize() {
 
 bool MyGTKGLSceneWidget::on_configure_event(GdkEventConfigure* event) {
 	CALL_TRACE ; // This one runs mainly when GTK GL Widget is resized
-	// See boring_parts.cc. In short : gluPerspective(60.0, aspect, 0.1, 10.0);
+	// See boring_parts.cpp. In short : gluPerspective(60.0, aspect, 0.1, 10.0);
 	return updateGLProjectionMatrix(get_gl_context(), get_gl_window(), get_width(), get_height());
 }
 
diff --git a/src/my_gtk_gl_scene_widget.hpp b/src/my_gtk_gl_scene_widget.hpp
index 212f346..b6fbe73 100644
--- a/src/my_gtk_gl_scene_widget.hpp
+++ b/src/my_gtk_gl_scene_widget.hpp
@@ -2,6 +2,7 @@
 #define MY_GTK_GL_SCENE_H
 
 #include "gtk_includes.hpp"
+#include "opencl_mesh_kit.hpp"
 
 // Class that will contain all the OpenGL logic for displaying the OpenCL computed data
 // Implementation is kept in gpudataviz.cc (I want to keep interesting code part in this file)
@@ -23,6 +24,7 @@ class MyGTKGLSceneWidget : public Gtk::DrawingArea, public Gtk::GL::Widget<MyGTK
 
 		bool do_mouse_logic(GdkEventType type, guint state, guint x, guint y);
 	private:
+		OpenCLMeshKit clKit;
 		struct camera_params { float rx; float ry; float tz; } camera;
 };
 
diff --git a/src/opencl_mesh_kit.cpp b/src/opencl_mesh_kit.cpp
new file mode 100644
index 0000000..a3eb6dc
--- /dev/null
+++ b/src/opencl_mesh_kit.cpp
@@ -0,0 +1,225 @@
+#include "opencl_mesh_kit.hpp"
+
+// TODO : print streamsdk::getOpenCLErrorCodeStr(res)
+#define CL_RETURN_VAL_IF_FAIL(val, expr) do { \
+	cl_int res=(expr); \
+	if ( res != CL_SUCCESS ) {	\
+		std::cerr << "file " << __FILE__ << ": line " << __LINE__ << " (" << __PRETTY_FUNCTION__ \
+			<< "): '" << "expr" << "' failed (return code : " << res << ")" << std::endl; \
+		return val; \
+	} \
+} while(0)
+
+cl_int OpenCLMeshKit::initCL(intptr_t gl_display, intptr_t gl_context, intptr_t gl_vbo, size_t meshWidth, size_t meshHeight, size_t groupSize) {
+	cl_uint id, numPlatforms;
+	cl_int res;
+	char pbuf[100];
+	cl_platform_id *platforms, platform;
+	bool usableDeviceFound=false;
+
+	this->meshWidth = meshWidth;
+	this->meshHeight = meshHeight;
+	this->groupSize = groupSize;
+
+	// Get platform count
+	CL_RETURN_VAL_IF_FAIL(10,
+		clGetPlatformIDs(0, NULL, &numPlatforms)
+	);
+	
+	std::cout << "Detected " << numPlatforms << " platform(s)" << std::endl;
+	if ( ! ( numPlatforms > 0 ) ) return 2;
+
+	// Allocate room for all platform IDs
+	platforms = new cl_platform_id[numPlatforms];
+
+	// Get platform IDs
+	CL_RETURN_VAL_IF_FAIL(11,
+		clGetPlatformIDs(numPlatforms, platforms, &numPlatforms)
+	);
+
+	// Enumerate platforms and grab informations
+	for(id=0;id<numPlatforms;id++) {
+		platform=platforms[id];
+
+		CL_RETURN_VAL_IF_FAIL(12,
+			clGetPlatformInfo(platform, CL_PLATFORM_VENDOR, sizeof(pbuf), pbuf, NULL)
+		);
+		std::cout << "Platform " << id << " : " << pbuf << std::endl;
+
+		// Dynamically get the function pointer for clGetGLConetextInfoKHR
+		clGetGLContextInfoKHR_fn clGetGLContextInfoKHR_proc = (clGetGLContextInfoKHR_fn) clGetExtensionFunctionAddressForPlatform(platform, "clGetGLContextInfoKHR");
+		if (!clGetGLContextInfoKHR_proc) {
+			std::cerr << "clGetExtensionFunctionAddressForPlatform(platform, clGetGLContextInfoKHR) failed" << std::endl;
+			continue;
+		}
+
+		// Try to get the device corresponding to the GL context/display on this platform
+		cl_context_properties cpsGL[] = {
+			CL_CONTEXT_PLATFORM, (cl_context_properties)platform,
+			CL_GLX_DISPLAY_KHR, gl_display,
+			CL_GL_CONTEXT_KHR, gl_context,
+			0
+		};
+
+		std::cout << "cl_context_properties cpsGL :" << std::endl;
+		std::cout << "\tCL_CONTEXT_PLATFORM :" << (void *)cpsGL[1] << std::endl;
+		std::cout << "\tCL_GLX_DISPLAY_KHR :"  << (void *)cpsGL[3] << std::endl;
+		std::cout << "\tCL_GL_CONTEXT_KHR :"   << (void *)cpsGL[5] << std::endl;
+
+    size_t deviceSize=0;
+    // get deviceSize (should be 1*sizeof(cl_device_id) with CL_CURRENT_DEVICE_FOR_GL_CONTEXT_KHR)
+    res=clGetGLContextInfoKHR_proc(cpsGL,CL_CURRENT_DEVICE_FOR_GL_CONTEXT_KHR,0,NULL,&deviceSize);
+    if ( res!=CL_SUCCESS || deviceSize!=1*sizeof(cl_device_id)) {
+      std::cerr << "clGetGLContextInfoKHR_proc(cpsGL,CL_CURRENT_DEVICE_FOR_GL_CONTEXT_KHR,0,...) failed" << std::endl;
+			std::cerr << " (return code : " << res << ")" << std::endl;
+      continue;
+    }
+
+		cl_dev=0;
+		res=clGetGLContextInfoKHR_proc(cpsGL,CL_CURRENT_DEVICE_FOR_GL_CONTEXT_KHR,deviceSize,&cl_dev,NULL);
+		if ( res!=CL_SUCCESS || cl_dev==0 ) {
+			std::cerr << "clGetGLContextInfoKHR_proc(cpsGL,CL_CURRENT_DEVICE_FOR_GL_CONTEXT_KHR," << deviceSize << ",...) failed" << std::endl;
+			std::cerr << " (return code : " << res << ")" << std::endl;
+			continue;
+		}
+
+		std::cout << "cl_device :" << (void *)cl_dev << std::endl;
+
+		cl_ctx = clCreateContext(cpsGL,1,&cl_dev,0,0,&res);
+		if ( res!=CL_SUCCESS ) {
+			std::cerr << "clCreateContext() failed" << std::endl; 
+			std::cerr << " (return code : " << res << ")" << std::endl;
+			continue;
+		}
+
+		cl_cq = clCreateCommandQueue(cl_ctx,cl_dev,0,&res);
+		if ( res!=CL_SUCCESS ) {
+			std::cerr << "clCreateCommandQueue() failed" << std::endl; 
+			std::cerr << " (return code : " << res << ")" << std::endl;
+			continue;
+		}
+
+		usableDeviceFound=true;
+		break;
+	}
+
+	if (! usableDeviceFound) {
+		std::cerr << "No OpenCL device has been successfully initialized" << std::endl;
+		return 13;
+	}
+
+	cl_vbo = clCreateFromGLBuffer(cl_ctx, CL_MEM_WRITE_ONLY, gl_vbo, &res);
+  if ( res!=CL_SUCCESS ) {
+		std::cerr << "clCreateFromGLBuffer(..., gl_vbo, &res) failed" << std::endl;
+		return 14;
+	}
+
+	std::cout << "OpenCL initialization done." << std::endl;
+	return 0;
+	
+}
+
+cl_int OpenCLMeshKit::compileKernels(std::list<std::string> names, const char source[], size_t sourceLen) {
+	cl_int res=0;
+
+	const char *p_source=source;
+
+	cl_program program = clCreateProgramWithSource(cl_ctx, 1, &p_source, &sourceLen,&res);
+	if ( res!=CL_SUCCESS ) {
+		std::cerr << "Failed to clCreateProgramWithSource(<source of zero_z kernel>)" << std::endl;
+		return 21;
+	}
+
+	res = clBuildProgram(program, 1, &cl_dev, "", NULL, NULL);
+  if ( res!=CL_SUCCESS ) {
+    std::cerr << "Failed to clBuildProgram()"  << std::endl;
+		return 22;
+	}
+
+	for (std::list<std::string>::iterator ii = names.begin(); ii != names.end(); ++ii) {
+		std::string kName = (*ii);
+		char *kNameZTS = new char[kName.length()+1];
+		std::strcpy(kNameZTS, kName.c_str());
+
+		cl_kernel kernel = clCreateKernel(program,kNameZTS,&res);
+		delete [] kNameZTS;
+
+		if ( res!=CL_SUCCESS ) {
+			std::cerr << "Failed to clCreateKernel(program,\"" << kName << "\",&res);" << std::endl;
+			return 23;
+		}
+		kernels[kName]=kernel;
+	}
+
+	return 0;
+}
+
+cl_int OpenCLMeshKit::execKernel(std::string kernelName) {
+	/*TODO
+
+cl_int execKernel(cl_context cl_ctx, cl_command_queue commandQueue, cl_kernel kernel, size_t mesh_width, size_t mesh_height, size_t group_size, cl_mem cl_vbo, float time) {
+
+  cl_int res;
+  cl_event  eventND[1];
+
+  // Set local and global work group sizes
+  size_t globalWorkSize[2], localWorkSize[2];
+  globalWorkSize[0]=mesh_width;
+  globalWorkSize[1]=mesh_height;
+  localWorkSize[0]=group_size;
+  localWorkSize[1]=1;
+
+  res=clEnqueueAcquireGLObjects(commandQueue, 1, &cl_vbo, 0, 0, NULL);
+
+  res=clSetKernelArg(kernel, 0, sizeof(cl_mem), &cl_vbo); // float4 *pos
+  res=clSetKernelArg(kernel, 1, sizeof(cl_uint), (void *)&mesh_width);
+  res=clSetKernelArg(kernel, 2, sizeof(cl_uint), (void *)&mesh_height);
+  res=clSetKernelArg(kernel, 3, sizeof(float), &time);
+
+  // Execute kernel on given device
+  res=clEnqueueNDRangeKernel(commandQueue, kernel, 2, NULL, globalWorkSize, localWorkSize, 0, NULL, eventND);
+  if ( res != CL_SUCCESS ) {
+    fputs("Failed to clEnqueueNDRangeKernel()\n", stderr);
+    return 1;
+  }
+
+  res=clFlush(commandQueue);
+  res=clWaitForEvents(1,eventND); //XXX: SimpleGL utilise une attente active, pourquoi ?
+  res=clReleaseEvent(eventND[0]);
+  res=clEnqueueReleaseGLObjects(commandQueue, 1, &cl_vbo, 0, 0, 0);
+  res=clFinish(commandQueue);
+
+  return CL_SUCCESS;
+}
+*/
+	return 0;
+}
+
+void OpenCLMeshKit::releaseKernels() {
+	for (std::map<std::string,cl_kernel>::iterator ii = kernels.begin(); ii != kernels.end(); ++ii ) {
+		clReleaseKernel((*ii).second);
+	}
+	kernels.clear();
+}
+
+cl_int OpenCLMeshKit::resetVBO() {
+	cl_int res;
+	std::map<std::string, cl_kernel> user_kernels=kernels;
+
+	std::list<std::string> n; n.push_back("zero_z");
+	res = compileKernels(n, kernel_src_zero_z, sizeof(kernel_src_zero_z));
+
+	if(res==0) res = execKernel("zero_z");	
+	
+	releaseKernels();
+	kernels=user_kernels;
+
+	return res;
+}
+
+void OpenCLMeshKit::setGroupSize(size_t groupSize) {
+	this->groupSize=groupSize;
+}
+
+OpenCLMeshKit::~OpenCLMeshKit() { }
+
diff --git a/src/opencl_mesh_kit.hpp b/src/opencl_mesh_kit.hpp
new file mode 100644
index 0000000..aa3a013
--- /dev/null
+++ b/src/opencl_mesh_kit.hpp
@@ -0,0 +1,60 @@
+#include <CL/opencl.h>
+#include <iostream>
+#include <map>
+#include <list>
+#include <cstring>
+
+#ifndef STRINGIFY
+  #define STRINGIFY(x) #x
+#endif
+
+class OpenCLMeshKit
+{
+	public:
+		//RAII is violated but it is really triky to do differently
+		cl_int initCL(intptr_t gl_display, intptr_t gl_context, intptr_t gl_vbo, size_t meshWidth, size_t meshHeight, size_t groupSize);
+		cl_int compileKernels(std::list<std::string> names, const char source[], size_t sourceLen);
+		cl_int execKernel(std::string kernelName);
+		void releaseKernels();
+		void setGroupSize(size_t groupSize);
+
+		// Quick and dirty function to initialize a test mesh
+		cl_int resetVBO();
+
+		virtual ~OpenCLMeshKit();
+
+	protected:
+		size_t meshWidth;
+		size_t meshHeight;
+		size_t groupSize;
+
+		cl_context cl_ctx;
+		cl_device_id cl_dev;
+		cl_command_queue cl_cq;
+		cl_mem cl_vbo;
+
+		std::map<std::string, cl_kernel> kernels;
+};
+
+/* Kernel for resetVBO() 
+To write your own kernels, take this one a make the calculus you want for z variable staying in [-0.5;0.5] if you want everything a 1*1*1 cube */
+const char kernel_src_zero_z[]=STRINGIFY(
+
+	__kernel void zero_z(__global float4 *pos, unsigned int width, unsigned int height, float time) {
+		unsigned int nx = get_global_id(0);
+		unsigned int ny = get_global_id(1);
+		/* calculate uv coordinates of the mesh point [0.0;1.0] */
+		float u = nx / (float) width;
+		float v = ny / (float) height;
+		/* calculate centered coordinates [-0.5;0.5] */
+		float x = (u*2-1)/2;
+		float y = (v*2-1)/2;
+		/* We only use normalized quaterinons here */
+		float w = 1.0f;
+		/* Calculate the desirated value of the mesh point */
+		float z = 0.0f;
+		/* Write output vertex (centered) */
+		pos[ny*width+nx] = (float4)(x, y, z, w);
+	}
+
+);