From 35e25937ad05e409340e7cd356c3ce1a45a5a3f9 Mon Sep 17 00:00:00 2001 From: Ludovic Pouzenc 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 #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 glCtx, Glib::RefPtr 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 //#include -int initLibs(); - #ifdef HAS_OPENCL -#include -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 +#include //#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 &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 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::WidgetmeshWidth = 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 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()" << 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::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::iterator ii = kernels.begin(); ii != kernels.end(); ++ii ) { + clReleaseKernel((*ii).second); + } + kernels.clear(); +} + +cl_int OpenCLMeshKit::resetVBO() { + cl_int res; + std::map user_kernels=kernels; + + std::list 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 +#include +#include +#include +#include + +#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 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 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