| Trees | Indices | Help |
|---|
|
|
1 #!/usr/bin/env python
2
3 #! This file is a literate Python program. You can compile the documentation
4 #! using mylit (http://pypi.python.org/pypi/mylit/).
5 ## title = "glitter Example: OpenCL"
6 ## stylesheet = "pygments_style.css"
7
8 # <h1><i>glitter</i> Example: OpenCL</h1>
9
10 # <h2>Summary</h2>
11
12 # This program will open a GLUT window and render an animated particle
13 # simulation. The particle simulation is computed in OpenCL, and the results
14 # are written directly to an OpenGL array buffer.
15
16 # <img src="opencl.png">
17
18 # <h2>Front matter</h2>
19
20 # <h3>Module docstring</h3>
21
22 # The module docstring is used as a description of this example in the
23 # generated documentation:
24 """Basic OpenGL/OpenCL interoperability example.
25
26 @author: Stephan Wenger
27 @date: 2012-02-29
28 """
29
30 # <h3>Imports</h3>
31
32 # We use numpy for creating the particle data:
33 import numpy
34
35 # <a href="http://mathema.tician.de/software/pyopencl">PyOpenCL</a> (<a
36 # href="http://documen.tician.de/pyopencl/">documentation</a>, <a
37 # href="http://pypi.python.org/pypi/pyopencl">PyPi</a>, <a
38 # href="http://sourceforge.net/projects/pyopencl/">download</a>) is a
39 # convenient Python wrapper for OpenCL:
40 import pyopencl as cl
41
42 # For creating an OpenCL context that shares data with OpenGL, we need
43 # <code>get_gl_sharing_context_properties</code>:
44 from pyopencl.tools import get_gl_sharing_context_properties
45
46 # We can usually import classes and functions contained in <i>glitter</i>
47 # submodules directly from glitter:
48 from glitter import ArrayBuffer, VertexArray, get_default_program
49
50 # Modules with external dependencies other than numpy, such as platform
51 # dependent parts like methods for the generation of an OpenGL context,
52 # however, have to be imported from their respective submodules:
53 from glitter.contexts.glut import GlutWindow, main_loop
54
55 # <h2>OpenCL</h2>
56
57 # <h3>OpenCL kernel</h3>
58
59 # The OpenCL kernel (adapted from a <a
60 # href="http://www.cmsoft.com.br/index.php?option=com_content&view=category&layout=blog&id=99&Itemid=150">CMSoft
61 # tutorial</a>) computes a single iteration of a simple Eulerian particle
62 # simulation. Kernel code can be loaded from a file or defined inline as a
63 # Python string:
64 kernel_code = """
65 __kernel void animate(__global float4* positions,
66 __global float4* colors,
67 __global float4* velocities,
68 __global float4* initial_positions,
69 __global float4* initial_velocities,
70 float dt) {
71 unsigned int i = get_global_id(0);
72 float4 p = positions[i];
73 float4 v = velocities[i];
74 float life = velocities[i].w;
75
76 life -= dt;
77 if (life <= 0.0) {
78 p = initial_positions[i];
79 v = initial_velocities[i];
80 life = 1.0;
81 }
82
83 v.z -= 9.8 * dt;
84 p.xyz += v.xyz * dt;
85 v.w = life;
86
87 positions[i] = p;
88 velocities[i] = v;
89 colors[i].w = life;
90 }
91 """
92
93 # <h3>OpenCL interaction</h3>
94
95 # The <code>CLCode</code> class will comprise all OpenCL interaction:
97 # <h4>Initialization</h4>
98 # The constructor receives OpenGL buffers for positions and colors as well
99 # as a numpy array of velocities and a timestep for the simulation:
101 # First, we have to initialize the OpenCL context. That means we have
102 # to get a list of available platforms and select one:
103 platform = cl.get_platforms()[0]
104 # Then, we can create a context. Passing
105 # <code>get_gl_sharing_context_properties()</code> as a property
106 # ensures that we share state with the active OpenGL context:
107 self.ctx = cl.Context(properties=[(cl.context_properties.PLATFORM, platform)] +
108 get_gl_sharing_context_properties(), devices=None)
109 # A command queue is necessary for serializing OpenCL commands:
110 self.queue = cl.CommandQueue(self.ctx)
111 # Finally, we can compile the kernel:
112 self.program = cl.Program(self.ctx, kernel_code).build()
113
114 # The constructor parameters are stored for later use:
115 self.gl_positions = gl_positions
116 self.gl_colors = gl_colors
117 self.velocities = velocities
118 # The <code>dt</code> value will later be passed to an OpenCL kernel as
119 # a 32-bit float. We therefore wrap it in a numpy <code>float32</code>
120 # object:
121 self.dt = numpy.float32(dt)
122
123 # Next, we generate OpenCL buffers. The positions and colors are
124 # contained in OpenGL buffers, which we wrap in PyOpenCL's
125 # <code>GLBuffer<(/code> class:
126 self.cl_positions = cl.GLBuffer(self.ctx, cl.mem_flags.READ_WRITE, self.gl_positions._id)
127 self.cl_colors = cl.GLBuffer(self.ctx, cl.mem_flags.READ_WRITE, self.gl_colors._id)
128 # Note how we had to extract the <code>_id</code>s from the
129 # <code>ArrayBuffer</code> objects. In pure <i>glitter</i> code, you
130 # should never (have to) access this value; however for interaction
131 # with other OpenGL-related libraries, this cannot always be avoided.
132
133 # The velocities are given as a numpy array, which is simply uploaded
134 # into a new OpenCL <code>Buffer</code> object along with the initial
135 # values of the positions and colors:
136 self.cl_velocities = cl.Buffer(self.ctx, cl.mem_flags.READ_ONLY |
137 cl.mem_flags.COPY_HOST_PTR, hostbuf=velocities)
138 self.cl_initial_positions = cl.Buffer(self.ctx, cl.mem_flags.READ_ONLY |
139 cl.mem_flags.COPY_HOST_PTR, hostbuf=self.gl_positions.data)
140 self.cl_initial_velocities = cl.Buffer(self.ctx, cl.mem_flags.READ_ONLY |
141 cl.mem_flags.COPY_HOST_PTR, hostbuf=self.velocities)
142
143 # <h4>Execution</h4>
144 # The <code>execute</code> function executes the OpenCL kernel several
145 # times in a row:
147 # First, we have to make sure that OpenGL is done using the buffer objects:
148 cl.enqueue_acquire_gl_objects(self.queue, [self.cl_positions, self.cl_colors])
149 # Now, we can safely call the kernel. Its arguments are buffer objects:
150 args = (self.cl_positions, self.cl_colors, self.cl_velocities,
151 self.cl_initial_positions, self.cl_initial_velocities, self.dt)
152 # The kernel will be executed several times with a small step size.
153 # This increases the accuracy with respect to a single step with a
154 # large step size. However, it is not necessary to display all the
155 # intermediate results.
156 for i in xrange(0, sub_intervals):
157 # In each step, the <code>animate</code> kernel function is called.
158 # Its arguments are the queue object that schedules its execution,
159 # the global and local block sizes, and any arguments that will be
160 # passed to the actual kernel.
161 self.program.animate(self.queue, [len(self.gl_positions)], None, *args)
162 # Finally, we allow OpenGL to access the buffers again:
163 cl.enqueue_release_gl_objects(self.queue, [self.cl_positions, self.cl_colors])
164
165 # <h2>Main class</h2>
166
167 # We wrap all the OpenGL interaction in a class. The class will contain an
168 # <code>__init__()</code> method to set up all OpenGL objects, any required
169 # callback methods, as well as a <code>run()</code> method to trigger execution
170 # of the GLUT main loop.
172 # <h3>Initialization</h3>
173
174 # When a <code>OpenCLExample</code> instance is created, we need to
175 # initialize a few OpenGL objects.
177 # First, we create a window; this also creates an OpenGL context.
178 self.window = GlutWindow(double=True, alpha=True, depth=True)
179
180 # Then, we set the GLUT display and keyboard callback functions which
181 # will be defined later.
182 self.window.display_callback = self.display
183 self.window.keyboard_callback = self.keyboard
184
185 # Here, we generate numpy arrays to hold the positions, colors, and
186 # velocities of the particles:
187 num = 200000
188 positions = numpy.empty((num, 4), dtype=numpy.float32)
189 colors = numpy.empty((num, 4), dtype=numpy.float32)
190 velocities = numpy.empty((num, 4), dtype=numpy.float32)
191
192 # So far, the array contents are undefined. We have to initialize them with meaningful values:
193 positions[:, 0] = numpy.sin(numpy.arange(0, num) * 2 * numpy.pi / num) * (numpy.random.random_sample((num,)) / 3 + 0.2)
194 positions[:, 1] = numpy.cos(numpy.arange(0, num) * 2 * numpy.pi / num) * (numpy.random.random_sample((num,)) / 3 + 0.2)
195 positions[:, 2:] = 0, 1
196 colors[:] = 0, 1, 0, 1
197 velocities[:, :2] = 2 * positions[:, :2]
198 velocities[:, 2] = 3
199 velocities[:, 3] = numpy.random.random_sample((num,))
200
201 # Instead of simply generating a vertex array from the position and color
202 # data, we first generate array buffers for them:
203 gl_positions = ArrayBuffer(data=positions, usage="DYNAMIC_DRAW")
204 gl_colors = ArrayBuffer(data=colors, usage="DYNAMIC_DRAW")
205 # These array buffers will later also be used by OpenCL. We do not need to
206 # wrap <code>velocities</code> in this way, as it will only be used by
207 # OpenCL and can be wrapped in an OpenCL buffer directly.
208
209 # We now create a vertex array that will pass the position and color data
210 # to the shader. The vertex array constructor accepts
211 # <code>ArrayBuffer</code> instances:
212 self.vao = VertexArray(gl_positions, gl_colors)
213
214 # In the OpenGL core profile, there is no such thing as a "standard pipeline"
215 # any more. We use the minimalistic <code>defaultpipeline</code> from the
216 # <code>glitter.convenience</code> module to create a shader program instead:
217 self.shader = get_default_program()
218
219 # Here, we create the <code>CLCode</code> object that manages OpenCL
220 # interaction. It is passed the OpenGL buffer objects as well as a numpy
221 # array of velocities.
222 self.clcode = CLCode(gl_positions, gl_colors, velocities)
223
224 # <h3>Callback functions</h3>
225
226 # <h4>Display function</h4>
227
228 # Here we define the display function. It will be called by GLUT whenever the
229 # screen has to be redrawn.
231 # First we clear the default framebuffer:
232 self.window.clear()
233
234 # To draw the vertex array, we use:
235 self.vao.draw()
236
237 # After all rendering commands have been issued, we swap the back buffer to
238 # the front, making the rendered image visible all at once:
239 self.window.swap_buffers()
240
241 # <h4>Timer function</h4>
242
243 # The animation is controlled by a GLUT timer. The timer callback animates the
244 # particle system, schedules the next timer event, and causes a screen redraw:
246 # We first tell an instance of the <code>CLCode</code> class to execute the
247 # OpenCL kernel:
248 self.clcode.execute(10)
249
250 # The following line schedules the next timer event to execute after one millisecond.
251 self.window.add_timer(1, self.timer)
252
253 # Finally, we tell GLUT to redraw the screen.
254 self.window.post_redisplay()
255
256 # <h4>Keyboard function</h4>
257
258 # To further illustrate the concept of GLUT callbacks, here's a keyboard
259 # handler that will simply make the program exit when any key is pressed:
262
263 # <h3>Running</h3>
264
265 # We will call the <code>run()</code> method later to run the OpenGL code.
267 # To start the animation, we call the timer once; all subsequent timer
268 # calls will be scheduled by the timer function itself.
269 self.timer()
270
271 # The default program is bound by using a <code>with</code> statement. At
272 # the same time, we can pass in additional uniform variables, such as the
273 # modelview matrix:
274 with self.shader(modelview_matrix=((1, 0, 0, 0), (0, 0, 1, 0), (0, 1, 0, 0), (0, 0, 0, 2))):
275 # With the shader bound, we enter the GLUT main loop.
276 main_loop()
277
278 # When the main loop exits, control is handed back to the script,
279 # unless <code>SystemExit</code> has been raised by the keyboard
280 # handler.
281
282 # <h2>Main section</h2>
283
284 # Finally, if this program is being run from the command line, we instanciate
285 # the main class and run it.
286 if __name__ == "__main__":
287 OpenCLExample().run()
288
| Trees | Indices | Help |
|---|
| Generated by Epydoc 3.0.1 on Fri Mar 16 17:56:05 2012 | http://epydoc.sourceforge.net |