Chouquette's blog

VLMC 0.1.0 released !

Hugo — Wed, 14 Apr 2010 09:47:00 +0100

Good news everyone !

The first pre-pre-...-pre-alpha of VLMC just has been released for win32. Please keep in mind that this is realy just about testing, and is not intented to be fully functionnal. Many bugs can occur, even crash, and using it may lead to an early end of the world.

If you're willing to test it though, you can find it there : ftp://ftp.videolan.org/pub/videolan...

Please, and I really meen it : please keep in mind that we're currently only two developpers, and that vlmc is developped in our free time. Before yelling at us because vlmc didn't work the way you expected it to, think twice (and maybe more). Instead, if you find a bug, even the smallest one, if you feel something needs improvements, if you think of a nice feature, or if you're willing to help us, please do ! Open up tickets (www.vlmc.org), send us patches, come and discuss with us on IRC :)

Have as many fun as possible testing VLMC !

VLMC: More infos on release process

Hugo — Wed, 17 Mar 2010 22:10:00 +0000

As this seems to bother many people, here are some update on the VLMC release :)

As VLMC was showed on the Solutions Linux expo (in Paris), we had time to chat with some people interested in VLMC. It was great having feedback, and allow us to spot many bugs, that have mostly been fixed :)

We decided today that it was pointless releasing VLMC for linux until VLC 1.1 has been release. So we will wait until then. We're going to release for win32 soon, as we have fixed the last specific bug tonight (unless we found another in the next days) For those of you who follow the git, you may have noticed that a 0.0.1-bugfix branch has been created a month ago. This branch will be droped, as it became pointless as time goes. In the last weeks, many parts of the code has been rewritten, and since we didn't release, and had time to fix the bug introduced by this big refactoring, there's no point keeping this old code up. Here's some of the things that has been fixed/rewritten :

Settings widget are now autogenerated.
Settings are now categorized, and variables must be created before using.
Media list has been unified, and useless classes have been removed
Synchronisation problem are fixed.
Memory usage has been reduced.
Some GUI bugs have been fixed
Import has been mostly rewritten, and is more efficient now
Win32 fixes
Many crash are now gone
Many more cleanups

In the last two weeks, I'd say i dropped from 1000 to 2000 lines of code. There's still many things to do, bug the code is gettings cleaner and cleaner everyday.

I hope the win32 release will come up in the next days. The binary release will come later, along with VLC 1.1

For those of you who use gentoo or archlinux, an ebuild/aur package is already available. In any case, this is the page to watch : http://trac.videolan.org/vlmc/wiki/Downloads

VLMC, On the way to the release

Hugo — Mon, 15 Mar 2010 11:02:00 +0000

Hi folks,

Just a quick post to inform you of the current VLMC state...

We've been working hard to fix the latest bugs, and to clean the code up. We still have a few problem under win32... some crashes, and most of all, the project preview doesn't work. We were planning on releasing yesterday, but it seems that we'll be late, again... We're doing our best to fix as many bug as possible, as quick as possible, but that's not easy, as VLMC team counts only 2 devs right now. Contributors are wanted and welcome !

About OSX, we can't test, as nobody on the team owns a mac. If you're a macosx developer, and you're willing to spend some time helping us out, you're more than welcome.

As you may have noticed, our website is back online. It's currently just a bugtracker, but it will become a real website, eventually. Please submit every bug you come across !

I really hope next post will be to announce a release :)

What about VLMC ?

Hugo — Tue, 09 Mar 2010 17:35:00 +0000

Hi everyone,

Long time since last post... at this time, VLMC wasn't abble to render anything, and was pretty much all about loading videos... and doing nothing with it....So what now ?

Many progress has been made since this time ! Right now, we're able to do some basic video editing. Cutting, moving, trimming, resizing, extracting parts from a rush... and many more to come hopefully ! Why didn't we talked much in the last two month ? Well... first VLMC was our end of studies project... we had to stop coding, and prepare for the boring stuff... oral defense, papers... but that's over. Then, we had a "small" problem with our server... by small, i mean that it completely burned down, which was pretty annoying...

We're currently moving to the VideoLAN infrastructure. You can already fetch the new git at : git://git.videolan.org/vlmc.git or if you feel like it, there's the RSS feed : http://git.videolan.org/?p=vlmc.git;a=rss Mailing list should be back by the end of the week, along with the bugtracker.

As the school project ended, the VLMC team was pretty much shrinked to a 2 or 3 developpers team... we definitely need some contributors ! Developers, translators, testers, advices, anything can be of assistance ! VLMC is now a fully open source project, and definitely needs your help !

