Numba and OpenGL Interop

17 views
Skip to first unread message

Chris Uchytil

unread,
Aug 15, 2016, 8:26:35 PM8/15/16
to Numba Public Discussion - Public
Apologies for all the questions asked, I've almost finished the project so this should be the last one. Does Numba support OpenGL Interoperability? I am having trouble figuring out how to, or if it is possible, to convert these two functions from C++ to Python. I get all the OpenGL stuff, just not sure how to establish PBO's/Textures for access in OpenGL. I found an example for PyCUDA by Peter Berrington (https://wiki.tiker.net/PyCuda/Examples/GlInterop) that has helped me a lot but I can't find anything for Numba.

// texture and pixel objects
GLuint pbo = 0;
GLuint tex = 0
struct cudaGraphicsResource *cuda_pbo_resource;

void render() {
unchar4 *d_out = 0;
cudaGrarphicsMapResources(1, &cuda_pbo_resource, 0);
cudaGraphicsResourceGetMappedPointer((void**)&d_out, NULL,
cuda_pbo_resource);
kernelLauncher(d_out, W, H, loc);
cudaGraphicsUnmapResources(1, &cuda_pbo_resource, 0);
}

void initPixelBuffer() {
glGenBuffers(1, &pbo);
glBindBuffer(GL_PIXEL_UNPACK_BUFFER, pbo);
glBufferData(GL_PIXEL_UNPACK_BUFFER, 4*W*H*sizeof(GLubyte), 0,
GL_STREAM_DRAW);
glGenTextures(1, &tex);
glBindTexture(GL_TEXTURE_2D, tex);
glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST);
cudaGraphicsGLRegisterBuffer(&cuda_pbo_resource, pbo,
cudaGraphicsMapFlagsWriteDiscard);
}

Stanley Seibert

unread,
Aug 16, 2016, 11:31:15 AM8/16/16
to Numba Public Discussion - Public
Numba has no mechanism for OpenGL interop, so I think PyCUDA is probably better for your application.

--
You received this message because you are subscribed to the Google Groups "Numba Public Discussion - Public" group.
To unsubscribe from this group and stop receiving emails from it, send an email to numba-users+unsubscribe@continuum.io.
To post to this group, send email to numba...@continuum.io.
To view this discussion on the web visit https://groups.google.com/a/continuum.io/d/msgid/numba-users/add4eb0b-89f5-4cb6-8e9a-199f094db104%40continuum.io.
For more options, visit https://groups.google.com/a/continuum.io/d/optout.

Chris Uchytil

unread,
Aug 16, 2016, 11:40:06 AM8/16/16
to numba...@continuum.io

Well shoot. Ok


To view this discussion on the web visit https://groups.google.com/a/continuum.io/d/msgid/numba-users/CADv3RKS3w2pXnwnjQiNYfa_98aVRAzUuXCVZ4yD4vbqFRUVWcg%40mail.gmail.com.

Siu Kwan Lam

unread,
Aug 16, 2016, 2:12:48 PM8/16/16
to numba...@continuum.io
There's a hack if you don't mind.  You can make numba, pycuda and opengl work together.  I have modified the pycuda interop sample here: https://gist.github.com/sklam/2ff89e40721d1f1a007449f02aee3990

Key modifications: 


2) Write the kernel: https://gist.github.com/sklam/2ff89e40721d1f1a007449f02aee3990#file-glinterop-py-L321-L331.  I have to force compilation here to make numba own the cuda context.

3) Call the kernel with device array from external memory: https://gist.github.com/sklam/2ff89e40721d1f1a007449f02aee3990#file-glinterop-py-L178-L191


Well shoot. Ok


To unsubscribe from this group and stop receiving emails from it, send an email to numba-users...@continuum.io.

--
You received this message because you are subscribed to the Google Groups "Numba Public Discussion - Public" group.
To unsubscribe from this group and stop receiving emails from it, send an email to numba-users...@continuum.io.

--
You received this message because you are subscribed to the Google Groups "Numba Public Discussion - Public" group.
To unsubscribe from this group and stop receiving emails from it, send an email to numba-users...@continuum.io.

