OpenCL GPGPU real-time ray tracing

OpenCL is the new open standard for GPGPU (general-purpose computing on graphics processing units) programming. It supports both graphics cards and multi-core processors with the same C99-based source code. I've re-written the NVIDIA CUDA version of my real-time ray tracer using OpenCL, and run it on an ATI/AMD Radeon HD 5870 GPU. For comparison I also ran it in CPU mode on the Intel i7-980X (6 core, hyperthreaded) processor.



The OpenCL CPU version provided similar results to the pthreads-generic version tested earlier. With 10 objects in the scene the GPU performance was only slightly better than the CPU version. But as the number of objects increased the GPU version maintained performance much better than the CPU version. This indicates that with low object counts the actual work done by each GPU core is small relative to the setup and copying results back overheads. Yes, a realistic ray tracer would be culling the scene graph so fewer objects would be evaluated for each point. And the GPU version ran out of constant memory above 800 objects. But this is much better scaling and performance than the NVIDIA CUDA test I did a few years ago.

Hyper-Threading

I've re-run my real-time ray tracing and Go ACO TSP programs on a 6 core processor with hyper-threading. The Intel i7-980X supports two hardware threads per core, so the OS sees it as a 12-core processor.

Here hyper-threading has improved performance for both programs in all configurations. At one core both programs were sped up 28%. At 6 cores ray tracing was sped up 9% versus non-HT, while ACO TSP was sped up 18% versus non-HT. The higher speedup for ACO TSP is expected as it matches the general speedup from increased cores.

Both programs were run with 24 threads in all test configurations. Having more program threads than cores*hardware_threads_per_core is key. Otherwise the OS may schedule multiple threads on the same core while leaving others idle. Linux is hyper-theading aware and attempts to avoid this scheduling problem, but doesn't always get it right. In testing with 6 threads and 6 cores, I saw 28% decreases in performance for both programs when hyper-threading was enabled.

The Intel i7-980X also supports "Turbo Boost", where one core is automatically overclocked when other cores are idle. Turbo boost was disabled for the tests above. With one core and HT disabled, Turbo boost provides a speedup of 8% for both programs. Surprisingly, with six cores and HT enabled (so all cores should be fully utilized) Turbo boost provides a speedup of 5% for both programs.

The performance of Go has also improved relative to the single-threaded C version of ACO TSP. For 24 ants the C version finishes in 60 seconds, so on one core Go is only 1.7x slower than C. At 2 cores Go has exceeded the performance of the single-threaded C version. Previously Go was 2.8x slower than C on one core; I don't know whether this improvement is due to an improved Go compiler, the move to 64-bit compilers and OS, or architectural differences between the Intel Core 2 and i7-980X.

Multi-Core Ant Colony Optimization for TSP in Go

Go is a new statically typed, garbage collected, concurrent programming language. It is native compiled and offers near-C performance. Like Erlang it uses a CSP model of lightweight processes ("goroutines") communicating via message passing. Messages can be synchronous (unbuffered) or asynchronous. Go supports multi-core processors.

On my usual ACO TSP test case Go is the best performing language so far. On a single core it was 2.8x slower than C, beating the previous best Shedskin Python. Go's speedup from 1 to 2 cores was 1.9, matching previous best Erlang. With 4 cores Go exceeds C's single core performance -- the first tested language to achieve this goal.



Go's lightweight processes are scheduled onto OS threads by the Go runtime system. This allows the system to support many more processes. I was able to run ACO TSP with 100,000 ant processes. The runtime system scheduler appears to treat these processes somewhat like futures, lazily scheduling them when another process is blocking on receipt of a message from a shared channel. This is inefficient for the type of parallel processing used in ACO (many seconds of parallel computation ending with a single message send), as only a single process (and thus a single core) is active at a time. This is a known issue, and the simple solution suggested on the golang-nuts mailing list is to add runtime.Gosched() calls to yield the processor. For my case this was insufficient, and adding additional heartbeat messages to each process helped force maximal multi-core usage.

See here for code and more details.

CUDA GPGPU on MacBook Pro Laptop

I've ported my pthread generic and CUDA GPGPU real-time ray tracers to Mac OS X running on a MacBook Pro laptop. It has a 2.4GHz Intel Core 2 Duo processor and a NVIDIA GeForce 9600M GT graphics processor. The Performance for 1 and 2 software threads is comparable to my older Linux desktop, and the CUDA is within a factor of 3 of the desktop graphics card.

Threads    Frames Per Second
1           8.5
2          17.4
GPU        16.7

Citations and Uses of Ray Tracer Code

The real-time ray tracing code I wrote for PS3 Cell and CUDA GPGPU is being used by others. David Yuen's group at the Department of Geology and Geophysics and Minnesota Supercomputing Institute at the University of Minnesota used my PS3 code as the starting point for tsunami modeling. From their IEEE Vis 2007 poster session Use of Ray Tracing Techniques on Tsunami Simulation Data with the PlayStation® 3:

"This poster will discuss how we have used the PlayStation® 3’s cell processor to implement and employ parallelized visualization, through the form of ray tracing techniques, on tsunami data. By adapting Eric Rollins’ ray tracing package, we were able to implement his ray casting techniques on triangles. This allowed us to triangulate our tsunami datasets, stored as [x, y, height] data, and feed the triangulated data into the program. We rendered 4000 triangle objects from the data and applied a colormap based on each triangle’s height value. Finally, we displayed the generated image."

They also published a paper Experiments in scientific computation on the PlayStation 3 in the Springer journal Visual Geosciences. This PowerPoint presentation shows images from my and their work.

Other students are trying CUDA GPGPU ray tracing. Michael Allgyer (DOC) cited my work in his masters thesis proposal. Umut Erturk asked me about PS3 versus GPGPU for his thesis. Others are trying to get my CUDA code working on windows.

C baseline for ACO TSP

In my previous post on the Multi-Core Problem I discussed comparing the performance of multi-core-friendly languages to the baseline performance of single-threaded C code. The test case I have been using is Ant Colony Optimization for TSP. In a recent post Leonardo Maffi has rewritten my ACO TSP code in C. This provides a baseline for comparisons:

           Time   Multiple   Speedup    
        in secs       of C    (1->2)          
C           0.4          1       n/a
Scala      10.0         25       1.7
Haskell    22.0         55       1.4
Erlang     93.0        233       1.9

gcc (GCC) 4.2.4 (Debian 4.2.4-1).
Compiled gcc -Wall -O3 -s ant_c.c -o ant_c
Still 1000 iterations, 200 cities.
Times for other languages from here .
Initial speedups from here (Scala interpolated), and do not model decay at higher numbers of cores.

The Multi-Core Problem

"Since 2002, the limits of power, available instruction-level parallelism, and long memory latency have slowed uniprocessor performance [growth] recently, to about 20% per year [down from 52% per year]." Hennessy and Patterson, Computer Architecture: A Quantitative Approach, Fourth Edition (2007), p.3.

Post-2002 the Moore's Law growth in transistor count has been increasingly devoted to adding additional cores to CPUs, rather than trying to improve single-core performance. In the prior decade single-threaded C/C++ applications were sped up "for free" with successive CPU generations. This free ride has ended, and left programmers with a choice. Either they can rewrite their single-threaded C/C++ applications to be multi-threaded (either explicitly, with pthreads, or more implicitly, say with OpenMP), or they can rewrite their applications in a different more multi-core-friendly programming language.

In choosing an alternative language two important related factors must be examined. First, how much of a single-threaded performance hit is taken moving to the new language? If the new language is say 40 times slower than C/C++, you will need a 40-core CPU to get back to where you started (assuming perfect linear speedup). Second, how much worse than linear is the actual speedup? If the speedup starts out at 1.5 with 2 cores, and drops below 1 at 8 cores, the new language will probably never catch single-threaded C/C++.

Multi-Core Ant Colony Optimization for TSP in Scala

I've added a new page Multi-Core Ant Colony Optimization for TSP in Scala.

Scala is a hybrid functional and object-oriented programming language. It is based on the Java JVM, so it has available Just-in-Time (JIT) native compilation and native thread support. It has access to the wide variety of class libraries available for Java.

Like Erlang, Scala supports an Actor-style message-passing concurrency model. More traditional thread-based concurrency with locks and shared memory is also available. Unlike Erlang and Haskell, Scala values are not immutable. Like SML and Haskell, typing is static with type inference.



Erlang and Alice SML omitted from graph. Note change from previous tests: 1000 interations, instead of 100. Still 200 cities.

Cores  Erlang  Haskell  Scala  OCaml  MLton  Alice
                                        SML    SML
    1     372       86   41.6           6.0    676
    2     192       63
    3     141       58
    4      98       54   14.6    1.6

All times in seconds.

If anyone knows how to reduce the number of cores used by the JVM, I'll fill in the rest of the table for Scala.

Functional Languages on Quad-Core (revised)

Yesterday's post generated lots of suggestions. I've revised the Erlang code to use dict mapping and updating functions. I didn't use the new Erlang arrays since they are one-dimensional and I need two dimensional arrays. Here are the new results:

Cores   Erlang   Haskell   SML
1       38.4     8.6       0.76
2       19.6     6.1
3       14.5     5.5
4       10.3     5.1

Here is the revised Erlang source code, as well as the Haskell quad-core source code used in tests. The MLton Standard ML implementation is still single-core, as MLton does not have SMP support (see comments below).

Update: Added new post with Scala implementation.