Right now, I'm currently rewriting many parts of the code, that needed to be cleaned before anything else happen (you don't build you house over a swamp don't you ? ;) ) This should be done by the end of the week. After that, a small bug in the file rendering, and a crash when moving clips in the timeline prevents us from releasing a developer's preview. Once fixed, release will be on the way. Don't forget that we require a stable VLC version before packaging anything (unless we build it statically... but no.) The VLC 1.1 should be out in a few weeks now. Please be patient ! :)

Here you can find the latest demo video :

See you soon on #vlmc @ freenode !

GSoC : Day 1

Hugo — Mon, 25 May 2009 08:12:00 +0100

Here we start :)

This is the first real GSoC post, hurray \o/

For this first day, I'll need to code a quick render preview widget.

I already started to split the VLMC's RenderWidget into 2 separate classes : one for the clips/media preview, one for the timeline/workflow render, but the second isn't implemented yet, this will be my first day job. I'll post along the day, as the code is written :)

In order to provide a little more genericity and flexibility to the code, i'll add an generic class on the top of the two described above.

First surprise, the Q_DISABLE_COPY macro disable the default ctor... this is quite normal since this maccro adds a private copy ctor, and as soon as the compiler sees a user defined constructor, it disable every default ctor... so normal but surprising :p

After a quick code fix, it works as good as the old version :)

I was kinda scared by the idea of having pure virtual slots, but this works just fine. I really do love Qt more and more every day :o

So now, VLMC has a GenericPreviewWidget... it just needs to be used (as the generic type) a little more, and moreover, it needs a RenderPreviewWidget implementation :)

10:33AM : VLMC now has an "independent" GUI for its two renderers. Code has been pushed to GIT, and now the fun part begin, with the render preview implementation, which depends on in-vmem codec, freshly fixed by etix :)

01:00Pm : invmem isn't as simple as I expected :(
At the moment, it just displays a white screen, and it seems that the clip workflow that wasn't tested isn't functionnal... Hopefully this will be solved soon :/

01:05Pm : woops, I just didn't remember how to use memset anymore... shame on me, the problem definitly doesn't come from the invmem codec...
Next step will be to export what the vout:vmem save into the memory, and see if that's correct...

2:15Pm : Damn it.... each time a graphic window pops over VLC, my X-server crashes... I hope that the recent Xorg update (1.5.3-r5 -> 1.5.3-r6 on gentoo is to blame, but I don't feel like it :/ )
It reminds me of a bug I came across last year, each time a graphic window would cover an opengl window, X was completly frozen... the only solution was to reboot the computer, and it looked like the intel graphics drivers were to blame... i *really* do hope that this won't be the case here :/
This hasn't been solved yet... that's sad :/

3:30Pm : Hurray it works !
We now have a basic clip workflow that output a clip to a video buffer, and a main workflow that asks the first track to output the first clip output, and output it through the render preview widget, using the invmem codec. Basically, it just reads the video at the moment, but now our primary workflow is functionnal.
Next step : we should use Clip instead of Media. A clip is a representation of a media inside the timeline, and can cover just a part of it, or the entire media. Inside of the timeline, we must use Clip, but we use Media at the moment.
This implies that we have to handle at least :

Videos that doesn't start at "0"
Videos that ends (this case hasn't been tested yet, funny isn't it ?)
Videos that ends before the real end of the video
Generating a Clip from the Media, using the GUI.
Probably many other things :)

Recent VLC patches

Hugo — Mon, 11 May 2009 09:35:00 +0100

Hi,

This is a follow up from the previous post, and a little more :)

Regarding the libvlc_media_player_play() race condition, a patch just has been commited by fenrir.
With this new patch, VLMC moved to a unstable-creeping software to something a lot more stable. No more freeze when loading a video, which is good :)

In addition, a tiny patch has just been commit, this is really not a big deal, but I really wonder why this wasn't like this in the first place :o A filename must be const...
As this wasn't the case, we ended up doing things like this :

QTemporaryFile tmp;
tmp.setAutoRemove( false );
tmp.open();
strncpy(m_tmpSnapshotFilename, tmp.fileName().toStdString().c_str(), tmp.fileName().length());
m_mediaPlayer->takeSnapshot( m_tmpSnapshotFilename, 32, 32 );

Since the takeSnapshot parameter cannot be const, we have to copy it and pass it as a non const parameter... sigh :( However, this is fixed :)

Patch reference

Using VLC with Qt-Creator

Hugo — Thu, 30 Apr 2009 16:32:00 +0100

QtCreator is great, really really great, so why don't using it when browsing/editing/patching vlc source code ? This article is based on the original Jean-Baptiste Kempf post, in the vlc-devel mailing list.

I feel like I'm gonna have to look into the VLC source code very soon...
A "problem" I came across is that once you've launch the playback with libvlc_media_player_play, you have to let the first frame render before accessing the "vlc message pump" again (ie the Control_Pop and co).