To post to this group, send email to numba...@continuum.io.
--
Siu Kwan Lam
Software Engineer
Continuum Analytics

Chris Uchytil

unread,
Aug 16, 2016, 2:16:03 PM8/16/16
to numba...@continuum.io

Wow, this is really cool. I'll mess around with the code later today. Do you know if it takes a hit to calculate speed in any significant way?


Well shoot. Ok


To unsubscribe from this group and stop receiving emails from it, send an email to numba-users+unsubscribe@continuum.io.

--
You received this message because you are subscribed to the Google Groups "Numba Public Discussion - Public" group.
To unsubscribe from this group and stop receiving emails from it, send an email to numba-users+unsubscribe@continuum.io.

--
You received this message because you are subscribed to the Google Groups "Numba Public Discussion - Public" group.
To unsubscribe from this group and stop receiving emails from it, send an email to numba-users+unsubscribe@continuum.io.
--
Siu Kwan Lam
Software Engineer
Continuum Analytics

--
You received this message because you are subscribed to the Google Groups "Numba Public Discussion - Public" group.
To unsubscribe from this group and stop receiving emails from it, send an email to numba-users+unsubscribe@continuum.io.

To post to this group, send email to numba...@continuum.io.
To view this discussion on the web visit https://groups.google.com/a/continuum.io/d/msgid/numba-users/CAHjo7%3DV54JXX2iRV8X2mH%2BXJyu-Wgzf8O1oG6c%2BuwC4zRv-%2BYw%40mail.gmail.com.

Siu Kwan Lam

unread,
Aug 16, 2016, 2:22:05 PM8/16/16
to numba...@continuum.io
I got the same framerate, 21fps, on my macbook pro.  So, no noticeable difference.


Well shoot. Ok


To unsubscribe from this group and stop receiving emails from it, send an email to numba-users...@continuum.io.

--
You received this message because you are subscribed to the Google Groups "Numba Public Discussion - Public" group.
To unsubscribe from this group and stop receiving emails from it, send an email to numba-users...@continuum.io.

--
You received this message because you are subscribed to the Google Groups "Numba Public Discussion - Public" group.
To unsubscribe from this group and stop receiving emails from it, send an email to numba-users...@continuum.io.
--
Siu Kwan Lam
Software Engineer
Continuum Analytics

--
You received this message because you are subscribed to the Google Groups "Numba Public Discussion - Public" group.
To unsubscribe from this group and stop receiving emails from it, send an email to numba-users...@continuum.io.

To post to this group, send email to numba...@continuum.io.
To view this discussion on the web visit https://groups.google.com/a/continuum.io/d/msgid/numba-users/CAHjo7%3DV54JXX2iRV8X2mH%2BXJyu-Wgzf8O1oG6c%2BuwC4zRv-%2BYw%40mail.gmail.com.

For more options, visit https://groups.google.com/a/continuum.io/d/optout.

--
You received this message because you are subscribed to the Google Groups "Numba Public Discussion - Public" group.
To unsubscribe from this group and stop receiving emails from it, send an email to numba-users...@continuum.io.

To post to this group, send email to numba...@continuum.io.
To view this discussion on the web visit https://groups.google.com/a/continuum.io/d/msgid/numba-users/CAAQHW3kEQf_y56q47Jqw0V6M68W5pevEP3y_qLVe1MsAC0aHmw%40mail.gmail.com.

For more options, visit https://groups.google.com/a/continuum.io/d/optout.

Chris Uchytil

unread,
Aug 16, 2016, 2:32:30 PM8/16/16
to Numba Public Discussion - Public
I can't express how cool this is. Thank you SOO much.

Chris Uchytil

unread,
Aug 16, 2016, 5:50:24 PM8/16/16
to Numba Public Discussion - Public
Hey Siu, if you have time let me know if you can look this over. In the link is the C++ code which I am trying to copy and Python code. The Python code is close to being complete but in the main function there are two things that I don't think are working. There is an issue with the PBO mapping. It is telling me I cannot map an None type. Thanks so much.

Chris Uchytil

unread,
Aug 17, 2016, 10:03:01 PM8/17/16
to Numba Public Discussion - Public
So I've made some progress. On both of my attempts I am getting an error that states

