Pedro Leite

words – an empty shell, a lack of meaning

OpenStereo Plugin using CUDA

Of course I wouldn’t let OpenStereo apart from my CUDA studies. I’ve implemented an OpenStereo plugin that uses CUDA to do the fusion process. Below is the preliminary kernel code:

#include <stdlib.h>
#include <stdio.h>
#include <math.h>


#define BLOCK_SIZE 16


__global__ void fuse_images(float* left_tex, float* right_tex, float* stereo_tex,
float* matrix_left, float* matrix_right, int width, int height) {
int ix = __mul24(blockIdx.x, BLOCK_SIZE) + threadIdx.x;
int iy = __mul24(blockIdx.y, BLOCK_SIZE) + threadIdx.y;
int i = 0, j = 0, k = 0;
float c = 0.0f;

if (ix >= width || iy >= height)
return;


// rows
for (; i < 3; i++) {
// columns
k = (ix + __mul24(iy, width)) * 3;
for (j = 0; j < 3; j++) {
c += matrix_left[j + __mul24(3, i)] * left_tex[k + j];
c += matrix_right[j + __mul24(3, i)] * right_tex[k + j];
}
stereo_tex[k + i] = c;
c = 0.0f;
}
}

That function runs on the device (GPU) while the host side does the OpenGL stuff, besides calling the kernel code. A call to that function is done this way:

// executes the kernel
fuse_images<<< grid, threads >>>(cuda_left_tex, cuda_right_tex, cuda_stereo_tex,
getCudaMatrixFromMethod(method, isLeft), getCudaMatrixFromMethod(method, !isLeft), width, height);

I need to test this plugin in a CUDA device, since its emulation is working almost perfectly.
P.S.: Yeah, posting source code in wordpress sucks =/

Advertisements

October 30, 2007 Posted by | cuda, development, open source, opengl, openstereo, stereoscopy | 1 Comment

OpenStereo Demonstration Video

This is a short demonstration video about the OpenStereo library in action.

http://www.youtube.com/watch?v=g-xrMplAE74

P.S.: I know that time for titles was short… My bad!

July 26, 2007 Posted by | nehe, open source, openstereo, screencast, stereoscopy | Leave a comment

NeHe’s Lessons

Suddenly this will of port every single lesson at NeHe’s site to OpenStereo (compiling under Linux and Windows) took me. Comparing the results obtained against the original ones can lead me to fix some undiscovered bugs in OpenStereo. So, let the porting marathon begin!

P.S.: Obviously, I’ll contact the GameDev.net team to release the ported source code.

July 26, 2007 Posted by | development, linux, nehe, open source, opengl, openstereo, stereoscopy, windows | Leave a comment

OpenStereo 0.1.0

Today, I decided to release the first version of OpenStereo under the LGPL v2 license (maybe I’ll consider the LGPL v3 after reading it ;). OpenStereo was my graduation thesis, or graduation conclusion work, if this suits you well and now is my open source project. It was presented on October 2006, but only now I decided to release it, even it meant to be an open source project.

Starting from 0.1.0, the OpenStereo library provides multi-platform support for development of real time stereoscopic applications. It’s based on a plugin subsystem, which abstracts the underlying graphics system. At the current time, only OpenGL applications are supported, but one plugin subsystem’s advantage is that a DirectX plugin can be written, enabling DirectX applications to be developed with and/or ported to OpenStereo. Of course, this is a future work. What I have on 0.1.0 is, though, a fragile not so well defined architecture, once it appears more like a plugin subsystem, than a stereoscopic library. Indeed, the libSL only does a good plugin management, but as time goes by, such architecture will evolve into a more robust one. Fortunately, it works! Applications can be more immersible than ever, all of this without requiring the user to buy expensive graphics cards or equipment like HMD glasses (indeed it’s necessary that the user have, at least, a red-blue anaglyph glass).

For a quick preview of what OpenStereo can do, let’s take the NeHe’s Lesson 26. This lesson teaches how to create “EXTREMELLY realistic reflections”, as the text states. The following screenshot shows the application running in 3D mode, the way it was meant to be:

NeHe's Lesson 26

First of all, some modifications were made on the code, so the application code can be compiled on both Windows and Linux environments. OpenStereo was embedded on the code to gather the function that draws the scene, thus allowing left and right pairs image generation (OpenStereo uses the anaglyph method to generate stereoscopic output). Camera values were correctly configured, including support to update eye separation and anaglyph method values with keyboard. After all these modifications (simple ones), the final result can be seen on the following screenshot:

NeHe's Lesson 26 ported to OpenStereo

Any application that uses OpenStereo will suffer from a performance penalty, mainly because OpenStereo needs to 1) render the scene twice, generating the left and right images and 2) fuse and present them to the screen. The fusion part generally depends on the selected plugin (this might be OpenStereo’s job and not plugin’s one).

The OpenStereo source code comes along with three plugins: one that fuses the images at the CPU level, gathering left and right images through glReadPixels; one that fuses images using the Cg Toolkit library (this dependency will be soon removed, using only OpenGL 2.0 and the ARB extensions to apply a pixel shader) at the GPU level, where the left and right images are gathered through glTexCopySubImage2D; and finally one that uses the same fusion process defined by the previous plugin, only that the scene is rendered directly into a framebuffer object, through the framebuffer object extension.

Some work must be done to ensure that the architecture will generate the stereoscopic output at the library level, not at the plugin level. Also, new plugins to DirectX (or this could be done on the architecture?) must be developed, providing support for a wide range of applications and platforms (hardware and software).

The OpenStereo source code will soon (this week) be available at OpenStereo’s Google Code project site.

July 26, 2007 Posted by | directx, linux, nehe, open source, opengl, openstereo, portable, release, screenshot, stereoscopy, windows | Leave a comment