It would be nice if I could provide a patch to look at the control stack *before* anything is done... however, I will have to browse some pieces of the VLC source code, and I'm kinda sick of greping and filtering the results by hand... Since Qt-Creator can be used with VLC, this is a good occasion to set it up !

Step 1: Launch Qt-Creator :D
Step 2: Create a new session (that can be seen as a sort of visual studio solution) named VLC (or whatever you wan't actually) => File -> session -> session manager -> add. The new session is set as the current one.
Step 3: If you haven't already, create a build subdir into your VLC directory :

$> mkdir build && cd build && ../bootstrap && ../configure and wait a few minutes :)

Step 4: Import the VLC project : File -> new ; Project/Import of Makefile-based Project ; Select your VLC dir as the project path, and name your newly imported project VLC
Step 5: close Qt-Creator
Step 6: Download those to files : http://people.videolan.org/~jb/QtCreator/VLC.files and http://people.videolan.org/~jb/QtCreator/VLC.includes inside your VLC dir (you should have to overwrite the ones Qt-Creator just generated
Step 7: Relaunch Qt-Creator
Step 8: Reload your last session (Should be the one you just created)
Step 9: Setup the build configuration : go to the big "project" button on the left, and adjust the field "Build directory" to your vlc build dir.
Go to the "run settings" tab, and create a new run configuration by clicking the blue "+" button. Name: VLC, Excutable: /path/to/your/vlc/build/vlc

j-b advised to add "-lqt" as the only argument, everytime I try, vlc just launched and quit, so I left the "argument" field empty.

Step 10: enjoy the "ctrl + click" on a symbol which brings you to its definition ! w00t !!

Original post by Jean-Baptiste Kempf

Getting into VLC source code

Hugo — Thu, 30 Apr 2009 10:54:00 +0100

As VLMC is built uppon VLC, I'll frequently have to take a closer look at the source code... At the moment, we have an annoying bug : when we ask for a media player to play a media, it's supposed to fire a "libvlc_MediaPlayerPlaying" somehow (at least it works for most of the media), but it doesn't with a specific mpeg file... let's try understanding why o//

First, we'll have to locate the file in which libvlc_media_player_play is declared.
Like most of the media_player functions, this one is located in "src/control/media_player.c"
Let's start by analysing what is does :