Failed at nopython (nopython mode backend)
Internal error:
Exception: Require NRT

So it is failing in the Kernel. I think it has something to do with the array I am passing in but I'm not 100% sure.

Siu Kwan Lam

unread,
Aug 18, 2016, 12:26:24 PM8/18/16
to Numba Public Discussion - Public
That's likely due to array operations that require dynamic allocation.  Check the ndim of the arrays.  Providing the function signature will help ensuring the argument types.

--
You received this message because you are subscribed to the Google Groups "Numba Public Discussion - Public" group.
To unsubscribe from this group and stop receiving emails from it, send an email to numba-users...@continuum.io.
To post to this group, send email to numba...@continuum.io.
To view this discussion on the web visit https://groups.google.com/a/continuum.io/d/msgid/numba-users/0bdcd1ab-0c37-4ddf-bf6f-2eea4449e0f4%40continuum.io.

For more options, visit https://groups.google.com/a/continuum.io/d/optout.

Chris Uchytil

unread,
Aug 22, 2016, 2:32:50 PM8/22/16
to Numba Public Discussion - Public
Here is my code so far. I think I am very close but it am getting a new error in the kernel.

from OpenGL.GL import *
from OpenGL.GLUT import *
from OpenGL.GLU import *
from OpenGL.GL.ARB.vertex_buffer_object import *
from OpenGL.GL.ARB.pixel_buffer_object import *

import numpy as np
import sys
import pycuda.gl as cuda_gl
import pycuda.driver as cuda_driver

import math
from numba import jit, cuda as nbcuda

W = 600
H = 600
loc = [W/2, H/2]
DELTA = 5
dragMode = False

TX = 32
TY = 32

pbo, tex, pycuda_pbo, distanceKernel = [None] * 4

class ExternalMemory(object):
"""
Provide an externally managed memory.

Interface requirement: __cuda_memory__, device_ctypes_pointer, _cuda_memize_
"""
__cuda_memory__ = True

def __init__(self, ptr, size):
self.device_ctypes_pointer = ctypes.c_void_p(ptr)
self._cuda_memsize_ = size

def render(): #Process
pbo_mapping = pycuda_pbo.map()

source_ptr = ExternalMemory(pbo_mapping.device_ptr(), W*H * 4) #Change shape to W*H when dtype=cl_array.vec.uchar4, remove 4

d_out = nbcuda.devicearray.DeviceNDArray(shape = W*H * 4,
strides = (1,),
dtype = np.dtype('uint8'), #change to dtype=cl_array.vec.uchar4 when supported and change kernel function
gpu_data = source_ptr)

blockSize = (TX, TY)
gridSize = ((W + TX - 1)/TX, (H + TY - 1)/TY)
distanceKernel[gridSize, blockSize](d_out, W, H, loc)

cuda_driver.Context.synchronize()

pbo_mapping.unmap()

def process(): #Process Image
global pycuda_pbo, pbo

assert pbo is not None

pycuda_pbo.unregister()

glBindBufferARB(GL_PIXEL_PACK_BUFFER_ARB, long(pbo))

glReadPixels(0,                #start x
0,                #start y
W,                #end   x
H,                #end   y
GL_RGBA,          #format
GL_UNSIGNED_BYTE, #output type
ctypes.c_void_p(0))

pycuda_pbo = cuda_gl.BufferObject(long(pbo))

render()

glBindTexture(GL_TEXTURE_2D, tex)

def drawTexture():
glEnable(GL_TEXTURE_2D)
glBegin(GL_TEXTURE_2D)
glBegin(GL_QUADS)
glTexCoord2f(0.0, 0.0)
glVertex2f(0,0)
glTexCoord2f(0.0, 1.0)
glVertex2f(0,H)
glTexCoord2f(1.0, 1.0)
glVertex2f(W,H)
glTexCoord2f(1.0, 0.0)
glVertex2f(W,0)
glEnd()
glDisable(GL_TEXTURE_2D)

def display():
process()
drawTexture()
glutSwapBuffers()

def create_PBO():
global pbo, pycuda_pbo

