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/opencl_mesh_kit.hpp | 60 +++++++++++++++++++++++++++++++++++++++++++++++++ 1 file changed, 60 insertions(+) create mode 100644 src/opencl_mesh_kit.hpp (limited to 'src/opencl_mesh_kit.hpp') 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