Releasing the input thread (which is the thread that handle the media playback) in case a media was already playing.
Creating a new input thread
Adding a bunch of rendering related variables (we'll certainly have to go back to variables later)
Adding 3 callbacks : input_seekable_changed, input_pausable_changed, input_event_changed

As far as i know, everything that's related to the media playback is managed by the input_thread, so let's take a closer look at the input_CreateThread() function...

This new function is located in "src/input/input.c". It's actually some kind of a wrapper on the "__input_CreateThreadExtended" function, which is located 10lines below. The only thing it does it to create a new thread, and launch it wich the "Run" function, which is located in the same file.

This Run function will call : Init(), MainLoop(), and finally, End(). At this time, I really don't care about End(), so I'll focus on the Init and MainLoop functions, which are, once again, located in the same file.

If I'm not mistaking, VLC worlflow consists of getting an elementary stream from a demuxer (among many other things) The init function creates the elementary streams (ES), creates 2 vlc-variables containing the bit-rate and the sample-rate.
Once this is done, it changes the media state to "Opening", and then call the InputSourceInit function.

This function is in charge of selecting the appropriate demux, or eventually selecting one that has been set by the command line. Selecting the demux is done by calling the demux_New() function. The selected demux is used to get a bunch of informations about the media, such as "can-seek", "can-pause"... and do many other things that I don't really understand at the moment :(
Let's go back to the Init function, and see what is does next...

Basically, it gets information about the file, init some stuff such as subtitles, and then changes the state to "Playing", which is what isn't always working in vlmc... I'm pretty confident that our problem is in this state change...

Even if this in kind of dirty, i'm going to add debug code to see what the hell happens in this function.

I'm going to add debug into the input_ChangeState, and to be sure that our problem isn't event related (and it seems like it, since the playback continues when we display the VLC debug lines...), into the input_SendEventState, which is located in the "src/input/event.c" file.

The "difficult" thing is that the input_SendEventState uses vlc-variables, which is a really nice thing, but something that's hard to apprehend (or at least it was for me). variable value is changed that way :

var_Change( p_input, "state", VLC_VAR_SETVALUE, &val, NULL );

When the "state" var is changed, it triggers a callback... but which one ?! Let's call or friend grep : grep --color -HRn "\"state\"" src/* Which highlight this line in particular : src/input/var.c:87: CALLBACK( "state", StateCallback ), Let's take a closer look... this callback calls "input_ControlPush( p_input, INPUT_CONTROL_SET_STATE, &newval );", provided that the new state is PLAYING_S, or PAUSE_S.

Good news everyone ! The state that was set is indeed PLAYING_S. We can keep going...

input_ControlPush takes us back into input.c.

What's the control inside vlc ?? The VLC playback thread has an array of struct input_control_t, that are executed by the thread, from the MainLoop(). This function will basically push a new control to be executed later... so this code is definitely asynchronous (did you still doubt it ?!)

Ok no worries, let's go to the MainLoop code, and see how this control stack is handled...

MainLoop() is located in input.c (big file ain't it ? :p). At every iteration, it clears the control stack by calling ControlPop(), until the stack is empty. The poped value is handled by the Control() function.

Control function is basically a good old fashioned switch/case, in which we're going to seek for the control we pass. We finally come down to this piece of code :

case INPUT_CONTROL_SET_STATE:
    if( ( val.i_int == PLAYING_S && p_input->p->i_state == PAUSE_S ) ||
        ( val.i_int == PAUSE_S && p_input->p->i_state == PAUSE_S ) )
    {
        ControlUnpause( p_input, i_control_date );
 
        b_force_update = true;
    }
    else if( val.i_int == PAUSE_S && p_input->p->i_state == PLAYING_S /* &&
              p_input->p->b_can_pause */ )
    {
        ControlPause( p_input, i_control_date );
 
        b_force_update = true;
    }
    else if( val.i_int == PAUSE_S && !p_input->p->b_can_pause && 0 )
    {
        b_force_update = true;
 
        /* Correct "state" value */
        input_ChangeState( p_input, p_input->p->i_state );
    }
    else if( val.i_int != PLAYING_S && val.i_int != PAUSE_S )
    {
        msg_Err( p_input, "invalid state in INPUT_CONTROL_SET_STATE" );
    }
    break;

This is were we're going to put our second piece of debug. Now let's build'n pray !

As expected, the second debug output isn't displayed... just have to figure out why :) Since this will consist only of putting debug text here and there, I'll skip the details, and only explain the problem :p

In the end, the problem seems to be a race condition inside the libvlc_media_player_play() function code... More precisely :

void libvlc_media_player_play( libvlc_media_player_t *p_mi,
                                 libvlc_exception_t *p_e )
{
    input_thread_t * p_input_thread;
 
    if( (p_input_thread = libvlc_get_input_thread( p_mi, p_e )) )
    {
        //stoping the thread, releasing the old media...
    }
 
    //Some stuff
 
    //Creates the thread, thats spawn onto the Run() function, which will call Init() (which basically send the event)
    p_mi->p_input_thread = input_CreateThread(
            p_mi->p_libvlc_instance->p_libvlc_int, p_mi->p_md->p_input_item );
 
    if( !p_mi->p_input_thread )
    {
        vlc_mutex_unlock( &p_mi->object_lock );
        return;
    }
 
    p_input_thread = p_mi->p_input_thread;
 
    //Another block of stuff
 
    //Setting the callbacks	
    var_AddCallback( p_input_thread, "can-seek", input_seekable_changed, p_mi );
    var_AddCallback( p_input_thread, "can-pause", input_pausable_changed, p_mi );
    var_AddCallback( p_input_thread, "intf-event", input_event_changed, p_mi );
 
    vlc_mutex_unlock( &p_mi->object_lock );
}

In all my tests, the call to var_AddCallback was made *after* the Init function completed, which is baaaaad, since no event can be fired.

Anyway, a mail has been sent to the vlc-devel mailing list in order to get some instructions on how to fix this (I don't feel like fixing it on my own without asking for advice, since it's somehow a big big part of the libvlc, and it touches some critical part of vlc, IE. the input thread), This should be fixed very soon :)

This is a one-shot post, please forgive me if some incoherence appears ;)

GSoC : Let's get started !

Hugo — Mon, 27 Apr 2009 07:27:00 +0100

Hi everyone !

It's been a looooong time since I've post anything arround... but this is about to stop, since I'm going to participate to the 2009 edition of the Google summer of code, for the VideoLAN project, more preciselly for the vlmc project (www.vlmc.org)

Just a few words to present the project : The aim is to provide a simple but powerfull video editing software, based on vlc.

At this time, it doesn't do much, just video loading, getting a snapshot and some meta informations, and a basic timeline that doesn't do much either.

In the next days, we will have to enable a basic workflow, so that we will be able to concatenate two videos. Once this is done, the fun part will begin ! My work during this gsoc will consist of conceiving an effects/transition API, add some functionnality to the timeline widget (in order to reflect whan can be done with the API), and probably many things that hasn't be though of right now :p

I'll try to keep the blog updated, as the features appears on the project :)

By the way, a big thanks to VideoLan association president Jean Baptiste Kempf, and to my mentor and friend, Ludovic Fauvet :) Some kisses to the vlmc team, which will go on vacation this summer, while I'll be coding :D

