bouliiii's blog

Sauerbraten: a game and an engine that rock

2011-12-18T02:16:00.000-08:00

Quick post about a fantastic game and engine: Sauerbraten (you may find it here)
If you are tired by design pattern BS, static type introspection with SFINAE joke, over-design of everything just to print a string... Sauerbraten is for you.

This is one of the code bases that recently impress me the most. Just straightforward and fantastic code. Everything is packed in 70,000 lines of code. Very nice custom (and really fast) UDP network code (ENet), insanely small scripting engine, straight-to-point rendering code, very cool AI, superb UI (3d is even supported), simple but powerful shader system...

Look at the code really and you will see what brutal and brilliant coding style means.

Global variables and cvar in tasking systems...

2011-12-05T22:30:00.000-08:00

Usually it is just better to avoid global variables but there is something I really want in my small video game engine is a-la-quake console variables (cvar). They are handy and you can tweak your system in a very easy way.

Problem if you have a fully distributed system changing a cvar from the console is just terrible. If you do some frame-overlapping (for example), you may completely screw your system (saying you decrease / increase the maximum number of particles on screen...). Well, cvars simply become nasty race conditions.

Initially, I wanted to do something complicated ie some copy-on-write stuff. You have a fully reference counted copy of your cvars and when you modify them you copy them into a new reference counted state. Then, instead of reading your variables from the global state, you fetch them from your copy. As long as the copy is valid (ie no modification is done), you can keep it. So it is not that expensive.

Problem is that it is just uber-complicated and your cvar are not global anymore and therefore not that handy. Too bad. The idea I finally decided to implement (not done yet) is to lock the tasking system when you change a global cvar. Basically, you stop all the other threads: you simply need to make them sleep after they run one given task and once everyone except the locking thread is sleeping, you modify the variable.

So it is something like:
TaskingSystemLock();
cvar = ...
TaskingSystemUnlock();

It is brutal and super expensive, but you don't care since this is really rare anyway. From the other thread perspectives (the ones that go to sleep), you absolutely know that nothing is going to happen while you run a task since the lock does not cross the task boundary.

So, for the usual path (ie you just read the cvar), you do not need any mutual exclusion since you know nothing is going to happen inside the task itself.

It is going to take some time to have a real life test case in point-frag but the on-going implementation in yaTS seems to be straightforward.

EDIT: Note also that several threads can request a lock simultaneously. This is not a problem and this is also properly serialized (like any other lock). Only thing is that all variables that may be modified globally have to be reloaded after the lock.

Quick and dirty compiler tests

2011-12-02T01:43:00.000-08:00

With point-frag, I am trying to support as many compilers and as many platform (well, Windows, Linux and soon Mac...) as I can. I hope to be able to do a benchmark from it at some point. I anyway quickly compared GCC, ICC and Clang on Linux with my small ray tracing test.
Here the numbers for my nehalem i7 (4 cores / 8 threads, triple channel memory, basically the high end chip 3 years ago). Code is compiled for SSE2 and does not use anything more recent.

GCC (4.6.2)
ray packet: 90 million rays/s
single rays: 9,9 million rays.s

ICC (12.1.0)
ray packet: 95.6 million rays/s
single rays: 10.6 million rays.s

CLANG (3.1 compiled from svn)
ray packet: 94.2 million rays/s
single rays: 9.9 million rays.s

With no surprise, the Intel compiler is slightly faster on a Intel processor than the other compilers. More surprisingly, the speed difference is quite small. Two years ago, GCC / ICC difference was about 15% on the same machine. Latest clang also outputs some fast code.

Good job open source compiler guys!

ompf.org is dead. Long live ompf!

2011-11-30T23:17:00.000-08:00

The best ray tracing / graphics related forum in the world ompf.org is down.
Fortunately, Jacco Bikker sets up a new one here:

http://igad2.nhtv.nl/ompf2

Cheers,
Ben

Ray traced occlusion culling in point-frag