Functional Languages on Quad-Core

In an earlier series of articles ending with Multi-Core Ant Colony Optimization for TSP in Erlang I evaluated the performance of the Standard ML, Haskell, and Erlang functional programming languages using a test problem of solving the traveling salesman problem using ant colony optimization. At the time I only had available a dual-core machine. I have re-run the tests using an Intel quad-core machine and the latest language versions available for Debian Linux.



The results match my previous summary: Erlang is much slower than the others. Erlang is a bytecode interpreted virtual machine while the others are native code compilers. Erlang is also dynamically typed while the others are statically typed. Like Haskell, Erlang has immutable (single assignment) variables which force lots of copying. Unlike Standard ML and Haskell, no array type is provided so I had to inefficiently implement them using dictionaries. In Erlang each process (thread) does have a separate heap, so there is little interference between processes due to memory allocation or garbage collection. Near perfect speedup is achieved.



The CPU utilization graph (Erlang, then Haskell, then Standard ML) does show the lack of interference between Erlang threads which makes its speedup possible. Haskell achieves little speedup at higher numbers of cores.

Update: thanks, responses to comments:

As discussed at the end of an earlier posting, based on email feedback I did try HiPE compilation with Erlang 5.5 on AMD. It provided no significant improvement, and was not compatible with multi-core operation. This time I again tried HiPE compilation with Erlang 5.6.2, and get the following with erl:

$erl
Erlang (BEAM) emulator version 5.6.2 [source] [smp:4] 
  [async-threads:0] [kernel-poll:false]

Eshell V5.6.2  (abort with ^G)
1>  c(ant,[native]).
./ant.erl:none: Warning: this system is not configured for
  native-code compilation.
{ok,ant}
2> 

Same for erlc:

$erlc +native ant.erl
./ant.erl:none: Warning: this system is not configured for 
  native-code compilation.

So it appears the latest i386 Erlang available from Debian Linux apt-get doesn't support native compilation, probably because they chose to enable SMP support instead.

Erlang arrays are new in the latest release. Thanks for the info, I'll try them out.

Yes, MLton SML with one core was the fastest. Older (pre-quad-core) versions of the code are available linked from the various ACO TSP pages and here: Standard ML, Haskell, Erlang.

Actually the next language I will probably try is Scala. Based on the Java JVM, it may inherit JIT native compilation and good SMP thread support. It also supports an Actor concurrency model like Erlang. Apparently Caml dialects don't support SMP. Of course for many problems (like ACO TSP :-) you are better off simply starting N different fast single-threaded instances of the program and collecting the results in Ruby or some-such. I am really looking for instances where using a SMP-capable language is a clear win over single-threaded C++. I haven't written a single-threaded STL C++ version of ACO TSP; I expect it could easily crush Mlton SML once you finally got the bugs out...

Here are the raw numbers, in seconds:

Cores   Erlang   Haskell   SML
1       87.0     8.6       0.76
2       46.3     6.1
3       34.6     5.5
4       26.2     5.1

When I increase the iterations in Haskell by 5x the runtime goes up by 5x (43 secs with 1 core, 26 secs with 4 cores), still showing the same pattern of decreasing improvement with increased cores.

Update 2: see new post for revised Erlang code and results.

Update 3:
Turns out Erlang HiPE is available in Debian as a separate package (erlang-base-hipe), and is now SMP compatible. But the improvement is only 5%, or 11% if you also native compile dict.erl.

Progress in GPGPU Ray Tracing

I've been getting traffic and questions about my pages on Cell and CUDA ray tracing. Saarland University has been doing some interesting work on Cell and CUDA ray tracing including this paper: Realtime Ray Tracing on the GPU with BVH-based Packet Traversal.

OpenMP

Intel C++ supports OpenMP, a standardized API for shared-memory multiprocessing in C++. By adding a simple #pragma loops can be split up to automatically run in parallel on multiple cores. I added a simple change to the per-screen-line rendering loop of my generic ray tracer:

    int screenX;