Cuda... episode 2

Hugo — Sun, 05 Oct 2008 15:16:00 +0100

Le premier programme en cuda, c'est fait, maintenant il serait bon de ne pas avoir a se prendre la tete avec les allocations memoires, ou tout du moins les simplifier un peu...

Ha aufait je vais ecrire en francais a partir de cet article, apres tout il n'y a pas d'articles en francais sur cuda (en tout cas pas la derniere fois que j'ai regarde :p) et puis ca me permettra surement d'etre plus clair.

Donc, personnellement les cudaMemcpy, cudaFree etc, ca va 5minutes, mais ca va me gonfler de devoir taper cudaMemcpyHostToDevice a longueur de temps, donc je me suis code une petite classe pour faire officer de "container" cuda. Il n'y aura rien d'extraordinaire la dedans, juste une petite classe fonctionnelle :)

Un rapide "cahier des charges", la classe doit pouvoir :

Envoyer des donnees vers la carte graphique
S'occuper seule comme une grande d'allouer la memoire dans la memoire de la carte graphique
liberer cet espace
Permettre de recuperer le pointeur vers le device pour pouvoir appeler les kernels
Pouvoir fonctionner avec n'importe quel type
Recuperer les donnees traitees par la CG

Ce qui nous donne la declaration de classe suivante :

template <typename T>
class		CudaContainer
{
public:
	CudaContainer(T* var, size_t size);
	CudaContainer(T& var);
	CudaContainer(size_t size);
	~CudaContainer();
	void		sendToDevice();
	void		getFromDevice();
	T*		getHost()
	T*		getDevice()
private:
       //probably more stuff here
};

ce qui une fois implemente, donne le fichier fourni en piece jointe de ce billet :)

Je n'ai pas commente abusivement, je pense que les noms de methodes et des fonctions de l'API cuda sont assez clairs ;)

Et voila le main du billet precedent, retouche pour fonctionner avec ce "container" :

int		main()
{
	const int NbElems = 100;
	float *toSquare = new float[NbElems];
	for (int i = 0; i < NbElems; ++i)
		toSquare[i] = (float)i;
 
	CudaContainer<float>		c(toSquare, sizeof(float) * NbElems);
	c.sendToDevice();
	//Cuda Kernel execution
	{
		dim3 blockSize(5, 5);
		const int nbThreads = blockSize.x * blockSize.y;
		int nbBlocks = NbElems / nbThreads + (NbElems % nbThreads == 0 ? 0 : 1);
		SquareArray<<<nbBlocks, blockSize>>>(c.getDevice(), NbElems);
	}
	c.getFromDevice();
	for (int i = 0; i < NbElems; ++i)
		std::cout << i << " => " << toSquare[i] << "\n";
	std::cout.flush();
}

Ce qui, a mes yeux, est toujours un peu plus agreable a lire/utiliser, qu'avec les fonctions initiales :)

Shader loading class

Hugo — Wed, 01 Oct 2008 22:52:00 +0100

For a little project I'll probably speak about here someday, I wrote a shader loading class to ease my pain, as i was discovering shaders...

I didn't wan't to bother loading shaders to try them, so I decided to write a class to do this for me once and for all.

Just as a reminder, a shader is a piece of code that will be executed by your graphical controller, instead of what's called the fixed function pipeline, which consists of a set of instructions for the vertice (object space to screen space, basic lighting), and a set of instructions for the fragments -aka pixels- (applying textures, fog, color blending).

You can write your shader in GLSL (openGL Shading Language), HLSL (High Level Shader Language, for DirectX), and CG (C for Graphics), which is the nvidia language, that will be compiled in binary code that can match with the ones generated by GLSL or HLSL compiler (but it does fail to sometimes :p )

The language i used at this time was GLSL, because it seemed to me that it would be easier to apprehend than HLSL, and I didn't really knew what CG was about...

This is basic c++, no killer language feature here, no pownage algorithm. Source can be found at the bottom.

It can be used like this :

Shaders* shad = new Shaders();
shad->setPixelShader("lightPS.glsl");
shad->setVertexShader("lightVS.glsl");
shad->makeAll();
 
//......
 
shad->enable();
glBegin(GL_QUADS);
glColor3ub(0,0,255); //face bleue
glNormal3f(0.f, 1.f, 0.f);
glVertex3d(-10, -1, 10);
glNormal3f(0.f, 1.f, 0.f);
glVertex3d(-10, -1, -10);
glNormal3f(0.f, 1.f, 0.f);
glVertex3d(10,	-1, -10);
glNormal3f(0.f, 1.f, 0.f);
glVertex3d(10,	-1, 10);
 
glEnd();
shad->disable();

and here is the test lighting shader i used for test :

Vertex shader :

varying vec3	normal;
varying vec3	vl;
 