2011-11-28T23:26:00.000-08:00

Arabic city (culled by the technique described here. The model can be found here)

Hello all,
I spent a decent amount of time implementing an occlusion culling technique for the video game engine I am developing: point-frag.

Various techniques
Initially, I thought hard about the different techniques to go beyond the good old frustum technique. First idea that came up was to use occlusion queries and implement CHC or CHC++ algorithms. Both sound good but they are pretty complicated. The first one does a bad job at reordering draw calls to decrease state changes and still badly suffers from stalls while the other one (the better version) is a combination of so many techniques (tighter bounding volumes, multi-queries, batching of already processed nodes, use of heuristics) that it already dissuades me to implement it. Both articles also show that this is not that easy to achieve something close to the optimal solution.

Other idea is the classical precomputed potential visible sets (see the work done by Seth Teller here). This was the technical solution used for Quake (BSP+PVS) and it was the direct implementation of Seth Teller's PhD. The solution is nice but I did not want anything that was too much time consuming like a preprocess since I would like to have a rendering pipeline as stupid as possible (i.e. rendering something which is as much as possible done at run-time and does not depend on anything precomputed (it will be also true for lighting BTW))

Software z-buffer
This leads me to look for something simpler: basically the solution implemented in Ps3 Warhawk, ie a software occlusion technique using a CPU-computed z buffer. This is basically the solution implemented in Battlefields games. You have a lot of nice details here. Note that Killzone 3 also implements a similar solution. All PDF from Guerilla are here.

All these games basically consider a set of occluders usually provided by the artists. They implement a software tiled based rasterizer. Basically:

You bin the triangles per tile. Many threads simply traverse each a sub-set of all occluders and find the set of tiles touched by each of them. You then maintain the list of occluders per tile. Note that it seems that KZ3 does not multi-thread this part
Then, a bunch of tasks rasterizes the occluders in each tile. Since tiles do not influence each other, this is pretty straightforward

HiZ buffer (highest level) as computed

Once done, you have a z buffer that you can query to know if one given object is visible or not. Usually (this is what I did at least), you take its bounding box, clamp it onto the screen and consider the zmin value (ie the distance to the closest point in the bounding box).

The technique is cool, simple and reasonably fast. There is no GPU dependency, you don't have to struggle with your drivers, everything is perfect.

Except that I don't think rasterization is the perfect tool here :-), at least for me

Ray tracing to compute the z-buffer
This is not supposed to be a rasterization vs ray tracing debate. The thing is that we speak here about software rasterization and software ray tracing. There is no debate at least right now for me with hardware rasterization. GPU rasterizers are fantastic and beautifully complex machines and GPU themselves are fantastic throughput processors.

Small buffers are better with ray tracing
Software rasterizers are another story. If you look at KZ3 numbers, they are kind of bad. They need 2.5ms to complete the rasterization with 1500 triangles for 640x360 resolution. Battlefield uses even smaller z-buffer (256x114).
(Note that a good old G80 is able to 200,000 triangles in 1ms in 1Kx1K)
This is my first problem. For such a small buffer, rasterization is already bad. Rasterization plays nicer with big resolution in particular to offset the initial triangle setup cost. The smaller the screen, the more the setup steps are costly.
With ray tracing, there is no such overhead

Dynamic occluders are not that a problem with few triangles
For dynamic scenes, ray tracing is not really different from tiled based rasterizer. Indeed, a 3D spatial binning of the triangle before building your bounding volume hierarchy (BVH) is not really different from a 2D binning for tiled based rasterization. Actually, 3D binning is even simpler since there is no perspective transformation at all to do on the vertices. You just bin the triangles in a grid. Once done, the BVH build is super fast. See Ingo's paper here for example. It is really fast.

