blob: 29e8d64b6e86812229b5e52c7fbd05e76b6bc350 (
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
|
#include <CL/opencl.h>
#include <iostream>
#include <cstring>
#include <map>
#include <time.h>
#ifndef STRINGIFY
#define STRINGIFY(x) #x
#endif
class OpenCLMeshKit
{
public:
//RAII is violated but it is really tricky 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(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 [-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) {
unsigned int nx = get_global_id(0);
unsigned int ny = get_global_id(1);
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;
}
);
|