void	main(void)
{
	vec4	light = vec4(0.0, 1.f, 0.0, 1.0);
 
	/// Putting the normal in the world space
	normal = gl_NormalMatrix * gl_Normal;
       /// Calculating the vertex position in world space
	vec3 point = gl_ModelViewMatrix * gl_Vertex;
       /// calculating the vertex-light ray
	vl = (gl_ModelViewMatrix * light).xyz - point;
       /// Calculating the vertex position in the projection space
	gl_Position = gl_ModelViewProjectionMatrix * gl_Vertex;
       /// Passing the color of the vertex to the fragment shader
	gl_FrontColor = gl_Color;
}

Fragment/Pixel shader :

varying vec3	normal;
varying vec3	vl;
 
void	main(void)
{
        /// normalizing both vectors to have a correct dot product result
	vec3 n = normalize(normal);
	vec3 lvect = normalize(vl);
       /// We get the angle between both vectors by using dot product, and clamping it between 0 and 1 to be sure
	float factor = clamp(dot(n, lvect), 0.0, 1.0);
       /// Finally, this angle is used to tweak the color, simulating the light
	gl_FragColor = gl_Color * factor;
}

with proper opengl initialisation and render, here is what we get : Et voila ! o//

First steps in CUDA

Hugo — Thu, 18 Sep 2008 23:42:00 +0100

I really want to throw myself in CUDA, so here I go. The first step will be to create a simple project, that runs under visual studio.

First, I just want to have a working environment that works with cuda. Since I don't want to fight with VS to handle cuda, i just searched for a VS2005 (sigh, cuda doesn't support VS2008 yet :'( ) project template. I found this one, which was the google first guess : http://forums.nvidia.com/index.php?showtopic=69183 It provides the Debug/Release/EmuDebug/EmuRelease modes, and the appropriated paths for compiling, which is exactly what i was looking for.

Good surprise, i works perfectly. The provided sample compiles without any issue, but it won't launch itself, because it can't find "cutild64.dll". It seems that the SDK doesn't set the PATH so the system can find cuda's dlls... I don't want my cuda's dll to be somewhere in \system32 or anywhere else than the cuda sdk, so i will upgrade my env variables. Go to the environment variable setting pannel (win key + pause key => advanced => env variables) For the moment I will only add debug DLLs, in order to try. I add a CUDA_DLL_PATH var that is equal to "%NVSDKCUDA_ROOT%\bin\win64\Debug" and added the following to my Path variable : ";%CUDA_DLL_PATH%"

After relaunching visual studio, it worked verry well. So let's add the release path, (before i forget to :D ) by adding "%NVSDKCUDA_ROOT%\bin\win64\Debug" to your newly created CUDA_DLL_PATH From now, I can compile and then launch CUDA project, after associating .cu files to visual c++, it's finally ok to start coding :)

Now lets go into cuda itself...

In order to launch code on the GPU, CUDA provides "Kernels". These are function that you can call from the CPU, but that will be run on the GPU. These functions are declared by the "global" keyword.

The "device" keywords means that the function will be callable from the GPU only, and run on it. The "host" keyword means that the function will be callable from the CPU only, and run on it. "host" is the default function qualifier. If you don't specify one, it will be host by default.

GPUs are highly parallel architectures, a high number of threads will run through your code, but you have to specify this number of threads. This is done when calling the kernel :

kernelFunctionName<<<blocksNumber, blockSize>(params);

It allows you to set your thread blocks parameters. A thread block is... well it's a block containing threads :D These informations are contained in a cuda special type called "dim3" which is basically an array of 3 integers.

If you set your blockSize to "dim3(5, 5)" you will have 5 * 5 ( * 1, for the implicit Z dimension) = 25 threads.

Then you can choose to create multiple threads blocks, by doing the same for the first kernel parameter (blocksNumber).

A set of threads blocks is called a grid.

If you choose dim3(2, 2) you will have 4 blocks of 25threads running through your code. Optimally choosing the right number of blocks and threads probably requires a high knowledge of the cuda doc, mostly the chapter 5, which is focused on optimisation.

But if you call your kernel, you must have a way to get the results of your computation back. This is done by allocating memory on the device, copying data to it, and then copy the results back in the main memory. CUDA provides a set of functions for this purpose. The one I will use are :

cudaMalloc(void** ptrToAlloc, size_t sizeToAlloc);
cudaMemcpy(void* destination, void* source, size_t sizeToCopy, cudaMemcpyKind mode); //cudaMemcpyKind is an enum to set the copy mode (host to device, device to host...)
cudaFree(void* ptrToFree);

For a test, i will write a little main.cu which will square every elements in an array. :

First, the main itself. It consists of classic c/c++