Ray packets are super fast
This is the key. Once you have a hierarchy over your occluders, you just create SIMD packets of rays and you traverse the BVH. This is damn fast. point-frag includes a complete ray packet traversal (and also single ray traversal for future lighting ideas I want to play with). This is just super fast. For Fairy Forest model (see here), I easily achieve 95 millions rays/sec on a 2600K and even 35 millions rays/sec on a 4 cores Fusion AMD laptop. Once again, ray tracing in that case is really simple. There is no perspective transform, no clipping. Just purely simple 3D computations. No corner case.

To be honest, I also implemented in the past a highly optimized tile based rasterizer on CPU. For Fairy Forest (about 180,000 triangles), on 1024x1024 it was only 30% faster that the highly optimized ray packet traversal. However, on 256x256, the ray packet tracer was 80% faster than the rasterizer with a much much simpler code. If you are interested, the DynBVH article by Ingo (Wald) is here. It is a beautiful and simple algorithm. Really cool.

Ray traced occlusion culling is the only valid solution when you do NOT have artists
Yep. This is my case right now. I only have me and google sketchup :-). So, I basically need to compute occlusions on the initial data. And once again, triangle setups and clipping is too expensive if you have 800,000 occluders (ie triangles). Ray tracing does not really care and this is really good.

Ray tracing acceleration structures can be small
This is a usual problem with ray tracing. BVH + triangle soups are huge. Instead of triangle soup, you can use indexed face sets. This is better but still you duplicate your data. Fortunately, at least for static scenes, you can use a bounding volume hierarchy (BVH) as a compression scheme. Manfred (Ernst) and I, when I was still at Intel Labs, indeed demonstrated a simple quantization + compression technique that uses the BVH itself to hierarchically encode a complete geometry. The very good things are:

The decompression algorithm is so fast that you do not need to have a cache. You just decompress the structure on the fly at the same time you traverse it.
The compression rates are very good. About 5x compared to an indexed face set (index buffer + vertex buffer + BVH) and, if I remember properly, about 8x compared to a triangle soup + BVH

The article is here.

For all these reasons and limitations, I therefore implemented a ray traced z-buffer in point frag. You can actually see the code in the repository and even try it if you want to (be aware, I only tested the code on a very small set of machines and this requires OGL 3.3).

Timing are pretty good. Computing a 256x128 z buffer over 400,000 triangles and culling 1,500 boxes takes 1ms on my nehalem machine. The ray tracing part is multi-threaded but not the culling part (yet). Both are however decently optimized. Finally, the overall rendering structure is simplistic today but I have to start somewhere.

That's all folks :-)

EDIT: note that if you use SPUs and you only have 1,000ish triangles, using ray tracing is really appealing. You asynchronously upload the BVH (possibly compressed) entirely in the local storage and you traverse it at the speed of light :-)

Exponential Propagation while using Distributed Job Systems

2011-11-15T01:12:00.000-08:00

This is basically a common idiom with work-stealing like tasking systems. Problem with them is that you only have a local view of the complete system.

Typical opposite approach is a shared FIFO of tasks. You possibly have a lot of contention, the memory behaviour is really bad (compared to work-stealing which is really good) but you have an exact view of the running system at each time. The advantage for example is the way to handle priorities. You just make several shared priority queues and picking up the highest priority task is quite easy.

To overcome (partially) these kinds of problems, a usual idiom with work-stealing-like systems is to exponentially create information. People usually call that "nuclear reaction" :-)

Some examples:

Task sets with an atomic counter
This is the approach I chose for yaTS. A task set is just a task to run n times possibly with as many threads as you need. All the threads share the task set atomic counter. The problem is that you need to make it run as quickly as possible on the other threads. The technique is just to reschedule the task twice in its own queue to have this task stolen by other threads. Since all the other threads will also push the task back twice in their queues, you are sure that the task set will propagate exponentially across the queues

Task sets with ranges
This is the parallel_for approach used by TBB. Here you recursively subdivide the range [0,n-1] by splitting it into two parts. This creates a tree of tasks which once again propagate exponentially across the threads

