def create_PBO(w, h):
global pbo, pycuda_pbo, rpbo, rpycuda_pbo
num_texels = w*h
array = np.zeros((num_texels, 3),np.float32)
pbo = glGenBuffers(1)
glBindBuffer(GL_ARRAY_BUFFER, pbo)
glBufferData(GL_ARRAY_BUFFER, array, GL_DYNAMIC_DRAW)
glBindBuffer(GL_ARRAY_BUFFER, 0)
pycuda_pbo = cuda_gl.RegisteredBuffer(long(pbo))
rpbo = glGenBuffers(1)
glBindBuffer(GL_ARRAY_BUFFER, rpbo)
glBufferData(GL_ARRAY_BUFFER, array, GL_DYNAMIC_DRAW)
glBindBuffer(GL_ARRAY_BUFFER, 0)
rpycuda_pbo = cuda_gl.RegisteredBuffer(long(rpbo))
def __init__(self, **kwargs):
super().__init__(**kwargs)
self.camera.projection.update(near=0.1, far=10)
# we have to import them here, they require a running context
from pycuda import autoinit
from pycuda.gl import autoinit
import pycuda.gl as cuda_gl
# data
positions = np.random.random((self.N, 3)).astype("f4")
self.pbuffer = self.ctx.buffer(positions)
self.cbuffer = self.ctx.buffer(
reserve=4 * self.N * 4
) # self.N*4 floats (4 bytes per float)
# render program
self.render_prog = self.load_program("file.glsl")
# vao
self.vao = VAO(mode=mgl.POINTS)
self.vao.buffer(self.pbuffer, "3f", ["in_position"])
self.vao.buffer(self.cbuffer, "4f", ["in_color"])
# pycuda stuff
self.mod = SourceModule(
"""
__global__ void color_them(float4 *dest, int n, float time)
{
const int tid = blockIdx.x * blockDim.x + threadIdx.x;
if (tid >= n) {return;}
float v = (float)tid / (float)n;
dest[tid] = make_float4(__cosf(time*3.), __sinf(time*2.), __sinf(time), 1.);
}
""",
keep=True,
cache_dir="./cache",
)
self.buffer_cu = cuda_gl.RegisteredBuffer(self.cbuffer._glo)
self.color_them = self.mod.get_function("color_them")
def createPBO():
global gl_Tex
#Create texture which we use to display the result and bind to gl_Tex
glEnable(GL_TEXTURE_2D)
gl_Tex = glGenTextures(1)
glBindTexture(GL_TEXTURE_2D, gl_Tex)
glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_S, GL_CLAMP)
glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_T, GL_CLAMP)
glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_LINEAR)
glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_LINEAR)
glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA8, nWidth, nHeight, 0, GL_RGBA,
GL_UNSIGNED_BYTE, None)
#print "Texture Created"
# Create pixel buffer object and bind to gl_PBO. We store the data we want to
# plot in memory on the graphics card - in a "pixel buffer". We can then
# copy this to the texture defined above and send it to the screen
global gl_PBO, cuda_POB
gl_PBO = glGenBuffers(1)
glBindBuffer(GL_PIXEL_UNPACK_BUFFER, gl_PBO)
glBufferData(GL_PIXEL_UNPACK_BUFFER, nWidth * 4 * nHeight, None,
GL_STREAM_COPY)
cuda_POB = cuda_gl.RegisteredBuffer(long(gl_PBO))
def __init__(self,
K_ir,
K_rgb,
T_rel,
side=256,
units_per_voxel=7.8125,
mu=200.0):
self.K_ir = K_ir
self.K_rgb = K_rgb
self.T_rel = T_rel
self.side = side
self.units_per_voxel = units_per_voxel
self.mu = mu
# Process the module.
self.module = SourceModule(cuda_source)
self.update_reconstruction = self.module.get_function(
"update_reconstruction")
self.compute_depth = self.module.get_function("compute_depth")
self.compute_smooth_depth = self.module.get_function(
"compute_smooth_depth")
self.compute_tracking_matrices = self.module.get_function(
"compute_tracking_matrices")
self.measure = self.module.get_function("measure")
self.raycast = self.module.get_function("raycast")
self.depth_texture = self.module.get_texref("depth_texture")
self.smooth_depth_texture = self.module.get_texref(
"smooth_depth_texture")
self.F_texture = self.module.get_texref("F_texture")
self.depth_texture.set_filter_mode(drv.filter_mode.POINT)
self.smooth_depth_texture.set_filter_mode(drv.filter_mode.POINT)
self.F_texture.set_filter_mode(drv.filter_mode.LINEAR)
#print self.update_reconstruction.shared_size_bytes
#print self.update_reconstruction.num_regs
# Set the global constant matrices.
mod = self.module # For short access.
K, _ = mod.get_global("K")
invK, _ = mod.get_global("invK")
drv.memcpy_htod(K, np.float32(kc.K_ir.ravel()))
drv.memcpy_htod(invK, np.float32(la.inv(kc.K_ir).ravel()))
# Reserve GPU variables.
print drv.mem_get_info()
self.buffers = {}
self.T_gk = np.array(
[[1, 0, 0, 0], [0, 1, 0, 0], [0, 0, 1, -1000], [0, 0, 0, 1]],
dtype=np.float32)
self.T_gk_gpu = gpuarray.to_gpu(self.T_gk[:3])
self.Tgk1_k_gpu = gpuarray.to_gpu(np.eye(4, dtype=np.float32)[:3])
self.F_gpu = gpuarray.zeros((side, ) * 3, dtype=np.float32) - 1000
self.W_gpu = gpuarray.zeros((side, ) * 3, dtype=np.float32)
self.mask_gpu = gpuarray.zeros(480 * 640, dtype=np.float32)
self.smooth_depth_gpu = gpuarray.zeros((480, 640), dtype=np.float32)
for buffer_type in ['measure', 'raycast']:
buffers = gl.glGenBuffers(2)
for b in buffers:
gl.glBindBuffer(gl.GL_ARRAY_BUFFER, b)
gl.glBufferData(gl.GL_ARRAY_BUFFER, 640 * 480 * 12, None,
gl.GL_DYNAMIC_COPY)
buffers = map(lambda x: cudagl.RegisteredBuffer(int(x)), buffers)
self.buffers[buffer_type] = buffers
# Tracking data.
self.AA_gpu = gpuarray.zeros((640 * 480, 21), dtype=np.float32)
self.Ab_gpu = gpuarray.zeros((640 * 480, 6), dtype=np.float32)
self.omega_gpu = gpuarray.zeros(640 * 480, dtype=np.float32)
self.AA = np.empty(21, dtype=np.float32)
self.Ab = np.empty(6, dtype=np.float32)
self.active_tracking = False
print drv.mem_get_info()
self._prepare_F_texture()