data = np.zeros((W*H,4), np.uint8)

pbo = glGenBuffers(1)
glBindBuffer(GL_ARRAY_BUFFER, pbo)  #GL_PIXEL_UNPACK_BUFFER in C++ code
glBufferData(GL_ARRAY_BUFFER, data, GL_DYNAMIC_DRAW)  #GL_STREAM_DRAW in C++ code
glBindBuffer(GL_ARRAY_BUFFER, 0)
pycuda_pbo = cuda_gl.BufferObject(long(pbo))

def create_texture():
global tex
tex = glGenTextures(1)
glBindTexture(GL_TEXTURE_2D, tex)
glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST)
glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA, W, H, 0, GL_RGBA, GL_UNSIGNED_BYTE, None)

def exitfunc():
glBindBuffer(GL_ARRAY_BUFFER, long(pbo))
glDeleteBuffers(1, long(pbo));
glBindBuffer(GL_ARRAY_BUFFER, 0)
pbo = None

glDeleteTextures(tex);
tex = None

def keyboard(key, x, y):
if key == '\033': #\033 is escape key
exit()
elif key == 'a':
dragMode = not dragMode
elif key == 27:
exit()
glutPostRedisplay()

def mouseMove(x, y):
if dragMode == True:
loc[0] = x
loc[1] = y
glutPostRedisplay()

def mouseDrag(x, y):
if dragMode == False:
loc[0] = x
loc[1] = y
glutPostRedisplay()

def handleSpecialKeypress(key, x, y):
if key == GLUT_KEY_LEFT:
loc[0] -= DELTA
if key == GLUT_KEY_RIGHT:
loc[0] += DELTA
if key == GLUT_KEY_UP:
loc[1] -= DELTA
if key == GLUT_KEY_DOWN:
loc[1] += DELTA
glutPostRedisplay()

def printInstructions():
print "flashlight instructions"
print "a: toggle mouse tracking mode"
print "arrow keys: move ref location"
print "esc: close graphics window"


def main():
global cuda_gl, cuda_driver, distanceKernel

printInstructions();
glutInit(sys.argv)
glutInitDisplayMode(GLUT_RGBA | GLUT_DOUBLE)
glutCreateWindow("flashlight: distance image display app")
gluOrtho2D(0, W, H, 0)
glutDisplayFunc(display)
glutKeyboardFunc(keyboard)
glutSpecialFunc(handleSpecialKeypress)
glutPassiveMotionFunc(mouseMove)
glutMotionFunc(mouseDrag)

create_texture()

#Sets up GL interop
import pycuda.gl.autoinit
import pycuda.gl
cuda_gl = pycuda.gl
cuda_driver = pycuda.driver

@nbcuda.jit(device = True)
def clip(n):
if n > 255:
n = 255
elif n < 0:
n = 0
return n

@nbcuda.jit("(uint8[::1], int16, int16, float32[::1])")
def distanceKernel(d_out, w, h, pos):
c = nbcuda.blockIdx.x*nbcuda.blockDim.x + nbcuda.threadIdx.x
r = nbcuda.blockIdx.y*nbcuda.blockDim.y + nbcuda.threadIdx.y
i = (r*w + c) * 4  #Remove 4 later

if c >= w or r >= h:
return

d = math.sqrt((c - pos[0]) * (c - pos[0]) + (r - pos[1]) * (r - pos[1]))

intensity = clip(255 - int(d))

d_out[i] = intensity    #d_out[i][0] = intensity
d_out[i+1] = intensity  #d_out[i][1] = intensity
d_out[i+2] = 0          #d_out[i][2] = 0
d_out[i+3] = 255        #d_out[i][3] = 255

create_PBO()

glutMainLoop()
atexit(exitfunc)

if __name__ == "__main__":
main()
Message has been deleted

Chris Uchytil

unread,
Aug 31, 2016, 2:49:48 PM8/31/16
to Numba Public Discussion - Public
For reference for anyone else who finds this post. Here is the working code with Numba and OpenGL interop functionality.