(Note that the TBB approach is a bit nicer that the atomic version, since:

You do not have a lot of contention on one counter
since you push many tasks, you go back to the scheduler more often. This is better when a higher priority task just appears on your local queues
you can load balance in an automatic way by choosing when you stop to recurse)

Highly reactive and fast thread yields and wakeups
This is another issue with distributed queues: when does the thread need to sleep and when does it need to wake up? Once again, propagate information exponentially

The threads go to sleep when they do not find anything to do after n successive failed tries. Just use a condition wait for that.
A thread that creates a new task randomly tries to wake up two threads. Just use a condition signal for that. Either many tasks are going to be created, and since all the threads wake up possibly two other threads, very quickly, every one is working. Or, no tasks or very few tasks are created, and awake threads quickly return to sleep.

That is it. Ignoring some technical details to kill the threads and avoid dead locks, you can very quickly turn off and turn on your threads. Pretty cool to limit power consumption, to play nice with hyper-threading (it is really bad to spin for nothing while the other threads of your core are working for real) and to avoid any useless contention.

Point-frag: a distributed game engine

2011-11-06T00:12:00.000-07:00

These days and weeks, I am developing a small game engine in my spare time. I personally spent 18 months in the video game industry. It was pretty nice and overall video games are very interesting since the amount of technical challenges is huge: rendering of course, streaming, network, AI, a lot of system programming and so on...

So, after thinking about cool personnal projects I may do, I decided to play with a video game engine. The current code is here:
http://code.google.com/p/point-frag/
(do not try to run it right now, it is changing quickly and I am the only developer so it is unstable)

One of the first goals is to make it completely distributed: there is no heavy threads, no main thread or no renderer threads. Basically, *everything* is executing uniformly through the same tasking system (in that case, yaTS).

In all game engines I saw, several heavy threads usually exist. They are typically:
- main thread which is basically a loop
- renderer thread which run the OGL/DX/Gcm/whatever draw calls
- streaming threads
- audio thread
Usually many of them communicate with a dedicated command buffer with the main thread which is usually connected to them in a unidirectional or bidirectional way.

Beside that, all other threads are worker threads and usually communicate with a task interface (like C functions or C++ class or whatever). They are a resource that can be used at any time.

This basically makes the thing somehow cumbersome and we have to deal with several levels of parallelism in the system.

So, the main idea with point frag is to remove that and to have only worker threads:
- there is no main thread
- there is no game loop

Of course, typically you may need to pin some tasks to some threads (like something that does some OGL draw call), the system is flexible enough to run arbitrary task on arbitrary threads.

In the code you can see today, I simply display a model and use the tasking system to load the textures (asynchronously) and compress them into DXT1... Nothing spectacular but I must start somewhere.

However, the cool thing you may already see is the fact that the game loop is emulated by making frame_n spawn frame_{n+1}. This allows the system to continue in a continuation passing style way.

The good thing is that since there is no game loop, you can then design your complete game as a pipeline. The idea is to subdivide your frame into sub-tasks and see each sub-task as an element of a pipeline. This allows to have partial frame overlapping or if you duplicate the pipeline (so you have twice, three times a frame), you can even increase more the parallelism (with more latency).

Since everything is a task, it becomes easier to decrease (for example) the "renderer thread" (basically the one that has the OGL context) burden. Just subdivide the job to do to only have OGL calls in specialized OGL tasks. Everything else can go somewhere else.

Unfortunately, nothing is perfect and IOs are a problem. By IO, I mean any asynchronous requests handled by the system. When you read a file in a non-blocking way, you must do something to cover the latency. In yaTS, you can ask the tasking system to run anything. It is convenient.

However, the other task you call can also do an IO that can be much longer to complete than the IO before. This therefore blocks the calling task that performs the shorter IO.

Fortunately, you can handle this problem with a tasking system a la yaTS / TBB. Basically, it is possible to emulate co-routines with tasks and to yield capabilities

This is for another post (and that does not require assembly)
:-)

