diff --git a/examples/demo_meta_codepy.py b/examples/demo_meta_codepy.py index c080109b9dcfe45c16525db2eaa7709f9250b3a9..2ba293c5dfc3783f449b8bc6e0b060a90a4d0c3e 100644 --- a/examples/demo_meta_codepy.py +++ b/examples/demo_meta_codepy.py @@ -19,10 +19,10 @@ a_buf = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=a) b_buf = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=b) c_buf = cl.Buffer(ctx, mf.WRITE_ONLY, b.nbytes) -from codepy.cgen import FunctionBody, \ +from cgen import FunctionBody, \ FunctionDeclaration, Typedef, POD, Value, \ Pointer, Module, Block, Initializer, Assign, Const -from codepy.cgen.opencl import CLKernel, CLGlobal, \ +from cgen.opencl import CLKernel, CLGlobal, \ CLRequiredWorkGroupSize mod = Module([ @@ -33,14 +33,14 @@ mod = Module([ arg_decls=[CLGlobal(Pointer(Const(POD(dtype, name)))) for name in ["tgt", "op1", "op2"]]))), Block([ - Initializer(POD(numpy.int32, "idx"), + Initializer(POD(numpy.int32, "idx"), "get_local_id(0) + %d * get_group_id(0)" % (local_size*thread_strides)) ]+[ Assign( "tgt[idx+%d]" % (o*local_size), "op1[idx+%d] + op2[idx+%d]" % ( - o*local_size, + o*local_size, o*local_size)) for o in range(thread_strides)]))]) @@ -50,7 +50,7 @@ knl(queue, (local_size*macroblock_count,), (local_size,), c_buf, a_buf, b_buf) c = numpy.empty_like(a) -cl.enqueue_read_buffer(queue, c_buf, c).wait() +cl.enqueue_copy(queue, c, c_buf).wait() assert la.norm(c-(a+b)) == 0 diff --git a/examples/demo_meta_template.py b/examples/demo_meta_template.py index fc64934385b58c7ac6a2d5b72a5b4fb1327de688..a39e954221b94cd44876e3af42a3d0feca129849 100644 --- a/examples/demo_meta_template.py +++ b/examples/demo_meta_template.py @@ -23,8 +23,8 @@ from mako.template import Template tpl = Template(""" __kernel void add( - __global ${ type_name } *tgt, - __global const ${ type_name } *op1, + __global ${ type_name } *tgt, + __global const ${ type_name } *op1, __global const ${ type_name } *op2) { int idx = get_local_id(0) @@ -33,13 +33,13 @@ tpl = Template(""" % for i in range(thread_strides): <% offset = i*local_size %> - tgt[idx + ${ offset }] = - op1[idx + ${ offset }] + tgt[idx + ${ offset }] = + op1[idx + ${ offset }] + op2[idx + ${ offset } ]; % endfor }""") -rendered_tpl = tpl.render(type_name="float", +rendered_tpl = tpl.render(type_name="float", local_size=local_size, thread_strides=thread_strides) knl = cl.Program(ctx, str(rendered_tpl)).build().add @@ -48,6 +48,6 @@ knl(queue, (local_size*macroblock_count,), (local_size,), c_buf, a_buf, b_buf) c = numpy.empty_like(a) -cl.enqueue_read_buffer(queue, c_buf, c).wait() +cl.enqueue_copy(queue, c, c_buf).wait() assert la.norm(c-(a+b)) == 0 diff --git a/examples/gl_particle_animation.py b/examples/gl_particle_animation.py index 1d838a2a4a0884dc53f7d24e8319336c5b7ca3ee..c8ac9c20a461c4e307497a4430f4a832dead4f84 100644 --- a/examples/gl_particle_animation.py +++ b/examples/gl_particle_animation.py @@ -1,25 +1,27 @@ # Visualization of particles with gravity # Source: http://enja.org/2010/08/27/adventures-in-opencl-part-2-particles-with-opengl/ -import pyopencl as cl # OpenCL - GPU computing interface +import pyopencl as cl # OpenCL - GPU computing interface + mf = cl.mem_flags from pyopencl.tools import get_gl_sharing_context_properties -from OpenGL.GL import * # OpenGL - GPU rendering interface -from OpenGL.GLU import * # OpenGL tools (mipmaps, NURBS, perspective projection, shapes) -from OpenGL.GLUT import * # OpenGL tool to make a visualization window -from OpenGL.arrays import vbo -import numpy # Number tools -import sys # System tools (path, modules, maxint) +from OpenGL.GL import * # OpenGL - GPU rendering interface +from OpenGL.GLU import * # OpenGL tools (mipmaps, NURBS, perspective projection, shapes) +from OpenGL.GLUT import * # OpenGL tool to make a visualization window +from OpenGL.arrays import vbo +import numpy # Number tools +import sys # System tools (path, modules, maxint) width = 800 height = 600 num_particles = 100000 -time_step = .005 +time_step = 0.005 mouse_down = False -mouse_old = {'x': 0., 'y': 0.} -rotate = {'x': 0., 'y': 0., 'z': 0.} -translate = {'x': 0., 'y': 0., 'z': 0.} -initial_translate = {'x': 0., 'y': 0., 'z': -2.5} +mouse_old = {"x": 0.0, "y": 0.0} +rotate = {"x": 0.0, "y": 0.0, "z": 0.0} +translate = {"x": 0.0, "y": 0.0, "z": 0.0} +initial_translate = {"x": 0.0, "y": 0.0, "z": -2.5} + def glut_window(): glutInit(sys.argv) @@ -37,60 +39,79 @@ def glut_window(): glViewport(0, 0, width, height) glMatrixMode(GL_PROJECTION) glLoadIdentity() - gluPerspective(60., width / float(height), .1, 1000.) + gluPerspective(60.0, width / float(height), 0.1, 1000.0) + + return window - return(window) def initial_buffers(num_particles): np_position = numpy.ndarray((num_particles, 4), dtype=numpy.float32) np_color = numpy.ndarray((num_particles, 4), dtype=numpy.float32) np_velocity = numpy.ndarray((num_particles, 4), dtype=numpy.float32) - np_position[:,0] = numpy.sin(numpy.arange(0., num_particles) * 2.001 * numpy.pi / num_particles) - np_position[:,0] *= numpy.random.random_sample((num_particles,)) / 3. + .2 - np_position[:,1] = numpy.cos(numpy.arange(0., num_particles) * 2.001 * numpy.pi / num_particles) - np_position[:,1] *= numpy.random.random_sample((num_particles,)) / 3. + .2 - np_position[:,2] = 0. - np_position[:,3] = 1. - - np_color[:,:] = [1.,1.,1.,1.] # White particles - - np_velocity[:,0] = np_position[:,0] * 2. - np_velocity[:,1] = np_position[:,1] * 2. - np_velocity[:,2] = 3. - np_velocity[:,3] = numpy.random.random_sample((num_particles, )) - - gl_position = vbo.VBO(data=np_position, usage=GL_DYNAMIC_DRAW, target=GL_ARRAY_BUFFER) + np_position[:, 0] = numpy.sin( + numpy.arange(0.0, num_particles) * 2.001 * numpy.pi / num_particles + ) + np_position[:, 0] *= numpy.random.random_sample((num_particles,)) / 3.0 + 0.2 + np_position[:, 1] = numpy.cos( + numpy.arange(0.0, num_particles) * 2.001 * numpy.pi / num_particles + ) + np_position[:, 1] *= numpy.random.random_sample((num_particles,)) / 3.0 + 0.2 + np_position[:, 2] = 0.0 + np_position[:, 3] = 1.0 + + np_color[:, :] = [1.0, 1.0, 1.0, 1.0] # White particles + + np_velocity[:, 0] = np_position[:, 0] * 2.0 + np_velocity[:, 1] = np_position[:, 1] * 2.0 + np_velocity[:, 2] = 3.0 + np_velocity[:, 3] = numpy.random.random_sample((num_particles,)) + + gl_position = vbo.VBO( + data=np_position, usage=GL_DYNAMIC_DRAW, target=GL_ARRAY_BUFFER + ) gl_position.bind() gl_color = vbo.VBO(data=np_color, usage=GL_DYNAMIC_DRAW, target=GL_ARRAY_BUFFER) gl_color.bind() return (np_position, np_velocity, gl_position, gl_color) + def on_timer(t): glutTimerFunc(t, on_timer, t) glutPostRedisplay() + def on_key(*args): - if args[0] == '\033' or args[0] == 'q': + if args[0] == "\033" or args[0] == "q": sys.exit() + def on_click(button, state, x, y): - mouse_old['x'] = x - mouse_old['y'] = y + mouse_old["x"] = x + mouse_old["y"] = y + def on_mouse_move(x, y): - rotate['x'] += (y - mouse_old['y']) * .2 - rotate['y'] += (x - mouse_old['x']) * .2 + rotate["x"] += (y - mouse_old["y"]) * 0.2 + rotate["y"] += (x - mouse_old["x"]) * 0.2 + + mouse_old["x"] = x + mouse_old["y"] = y - mouse_old['x'] = x - mouse_old['y'] = y def on_display(): - """Render the particles""" + """Render the particles""" # Update or particle positions by calling the OpenCL kernel cl.enqueue_acquire_gl_objects(queue, [cl_gl_position, cl_gl_color]) - kernelargs = (cl_gl_position, cl_gl_color, cl_velocity, cl_start_position, cl_start_velocity, numpy.float32(time_step)) + kernelargs = ( + cl_gl_position, + cl_gl_color, + cl_velocity, + cl_start_position, + cl_start_velocity, + numpy.float32(time_step), + ) program.particle_fountain(queue, (num_particles,), None, *(kernelargs)) cl.enqueue_release_gl_objects(queue, [cl_gl_position, cl_gl_color]) queue.finish() @@ -101,11 +122,11 @@ def on_display(): glLoadIdentity() # Handle mouse transformations - glTranslatef(initial_translate['x'], initial_translate['y'], initial_translate['z']) - glRotatef(rotate['x'], 1, 0, 0) - glRotatef(rotate['y'], 0, 1, 0) #we switched around the axis so make this rotate_z - glTranslatef(translate['x'], translate['y'], translate['z']) - + glTranslatef(initial_translate["x"], initial_translate["y"], initial_translate["z"]) + glRotatef(rotate["x"], 1, 0, 0) + glRotatef(rotate["y"], 0, 1, 0) # we switched around the axis so make this rotate_z + glTranslatef(translate["x"], translate["y"], translate["z"]) + # Render the particles glEnable(GL_POINT_SMOOTH) glPointSize(2) @@ -130,17 +151,25 @@ def on_display(): glutSwapBuffers() + window = glut_window() (np_position, np_velocity, gl_position, gl_color) = initial_buffers(num_particles) platform = cl.get_platforms()[0] -context = cl.Context(properties=[(cl.context_properties.PLATFORM, platform)] + get_gl_sharing_context_properties()) +context = cl.Context( + properties=[(cl.context_properties.PLATFORM, platform)] + + get_gl_sharing_context_properties() +) queue = cl.CommandQueue(context) cl_velocity = cl.Buffer(context, mf.COPY_HOST_PTR, hostbuf=np_velocity) -cl_start_position = cl.Buffer(context, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=np_position) -cl_start_velocity = cl.Buffer(context, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=np_velocity) +cl_start_position = cl.Buffer( + context, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=np_position +) +cl_start_velocity = cl.Buffer( + context, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=np_velocity +) cl_gl_position = cl.GLBuffer(context, mf.READ_WRITE, int(gl_position)) cl_gl_color = cl.GLBuffer(context, mf.READ_WRITE, int(gl_color))