from OpenGL.GL import *
from OpenGL.GLUT import *
from OpenGL.GLU import *
from OpenGL.GL.ARB.vertex_buffer_object import *
from OpenGL.GL.ARB.pixel_buffer_object import *
 
import sys, numpy as np
import pycuda.gl as cuda_gl
import pycuda.driver as cuda_driver

#from interactions import *

import math
from numba import jit, cuda as nbcuda
import ctypes

W = 600
H = 600
loc = np.array([W/2, H/2], dtype = 'float32')
DELTA = 5
dragMode = False

TX = 32
TY = 32

pbo, tex, cuda_pbo_resource, distanceKernel = [None] * 4

class ExternalMemory(object):
"""
Provide an externally managed memory.

Interface requirement: __cuda_memory__, device_ctypes_pointer, _cuda_memize_
"""
__cuda_memory__ = True

def __init__(self, ptr, size):
self.device_ctypes_pointer = ctypes.c_void_p(ptr)
self._cuda_memsize_ = size

def render():
global cuda_pbo_resource, pbo, loc
assert pbo is not None

cuda_pbo_resource.unregister()

cuda_pbo_resource = cuda_gl.BufferObject(long(pbo))

pbo_mapping = cuda_pbo_resource.map()

source_ptr = ExternalMemory(pbo_mapping.device_ptr(), W*H*4) #Change shape to W*H when dtype=cl_array.vec.uchar4, remove 4

d_out = nbcuda.devicearray.DeviceNDArray(shape = W*H*4,
strides = (1,),
dtype = np.dtype('uint8'), #change to dtype=cl_array.vec.uchar4 when supported and change kernel function
gpu_data = source_ptr)

 #(   gridSize    ) (blockSize)
distanceKernel[((W + TX - 1)/TX, (H + TY - 1)/TY), (TX, TY)](d_out, W, H, loc)

cuda_driver.Context.synchronize()

pbo_mapping.unmap()

def drawTexture():
glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA, W, H, 0, GL_RGBA, GL_UNSIGNED_BYTE, None)
glEnable(GL_TEXTURE_2D)
glBegin(GL_QUADS)
glTexCoord2f(0.0, 0.0); glVertex2f(0,0)
glTexCoord2f(0.0, 1.0); glVertex2f(0,H)
glTexCoord2f(1.0, 1.0); glVertex2f(W,H)
glTexCoord2f(1.0, 0.0); glVertex2f(W,0)
glEnd()
glDisable(GL_TEXTURE_2D)

def display():
render()
drawTexture()
glutSwapBuffers()

def initPixelBuffer():
global pbo, cuda_pbo_resource, tex

data = np.zeros((W*H,4), dtype = 'uint8')

pbo = glGenBuffers(1)
glBindBuffer(GL_PIXEL_UNPACK_BUFFER, pbo)
glBufferData(GL_PIXEL_UNPACK_BUFFER, data, GL_STREAM_DRAW)
cuda_pbo_resource = cuda_gl.BufferObject(long(pbo)) #cudaGraphicsGLRegisterBuffer

tex = glGenTextures(1)
glBindTexture(GL_TEXTURE_2D, tex)
glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST)

def main():
global cuda_gl, cuda_driver, distanceKernel

printInstructions();
initGLUT()
gluOrtho2D(0, W, H, 0)
glutDisplayFunc(display)
glutKeyboardFunc(keyboard)
glutSpecialFunc(handleSpecialKeypress)
glutPassiveMotionFunc(mouseMove)
glutMotionFunc(mouseDrag)
#Sets up GL interop
import pycuda.gl.autoinit
import pycuda.gl
cuda_gl = pycuda.gl
cuda_driver = pycuda.driver

@nbcuda.jit(device = True)
def clip(n):
if n > 255:
n = 255
elif n < 0:
n = 0
return n

@nbcuda.jit("(uint8[::1], int32, int32, float32[::1])")
def distanceKernel(d_out, w, h, pos):
c = nbcuda.blockIdx.x*nbcuda.blockDim.x + nbcuda.threadIdx.x
r = nbcuda.blockIdx.y*nbcuda.blockDim.y + nbcuda.threadIdx.y
i = (r*w + c) * 4  #Remove 4 later