So, the no-heavy-threads approach may sound good but scheduling remains a damn problem. This is the good thing with thread over-subscription + blocking IOs: the system really helps you making the system progress in a relatively fair and time-shared environment. When you discard the system and do everything yourself, you are on your own

Anyway, point frag is my laboratory for experiments regarding tasking, rendering, procedural generation... Let's get some fun

To all HW vendors: do not thread your drivers

2011-11-02T12:53:00.000-07:00

This is kind of crazy. Many drivers are now spawning worker threads to do that or that task. Please make your driver with as few contention as you can but *really* stop spawning threads behind the developer back.

This just plays very badly with the application and now the developer has to deal with even more discrepancies from HW to HW and drivers to drivers possibly dealing with horrible thread over-subscription.

The only one that knows if the application needs to have a multi-threaded back-end (for creation resource for instance) is the user of the API.

Effect on data sharing between CPU cores

2011-10-11T21:45:00.000-07:00

Hello,

I made several benchmarks to stress the tasking system (yaTS) I recently developed and to see if it handles continuation / completion correctly.

Two tests actually spawn a binary tree of tasks.
You can see them in utests.cpp (CascadeNodeTask and NodeTask)

There is only one difference between both:

NodeTask. Here each node completes the root. This means that when a task just finishes to run, it decrements an atomic counter in the root task. When this counter becomes zero, the root is done
CascadeNodeTask. Here each node completes its parent. This basically means that the tasks finish in a cascade way (This is the classical and efficient way to do with work-stealing approach where the tasks are processed in depth-first order)

In (2), there is much less contentation than in (1) because in (1) the root cache line which contains the atomic is going to travel from cores to cores during the process.

This leads to interesting results on my i7 machine (4 cores / 8 threads)

NodeTask
1 thread == 237 ms
8 threads == 213 ms
Speed up == x1.1

CascadeNodeTask
1 thread == 237 ms
8 threads == 54 ms
Speed up == x4.4 (> 4 => Congratulations hyper-threading!)

This is basically the price to share.

(EDIT, also do not see that as a performance "study". This is just a random but interesting performance difference I saw while writing functional tests for yaTS code)

Disk asynchronous IO mess on Linux and libaio

2011-10-11T02:31:00.000-07:00

Damn there is something which is really imperfect with Linux: Asynchronous Disk IO. Basically, you want to issue a read or write and you want the function to return immediately.

Well, two classical solutions:

You open a file with O_NONBLOCK flag. No luck, the system does *not* need to open it as non-blocking actually. And after a test, the system actually blocks *anyway* when you open a disk file with O_NONBLOCK file. See here for a documentation
You use the portable posix AIO library like described here. Well, very bad thing, the implementation is done in user mode using pthread to emulate everything. Terrible for any high performing game like application where you do not want any kind of thread oversubscription

Looking for the perfect answer, there is finally a non-portable way to do exactly what I want. This is called libaio and this is just a thin layer above some Linux native system calls.

There is a quick and dirty example of libaio here.

So,

libaio does not spawn any thread behind your back
libaio is really asynchronous
libaio can batch some number IO jobs in one request

So, this is asynchronous IOs done right :-)

EDIT: but libaio does not seem to bufferize anything. I am wondering if there is a asynchronous buffered solution on Linux for disk IO even if in my case (streaming textures / models) unbuffered IOs should be OK. Well, that may require manual handling of buffer for big data you may not want to have fully on RAM before processing it. We will see :-)

yaTS - yet another Tasking System

2011-10-09T00:06:00.000-07:00

Hello all,

some real posts after a *long* time. Well, basically, I want to only post anything related to some piece of software I write and publish.

So, I think I will publish a small serie of posts related to a small library I wrote: yaTS which is here:

http://code.google.com/p/yats/

As I am extremely lazy for this first post, I will mostly take what I wrote in tasking.hpp :)

The idea here is to present what the tasking system does.

Basically, a "tasking system" offers the possibility to schedule and asynchronously run functions in shared memory "system threads". This is basically a thread pool.

However, yaTS tries to propose more in this API by letting the user:

Define *dependencies* between tasks
Setup priorities for each task ie a higher priority task will be more likely executed than a lower priority one
Setup affinities for each of them ie a task can be "pinned" on some specific hardware thread (typically useful when something depends on a context like an OpenGL context)

The core of this tasking system is a "Task". A task represents a function to call (to run) later. Each task can specify dependencies in two ways:

"Start dependencies" specified (see below) by Task::starts. Basically, to be able to start, a task must have all its start dependencies *ended*
"End dependencies" specified (see below) by Task::ends. In that case, tgo be able to finish, a task must have all its end dependencies *ended*

So, task1->starts(task2) means that task2 cannot start before task1 is ended
Also, task3->ends(task4) means that task4 cannot end before task3 is ended
Note that each task can only start one task and can only end one task

Specifying dependencies in that way allows the user to *dynamically* (ie during the task execution) create a direct acyclic graph of tasks (DAG). One may look at the unit tests to see how it basically works.

yaTS also classicaly implements a TaskSet which is a function which can be run n times (concurrently on any number of threads). TaskSet are a particularly efficient way to logically create n tasks in one chunk.

So, yaTS is somehow similar in TBB or other tasking systems.
However, I tried hard to make it small and I tried to have a API which is not too complicated and also reasonably powerful.

As you may see in utests.cpp, writing dynamic graphs of tasks is easy. You may use classical continuation-like tasking systems or wait for a task completion.

You may "pin" some task onto a particular HW threads. This will be useful when dealing with graphics API or anything with a context attached to a thread.

In the next posts, I would like to present a bit the current implementation details. Later, and as soon as the code is started, I would like to present how we can use this library to design a "game-loop-less" game engine.

Ben

Blog Resurrection Attempt

2011-09-06T03:07:00.000-07:00

Hello all,

This blog has been dead for a long time. I am trying to resurrect it :-)

I recently decided to work on a new personal project, a small video game engine. I basically extracted some code from the very nice Intel open source project, "Embree" (from Manfred Ernst and Sven Woop, outstanding guys I had the chance to work with in the past) you may find here:

http://software.intel.com/en-us/articles/embree-photo-realistic-ray-tracing-kernels/

Then, I quickly wrote some very basic pieces of code I threw here:

http://code.google.com/p/point-frag/

Nothing spectular actually.

Anyway, the first thing I would like to try is to build a "symetrical" multi-threading system. Basically, for several reasons (historical or because some hardware contexts are bound to one thread in particular) , games still use both "heavy" threads like the rendering thread or the main thread. The idea I would like to try is to build a tasking system where only worker threads exist (ie the main thread becomes a regular worker) and to replace every game / render loop by job pipelines.

Stay tuned :-)

Brute Force Tesselation and Quadratic Approximation of Subdivision Surfaces

2008-08-31T09:02:00.000-07:00

**Update** This post is somehow outdated and I just lost this small demo. Sorry for that.

Hello there, next post!
Here no relationship with ray tracing; on the contrary, it is all about rasterization and tesselation.
First, some screen shots of the demo provided with this post:

The standard mesh with no tesselation (2408 triangles)
___

A close up on the tesselated surface in wireframe mode (10,000,000 triangles). About 200 millions triangles per second are displayed on my low-end 8400GS (Here the frame rate is lower due to the wireframe mode slow on Nvidia)
__

A simple procedural displacement is applied on the tesselated mesh
__

This technique actually relies on several techniques developed by Tamy Boubekeur.

Patch Tesselation

The first thing to do is to tesselate the input mesh. To achieve this, we simply replace each triangle by a patch, i.e. a set of barycentric coordinate triangles. Then, instead of displaying each triangle, we replace each of them by an instanciated patch. In the demo, we therefore create two vertex streams. The first one contains all the data for each triangle and another one simply contains the barycentric coordinates of the triangles of a patch. Then, the instanciation ability of Direct3D allows us to repeat the patch along the surface of the object. For more details about the technique, may one refer to the Graphics Hardware '05 article by Tamy here. However, in this article, no Direct3D instanciation is used.

