summaryrefslogtreecommitdiff
path: root/src/opencl_mesh_kit.hpp
blob: ac3cf46c73438965b21203f1b8491b1174a2e346 (plain)
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
#include <CL/opencl.h>
#include <iostream>
#include <map>
#include <list>
#include <cstring>
#include <time.h>

#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, float karg_time);
		void releaseKernels();

		size_t getMeshWidth();
		size_t getMeshHeight();
		size_t getMeshItemCount();
		size_t getGroupSize();
		intptr_t getGLVBO();
		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;

		intptr_t gl_vbo; // Save this pointer for convinence (for data display code)

		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 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);
	}

);