1 # GL interoperability example, by Peter Berrington.
   2 # Draws a rotating teapot, using cuda to invert the RGB value
   3 # each frame
   4 
   5 from OpenGL.GL import *
   6 from OpenGL.GLUT import *
   7 from OpenGL.GLU import *
   8 from OpenGL.GL.ARB.vertex_buffer_object import *
   9 import numpy, sys, time
  10 import pycuda.driver as cuda_driver
  11 import pycuda.gl as cuda_gl
  12 from pycuda.compiler import SourceModule
  13 
  14 #this is all munged together from the CUDA SDK postprocessGL example.
  15 
  16 initial_size = 512,512
  17 current_size = initial_size
  18 animate = True
  19 enable_cuda = True
  20 window = None     # Number of the glut window.
  21 time_of_last_draw = 0.0
  22 time_of_last_titleupdate = 0.0
  23 frames_per_second = 0.0
  24 frame_counter = 0
  25 output_texture = None # pointer to offscreen render target
  26 (source_pbo, dest_pbo, cuda_module, invert,
  27  pycuda_source_pbo, pycuda_dest_pbo) = [None]*6
  28 heading,pitch,bank = [0.0]*3
  29 
  30 def create_PBOs(w,h):
  31     global source_pbo, dest_pbo, pycuda_source_pbo, pycuda_dest_pbo
  32     num_texels = w*h
  33     data = numpy.zeros((num_texels,4),numpy.uint8)
  34     source_pbo = glGenBuffers(1)
  35     glBindBuffer(GL_ARRAY_BUFFER, source_pbo)
  36     glBufferData(GL_ARRAY_BUFFER, data, GL_DYNAMIC_DRAW)
  37     glBindBuffer(GL_ARRAY_BUFFER, 0)
  38     pycuda_source_pbo = cuda_gl.BufferObject(long(source_pbo))
  39     dest_pbo = glGenBuffers(1)
  40     glBindBuffer(GL_ARRAY_BUFFER, dest_pbo)
  41     glBufferData(GL_ARRAY_BUFFER, data, GL_DYNAMIC_DRAW)
  42     glBindBuffer(GL_ARRAY_BUFFER, 0)
  43     pycuda_dest_pbo = cuda_gl.BufferObject(long(dest_pbo))
  44 
  45 def destroy_PBOs():
  46     global source_pbo, dest_pbo, pycuda_source_pbo, pycuda_dest_pbo
  47     for pbo in [source_pbo, dest_pbo]:
  48         glBindBuffer(GL_ARRAY_BUFFER, long(pbo))
  49         glDeleteBuffers(1, long(pbo));
  50         glBindBuffer(GL_ARRAY_BUFFER, 0)
  51     source_pbo,dest_pbo,pycuda_source_pbo,pycuda_dest_pbo = [None]*4
  52 
  53 def create_texture(w,h):
  54     global output_texture
  55     output_texture = glGenTextures(1)
  56     glBindTexture(GL_TEXTURE_2D, output_texture)
  57     # set basic parameters
  58     glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_S, GL_CLAMP_TO_EDGE)
  59     glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_T, GL_CLAMP_TO_EDGE)
  60     glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST)
  61     glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_NEAREST)
  62     # buffer data
  63     glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA,
  64                  w, h, 0, GL_RGBA, GL_UNSIGNED_BYTE, None)
  65 
  66 def destroy_texture():
  67     global output_texture
  68     glDeleteTextures(output_texture);
  69     output_texture = None
  70 
  71 def init_gl():
  72     Width, Height = current_size
  73     glClearColor(0.1, 0.1, 0.5, 1.0)
  74     glDisable(GL_DEPTH_TEST)
  75     glViewport(0, 0, Width, Height)
  76     glMatrixMode(GL_PROJECTION);
  77     glLoadIdentity();
  78     gluPerspective(60.0, Width/float(Height), 0.1, 10.0)
  79     glPolygonMode(GL_FRONT_AND_BACK, GL_FILL)
  80     glEnable(GL_LIGHT0)
  81     red   = ( 1.0, 0.1, 0.1, 1.0 )
  82     white = ( 1.0, 1.0, 1.0, 1.0 )
  83     glMaterialfv(GL_FRONT_AND_BACK, GL_DIFFUSE,  red  )
  84     glMaterialfv(GL_FRONT_AND_BACK, GL_SPECULAR, white)
  85     glMaterialf( GL_FRONT_AND_BACK, GL_SHININESS, 60.0)
  86 
  87 def resize(Width, Height):
  88     global current_size
  89     current_size = Width, Height
  90     glViewport(0, 0, Width, Height)        # Reset The Current Viewport And Perspective Transformation
  91     glMatrixMode(GL_PROJECTION)
  92     glLoadIdentity()
  93     gluPerspective(60.0, Width/float(Height), 0.1, 10.0)
  94 
  95 def do_tick():
  96     global time_of_last_titleupdate, frame_counter, frames_per_second
  97     if ((time.clock () * 1000.0) - time_of_last_titleupdate >= 1000.):
  98         frames_per_second = frame_counter                   # Save The FPS
  99         frame_counter = 0  # Reset The FPS Counter
 100         szTitle = "%d FPS" % (frames_per_second )
 101         glutSetWindowTitle ( szTitle )
 102         time_of_last_titleupdate = time.clock () * 1000.0
 103     frame_counter += 1
 104 
 105 # The function called whenever a key is pressed. Note the use of Python tuples to pass in: (key, x, y)
 106 def keyPressed(*args):
 107     global animate, enable_cuda
 108     # If escape is pressed, kill everything.
 109     if args[0] == '\033':
 110         print 'Closing..'
 111         destroy_PBOs()
 112         destroy_texture()
 113         exit()
 114     elif args[0] == 'a':
 115         print 'toggling animation'
 116         animate = not animate
 117     elif args[0] == 'e':
 118         print 'toggling cuda'
 119         enable_cuda = not enable_cuda
 120 
 121 def idle():
 122     global heading, pitch, bank
 123     if animate:
 124         heading += 0.2
 125         pitch   += 0.6
 126         bank    += 1.0
 127 
 128     glutPostRedisplay()
 129 
 130 def display():
 131     try:
 132         render_scene()
 133         if enable_cuda:
 134             process_image()
 135             display_image()
 136         glutSwapBuffers()
 137     except:
 138         from traceback import print_exc
 139         print_exc()
 140         from os import _exit
 141         _exit(0)
 142 
 143 def process(width, height):
 144     """ Use PyCuda """
 145     grid_dimensions   = (width//16,height//16)
 146 
 147     source_mapping = pycuda_source_pbo.map()
 148     dest_mapping   = pycuda_dest_pbo.map()
 149 
 150     invert.prepared_call(grid_dimensions, (16, 16, 1),
 151             source_mapping.device_ptr(),
 152             dest_mapping.device_ptr())
 153 
 154     cuda_driver.Context.synchronize()
 155 
 156     source_mapping.unmap()
 157     dest_mapping.unmap()
 158 
 159 def process_image():
 160     """ copy image and process using CUDA """
 161     global pycuda_source_pbo,source_pbo,current_size, dest_pbo
 162     image_width, image_height = current_size
 163     assert source_pbo is not None
 164 
 165     # tell cuda we are going to get into these buffers
 166     pycuda_source_pbo.unregister()
 167 
 168     # activate destination buffer
 169     glBindBufferARB(GL_PIXEL_PACK_BUFFER_ARB, long(source_pbo))
 170 
 171     # read data into pbo. note: use BGRA format for optimal performance
 172     import OpenGL.raw.GL as rawgl
 173 
 174     rawgl.glReadPixels(
 175              0,                  #start x
 176              0,                  #start y
 177              image_width,        #end   x
 178              image_height,       #end   y
 179              GL_BGRA,            #format
 180              GL_UNSIGNED_BYTE,   #output type
 181              ctypes.c_void_p(0))
 182 
 183     pycuda_source_pbo = cuda_gl.BufferObject(long(source_pbo))
 184 
 185     # run the Cuda kernel
 186     process(image_width, image_height)
 187     # blit convolved texture onto the screen
 188     # download texture from PBO
 189     glBindBuffer(GL_PIXEL_UNPACK_BUFFER_ARB, long(dest_pbo))
 190     glBindTexture(GL_TEXTURE_2D, output_texture)
 191 
 192     rawgl.glTexSubImage2D(GL_TEXTURE_2D, 0, 0, 0,
 193                     image_width, image_height,
 194                     GL_BGRA, GL_UNSIGNED_BYTE, ctypes.c_void_p(0))
 195 
 196 def display_image():
 197     """ render a screen sized quad """
 198     glDisable(GL_DEPTH_TEST)
 199     glDisable(GL_LIGHTING)
 200     glEnable(GL_TEXTURE_2D)
 201     glTexEnvf(GL_TEXTURE_ENV, GL_TEXTURE_ENV_MODE, GL_REPLACE)
 202     glMatrixMode(GL_PROJECTION)
 203     glPushMatrix()
 204     glLoadIdentity()
 205     glOrtho(-1.0, 1.0, -1.0, 1.0, -1.0, 1.0)
 206     glMatrixMode( GL_MODELVIEW)
 207     glLoadIdentity()
 208     glViewport(0, 0, current_size[0], current_size[1])
 209     glBegin(GL_QUADS)
 210     glTexCoord2f(0.0, 0.0)
 211     glVertex3f(-1.0, -1.0, 0.5)
 212     glTexCoord2f(1.0, 0.0)
 213     glVertex3f(1.0, -1.0, 0.5)
 214     glTexCoord2f(1.0, 1.0)
 215     glVertex3f(1.0, 1.0, 0.5)
 216     glTexCoord2f(0.0, 1.0)
 217     glVertex3f(-1.0, 1.0, 0.5)
 218     glEnd()
 219     glMatrixMode(GL_PROJECTION)
 220     glPopMatrix()
 221     glDisable(GL_TEXTURE_2D)
 222     glBindBuffer(GL_PIXEL_PACK_BUFFER_ARB, 0)
 223     glBindBuffer(GL_PIXEL_UNPACK_BUFFER_ARB, 0)
 224 
 225 
 226 def render_scene():
 227     glClear (GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT)# Clear Screen And Depth Buffer
 228     glMatrixMode(GL_MODELVIEW)
 229     glLoadIdentity ()      # Reset The Modelview Matrix
 230     glTranslatef(0.0, 0.0, -3.0);
 231     glRotatef(heading, 1.0, 0.0, 0.0)
 232     glRotatef(pitch  , 0.0, 1.0, 0.0)
 233     glRotatef(bank   , 0.0, 0.0, 1.0)
 234     glViewport(0, 0, current_size[0],current_size[1])
 235     glEnable(GL_LIGHTING)
 236     glEnable(GL_DEPTH_TEST)
 237     glDepthFunc(GL_LESS)
 238     glutSolidTeapot(1.0)
 239     do_tick()#just for fps display..
 240     return True
 241 
 242 def main():
 243     global window, cuda_module, cuda_gl, cuda_driver, invert
 244     glutInit(sys.argv)
 245     glutInitDisplayMode(GLUT_RGBA | GLUT_DOUBLE | GLUT_ALPHA | GLUT_DEPTH)
 246     glutInitWindowSize(*initial_size)
 247     glutInitWindowPosition(0, 0)
 248     window = glutCreateWindow("PyCuda GL Interop Example")
 249     glutDisplayFunc(display)
 250     glutIdleFunc(idle)
 251     glutReshapeFunc(resize)
 252     glutKeyboardFunc(keyPressed)
 253     glutSpecialFunc(keyPressed)
 254     init_gl()
 255 
 256     # create texture for blitting to screen
 257     create_texture(*initial_size)
 258 
 259     #setup pycuda gl interop
 260     import pycuda.gl.autoinit
 261     import pycuda.gl
 262     cuda_gl = pycuda.gl
 263     cuda_driver = pycuda.driver
 264 
 265     cuda_module = SourceModule("""
 266     __global__ void invert(unsigned char *source, unsigned char *dest)
 267     {
 268       int block_num        = blockIdx.x + blockIdx.y * gridDim.x;
 269       int thread_num       = threadIdx.y * blockDim.x + threadIdx.x;
 270       int threads_in_block = blockDim.x * blockDim.y;
 271       //Since the image is RGBA we multiply the index 4.
 272       //We'll only use the first 3 (RGB) channels though
 273       int idx              = 4 * (threads_in_block * block_num + thread_num);
 274       dest[idx  ] = 255 - source[idx  ];
 275       dest[idx+1] = 255 - source[idx+1];
 276       dest[idx+2] = 255 - source[idx+2];
 277     }
 278     """)
 279     invert = cuda_module.get_function("invert")
 280     # The argument "PP" indicates that the invert function will take two PBOs as arguments
 281     invert.prepare("PP")
 282 
 283     # create source and destination pixel buffer objects for processing
 284     create_PBOs(*initial_size)
 285 
 286     glutMainLoop()
 287 
 288 # Print message to console, and kick off the main to get it rolling.
 289 if __name__ == "__main__":
 290     print "Hit ESC key to quit, 'a' to toggle animation, and 'e' to toggle cuda"
 291     main()

PyCuda/Examples/GlInterop (last edited 2013-09-14 01:14:06 by AndreasKloeckner)