So, why is this method good? Actually, if the patches are large enough (i.e. if the tesselation rate is high), the pre-transform and post-transform caches used by the GPUs become more and more efficient so that we directly get the maximum theoretical performance of the rasterizers. For example, the frame rate does not change when you use 64 times more triangles. In fact, we can expect the same level of performance for the GPUs which will support tesselation units.without the need to use any hand-coded patches. In my 8400GS, I got 200 millions triangles per second but we may easily expect 1 billion triangles per second on DirectX11 GPUs.

By the way, I use in the demo a simple optimization of the patch using the very good and simple indexation optimizer, "Tipsify" (you may find the reference here). This reorders the patch indexation in linear time using only triangle / point adjacency while maintaining a very good level of performance (and this is much simpler and much faster that stripification algorithms like TriStrip or NvStrip).

Approximation of Subdivision Surfaces

The tesselation algorithm via the Direct3D instanciation is the core system to provide many triangles. However, the way triangles are generated (via tesselated patches) does not seem appropriate to the evaluation of subdivision surfaces since they are complex recursive algorithms. (By the way, one may find excellent references and details on subdivision surfaces in the awesomo tutorial located here). In the article, QAS (located here), Tamy proposes to replace the recursive evaluation of the subdivision surface by an analytical approximation. The principle is pretty simple: first, project the triangles vertices and the middle of every edge on the limit surface (i.e. the subdivision surface) and compute a quadratic Bezier triangle which fits these 6 points (the vertices and the edge middle points). This leads to a pretty fast and nice approximation of subdivsion surfaces.

About Adaptive Tesselation

In this article, Tamy proposes a very practical way to adaptively tesselate patches without cracks. This may lead for example to simple terrain rendering algorithms (if we ignore the possible memory problems which may require texture mip-mapping or page faulting systems). However, adaptive tesselation on current hardware (or rather APIs since Radeon already have tesselation units) may be less efficient for some cases since it requires to change the patches and therefore to queue more draw calls. However, with nicer APIs (like DX11), this will lead to super efficient solutions.

Some Details about Asset Production

What can we do with displacement? Actually, many tools already provide the assets. Using ZBrush, we already have the mesh + normal map + displacement map. Rendering this is simple: tesselate the mesh a subdivision surface approximation, displace it and light it correctly using the normal maps in the shader. Another super nice tool would to extend the mega-texture technique (the original article by Sylvain Lefebvre may be found here) to directly paint displacement on objects. However, the QAS technique may suffer from serious drawback when the artists decide to use creasing. In a modeler like Maya, the artist may work in subdivsion surface mode. He can view the tesselated mesh while edition the control mesh. He can then crease triangles, edges or vertices in any level of tesselation and then see the resulting subdisvion surface. For this kind of input, the QAS algorithm requires deep modification. Nevertheless, even if it seems necessary to provide ways to crease the geometry, none of the artists I work with produces the asset in this mode (A ZBrush approach + polygonal mode edition is generally the way they produce meshes).

The Demo

Here is the demo. This demo does not include displacement maps but I quickly hacked a stupid procedural displacement in the shader (that you may change). Using a displacement map does not change the performance on a 8800GT (due to the large number of texture units on it). On the 8600GT I have in office, the application becomes texture bound. I do not provide the complete source code yet since it is a bit in a dirty state. I will however try to post it ASAP.

Enjoy!

Cuda real time ray tracing - 100 millions ray/s?

2008-08-16T16:49:00.000-07:00

