summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
authorLudovic Pouzenc <ludovic@pouzenc.fr>2013-04-19 20:46:21 +0000
committerLudovic Pouzenc <ludovic@pouzenc.fr>2013-04-19 20:46:21 +0000
commitf88f52617be9966cddd3ec28d590704fb8a615eb (patch)
tree943ff25992e8e5721a28a18925f5acd634358ec1
parent151d0ff1af64b058f37610e61c6de7e7845416d2 (diff)
download2013-gpudataviz-f88f52617be9966cddd3ec28d590704fb8a615eb.tar.gz
2013-gpudataviz-f88f52617be9966cddd3ec28d590704fb8a615eb.tar.bz2
2013-gpudataviz-f88f52617be9966cddd3ec28d590704fb8a615eb.zip
Ajout option -Werror et affichage lorsque ca compile pas.origin/trunk
Amelioration des kernels en conséquence. Bugfix avec le sizeof(source) qui prennait un caractère de trop (le '\0') git-svn-id: file:///var/svn/2013-gpudataviz/trunk@31 371a6b4a-a258-45f8-9dcc-bdd82ce0ac9d
-rw-r--r--src/gpudataviz.cpp41
-rw-r--r--src/opencl_mesh_kit.cpp17
-rw-r--r--src/opencl_mesh_kit.hpp18
3 files changed, 44 insertions, 32 deletions
diff --git a/src/gpudataviz.cpp b/src/gpudataviz.cpp
index 1dc14f8..a3e577e 100644
--- a/src/gpudataviz.cpp
+++ b/src/gpudataviz.cpp
@@ -123,33 +123,38 @@ void MyGTKGLSceneWidget::on_realize() {
const char source[]=STRINGIFY(
/* This is OpenCL kernel code (Syntax like C but it's a different language) */
+
__kernel void water1(__global float4 *pos, unsigned int width, unsigned int height, float time) {
+ /* Getting current vertex indices (could be seen as a 2D array of float4) */
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 normalized coordinates [-1.0;1.0] */
- float x = u*2.0-1.0;
- float y = v*2.0-1.0;
- /* set some constants and calculate some intermediate values */
- float freq = 8.0 * 3.14;
- 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);
- /* 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) */
- pos[ny*width+nx] = (float4)(x, y, z, w);
+ /* A float4 vector that hold the output vertex */
+ float4 out;
+
+ /* Calculate centered mesh coordinates [-1.0;1.0]² */
+ out.x = nx / (float) width * 2.0f - 1.0f;
+ out.y = ny / (float) height * 2.0f - 1.0f;
+
+ /* Set some constants (should be preprocessor macros...) */
+ float freq = 8.0f;
+ float speed = 1.0f;
+ float amp = 1.0f / 10.0f; /* 0.1f does NOT works for me ! WTF !!! */
+
+ /* Calculate some intermediate values */
+ float dist = hypot(out.x,out.y); /* =sqrt(out.x*out.x+out.y*out.y); */
+
+ /* Calculate the desirated value of the mesh point z=f(x,y,t) */
+ out.z = amp * sinpi( freq * dist + speed * time ) / dist ;
+ out.w = 1.0f; /* We always use normalized quaterinons here */
+
+ pos[ny*width+nx] = out; /* Write output vertex */
}
);
// 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));
+ cl_res = this->clKit.compileKernels(/*knames,*/ source, sizeof(source)-1);
//knames.clear();
EXIT_IF_FAIL(6, cl_res==0);
diff --git a/src/opencl_mesh_kit.cpp b/src/opencl_mesh_kit.cpp
index 6e41c25..318a928 100644
--- a/src/opencl_mesh_kit.cpp
+++ b/src/opencl_mesh_kit.cpp
@@ -131,9 +131,22 @@ cl_int OpenCLMeshKit::compileKernels(const char source[], size_t sourceLen) {
return 21;
}
- res = clBuildProgram(program, 1, &cl_dev, "", NULL, NULL);
+ res = clBuildProgram(program, 1, &cl_dev, "-Werror", NULL, NULL);
if ( res!=CL_SUCCESS ) {
std::cerr << "Failed to clBuildProgram()" << std::endl;
+
+ // Shows the log
+ char* build_log;
+ size_t log_size;
+ // First call to know the proper size
+ clGetProgramBuildInfo(program, cl_dev, CL_PROGRAM_BUILD_LOG, 0, NULL, &log_size);
+ build_log = new char[log_size+1];
+ // Second call to get the log
+ clGetProgramBuildInfo(program, cl_dev, CL_PROGRAM_BUILD_LOG, log_size, build_log, NULL);
+ build_log[log_size] = '\0';
+ std::cerr << build_log << std::endl;
+ delete[] build_log;
+
return 22;
}
@@ -241,7 +254,7 @@ cl_int OpenCLMeshKit::resetVBO() {
cl_int res;
std::map<std::string, cl_kernel> user_kernels=kernels;
- res = compileKernels(kernel_src_zero_z, sizeof(kernel_src_zero_z));
+ res = compileKernels(kernel_src_zero_z, sizeof(kernel_src_zero_z)-1);
if(res==0) res = execKernel("zero_z", 0.0f);
diff --git a/src/opencl_mesh_kit.hpp b/src/opencl_mesh_kit.hpp
index 5efd30e..29e8d64 100644
--- a/src/opencl_mesh_kit.hpp
+++ b/src/opencl_mesh_kit.hpp
@@ -52,18 +52,12 @@ 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 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 */
- float z = 0.0f;
- /* We only use normalized quaterinons here */
- float w = 1.0f;
- /* Write output vertex (centered) */
- pos[ny*width+nx] = (float4)(x, y, z, w);
+ float4 out;
+ out.x = nx / (float) width * 2.0f - 1.0f;
+ out.y = ny / (float) height * 2.0f - 1.0f;
+ out.z = 0.0f;
+ out.w = 1.0f;
+ pos[ny*width+nx] = out;
}
);