From 35e25937ad05e409340e7cd356c3ce1a45a5a3f9 Mon Sep 17 00:00:00 2001
From: Ludovic Pouzenc <ludovic@pouzenc.fr>
Date: Sun, 3 Mar 2013 20:27:37 +0000
Subject: 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
---
src/boring_parts.cpp | 122 ----------------------
src/boring_parts.hpp | 5 +-
src/compil.sh | 2 +-
src/gpudataviz.cpp | 54 +++++-----
src/my_gtk_gl_scene_widget.hpp | 2 +
src/opencl_mesh_kit.cpp | 225 +++++++++++++++++++++++++++++++++++++++++
src/opencl_mesh_kit.hpp | 60 +++++++++++
7 files changed, 313 insertions(+), 157 deletions(-)
create mode 100644 src/opencl_mesh_kit.cpp
create mode 100644 src/opencl_mesh_kit.hpp
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);
+ }
+
+);
--
cgit v1.2.3