**Update** This post is somehow outdated. The code of Radius-Cuda may be found on the internet I think. Look at ompf.org or similar places. On my side, I lost the code. Sorry :-(

Ray Tracing!
There is a big buzz around these words nowadays and flamewars are everywhere. What I can however say is that I spent a long time writing ray tracers, Monte-Carlo renderers and: it is fun!

I recently decided to write a ray tracer using Cuda. There were two reasons for this. The first one is that I never wrote one line of Cuda. The second one is that ray tracing was said to be slow on G80 graphics cards due the heavy use of caches and stacks during the accelaration structure traversal generally required by ray tracers.

As usual, I first fetched all the articles I could find on the net, downloaded Cuda, installed it, runned some samples and began to code the thing. As I needed a minimalist ray tracing framework, I decided to drop my old ray tracer, yaCORT, and installed Radius, a small but pretty fast ray tracer developed by Thierry Berger-Perrin. As I had moreover already developed a Bounding Volume Hierarchy on it, it was perfect to inject a Cuda ray tracer with it.

I therefore decided to write a kd-tree traversal and first results were disappointing. Fortunately, with the help of Brian Budge (who posts as me, on the uber-excellent ray tracing forum ompf.org), I managed to get quite good performance.

Summary of the implementation

a stack per ray: no need to synchronize rays using a lot of __syncthreads(). The cost is that the stack will be necessary on local memory (slow) but we do not need anymore to handle ray packets or things like that as done with shared stacks or SIMD ray tracing. The Cuda compiler handle itself the SIMD packing and the branches in the code.
Everything is in textures to cache the access to the kd-tree nodes, the triangles and so on. This is much simpler and faster than coalesing accesses to global memory since no synchro are needed.
KILL ALL UNIONS. Unions are bad even in C (and forbidden in C++ as soon as there is a : constructor). This command goes with this one: KILL ALL ARRAYS unless you put it on shared memory. The cuda compiler seems to put stack allocated array (stg like float my_array[4]) on local memory and accesses will just kill your app. It may work with indirection known at compile time but otherwise, this is just BAD.
Use dirty switches. In my implemention, I need to perform circular permutations on the coordinates (due to the kd-tree traversal and the Baduel / Wald ray-tri intersection). I tested two approaches: an shared array strategy with a simple modulo computation where all the rays are put in shared memory, and a coding horror switch permutation (see the code). The second one was faster. The reason seems to be that in this case, a register allocation (allocating rays on the stack requires 6 more registers) is more efficient that a shared array allocation.

Performance

The performance are quite good. I am currently developing on a laptop with the low end GeForce 8400GS but made some preliminary tests on a 8800 GT. With the last optimizations I made, I think that between 12 millions rays / s and 40 millions rays/s may be computed on a GeForce 8800GT on the demo I give in the code. I will test soon the code on a GTX280 and without being too optimistic, I think that 100 millions rays / sec may be generated on it. This may be one of the fastest ray tracers wrote on a single processor.
The other good point is that the memory footprint is much smaller that a CPU ray tracer. Since the GPU is better at computing than at traversing large structures, a small kd-tree has the best performance. For 180,000 triangles, only 5MB are required for the kd-tree.

About Cuda

After the first deceptions due to my poor knowledge of the API and the compiler behaviour, I think that Cuda is a nice API, easy to use. G80 architectures are quite impressive and for a non-friendly application like ray tracing, speed is pretty good with no need to write one line of SIMD code. The bad point is that there is no way to have a real low-level access. After this nice experience, I am really eager to see what the upcoming Larrabee will be able to do with ray tracing. The good point is that Intel will certaintly provide both a high level API for fast prototyping (a la Cuda or the promising OpenCL) and a standard compiler / assembler since Larrabee cores are just x86s.

**Update** Source code is ... no more available on the previously posted link. I think you may find somewhere on ompf.org

Hello world

2008-07-29T15:37:00.000-07:00

Hello boys and girls.
I decided to open a Blog to keep on coding personal small projects, reading new stuffs and learning many things on computer science. This blog is therefore a pratical way for me to present small source codes or demos I am working on. I really hope you will appreciate it. By the way, you may still find my old home page from my PhD: there are some old stuffs on it which may be a bit interesting. The URL:
http://www710.univ-lyon1.fr/~bsegovia/