int		main()
{
        //setting the number of elements to compute
	const int NbElems = 100;
        //classic c++ allocation of the data to compute
	float *toSquare = new float[NbElems];
	for (int i = 0; i < NbElems; ++i)
		toSquare[i] = (float)i;
	float*	device = NULL;
        //allocating memory on the device
	cudaMalloc((void**)&device, sizeof(*toSquare) * NbElems);
        // and copying datas to this mevory space
	cudaMemcpy(device, toSquare, sizeof(*toSquare) * NbElems, cudaMemcpyHostToDevice);
	//Cuda Kernel execution
	{
                //we will have 5 * 5 threads blocks :
		dim3 blockSize(10, 10); //Z dimension is implicitly 1
		const int nbThreads = blockSize.x * blockSize.y;
               //we must have a grid containing enougth threads to compute every data, so the number of elements divided by the number of threads, and if it remains data to compute, we add one thread block
		int nbBlocks = NbElems / nbThreads + (NbElems % nbThreads == 0 ? 0 : 1);
               // and we launch the kernel !
		SquareArray<<<nbBlocks, blockSize>>>(device, NbElems);
	}
        // getting the results back
	cudaMemcpy(toSquare, device, sizeof(*toSquare) * NbElems, cudaMemcpyDeviceToHost);
       //and this doesn't need commentary i guess :p
	for (int i = 0; i < NbElems; ++i)
		std::cout << i << " => " << toSquare[i] << "\n";
	std::cout.flush();
       cudaFree(device);
       delete[] toSquare;
       return 0;
}

So we have our main code, but it misses a kernel to compile and run... Beside the global qualifier, a cuda function is exactly the same as in C :

__global__ void		SquareArray(float* in, int nb)
{
	int	id = ((blockIdx.x + blockIdx.y * blockDim.x) * (blockDim.x * blockDim.y))
				+ blockDim.x * threadIdx.y + threadIdx.x;
	if (id < nb)
		in[id] = in[id] * in[id];
}

The first line is just to know the array of the current array index this kernel has do compute. Knowing the thread and blocks size and id allow a cuda code to run lock free, which is kinda fast :p

Then there's just a "bounds" check, which, actually, isn't necessary. Without it the code will be executed just fine. Maybe there's a change in perfs, but I didn't pay attention to this yet.

Now it's complete, both main and kernel are done, compile and run, and you've a nice "array squarer" in cuda ! Hurray !

Now the next step will be to abstract the memory stuff in a class, to ease the pain :p

Calling contructor without parameter

Hugo — Sat, 13 Sep 2008 15:29:00 +0100

Maybe something like that already happens to you : You wan't to build an instance of your class without calling the constructor :

class   A
{
    private:
        A(){};
        ~A(){};
};
 
int main()
{
    A   a();
}

But this will compile on g++ without any problem... why is that ? Because you're not building an instance of your class, your declaring a function called "a", which will return an instance of A. If you pass parameters, g++ will understand this a a constructor call, but if you wan't to call a 0 parameter constructor, you will have to call your CTOR without parenthesis...

The almost useless c++ trick of the day

Hugo — Wed, 03 Sep 2008 23:37:00 +0100

When I learn c++, I remember my teacher telling : "A pure virtual method is a method without implementation", well that's not entirelly true... Actually, you can have a method beeing both pure virtual, and implementated. How is that possible ? well i must say that I don't know, since every assertions I ever made concerning how a compiler handles pure virtual methods was using the fact that a pure virtual was a null pointer in the class vtable... since you can have a body in it, it's obviously not null anonyme...

Anyway, implementing a pure virtual method is quite easy :

class  A
{
public :
	virtual   void pureVirtual() = 0
        {
            std::cout << "Me love otters" << std::endl;
        }
};

And this will compile without any problem !

The comportment of an implemented pure virtual method is quite simple, like any other pure virtual method, it prevents your class from beeing implemented. So here you can't have

A* a = new A();

it would print the usual error message : cannot instanciate abstract class. But what you can do, is call the implementation from a inherithed class of A, so here, you could do :

class B : A
{
public :
	virtual   void pureVirtual()
        {
            std::cout << "Not pure virtual anymore" << std::endl;
            A::pureVirtual();
        }
};
 
int main()
{
    B*b = new B();
    b->pureVirtual();    
}

which would print :

Not pure virtual anymore
Me love otters

And the worst part, is that even if a implemented pure virtual method seems a little bit of a non sense, not only it can be done, but it could be usefull, if you want to provide an interface with a default behavior, or probably many things else :)

Of debugging efficiency

Hugo — Sat, 30 Aug 2008 22:52:00 +0100

Important part of the developpment process is "wasted" in debugging. The most used technique (or at least, the technique I most observed at school) is the good old printf technique, or for c++ programmers std::cout/cerr, but there are a few drawbacks with this technique :