#ifdef USE_OMP
#pragma omp parallel for firstprivate(portPoint)
#endif
    for (screenX = 0;
         screenX < SCREEN_WIDTH;
         screenX++) {

With this change each pixel is rendered by a separate thread, with a maximum of 4 (the number of cores) running simultaneously. This use of threads is an alternative to the explicit pthread threading tested previously, where a separate thread was used per line of the screen. In this case OpenMP is simpler than explicit threading, though it yields slightly lower performance.

GCC 4.2 also supports OpenMP, but its performance is much worse than Intel.

Intel C++ Compiler

I recompiled my generic Ray Tracer code using the Intel C++ compiler. Using SSE3 auto-vectorization I obtained a 21% speedup over GCC.

Cell PPC Hardware Threads

The PowerPC unit in the PS3 Cell processor supports two hardware threads. Earlier I added pthread support to the generic ray tracer code. I have updated Real-Time Ray Tracing on the Playstation 3 Cell Processor with PPC performance measurements with pthreads and the IBM XLC compiler. Pthreads gives a 63% speedup, and XLC gives a 26% speedup, and together they give an 89% speedup:

CUDA GPGPU and Pthreads

I've added a new page, Real-Time Ray Tracing with NVIDIA CUDA GPGPU and Intel Quad-Core. I've created two new versions of the ray-tracer code, one utilizing NVIDIA CUDA, and another using pthreads to exercise multi-core CPUs. The CUDA code is fastest, but least realistic as a practical ray-tracer. Pthreads was certainly easiest -- a half hour change to pre-existing generic code. I think CUDA will perform much better on more constrained (in terms of control flow and random memory access) algorithms.

PS3 Cell SDK 2.1

I've been getting some compatibility questions about my PS3 Ray Tracer code, so I just upgraded my PS3 to Fedora Core 6 and Cell SDK 2.1. IBM changed the SPU thread API (to pthreads); I now provide code for both versions. They also now have a PPC native version of the IBM XLC compiler. It provides an over 50% performance improvement over GCC:

PS3 Cell Ray Tracing at MIT, GPGPU at UI

Earlier this year I wrote an article on Real-Time Ray Tracing on the Playstation 3 Cell Processor. MIT is offering a course on Multicore Programming on the PS3. One of the examples presented is the Blue Steel Ray Tracer.

Earlier I also wrote an article on Four-Dimensional Cellular Automata Acceleration Using GPGPU. The University of Illinois at Urbana-Champaign is offering a course on Programming Massively Parallel Processors featuring GPGPU programming on NVIDIA CUDA hardware.

Currently applications must be written specifically for the Cell, conventional multi-core CPUs, or GPGPU. Several organizations are working to develop a common programming model where the same application code could be compiled to run on any of these. One such system is from RapidMind. It's programming model is much like the use of OpenGL within a C++ program. Arrays are defined using RapidMind container types (like OpenGL graphical types), and then embedded programs are invoked on them (like OpenGL Shading Language).

Mail Bag

My web site gets a surprising amount of interesting email (beyond the usual headhunter spam and spam spam). This may be due to certain search engines regarding me as the world's foremost authority on four dimensional cellular automata, among other topics :-)

One email was from an author working on updating a book I had cited. I gave him and his assistant pointers on the latest I knew about emergent collective intelligence.

Another email was someone looking for a part name and source for a CPU chip socket in a photo from my robot project -- not for a robot or a CPU, so I have no idea how he found the page. I was able to help him.

The most amusing email was a 1337 haxor claiming to have broken the copy protection on one of my Palm games, so he could "fix" it (our backgammon didn't cheat, honest!). I replied with this link; he replied with this one. Touche.

Another person found my board game list, and was looking for a copy of the Avalon Hill Titan board game to play ($150 on e-bay!). Sorry, not selling & no leads.

I was also asked whether I have tried a SAN with more than one Coraid EtherDrive. I have not.

The latest email suggested improvements to Multi-Core Ant Colony Optimization for TSP in Erlang. At the time I had tried native compilation in Erlang, but didn't see any speedup and didn't mention it in the article. On Sun SPARC the correspondent was able to get significant speedup (2x +) on my application with native compilation after he separately native compiled the dict module. Turns out the native compiler only operates on a specific module, and you must native compile the system modules yourself. Otherwise the native speedup is lost in the translation between native and bytecode interpreted in calls between modules. I tried native compiling dict on my AMD system, and didn't get any improvement. I Also (re)discovered that native compilation is not yet compatible with SMP (multi-core) operation. The correspondent also got significant improvements by altering some of my data structure choices. I expect that all three ant algorithm implementations (Erlang, Haskell, and Standard ML (twice)) could be improved be re-visiting my data structure choices.

Software Transactional Memory

I saw an interesting tech talk recently by someone from Intel on their work on Software Transactional Memory. A question from the audience pointed out an interesting problem with exceptions (also discussed under "Implementation issues" in the Wikipedia article). In STM threads perform optimistic reads, and are retried at commit time if conflicts are detected. Because the reads are "dirty" inconsistent states may be seen, and null pointer or division by zero exceptions may be thrown. The application programmer must recognize this possibility and distinguish between transient (should be retried) and fatal errors. STM trades off the complexity of locking with the complexity of understanding the effects of inconsistent reads. Opposite to the suggestion in Wikipedia, perhaps the system could check for conflicts at the point of exception, and automatically retry if one is found. Of course STM still doesn't work for non-restartable operations, such as I/O. I find the ACID (Atomic, Consistent, Isolated, Durable) status of STM transactions debatable (and advocates don't claim to satisfy all of them).