if c >= w or r >= h:
return

d = math.sqrt((c - pos[0]) * (c - pos[0]) + (r - pos[1]) * (r - pos[1]))

intensity = clip(255 - d)

d_out[i] = intensity #d_out[i][0] = intensity
d_out[i+1] = intensity  #d_out[i][1] = intensity
d_out[i+2] = 0   #d_out[i][2] = 0
d_out[i+3] = 255 #d_out[i][3] = 255

initPixelBuffer()

glutMainLoop()
atexit(exitfunc)

def initGLUT():
glutInit(sys.argv)
glutInitDisplayMode(GLUT_RGBA | GLUT_DOUBLE)
glutCreateWindow("flashlight: distance image display app")

def exitfunc():
pbo_mapping.unmap()

glDeleteBuffers(1, long(pbo))
pbo = None

glDeleteTextures(tex)
tex = None

def keyboard(key, x, y):
global dragMode
if key == '\033': #\033 is escape key
exit()
elif key == 'a':
dragMode = not dragMode
elif key == 27:
exit()
glutPostRedisplay()

def mouseMove(x, y):
global dragMode, loc
if dragMode == True:
loc[0] = x
loc[1] = y
glutPostRedisplay()

def mouseDrag(x, y):
global dragMode, loc
if dragMode == False:
loc[0] = x
loc[1] = y
glutPostRedisplay()

def handleSpecialKeypress(key, x, y):
#global loc
if key == GLUT_KEY_LEFT:
loc[0] -= DELTA
if key == GLUT_KEY_RIGHT:
loc[0] += DELTA
if key == GLUT_KEY_UP:
loc[1] -= DELTA
if key == GLUT_KEY_DOWN:
loc[1] += DELTA
glutPostRedisplay()

def printInstructions():
print "flashlight instructions"
print "a: toggle mouse tracking mode"
print "arrow keys: move ref location"
print "esc: close graphics window"

Stanley Seibert

unread,
Aug 31, 2016, 2:51:46 PM8/31/16
to Numba Public Discussion - Public
Wow, I'm impressed you got this all working!  Would you be OK if we include some version of this in our documentation as an example?

--
You received this message because you are subscribed to the Google Groups "Numba Public Discussion - Public" group.
To unsubscribe from this group and stop receiving emails from it, send an email to numba-users+unsubscribe@continuum.io.

To post to this group, send email to numba...@continuum.io.
To view this discussion on the web visit https://groups.google.com/a/continuum.io/d/msgid/numba-users/1fa80025-1c11-47d6-9ce4-00afb4396294%40continuum.io.

Chris Uchytil

unread,
Aug 31, 2016, 2:57:27 PM8/31/16
to Numba Public Discussion - Public
Go right ahead. As a side note, something you might want to include in the documentation, I didn't include this in the code above because I am not sure if it is just my computer. Under the functions for mouseMove and mouseDrag, I multiplied the x and y values by 2 because for some reason, even though the window size is 600,600, my mouse coordinates were half of that, so the bottom right corner was 300,300 instead of 600,600.

Chris Uchytil

unread,
Aug 31, 2016, 3:03:24 PM8/31/16
to Numba Public Discussion - Public
Also Also, for further reference. If you don't want to type out your entire kernel code in your main function you can put it in a separate file and import the functions as long as the import statement is within the main function where the kernel would go. Something like this.

def main():
global cuda_gl, cuda_driver, distanceKernel_var

printInstructions();
initGLUT()
gluOrtho2D(0, W, H, 0)
glutDisplayFunc(display)
glutKeyboardFunc(keyboard)
glutSpecialFunc(handleSpecialKeypress)
glutPassiveMotionFunc(mouseMove)
glutMotionFunc(mouseDrag)
#Sets up GL interop
import pycuda.gl.autoinit
import pycuda.gl
cuda_gl = pycuda.gl
cuda_driver = pycuda.driver

  from kernel import *

distanceKernel_var = distanceKernel #where distanceKernel is the name of the kernel function in the kernel.py file
initPixelBuffer()

glutMainLoop()
atexit(exitfunc)

Reply all
Reply to author
Forward
0 new messages