You need to recompile your code
Both printf and std::cout are buffering the data you want to print, and in some case, it can hide your memory allocation bugs, and your code will run with the debug stuff, but will crash without.
It's hard to get some informations about which function called this one, or about the entire callstack
It's annoying to get the process id, but mostly the thread id, in case multithreaded program developpment.

But like every problem, there's a solution, that isn't used enough considering the advantages it offers.

In visual studio, a breakpoint isn't just about pausing the execution when reached, it's a lot more, and I invite you to read the following arcticle

As you certainly already know, you can put breakpoints in your code bu pressing F9, but sometimes, you want to break only if something is equal to a specific value, or if the value has changed since the last call, or anything else. What I used to do was this sort of code :

void   randomFunction(int value)
{
  if (value == 1234)
     std::cout << "Breakpoint" << std::endl; //Here goes the breakpoint.
  //...
}

or something like that :

void  anotherFunction(int value)
{
  static int debugValue;
 
  if (value != debugValue)
    std::cout << "Changed value" << std::endl; //Another breakpoint
  debugValue = value;
  //...
}

And then put some breakpoints where they should be. But visual studio (since, i guess, 2005 version) allows you to do nice stuff with your breakpoints. You have a nice set of options if you right click on your breakpoint. You can set a breaking condition by clicking... condition.

Here is what you will get : (sorry for the french version of VS...)

Here, I want to break only if value == 0, so i just have to set this as my condition. You will notice the little "plus" icone of your breakpoint, which means this breakpoint is now conditionnal, and will only be triggered when your condition is satisfied.

If you want to check that your value hasn't change (or has changed), you can select "has changed". Your breakpoint will now be triggered when the comparaison is different from the last one. For an exemple, if i change my main to this :

int		main()
{
	testFunc1(1);
	testFunc1(0);
	testFunc1(0);
	testFunc1(2);
}

It will be trigged after the second call, since now value is really equal to zero, not the third, because value is still equal to zero, and the 4th and last time, because 2 isn't equal to zero, so the condition result changed from true to false

Another nice feature is the tracepoint, which is a breakpoint that doesn't break... in the break point right click menu, select "when hit" Once you're here, you will get this screen :

If you check "Print a message", you will be able to trace your code, like you can with printf, but this won't require code modification, so no need to build your code again. This is more appreciable when you want to add trace in a header file, containing your templates function for examples, and you don't want to rebuild your entire project. The huge drawback of this method is that is slows your program a lot

You can also break after a number of calls, only from a given thread id... this is a powerful tool :)

If you just discovered this feature, i hope you will use it a lot (but let's hope you won't have to ;) )

OpenCL

Hugo — Sun, 10 Aug 2008 01:06:00 +0100

Hi all, With almost 2 month late, i would like to point out the great news that OpenCL is :
Right now, the GP/GPU solutions are :

CUDA for NVidia hardware
CTM for AMD(ATI) hardware

which is sad for developpers, since we (and by we, I mostly mean I) don't want to have to use 2 APIs, and 2 languages, of different abstraction levels (CUDA is "high" level, where CTM is close to hardware (or to metal in that case \o/))
Considering NVidia goes well with CUDA, we could have bet on CTM death, but what's gonna happen is far better : OpenCL will born "soon", and will provide a unify way to use your GPU's power.
Gigaflop and mass parallel programming, here we are ! (as soon as specification and compiler will be out, which lets us a few years at least)

Let's wish good luck to the Khronos group which have to write the standards !

Hello World

Hugo — Wed, 06 Aug 2008 22:12:00 +0100

Hi all, Welcome on this blog. This isn't my first attempt, but i'll make my best to keep this one up to date. This won't be much of a "mylife" blog. More likely, this will be a technical blog.

But first of all, It would be best to introduce myself : I'm Hugo Beauzee-Luyssen, 21 years old, studying computer science in Epitech (Paris), and currently in internship at Ubisoft Paris, as a gameplay engineer. I've felt a need to share my developer experience since for a while, and here it is.

I'm currently working on a raytracer with a friend, this will be an excuse for writing a bunch of articles, but i'll write a new post a soon as i find a new trick, a new technique, or just something worth looking.

Even if i'll do my best, I apologize for any grammatical mistakes that could haunt my posts ;)

Ho and if you ask why mustelidae.fr, the answer is quite simple : everyone has a favorite debug word. In my school, some use "toto", others use "bite", you can even find some "chiche donne nous tout", but I prefer to use "loutre", which is the french word for "otter". Since I had to write course for 2nd year students, I used otters in every of my exemple, and so i was known as the man who loves otters. As you probably don't know (or maybe you do ?) otters are from the mustelidae animal familly (http://en.wikipedia.org/wiki/Mustelidae), and this is why my website is "mustelidae.fr" :)

I hope you'll enjoy reading the articles i'll post here, and event learn a few things !