Review + parsing du source pour trouver les noms des kernels à la vollée

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

6 files changed, 66 insertions, 54 deletions

diff --git a/src/README b/src/README
index f9570a2..00b0d0c 100644
--- a/src/README
+++ b/src/README
@@ -2,3 +2,5 @@
 # Compilation depedencies
 # Tested on Ubuntu 12.04 and 12.10
 # apt-get install pkg-config libgtkglextmm-x11-1.2-dev libgtkmm-2.4-dev
+#
+# Also needs an OpenCL-aware environnement. Could be AMDAPP SDK (ATI cards proprietary impl).
diff --git a/src/compil.sh b/src/compil.sh
index d913674..fa77686 100755
--- a/src/compil.sh
+++ b/src/compil.sh
@@ -1,8 +1,4 @@
 #!/bin/bash -e
-
-clear
-
-#set -x
 CXX="g++ -Wall -g"
 BUILD_PATH="../build"
 
@@ -29,7 +25,6 @@ else	echo "ERROR : pkg-config is unavailable or gtkglextmm-1.2 developpemnt file
 	echo " sudo apt-get install pkg-config libgtkmm-2.4-dev libgtkglextmm-x11-1.2-dev"
 	exit 1
 fi
-#set +x
 
 function build_cxx() {
 	echo "$PS4$CXX \$DEFINES \$INCLUDES -o $BUILD_PATH/$1 -c $2 \$LIBS"
diff --git a/src/gpudataviz.cpp b/src/gpudataviz.cpp
index 89adf40..1dc14f8 100644
--- a/src/gpudataviz.cpp
+++ b/src/gpudataviz.cpp
@@ -34,7 +34,7 @@ int main(int argc, char* argv[]) {
 	Gdk::GL::ConfigMode glMode = Gdk::GL::MODE_RGB | Gdk::GL::MODE_DEPTH;
 	EXIT_IF_FAIL(2, glconfig=Gdk::GL::Config::create(glMode) );
 
-	// Initialize the OpenGL scene widget (realization came later)
+	// Initialize the OpenGL scene widget (realization came later, no RAII)
 	MyGTKGLSceneWidget glScene(glconfig);
 
 	// Instantiate the GTK app (and realize glScene)
@@ -49,7 +49,9 @@ int main(int argc, char* argv[]) {
 
 // MyGTKGLSceneWidget implementation (extends a Gtk::DrawingArea with Gtk::GL::Widget)
 //  I want to keep all interesting code parts in this file, in natural reading order
-MyGTKGLSceneWidget::MyGTKGLSceneWidget(Glib::RefPtr<Gdk::GL::Config> &glconfig) {
+MyGTKGLSceneWidget::MyGTKGLSceneWidget(Glib::RefPtr<Gdk::GL::Config> &glconfig):
+	continuous_play(false), need_recompute(false), time(0.0f)
+{
 	CALL_TRACE;
 	set_gl_capability(glconfig);
 	Gdk::EventMask mask = Gdk::POINTER_MOTION_MASK | Gdk::BUTTON_MOTION_MASK \
@@ -57,15 +59,17 @@ MyGTKGLSceneWidget::MyGTKGLSceneWidget(Glib::RefPtr<Gdk::GL::Config> &glconfig)
 			      | Gdk::SCROLL_MASK;
 	set_events(mask); // The containing window should have those attributes too
 
-	// Some starting default values
-	this->camera.rx = -64.0f; this->camera.ry = -16.0f; this->camera.tz = -2.0f;
-	need_recompute=false; continuous_play=false; time=0.0f; 
+	// Camera is always looking the center of the mesh
+	// this->camera affects (only) the camera position
+	// rx,ry : the camera rotate around the mesh. Those vars are angles in degrees
+	// tz : distance between the center of the mesh and the camera (roughly)
+	this->camera.rx = -64.0f; this->camera.ry = 16.0f; this->camera.tz = 2.0f;
 }
 
 MyGTKGLSceneWidget::~MyGTKGLSceneWidget() { }
 
 void MyGTKGLSceneWidget::on_size_request(Gtk::Requisition* requisition) {
-	CALL_TRACE; // Technical stuff : GTK call this to ask the widget minimal size 
+	CALL_TRACE; // Technical stuff, GTK call this to ask for widget minimal size 
  	*requisition = Gtk::Requisition();
 	requisition->width = 320; requisition->height = 240;
 }
@@ -128,15 +132,13 @@ void MyGTKGLSceneWidget::on_realize() {
 			/* calculate centered normalized coordinates [-1.0;1.0] */
 			float x = u*2.0-1.0;
 			float y = v*2.0-1.0;
-			/* Calculate the desirated value of the mesh point */
+			/* set some constants and calculate some intermediate values */
 			float freq = 8.0 * 3.14;
-			float amp = 1.0 / 10.0; /* 0.1 does NOT works for me ! WTF !!! */
 			float speed = 1.0;
+			float amp = 1.0 / 10.0; /* 0.1 does NOT works for me ! WTF !!! */
 			float dist = sqrt(x*x+y*y);
-			float z;
-			//for(int i = 0; i < 40000; i++){
-			z = amp * sin( freq * dist - speed * time ) / dist ; 
-			//}
+			/* Calculate the desirated value of the mesh point */
+			float z = amp * sin( freq * dist + speed * time ) / dist ; 
 			/* We only use normalized quaterinons here */
 			float w = 1.0;
 			/* Write output vertex (centered) */
@@ -145,10 +147,10 @@ void MyGTKGLSceneWidget::on_realize() {
 	);
 
 	// TODO : change API, use only one string and split it at each blanks
-	std::list<std::string> knames;
-	knames.push_back("water1");
-	cl_res = this->clKit.compileKernels(knames, source, sizeof(source));
-	knames.clear();
+	//std::list<std::string> knames;
+	//knames.push_back("water1");
+	cl_res = this->clKit.compileKernels(/*knames,*/ source, sizeof(source));
+	//knames.clear();
 	EXIT_IF_FAIL(6, cl_res==0);
 
 	glEnable(GL_BLEND);
@@ -168,7 +170,6 @@ bool MyGTKGLSceneWidget::on_configure_event(GdkEventConfigure* event) {
 	Glib::RefPtr<Gdk::GL::Window> glwindow = get_gl_window();
 
 	// *** OpenGL BEGIN ***
-	//FIXME could segfault if get_gl_window() has failed
 	if (!glwindow->gl_begin(get_gl_context())) {
 		std::cerr << "Oops : glwindow->gl_begin(get_gl_context())" << std::endl;
 		return false;
@@ -211,9 +212,9 @@ bool MyGTKGLSceneWidget::on_expose_event(GdkEventExpose* event) {
 	//Camera position update
 	glMatrixMode(GL_MODELVIEW);
 	glLoadIdentity();
-	glTranslatef(0.0, 0.0, this->camera.tz);
+	glTranslatef(0.0, 0.0, this->camera.tz * -1.0);
 	glRotatef(this->camera.rx, 1.0, 0.0, 0.0);
-	glRotatef(this->camera.ry, 0.0, 0.0, 1.0);
+	glRotatef(this->camera.ry, 0.0, 0.0, -1.0);
 
 	glClear(GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT);
 
@@ -229,7 +230,7 @@ bool MyGTKGLSceneWidget::on_expose_event(GdkEventExpose* event) {
 	// The comprehensible one 
 	//float c1=(1024*1024)/(m_w*m_h);	// coef 1 decreases with mesh point quantity
 	//float c2=(s_h*s_h)/(1024*1024);	// coef 2 increases with viewport pixel quantity
-	//float c3=(-2*-2)/(t_z*t_z);		// coef 3 decreases with mesh-camera distance
+	//float c3=(2*2)/(t_z*t_z);		// coef 3 decreases with mesh-camera distance
 	//float alpha=0.5*c1*c2*c3;		// Combine it all
 	//if (alpha < 0.01f) alpha = 0.01f;	// Prevent values outside acceptable range
 	//if (alpha > 1.0f) alpha = 1.0f;
@@ -349,7 +350,7 @@ bool MyGTKGLSceneWidget::do_mouse_logic(GdkEventType type, guint state, guint x,
 		gint dx = drag_x - x; // Delta movement (since last event)
 		gint dy = drag_y - y; // Not unsigned !
 		this->camera.rx -= mouse_sensivity * dy; // Camera position update
-		this->camera.ry -= mouse_sensivity * dx; // Yes dy for camera.rx, and -= operator :
+		this->camera.ry += mouse_sensivity * dx; // Yes dy for camera.rx, and -= operator :
 		// GTK mouse coords and OpenGL ones are not on the same coords system
 		drag_x = x; drag_y = y;
 		redraw=true;
@@ -358,8 +359,8 @@ bool MyGTKGLSceneWidget::do_mouse_logic(GdkEventType type, guint state, guint x,
 	// Camera zoom-in
 	MOUSE_WHEEL(GDK_SCROLL_UP, 0, 0) {
 		
-		if (this->camera.tz + 0.5f <= -0.5f) {
-			this->camera.tz += 0.5f;
+		if (this->camera.tz - 0.5f >= 0.5f) {
+			this->camera.tz -= 0.5f;
 			//std::cout << "camera.tz == " << this->camera.tz << std::endl;
 			redraw=true;
 		}
@@ -367,8 +368,8 @@ bool MyGTKGLSceneWidget::do_mouse_logic(GdkEventType type, guint state, guint x,
 
 	// Camera zoom-out
 	MOUSE_WHEEL(GDK_SCROLL_DOWN, 0, 0) {
-		if (this->camera.tz - 0.5f >= -9.0f) {
-			this->camera.tz -= 0.5f;
+		if (this->camera.tz + 0.5f <= 9.0f) {
+			this->camera.tz += 0.5f;
 			//std::cout << "camera.tz == " << this->camera.tz << std::endl;
 			redraw=true;
 		}
diff --git a/src/my_gtk_gl_scene_widget.hpp b/src/my_gtk_gl_scene_widget.hpp
index 1c7e305..62c8d16 100644
--- a/src/my_gtk_gl_scene_widget.hpp
+++ b/src/my_gtk_gl_scene_widget.hpp
@@ -32,8 +32,8 @@ class MyGTKGLSceneWidget : public Gtk::DrawingArea, public Gtk::GL::Widget<MyGTK
 	private:
 		OpenCLMeshKit clKit;
 		struct camera_params { float rx; float ry; float tz; } camera;
-		float time;
 		bool need_recompute;
+		float time;
 };
 
 #endif /*MY_GTK_GL_SCENE_H*/
diff --git a/src/opencl_mesh_kit.cpp b/src/opencl_mesh_kit.cpp
index 8ad18b3..6e41c25 100644
--- a/src/opencl_mesh_kit.cpp
+++ b/src/opencl_mesh_kit.cpp
@@ -120,7 +120,7 @@ cl_int OpenCLMeshKit::initCL(intptr_t gl_display, intptr_t gl_context, intptr_t
 	
 }
 
-cl_int OpenCLMeshKit::compileKernels(std::list<std::string> names, const char source[], size_t sourceLen) {
+cl_int OpenCLMeshKit::compileKernels(const char source[], size_t sourceLen) {
 	cl_int res=0;
 
 	const char *p_source=source;
@@ -132,27 +132,43 @@ cl_int OpenCLMeshKit::compileKernels(std::list<std::string> names, const char so
 	}
 
 	res = clBuildProgram(program, 1, &cl_dev, "", NULL, NULL);
-  if ( res!=CL_SUCCESS ) {
-    std::cerr << "Failed to clBuildProgram()"  << std::endl;
+	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;
+	char *p_word;
+	int state=0;
+	char *source2=strdup(source); // strtok will alter the source2 string
+	char *strtok_arg1=source2;
+
+	// Trivial parsing of source to find every kernel name and register them
+	res=CL_SUCCESS;
+	while ( res == CL_SUCCESS && ( p_word=strtok(strtok_arg1, "\n\r\t (") ) != NULL ) {
+		strtok_arg1=NULL; // strtok need it's first arg NULL after the first call
+		switch(state) {
+			case 0: // Searching "__kernel"
+				if ( strcmp(p_word, "__kernel")==0 ) {
+					state=1;
+				}
+				break;
+			case 1: // Skipping kernel return type (void)
+				state=2;
+				break;
+			case 2: // Grabbing kernel name and register it
+				cl_kernel kernel = clCreateKernel(program,p_word,&res);
+				if ( res!=CL_SUCCESS ) {
+					std::cerr << "Failed to clCreateKernel(program,\""
+					       	<< p_word << "\",&res);" << std::endl;
+				}
+				kernels[std::string(p_word)]=kernel;
+				state=0;
+				break;
 		}
-		kernels[kName]=kernel;
 	}
+	delete[] source2;
 
-	return 0;
+	return res;
 }
 
 cl_int OpenCLMeshKit::execKernel(std::string kernelName, float karg_time) {
@@ -225,8 +241,7 @@ 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));
+	res = compileKernels(kernel_src_zero_z, sizeof(kernel_src_zero_z));
 
 	if(res==0) res = execKernel("zero_z", 0.0f);	
 	
@@ -242,7 +257,7 @@ size_t OpenCLMeshKit::getMeshItemCount() { return this->meshWidth * this->meshHe
 size_t OpenCLMeshKit::getGroupSize() { return this->groupSize; }
 intptr_t OpenCLMeshKit::getGLVBO() { return this->gl_vbo; }
 
-void OpenCLMeshKit::setGroupSize(size_t groupSize) { 	this->groupSize=groupSize; }
+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
index ac3cf46..b18603b 100644
--- a/src/opencl_mesh_kit.hpp
+++ b/src/opencl_mesh_kit.hpp
@@ -1,8 +1,7 @@
 #include <CL/opencl.h>
 #include <iostream>
-#include <map>
-#include <list>
 #include <cstring>
+#include <map>
 #include <time.h>
 
 #ifndef STRINGIFY
@@ -14,7 +13,7 @@ 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 compileKernels(const char source[], size_t sourceLen);
 		cl_int execKernel(std::string kernelName, float karg_time);
 		void releaseKernels();
 
@@ -47,7 +46,7 @@ class OpenCLMeshKit
 };
 
 /* 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 */
+To write your own kernels, take this one a make the calculus you want for z variable staying in [-1.0;1.0] 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) {