Atom

PC CPU Task Parallelism Limits

2010-01-03T00:05:00.002-06:00

All numbers below are from a Windows XP + 2.39 GHz Intel Core 2 Duo laptop. Performance on other machines and configurations will vary (duh).

Some Numbers

Histogram of the measured clock cycles/task of one thread processing the same ~10K clock cycle task many times (in a loop) on the above laptop,

835330 ... 0-16K cycles 8005 ..... 16K-64K cycles 1178 ..... 64K-256K cycles 66 ....... 256K-1M cycles 20 ....... 1M-4M cycles 5 ........ greater than 4M cycles

This specific task is a fully ALU bound dummy task which spins doing pure math (all values in registers). Run time variability is a function of hyper-threading (minor effect) and operating system preemption (major effect). Time scale here is somewhere around 4 seconds of total run time (for 0.8M tasks), which could be 120 frames of a 30 Hz game on this laptop. Note the variability.

Preemption

A major problem for any task parallel system on the CPU is preemption at non-task boundaries. The net result is that tasks at random could easily get stalled for over 1 millisecond (happened 5 times in 4 seconds in the above example). Any tasks which depend on the result of those stalled tasks also get stalled.

The visible result is variable frame rates and a screwed over player/user experience when the application attempts to have more than a few task parallel dependencies per video frame.

Workarounds

The best solution is cooperative multitasking: tasks release the CPU when they are finished, programmers insure task run time is at the desired program response latency. As desktop operating systems have de-evolved over the years we have lost the ability to provide a correct solution.

Are there any workarounds?

UNDER-UTILIZE CPU - In order to always hit v-sync, the program can under-utilize the CPU such that the maximum amount of time taking to compute a frame is the minimum amount of time in the worst case preemption. Clearly preemption is not a bounded problem, a program will always have to accept some kind of frame drop.

DUPLICATE TASKS - With a painful case of over-engineering, one could build a task system with transactional memory, such that if a currently needed task dependency is not completed due to preemption, that a running thread could manually duplicate and run the dependent task.

Anything less painful?

Manual Self Preemption?

Seemed like a great idea: keep track of thread run time and yield execution to another task worker thread before the operating system would preempt the thread's time-slice. This way the program could simulate cooperative multitasking switching at task boundaries.

Epic fail on the XP Laptop!

Of the two ways to do this with fixed thread affinity: pair of threads yielding execution, or pair of threads blocking/releasing each other via signals, all have high overhead and do not fully solve the variability problem.

So anything less painful?

No!

DirectCompute 3D Fluid Simulation Tube

2009-12-29T20:07:00.002-06:00

Download Page for Jan Vlietinck's Simulation Demo

Random Pre-Christmas

2009-12-13T13:32:00.002-06:00

Tony Albrecht (SCEE): Pitfalls of Object Oriented Programming
One of the best I've seen on this topic. See more SCEE presentations on SCEE's Presentation Page.

CGW on Avatar Production

Unchared 2 Art on Z Brush Central
Thanks to @BrianKaris for this link.

PixelJunk Shooter
Is out now and is quite fun/awesome!

Parallel Beast

2009-12-08T22:48:00.005-06:00

Been too long since I've posted anything about parallel programming, considering my addiction with the topic. Guess I should also make this clear, that what I post here is my (sometimes radical) personal view and not necessarily representative of my employer.

Perpetual "Computational" Burnout

Burnout as in data moves around and around in a circle, but just cannot seem to put the power to any useful computation. Bill Dally's SC09: The Future of GPU Computing Presentation puts this to numbers on page 14,

"Moving a word across die = 10FMAs"
"Moving a word off chip = 20FMAs"

Data Routing and Work Scheduling Problem

Efficient parallel programming is effectively a data routing and work scheduling problem, balancing the cost of moving data and computation. Chip is effectively a static routed network, going parallel involves reappropriating chip area towards computation at the cost of something else. Massively parallel machines will continue to enable algorithms which take advantage of limiting data movement to scale in performance.

Scheduling is of critical importance because it insures data is reused instead of being moved to and from a holding area as a result of inefficient ordering of computation.

Memory Hierarchy

Furthermore as the balance of chip area enables more computation, the area devoted to holding data is used more for storing active computation rather than simply storing data. Area devoted to holding data includes registers, cache, and/or program managed local store.

The difference between a massively scalar processor and a massively parallel processor can be visualized by thinking of the on-chip memory hierarchy pyramid. At the tip of the massively scalar processor pyramid is the registers and at the bottom is L3. In contrast on a massively parallel processor the registers might be the widest part of the pyramid (inverted pyramid)!

The life time of the data in this holding area is greatly reduced as throughput increases. This places even more importance on scheduling, data must be reused in the ever smaller window of opportunity. The roll played by the memory hierarchy changes from holding data for a long time for use in much later computations, to more of a roll of smoothing out a short window of access patterns.

What this literally means is that to be efficiently parallel at a fine granularity, work must be explicitly scheduled group reusing data. Anything fine granularity and dynamic comes at a cost. The reason for vectorization and grouping vectors into blocks is part way to insure the tight scheduling required to group data access to common data.

Medium Granularity Scheduling: OS Hand of Perf Death

Preemption is a nightmare for the latency of task parallel programming dependencies. When the scheduler preempts your worker thread while it is in the middle of processing a task, dependent tasks wait on the rescheduling of that thread so the task can finish. Often that thread was just placed on the end of a queue, so there is going to be a good delay before that thread will run again. What is desired in this case is a cooperative or task aware scheduling, so that a thread will yield to another thread at a task boundary. Working around these system level problems causes a cascade of design complexities.

What I want personally in the future of Operating Systems and massively parallel machines is an OS which can allocate entire cores and memory to a application for exclusive access! Take your awful mess and get out of the way. Let me return to the days of microsecond and nanoseconds response time instead of the abysmal peaks of milliseconds of latency on scheduling.

Actually no need to wait, can do this right now.

Raw Programming a Parallel Beast

Hidden under the guise of a GPU is an awesome parallel machine, and a driver which provides chunks of exclusive access to the entire machine, and during those periods of exclusive access, no operating system or driver to screw anything up.

Getting Naked With the Parallel Beast!

Dead simple to play in CUDA and a GT200 series GPU: launch enough warps via 32xN sized CTAs to fill the machine just once using a Cx1 size grid. Where C is the number of CTAs to fill the machine. The result is N*C "hardware threads" (as in a traditional CPU) with 32-wide vector units, texture units, and unified memory! Use nvcc's "--maxrregcount" to set the maximum number of registers per warp such that C CTAs will correctly fill machine (and be careful with shared memory usage).

Use clock() to insure the warps exit at a task aware boundary at some Hz, then simply repeatedly relaunch the same kernel to continue the machine. Allocate a huge chunk of GPU memory to work within without fear that accessing that memory might ever incur a page fault. Use page-locked host memory to roll your own queue to communicate to and from the CPU for keyboard and mouse IO. Do a little OpenGL interop for drawing to the display, and you have a 100% software massively parallel machine to prototype or code on.

Scalar code in CUDA is easy, simply write C code and don't explicitly use "threadIdx.x" to have the code adapt to the lane of the 32-wide vector. Use a mix of scalar and vector code at will.

The launch kernel becomes the binary of the "operating system" and "application". Use "threadIdx.y + N*blockIdx.x" to obtain an index to the "hardware thread". CUDA 2.3 PTX ISA does NOT provide branch or call by register, workaround to provide a data dependent branch is a giant switch (GPU's actually a lot better at it than you think it is). Use global memory atomics to communicate between warps or CTAs.

Could run N*C instances of something like Contiki and do message passing through global memory. Or could test data structures for 100's of "cpu thread" scaling. Perfect for those who want to return to the glory days of 64KB but with 1Tflop and 1GB of memory ... 320x200 at 30 fps provides 0.5 Mflops per pixel per frame on a GTX 275. What's not to love about that!

Imagination is your limit.

Demo Tube 2

2009-12-08T00:52:00.002-06:00

PixelJunk Shooter Trailer

2009-12-03T13:07:00.002-06:00

PixelJunk Shooter trailer shows off more fluid effects and in-game awesomeness!

GigaVoxels Mandelbulb

2009-12-02T11:36:00.002-06:00

GigaBroccoli: The Mandelbulb into GigaVoxels Blog Post

Really cool work by Cyril, this time run-time on demand generation of 3D geometry into the GigaVoxels brick cache. See the blog link above for more info, videos, and screen shots.

PixelJunk Shooter Dec 10th!

2009-12-02T11:31:00.002-06:00

PixelJunk Shooter Thursday, December 10th

Real-Life "CFD Visualization"

2009-12-01T19:44:00.006-06:00

Longer exposure night shot of the sparks from a backyard bonfire visiting family on Thanksgiving. Thanks to my wife Kathryn for the photograph, and my nephew Joe for motivation required to take the time to burn wet firewood,

PixelJunk Shooter Youtubes

2009-12-01T13:06:00.003-06:00

Battlefield: Bad Company 2

2009-11-24T03:04:00.003-06:00

Been playing the PS3 beta because it is an awfully addictively fun game. Don't think I will ever want to go back to a game with a trivially static play area again, no matter how pretty the developer makes it. Baked global illumination, forget it, just not good enough any more. Nothing compares to the fun factor of being able to leave your mark, to be able to interact with your surroundings, and to do that in a tank!

The point of excellence in game design is when a game provides more than just the sum of canned content generated in production. Otherwise a game is simply an interactive movie: you play to see the next scene, and when the end is found, there is not anything more to it. Sure there is the market for the interactive novel, but games since the dawn of humans have been about enjoying the interaction between people in a system engineered to produce fun in the process. By this definition, BBC2 multi-player beta definitely captures the essence of what it is to be a game.

Other devs out there should be taking notes.

You've got greater than 1 Tflop to play with on high end GPUs in the PC space, over 100,000 times more processing power than many of us started programming with, and more to come in future years.

What are you going to do with that?

Many of you are simply going to increase your resolution, provide a more pretty interactive novel. Others are going to push forward with something new, provide gamers with an interactive experience which brings renewed energy into the industry, and learn to wield the scaling of the parallel machine for something beyond just graphics!

Real Time Global Illumination Using Temporal Coherence

2009-11-20T03:10:00.002-06:00

Martin Knecht has posted a video and his thesis on Real Time Global Illumination Using Temporal Coherence.

Time Flys and the Blog is Lonely

2009-11-19T18:39:00.005-06:00

Too busy to write but links of interesting stuff,

Johan Andersson's Parallel Futures of a Game Engine
Direct To Video on Deferred Rendering in Frameranger
Sander van Rossen's Progress on Virtual Texturing
Aras's Post on Deferred Shadow Maps
Stochastic Progressive Photon Mapping by T. Hachisuka and H. W. Jensen
Fermi spotted on Twitter
SC09 Papers
State-Based Scripting in Uncharted 2: Among Thieves
Quilez's Mandelblub

Demo Tube

2009-11-07T18:38:00.002-06:00

Link Soup

2009-11-01T18:48:00.002-06:00

meshula.net : Stone Soup - Thanks Nick for this post reminiscing of the exotic time before the common C compiler!

Real-time Parallel Hashing on the GPU - Neat paper, and poster at GTC. Builds upon Cuckoo hashing (using N>=2 hash functions instead of one, lookup requires at most a check of N places, insertion requires recursive eviction and insertion of the filled bin using the other hash functions), cuckoo hashing for N=3 hash functions can almost achieve 90% hash table occupancy. Paper presents parallel cuckoo hashing. First step is to use a high level hash function to divide the input into bins sized to fit into the local store, followed by parallel version of cuckoo hashing for the bin. Read the paper for more details.

Stochastic Progressive Photon Mapping - PPM extended to compute the average radiance over a region instead of a point.

Amortized Supersampling - Interesting, still problems with rapidly changing shading such as specular highlights, but great progress towards a complete solution.

Rudebox by Alcatrax

Jellyfish

Visited the Monterey Bay Aquarium the week before Halloween,

Random End of October

2009-10-28T22:56:00.002-05:00

Looks like Insomniac updated their website: Link to R&D Page. Going to have to finish Uncharted 2 before I get onto A Crack in Time however. Been meaning to watch Nürburgring 24 Hour Race in 3D.

GPU Technology Conference

2009-10-23T16:11:00.003-05:00

GPU Technology Conference Main Page | GTC Blog
Screen-casts plus audio recordings for the sessions are now starting to get posted. Look in the session catalog/calendar for the links.

1489% on 8-core machine

2009-10-20T10:14:00.002-05:00

No not 14.89x faster, but rather 14.89x slower!

False sharing is no fun from Joe Duffy's Weblog. Really like this post because of the simplicity as to which it shows the problem of sharing cache lines on the typical multi-core CPU.

If one were to assume that you'd always be exactly 14.89 times slower when sharing cache lines (clearly this would vary) across 8 other cores, and you wanted this slow down to only add 6.25% (or 1/16) time overhead to your code, you would want to only share cache lines about 0.45% of the time.

Actual number here is not the point (and wouldn't be correct anyway), but rather that on the CPU, parallel performance is found by doing a majority of computation in effectively isolated memory regions.

Naumachia

2009-10-13T10:19:00.003-05:00

Found this thanks to a twitter post from Jerome Liard (@blackjero) at Q-Games.

naumachia.aureasection.com

The "Kill Zone 2" effect (particle and transparent effects really make the engine) applied to 3D space combat. Impressive work from a team of just three developers!

Fermi

2009-09-30T23:56:00.004-05:00

NVIDIA Fermi Architecture Page

Download the white paper for some architectural details.

Personal View

Yeah the architecture is awesome! For HPC, double precision fused multiply add at 1/2 the single precision rate, ECC memory support, and 40-bit address space support! For compute, configurable shared memory/L1 split, function and data pointer support via a unified address space, faster global atomics, unified L2 and more...

Last Night in Wisconsin

2009-09-25T00:11:00.002-05:00

Yeah, last night in Wisconsin. Could talk about how I've actually physically moved out of Wisconsin and back to my Illinois Condo the weekend before September 11th when my lease was up, and how that has resulted in not having the ability to do any interesting CUDA and GPU computing programming for the past few weeks, but likely that would be too boring!

Prior Dormant Hobby

Had a 1969 911E which I started to build into a track car years ago and never finished. Sold off the car, trailer, and related stuff this Wednesday. Since I had this huge stalled project (full tear down and rebuild of the entire car), the racing bug was effectively squashed, the wife would have killed me if I bought another car without selling the old project. Actually I ended up doing just that! Made matters worse that the 911E body and trailer was in her parents barn. Also bad that I took her WRX to the track a few times and destroyed the rotors which I didn't fix for a year which ultimately resulted in her taking the car in to get the brakes done locally. Kathryn, you really are the best wife ever! Yes, she reads this blog once and a while ... I'll get an email saying, "I actually understood that post" ;)

With the project car gone, and a very track friendly daily driver to replace it, the racing bug is biting back hard.

Found that video via this rx7club forum post while looking for someone who is using a G-Force Transmission dog box in a road racing track car.

Need to do the brakes, brake lines, harness, and more first, but a G-Force dog box conversion of the Tremec T-56 in my car, is on the list long before adding more HP. Yes it is still going to be my daily driver, no the sound doesn't bother me, and yes one of my old daily drivers didn't have heat, air conditioning, or an interior (so safe to classify me as crazy with respect to what I drive).

Looks like Thunder Hill Raceway is less than a 3 hour drive from where I will be living...

Function 09: Behind Elevated

2009-09-24T22:52:00.002-05:00

Function 09 - Behind Elevated

Pixel Junk Shooter: TGS 09 Gameplay!

2009-09-24T22:35:00.002-05:00

More videos and images can be found on Gigazine

R800

2009-09-23T14:41:00.009-05:00

R800 info is out, architecture looks similar to R700. Hard to get a clear answer on if texture filtering has indeed moved completely to the SIMD units (EDIT: perhaps not, see new comments/edits). Also have not seen any definitive answers with regards to ~~triangle setup performance and~~ info on append/consume performance. HD 5870 specs appear to be as follows,

EDIT: hardware.fr tests show same triangle setup rate as R700. No parallel triangle setup, but perhaps double samples rasterized per triangle per clock. Very interesting indeed!

EDIT: ATI's Dave Baumann on B3D post: "Texture interpolators have been removed from the design and is done on the shader core. In general we are seeing this as a performance improvement - its also the reason why one of the Vantage feature test gets a disproportionate increase over the previous gen."

OVERALL - 2.7 Tflops single precision. - 544 Gflops double precision. - 153.6 GB/sec bandwidth. - 20 SIMD cores. TEXTURE - 272 Gtex/sec 32-bit pixel unfiltered samples. - 68 Gtex/sec bilinear filtering. - 1 TB/sec L1 texture cache. - 16KB TEX L1 per SIMD core (320KB total). - Texture units able to read compressed AA color buffers. - Texture filtering done in SIMD units (according to AnandTech)??? COMPUTE - 8KB extra compute L1 per SIMD core (160KB total). - 32KB local data store (640KB total). - 64KB global data store. MC/RBE/ROP - 435 GB/sec L2 cache bandwidth. - 128KB L2 cache per memory controller (512KB total). - 4 64-bit memory controllers (MCs). - Render Back-Ends (RBEs) can process 32 pixel/clk. - RBEs look to be divided across the 4 MCs. - Fast color clears. SIMD - EDIT: vertex attribute interpolation done in SIMD units. - EDIT: Support for full speed 24-bit integer math. - Dedicated Sum of Absolute Differences instruction. - Faster Dot Product instruction. - Required support for DX11 instructions. - Full speed denormals?

Bad Industry Humor: To 12 MPix and Beyond!

2009-09-21T21:56:00.008-05:00

Second post in the awfully-bad-industry-humor or would-have-been-better-off-sleeping catagory,

"At first we thought it would be a great idea, the new ATI GPU was to have something like 10x the bandwidth and 20x the compute power of the Xbox 360. So what was working at 720p on the 360 should work about the same on a six 1080p display array off the new GPU. Napkin math said peak capacity amortized to around 8 Kflops/pixel and 800 bytes/pixel bandwidth at 720p/30fps on the 360 and about the same targets on the 12 MPix (multi-monitor) output of the new GPU.

The project would be simple ... (just use nearly the same engine, add tessellation and up the texture quality) ... or so we thought.

Early on, decided to go with virtual texturing. With deferred shading, needed about 380 MB alone for the G-buffer (2xMSAA 16B/pixel). Another minimum of 1 GB would be needed for four layers of a compressed 16Kx16K virtual texture. This virtual texture would provide 21 four layer texels per screen pixel (a great ratio).

That is when the problems started.

Artists used to work with 2Kx2K source stamp textures at maximum size. Now they were required to work at 8Kx8K just to be safe. This required an upgrade of all the artists machines to 64-bit Photoshop CS4 to support enough address space so they could have more than 16 layers active in Photoshop without swapping to disk (we use 16-bit per channel source images). Had to get rid of all the Mac machines because 64-bit wasn't supported on the Mac with Photoshop.

Old source material just would not cut it with the new 64 Mpix texture resolution requirement. We had to purchase a few 40 Mpix Phase One digital camera backs and special medium format camera gear just to gather good source photos (at a cost of about $30-40K US per each group of camera/back/lenses/etc). Also had to hire a special photographer to help train the artists on the new gear (teach details like how to choose the proper aperture for sharp photos, etc).

Furthermore our publisher would only allow us to use four dual-layer Blue-Ray discs, which is a measly 200 GB of data. The art department was outraged, what is the point of having 1 to 2TB desktop drives when the game can only use 200 GB of space? Had something to do with the publisher saying that 12x speed BR drives would require a little over a one hour install to the HDD, and that was the upper maximum time that the average user would stand when every 16 to 20 minutes they had to put in another BR disk to complete the install.

And don't get me started on the required upgrade to our in-house GI farm, storage solution, and network ... lets just say re-baking the lighting and virtual texture was almost measured in weeks not hours.

In the end the project got canned and we ran out of money, going for 10x the resolution and quality on 2x the budget of a console title just was not possible yet. We never should have bet on technology that was going to only sustain or reduce the bandwidth (or compute) to pixel ratio ... or worse yet, require hanging an array of monitors from the ceiling.

Later we heard of another post DX11 launch title going the opposite direction, doing something VGA retro, real-time photon mapping at 320x200 at 30Hz with motion blur and a real camera lens model. The graphics engine was called something like The Turing Engine because most viewers could not tell the difference between the real-time 320x200 rendering and a VHS tape of similar real life source video. Something like 2 Mflops/pixel and 100 KB/pixel of bandwidth at 30 fps, damn, why didn't I think of that!"

If for a second you took any of that seriously, you really need to fill up on coffee before hitting the early morning blog run!

CEDEC 2009: Anti-Downsized Buffer Artifacts

2009-09-18T02:10:00.003-05:00

CEDEC 2009: Anti-Downsized Buffer Artifacts
"The page you requested was too large to translate."
Anyone have a good summary?

Food-For-Thought Supermarket Aisle

2009-09-17T21:16:00.005-05:00

This post is to keep the twitter peanut gallery happy ;) Thanks Ivan-Assen Ivanov for the inspiration, sure was a lot more fun then doing something productive like continuing my search for a place to live when I move...

Starting With Some Simple Examples

Q: If you are playing chess with a distant opponent via snail mail, and want to play the game faster than the US postal service can deliver messages, how do you do it?

A: You send not only your next move, but also your response to all of your opponent's possible next moves, under agreement neither player will cheat.

In the above case, latency is the bottleneck, compute time is abused with redundant work to solve the latency problem.

Q: A pre-gunpowder age general wants to communicate a real-time updates of highly complex battle field commands to distant field officers but only has a limited set of flags, how is this done?

A: General and field officers pre-agree upon a complex battle plan option (or meaning) per flag.

Bandwidth is the bottleneck, solved by taking advantage of local memory using the ultimate form of compression.

Q: A manager has 12 phone meetings in a given very busy day, how is this handled efficiently?

A: A secretary pre-schedules each phone meeting.

This is a rough example of the problem of task preemptive scheduling (interrupted phone calls) when there are dependencies (communication) between nodes.

Humans Are Natural Parallel Programmers

Ironically, humans as a group are great at solving problems with distributed memory, as individuals have only a subset of the knowledge stored in the continuation of the species. Many of us, myself included, couldn't possibly remember everything we need to know on a daily basis. In fact, often we get better at remembering where we can find information (indexes or addresses) than we are at remembering the information itself. Humans organize into hierarchies to solve problems, as well as pipeline processes, think about manufacturing (thanks Mike Acton), supply chain management, etc.

The point I am getting at here, is that going parallel is natural and is easy to think about when grounding the problem in terms of something real.

Human Species Machine

How about something more crazy...

Massively parallel programming gets rather interesting when thinking about distributed memory, non-coherent redundancy, the problem of node failure, message passing, latency, and compression, all together.

One fantastic quality of humans is that we all have different memories of similar events or ideas, a kind of distributed non-coherent redundancy. The memory of an event (or idea) is physically distributed across many nodes, with great overlap in the "important" data and less overlap in details which are more important to individual nodes.

People live and die (node failure) and pass on ideas though ripples left in the fabric of this non-coherent distributed memory. Human language and culture provides a huge distributed context as to which to compress data for message passing. A single phrase can draw upon this giant context for meaning.

Stability Without Coherent Memory

Simple patterns can serve to generate intrinsic randomness (divergence), or can serve to maintain and safely transmit information (convergence). Digital computation is a great example of this (given the underlining "analog" flow of electrons).

With the human species machine, constant message passing serves to shape and change the distributed memory. Ideas and information become stable when replicated across many nodes. The more parallel the machine, and the more non-coherent, distributed, and redundant the memory gets, the less a single node failure matters. In fact, given the amount of message passing, the information is constantly in flux at the scale of individual nodes, but as a global whole, a stable idea is highly coherent collectively.

BTW, if you are still reading this post, I would highly suggest reading Stephen Wolfram's New Kind of Science (at least twice)!

NVIDIA and Moving

2009-09-16T07:37:00.005-05:00

Moving out to Silicon Valley soon to start a new job at NVIDIA!

Clearly the blog will be changing, should be obvious what I will NOT be writing about. Also more excited about what I will be working on at my new job than anything I would have worked on at home. Future "at home / free time" focus will likely be shifting towards actual technical papers or presentations or demos.

Soon Farewell to Wisconsin

Had a great time working many of the people at Humanhead: Brian, Jimmy, Ash, Giovanni, Michael, Norm, and more.

Also have a lot of great memories of Madison and Wisconsin: Lao Laan Xang (favorite local restaurant), Ford's Gym, Devils Lake (hiking, climbing), Blue Mound, Farmer's Market on the Square, EAA Airshow, Apostle Islands, and Road America (race track).

Did not even scratch the surface of everything I could have done when in Wisconsin. For instance really did not take advantage of the good proximity to two great road courses: Road America and Black Hawk Farms. Last Road America visit warped the rotors on the WRX (which is ironically my wife's daily driver). Did not get around to either replacing the rotors (done) or getting myself a track friendly car (also done) until it was too late (too busy, and now moving). Hopefully will rectify this problem at Infineon Raceway or Laguna Seca after I move.

Quite "yes we can"

2009-09-15T02:33:00.001-05:00

Quite "yes we can", 1st at Chaos Constructions 2009 combined 4k compo.

Selling My Truck

2009-09-13T12:53:00.004-05:00

In case anyone in the Chicago area is interested (the truck is garaged in Westmont IL),

Kelley Blue Book Value: $5210
Selling For: $5000 or best offer

Contact: Timothy Farrar
Phone: 630-430-6656
Gmail: farrarfocus

Had this truck custom made from the factory with special options like the close-ratio New Venture Gear transmission with single shift rail rated to 300 ft lbf of torque, and the Steel wheels (actually durable compared to crappy Aluminum). Great for towing, great for weekend trips, fun to drive. Was not planning on ever selling this truck, but moving across the country in two weeks and I have too many vehicles (truck, project car, and trailer need to go).

Chevy S-10 2WD Extended Cab

Original Owner
4.3 L V6
Heavy duty 5 speed manual transmission option
Limited slip differential
Bed liner and topper
Hitch
98K highway miles
New tires
Heavy duty suspension option

Specific Options

G80 : AXLE POSITRACTION, LIMITED SLIP GU4 : AXLE REAR, 3.08 RATIO K60 : GENERATOR, 100 AMP L35 : ENGINE, GAS, 6 CYL, 4.3L, CPI, V6, 90 DEG M50 : TRANSMISSION, MAN 5 SPD, NVG PH1 : WHEEL, 15 X 7, STEEL QCE : TIRE ALL, P205/75R15/N BL R/PE ST TL ALS Z85 : CHASSIS PACKAGE, HIGH PAYLOAD FIRM RIDE

Note to Self

S10 Forum Ad

DirectX SDK August 2009

2009-09-12T11:32:00.001-05:00

DirectX SDK August 2009

Hotchips 2009 OpenCL Presentations

2009-09-12T11:27:00.002-05:00

Hotships 2009 OpenCL Presentations

LLVM - Block Implementation Specification

2009-09-09T20:34:00.002-05:00

http://clang.llvm.org/docs/BlockImplementation.txt
http://clang.llvm.org/docs/BlockLanguageSpec.txt

Interesting read for anyone looking to get a better idea on GCD low level details. Have not yet had the time to really look at GCD in enough detail yet myself. However, it is very exciting for Apple to take a step in the right direction and both (1) make the OS smaller and faster, and (2) start to address issues of parallel programming.

L1 Misses

2009-09-09T20:08:00.003-05:00

Dark Shikari has a really great blog post talking about L1 misses. Quoting from that post,

"And yet in many cases–such as in x264–much more time is wasted on L1 misses than L2 misses.

The AMD processor documentation says that the L2->L1 prefetcher is not strided, and tests on Intel chips suggest the same. This means that if we are performing, for example, an access of a block of image data that is in L2 but not L1 cache, every single line of data will cause an L1 cache miss. The benchmarks seem to agree; the first chroma motion compensation during qpel in x264 takes more than twice as long as the others!"

This is really a fantastic example of some of the less talked about limitations of using CPU caches and processors optimized for low latency serial computation. Will be interesting to see if this becomes more or less of a problem as CPU style architectures have ever larger vector units and end up using L1 more as a virtual register file.

Erlang

2009-09-09T19:22:00.002-05:00

Never really looked at Erlang, so just watched Erlang : Software for a Concurrent World to get a general overview. Lots of interesting bits in that presentation. One thing I think might be a possible weakness of Erlang is the idea of "location transparency". Seems as if when going parallel with message passing, that computation location and data location becomes a critical factor in algorithm design (perhaps less critical when nodes are interconnected by high-radix routers). That presentation referenced Tilera Tile64 hardware, which I had not seen before. Interesting in how the above Erlang presentation talks about how shared memory doesn't scale, and yet the Tilera hardware dedicates area to a giant shared L2. That last link from Ars references The Raw Project, which I also had never seen ... starting to feel like I have been living in a vacuum! Anyway, for those interested in what others are doing with parallel systems and programming, lots of interesting reading is only a google away from many of the above topics.

Early September Link Dump

2009-09-03T01:20:00.005-05:00

Still no time for anything but links and videos...

GDC 2009 Europe and Canada
HPG 2009
Micro Rendering
Grand Central Dispatch

Krakatoa particle renderings, links to these videos were posted on Twitter.

Mid-August Core Dump

2009-08-19T22:17:00.004-05:00

Been busy... too much interesting stuff, too little time.

Logarithmic Z Buffer

Journal of Lethargic Programmers Post thanks to tweet from Jerome at Q-Games (@blackjero). Provides large Z-buffer range in 24-bits at the expense of destroying correct interpolation, might be something to revisit with DX11 tessellation (tessellate to the point where interpolation artifacts are not a problem).

Assembly 2009 Seminars

Direct link to all. Watched the Technology Behind Love presentation. Interesting tech, smart use of OpenGL buffer objects, etc, but I'm not wild about the min filter as a post process (looks better with those thin fine details).

Carmack Keynote

Watch vid here. Refreshing to hear someone just speak their mind and not care, even if I don't agree with all of it. Not knowing people on the bottom floor, was priceless! Seems like they are taking a big hit with a mess of single thread OO game code forcing adoption of full frame of latency on all parallel job results. Sounds awfully painful to me.

Best way forward IMO is to go as wide (parallel) as possible with everything including game code! Tight pipeline of independent jobs ordered by common sets of a small number of dependencies, enables low enough latency to make use of parallel results in-frame. If one is doing this right, jobs won't look like this (example with 5 cores),

AAA.BBBCDD.FFFF. AAA.BB.CDE.FFFF. AA..BB.CDEEFFF.. AA..B..CDE.FFFF. GGGGGGGGGGG.HIIJ

Note the bubbles in the pipeline, and the clear sync points between groups of jobs, and one large serial operation (G) which isn't latency friendly. Instead the ideal result would look something like this,

AABCDDEEFIJJ AAACDDEEFIJJ ABBCCDEEGIIJ AABBCDDEEEIJ ABBCCDEEHIIJ

Note how enough independent work is interleaved to handle run-time variation. For example job B can start before A is finished (D might be the first job depending on results from A). Also note the task parallel G and H, as well as how jobs go as wide as possible to keep the latency down! Dead simple, dead fast!

I know, can be tough (and impossible) for some to do given the constrains of legacy code base. Plus integration of 3rd party middleware can be a real problem (and with regards to this, clearly I'm not a fan of SPURS). But sooner or later, everything will need to work this way to scale on future hardware.

Light Propagation Volumes in CryEngine 3

Grab the presentation off the Crytek Technology Page. Really appreciate the PS3 and 360 run-time numbers. Impressive in the trade-offs made for real-time performance. Locking the volume to world grid points (CCSM style) is a great idea. This hierarchical irradiance volume method might very well just be the beginning of all sorts of interesting approximations ... just wait until someone revisits this with DX11!

GDC Europe: Cevat on The Future of Gaming Graphics

Sweeney HPG Keynote

Slides here. Like functional programming, but NOT into either software transactional memory and cache coherency. Also the cost slide is way misleading: STMD shader programming model is one of the most productive and easy programming models ever, traditional C/C++ and vector intrinsics is many times more costly given a programmer who actually know how to solve problems using both methods! Lets not forget that the STMD shader/kernel model has been critical for the rapid evolution of the GPU: everything scales, everything is portable across quite different implementations!

Giga Voxels

Siggraph 2009 Slides. It was nice to briefly talk to Cyril at Siggraph. Cyril keep up the great work!

Siggraph 2009: Thursday - Frostbite

2009-08-13T09:28:00.002-05:00

Beyond Programmable Shading I Talk,
Parallel Graphics in Frostbite - Current & Future by Johan Andersson

The deferred lighting part of the presentation provides a glimpse into the types of things possible with DX11. Running the light intersection step in the same kernel as the lighting is an excellent example of task parallel computation in a data parallel kernel. Note this has the advantage of keeping intermediate data on-chip!

Optimization

In the presentation, the depth min/max computation and the parallel lighting tests, are good examples of the types of problems one will need to solve using group share memory (or shared memory for those in CUDA). At some point one will want to optimize the interactions between threads communicating through this shared memory.

Below I've included some simple CUDA test results on a GTX 275 which provides an example on how to go about optimizing shared memory interaction between threads.

Note, I don't have details on DX11 hardware, so the logic below may or may not carry over to actual DX11 hardware!

MIN/MAX REDUCTION TESTS

Numbers below represent overall run time in GPU clock cycles as reported by the CUDA clock() function (took the minimum of a few tests since run-time even on the GPU varies), so lower is better. The kernel returns the GPU clock() at the end point in the kernel, so I use a CPU side min/max reduction to get a rough idea of GPU time. The results are good for general trends, but one would have to test each in a physical application to really know (and to fully verify actual results, which I only did roughly via eyeballing a few blocks).

32x8 block size (256 threads, width matches CUDA warp length).
32x32 block grid (1024x256 threads).

CLOCKS .... TEST ---------------- 88286 ..... no min/max reduction, but syncthreads (overhead of test) 113880 .... non-atomic log-time horz only 174712 .... non-atomic log-time horz, with atomic min/max vertical 338372 .... non-atomic linear-time horz only 397648 .... non-atomic linear-time horz, with atomic min/max vert 408048 .... non-atomic linear-time horz, with non-atomic linear-time vert 1028708 ... full colliding atomic min/max

I'm going to stress this again: it is important to note here that run time does vary (the difference between atomic and non-atomic linear-time vertical reductions should be logically the same), and the test reporting itself can effect run times (recording GPU clock cycles requires writing results back to memory, uses registers, as well as has some ALU overhead, for example a 64-bit to 32-bit conversion if I remember right). The only thing gathered from these numbers should be general trends, programmers will have to test in an actual application to see side effects of varring shared memory requirements and register usage!

The "full colliding atomic min/max" represents the baseline case in the presentation. It was neat to see even this initial "not production tested or optimized" (see page 21) deferred lighting demo running with 1000 very large intersecting lights running quite a bit faster than the needs for a real-time game...

With the tests, I did not try all the possibilities (haven't had the time), for example doing the vertical stages first, and also I did not try a log-time vertical reduction (because it would have required lots of syncthreads).

TERMS AND RESULTS

horizontal - means across threads in a warp
vertical - means across warps

Trends verify that colliding shared memory atomics are indeed serialized on NVidia's GT200 chipset (it is only logical that they would be). Best performance I found was a non-atomic log-time parallel reduction horizontally (5 steps for a 32-wide block, each step writes results back to shared memory for use in the next step), followed by a barrier and then a vertical reduction. Knowing the hardware, it should be possible to do the vertical reduction steps without atomic operations, and without group share scatter. Without knowing the DX11 in detail, I can only guess that this path would work in CS4 and also be the fast path in on CS5 hardware.

Stochastic Visibility in Distorted Fisheye

2009-08-12T17:22:00.002-05:00

Talked about this before, "Screenshot would have been included if computer was alive". Finally got around to doing a few screen grabs.

Quick refresher, the engine is a modified version of the one in this video. The screens below do NOT show a rendered view, instead they show a direct rendering of the {x,y,z} coordinates (modified by a transform to color them) of the nodes themselves. Yes the scene tree is physically stored in the distorted fisheye projection, parent nodes in the hierarchy show up as the salt and pepper noise in the screen grabs.

The persistent scene is 1024x1024 or 1M nodes (minus the edges of the fisheye projection) of which their 8M children are processed each frame. I bumped up the number of roots in the scene tree to 64K (from the 4K I mentioned in the previous blog post).

It was awesome to see the stochastic visibility algorithm work on all sorts of horrid cases like insane amounts of object to object overlap (overlap being both from rooted objects overlapping, as well as the root's tree nodes self overlapping).

I've been distracted with the CUDA version, so I haven't had time to try the collision improvement. I am keeping this alive however, just in case I get surprised by the triangle rate on DX11 hardware...

View of 64K nodes from outside the structure,

BTW, background color is a function of eye position,

Inside,

View of nodes self intersecting from outside,

And from inside the structure (note the corners of the distorted fisheye are not part of the image, this is a 360 degree view),

Other shots also with different fractal rules (different amounts of overlap, etc),

GPU Address Space Mapping II

2009-08-12T11:28:00.003-05:00

This is somewhat a continuation of the topics from this prior post.

Decided to look at the issues of GPU memory from a different perspective, from the perspective of the packets of information which need to be routed around on chip. I'm using the CUDA Compute 1.3 model as a rough guide, and building an oversimplified hypothetical GPU model which probably is different from the actual hardware.

The aim of this post is to cover some of the reasons GPUs have a different architecture than CPUs, and to provide a natural intuition on how to go about solving problems on the GPU.

Hypothetical GPU Model

The GPU modeled in this post has a bunch of independent ALU units connected to a bunch of independent MC (memory controllers) via a very fast on-chip interconnect network. If there are N MCs then in this example, global memory is banked N-way, so that memory requests can be load-balanced between the MCs. The individual bank width is large (perhaps 32 bytes) to amortize the bits used for address per read/write transaction (for efficient on-chip routing and GDDR access).

INTERCONNECT PACKETS FOR GLOBAL MEMORY ACCESS

Since Compute 1.3 supports 32B (B=byte) segments, the hypothetical hardware model routes 32B data segments to and from the MCs. Each routed data packet also includes extra bits such as segment address, type of transaction (read, write, or atomic operation), mask bits (to mask words of the operation), packet routing header, etc.

Note I'm working under the assumption that global atomic ALU ops are done at the MC (it might not be, or it might be done at the ROP in real hardware).

Using this model, larger segment transactions (and atomics which have 2 operands, or work on 64-bit values), require multiple packets to service a single request. Below is a rough idea of packets,

transaction ........ {send} -- {receive} -------------------- ------------------- write .............. {extra bits, 32B segment} -- {} read ............... {extra bits} -- {return extra bits, 32B segment} atomic ............. {extra bits, 32B atomic operand(s)} -- {} atomic w/ return ... {extra bits, 32B atomic operand(s)} -- {extra, 32B old values}

FINE VS COARSE ROUTING

Under the simple model above it becomes easy to see why fine granularity scatter/gather is not efficient.

1. Want to stick to large packets to amortize data payload per extra bits. Having an address per word of data would likely cut the amount of data in a fixed size packet in half.

2. Even if the packet did include an address per data word, the packet itself is routed to one MC. Given memory is banked to MCs, scatter/gather to random addresses (even if this random set of addresses contains good data locality) would likely need to go to different MCs, again cutting the amount of data per packet.

Given a wide interconnect bus (large packets), an attempt to solve (2) would require coalescing buffers per MC on each ALU unit (buffer space to collect enough requests per MC to fill a large packet). As the number of MCs increase the coalescing buffer space per MC would likely decrease (unfriendly scaling for a set buffer size). Fixed function buffer space would likely trade area from ALU, local store or register area.

Note here that if packets switch from being routed to/from MCs to ALUs (which might be the case if the GPU had hardware support for queues for a programmable pipeline), the coalescing buffer problem remains. Attempting to work around this problem might require either a smaller packet size, or some kind of hierarchical network (more on this below in GPU Problem Solving).

TEXTURES, FRAMEBUFFERS, AND FIXED FUNCTION

Texture units provide fine granularity gather. Continuing the above model, the compromise of having a huge amount of addressing/extra bits per result payload, is partly mitigated by only locally routing requests to a single cluster shared texture unit per ALU. However the cost of the texture units in area is relatively large.

Framebuffer ROP units provide fine granularity scatter. The compromise of high address/extra to payload data ratio is partly mitigated by only routing requests to a single ROP for a full tile of data computed at one ALU unit. The cost of ROP units in area is also relatively large!

See the common pattern, fine granularity gather/scatter, but only efficient for tight data locality. Both solve the network problem by routing to/from one location. For the ROPs, having an address mapping which maps 2D tiles to ROPs key.

ROPs in conjunction with fixed function raster hardware, also provide a combined atomic operation (min or max) conditional scatter of a packet of data (think in terms of multiple render targets with Z and stencil). This is a tremendously powerful construct which is very hard to emulate efficiently in compute.

TANGENT ON DX11

A good amount of chip area is used for fixed function, might as well use both fixed function and compute to maximize what can be done with the GPUs of the next generation. DX11 API offers some interesting possibilities depending on what decisions IVHs made with the hardware (assuming I'm understanding DX11 correctly)...

Two possibilities include,

1. Higher fixed function triangle setup throughput. The fast path with the low to mid-end cards of the DX10 generation for data scatter was to do point scatter with all the computation done in the vertex shader and use the fixed function hardware to scatter via the ROP units. Towards higher end cards, triangle rates did NOT scale enough to make this the fast path.

2. RT UAVs. Render Target Unordered Access Views could be an interesting construct to play with, if this path accesses memory via a tiled mapping (possibly backed by a little cache). Meaning in theory RT UAVs might support data local scatter/gather better than via global memory access. Better could mean either lower bandwidth or lower ALU work (which would be required to simulate scatter via compute methods).

Going to have to wait for DX11 cards to go public to know for sure.

Essence of the GPU

MAKING THE COMMON CASE FAST

For a program to scale, the majority of work must be independently parallel.

1. The CPU does not provide an avenue to solve problems which require high parallel throughput.

2. However the GPU provides both parallel throughput and the option to trade parallel throughput for emulation of non-parallel constructs.

With high parallel throughput, other less used constructs (less used because they would not scale, for example memory coherency, fine granularity random access, etc) can be emulated in software on parallel hardware. This also provides the programmer ability to use extra program specific knowledge to tailor a better (than general purpose CPU) solution to the problem.

THE LIE THAT GPU PROGRAMMING IS INHERENTLY HARD

Had the situation been reversed, and a majority of programmers were programming on highly parallel machines, given a serial machine, the common programmer would have a really hard time solving problems!

Intuition on GPU Problem Solving

BUILD A SCALABLE NETWORK, AND YOU HAVE THE SOLUTION TO YOUR PROBLEM!

Best technique I've found is to describe the problem in terms of a network (or graph) of transforms on the data. Common problems, such as the following, all have a similar solution.

1. Data transfer between nodes in the network is too fine grained.

2. Need to scatter to too many nodes, or gather from too many nodes.

3. The node requires too much local store space.

The solution is to build some kind of hierarchical network and divide the problem into smaller parts. This type of problem solving is the parallel version of the divide and conquer method common to solving classic serial CPU (problems like sorting).

Parallel machines force you to factor your problem. How the problem is factored (or how the cost of computation or communication is amortized) can be seen directly in the network (or graph) that is the solution to the problem!

With serial CPU programming, for the most part, programmers do not even think about the cost of gathering data to process because of the giant fixed function CPU and memory hierarchy. With a parallel machine, the cost to gather data to process is the cost to route that data through the network. With current GPUs, the data transferred between distant nodes in the network (or graph) gets transferred to and from global memory. Near nodes (threads) communicate through a local store (small shared memory).

LINK TO BRAIDED PARALLELISM

Given the ability for the hierarchical elements of the network to scale in width (amount of parallelism) and height (of a reduction tree for example), then the solution can scale to either problem size or machine width.

Two core options for the network/graph (can mix/match these),

1. Fixed routing, jobs are draw calls (actual draws or compute invocations). For a fixed graph, run time can be bound to a relatively fixed cost (easy to meet frame rate requirements for example).

2. Braided parallelism via variable routing, uber-kernel with "locked" threads. Enables the network/graph to adapt at run-time to the varying needs of the program.

In Closing

HYPOTHETICAL GPU MODEL VS REAL HARDWARE

Quick CUDA tests on a GTX 275 of a simple non-atomic global read-increment-write kernel, with varying amounts of possible segment bank collisions show quite different run times. Performance when all blocks access the same segment is many times worse than when blocks access different memory segments. This result is predicted well by the hypothetical GPU model described at the start of this post.

Note actual global memory performance on the GTX 275 is quite complex. For example, when running more than one block per multiprocessor, in at least one test it appeared that it is more important that the blocks access different segment banks, than if the individual warps in a given block access different banks (suggests scheduling effects).

V-Ray RT and CUDA

2009-08-11T15:18:00.000-05:00

Video from Siggraph, CUDA V-Ray RT demo on a GTX 285, thanks to cgarchitect.

Macton’s Posterous

2009-08-10T00:50:00.003-05:00

Mike Acton: Recent Sketches on Concurrency ... almost lost this set of awesome presentations in the twitter sphere!

Assembly 2009

2009-08-10T00:33:00.002-05:00

Siggraph 2009: Thursday - State

2009-08-10T00:08:00.003-05:00

Continuing Beyond Programmable Shading topics...

GPU vs CPU State

TLB/etc on CPU provide virtual memory via address mapping. GPU state (such as buffers, textures and render targets) serves a similar purpose as the TLB/etc in terms of providing a special address mapping, however GPU state is more useful in that GPU state provides the information which enables an address mapping designed for efficient 2D and 3D data locality.

Not sure if GPUs will someday get virtual memory, I can see great arguments both ways.

CPU provides ISA which enables changing TLB, page tables, and other system state. In contrast GPUs provide a fixed function command buffer interface to change GPU state.

At some point the graphics API (and possibly the GPU) will need the ability to change GPU state without the CPU in order to keep on evolving...

Siggraph 2009: Thursday - Braided Parallelism II

2009-08-09T23:06:00.003-05:00

Continuing Beyond Programmable Shading topics...

Braided Parallelism II

Along with braided parallelism is the issue of programs with more variable run-times. Variable run-time programs implies the need to run more than one kind of program at the same time on the same core (perhaps through some future MPMD machine, or emulated MPMD via an uber-binary in a SPMD machine).

Core to both these issues is the possible problem of register allocation between running programs. Static allocation of chunks of registers to a program could become inefficient when simultaneously running multiple different programs which have different amounts of register occupancy or running programs which vastly vary in the number of live registers over the course of execution.

Not sure how much of a problem this will be in practice on typical usage cases.

The GPU wants to maximize the number of running programs which fit in the local register and shared-register space (also called group share or shared memory). This is important to maximize the amount of latency hiding capacity.

Larrabee skirts this issue with a smaller register file and a cache hierarchy. The program's stack tends to provide a good description of data locality (top of the stack is more likely to get used). I'm going to assume that register files are less area efficient than a block accessed local memory (because of the crossbar(s), etc). However having a coherent cache adds other complexities such as a TLB (TLB issues covered in another post).

Will be interesting to see how ATI and NVIDIA GPUs evolve with regards to this issue...

Siggraph 2009: Thursday - Issue of Coherent Cache

2009-08-09T15:46:00.004-05:00

Continuing Beyond Programmable Shading topics...

Coherent Caches

1. Given equal estimated {area, process, power}, where is the performance advantage of switching chip area from ALU to CACHE, and what is the right balance?

2. Given a performance advantage is found in (1), is the advantage from simply having a larger local store, or is cache coherency a critical factor?

3. How do algorithms which make use of cache coherency scale as GPUs continue to grow ever wider and more parallel?

Anyone else notice that a majority the examples Intel has shown on their rendering pipeline treats the cache as a program managed local store?

Siggraph 2009: Thursday - Braided Parallelism

2009-08-09T14:52:00.004-05:00

For me Thursday was a full day of Beyond Programmable Shading I and II, slides can be found here. There were a bunch of rather interesting talks and topics, and I'm going to attempt to elaborate on some of the core themes in a bunch of posts starting with this one...

Braided Parallelism

Scalability of a data and task parallel work mix is limited by data dependencies. It is the programmers responsibility to insure there is enough independent work to do between dependencies.

One important difference between task and data parallel work is that task parallel work cannot fill the entire width (all cores across the entire chip) of the machine at one time, and often data parallel work can.

Given this crucial observation, the real meaning of braided parallelism becomes clear: efficient braided parallelism implies that a machine can pipeline and execute (without idle ALU units) jobs which do NOT fill the entire parallel width of the machine.

Task parallel work likely involves data parallel elements, so perhaps a new term "task data parallel" should be used. As a machine gets wider, more and more work will likely fall into the task data parallel category.

A CPU like architecture is NOT needed for braided parallelism, and NVidia's OptiX ray engine is a good example of this. Seems safe to predict that future GPUs (and perhaps DX11 GPUs) will continue to get better at pipelining small jobs, and thus more efficient at braided parallelism.

Will be covering issues regarding the efficiency of braided parallelism on GPUs in later posts...

Siggraph 2009: Thursday - Id Tech 5

2009-08-09T14:41:00.003-05:00

Id Tech 5 Challenges

Highlights of the Beyond Programmable Shading presentation from Id.

Completely off topic, but I really like the sky-box seen in the screen shots. Great atmospheric diffusion, and a contrast range which photographers get through careful use of a polarizing filter to darken the sky in combination with proper exposure control (perhaps with a graduated neutral density filter) to keep the highlights from burning out too much.

VIRTUAL TEXTURING

The screen shots show how they map their UV texture space into the 128K x 128K texels. One can see a mix of world textures and character face textures in the virtual texture pages. Surely there is a relatively high importance in keeping good spacial locality for large multi-tile reads to minimize seek time.

They are not using tri-linear filtering, and instead use bi-linear only. In order to handle the LOD pop issue (caused by optical media seek latency going from a low resolution to high resolution tile), they first up-sample the lower mip, then continuously re-transcode (blend up-sampled tile with decoded tile from higher disc compression, then re-compress to DXT format and update the physical tile on the GPU). Should provide a visually seemless experience, in which detail smoothly blends in.

Over-subscription of the physical texture space is solved via adjusting a global LOD bias slider until the entire working set fits.

JOB SYSTEM

Their job system has job lists (which of which are fully independent) with signal and block on signal constructs to enforce scheduling dependencies. By design, job processing is given 1 frame of latency to complete.

They provided some approximate job costs,

2 ms - animation blending
2 ms - sorting for transparency
4 ms - collision detection
4 ms - obstacle avoidance
4 ms - misc
4 ms - audio
8 ms - virtual texturing
10 ms - rendering

RENDERING

They have sorted alpha blended foliage and particles. They also have a system for detail model generation which items such as rocks and pebbles (perhaps similar to a more advanced shrub system). This is an interesting blend between unique texturing, and instanced stuff to add tiny details.

Not in the presentation, but mentioned after the talk, they ultimately decided not to use Edge or SPU geometry processing. Seems as if the trade off was between higher performance (especially for a vertex bound smaller virtual texture feedback render target) at the cost of needing to keep vertex and index data CPU side for SPU processing, and likely they had free GPU memory (the physical texture for virtual texturing is limited by GPU texture size limits) but were pushing the limit on CPU memory.

GPU JOBS

They mentioned that they are starting to look into running some jobs GPU side, investing in some forward looking usage of the compute capacity of current and new hardware.

Siggraph 2009: Wednesday

2009-08-05T18:31:00.008-05:00

Alternative Rendering Pipelines in CUDA

Slides Off This Page

The first talk was about Ray tracing and REYES in CUDA. The second on OptiX.

There was one very important overwhelming theme NVidia presented which marks a distinct shift in how to program in CUDA (and get at the fast path on current hardware). The shift is the idea of persistent threads pulling work from queues and the usage of uber-kernels!

This has been hinted at prior to Siggraph in some GRAMPS discussion, and has been detailed in the Understanding the Efficiency of Ray Traversal on GPUs paper. The core problem with hardware block scheduling, is that one single long-running thread stalls the entire block. The solution to this problem is to issue only enough threads to fill the GPU once, and then pull work from queues (use work stealing to balance cores). Stalls can be limited to warp granularity, and warps can run separate "kernels" from one giant uber-kernel. Note is this effectively almost a form of software MPMD on a SPMD machine.

Cost of warp-coherent dynamic branching to different sub-kernels can be either linear or a log cost binary search (since CUDA doesn't currently have branch by register). Also register count for the uber-kernel is going to be the maximum register count used by any sub-kernel, but the advantage in ALU utilization from software scheduling greatly out-weighs the possible less efficient register usage for some sub-kernels!

The talk did very briefly cover the best way to fetch from a queue in CUDA (but not in enough detail IMO, likely because of a lack of time to go there). Anyway I've got a 2nd huge CUDA mega post in progress which covers many parallel programming constructs in detail which can be useful in the new age of alternative and conventional rendering pipelines which use "compute". Covering things like the performance of shared memory atomics (serialization cases, when to avoid them, etc), CS4 vs CS5, and more. This is going to have to wait for next week, but it is going to be a great follow up to details which are either too technical or too time consuming to place in these great Siggraph talks.

OptiX

OptiX is NVidia's new interactive ray tracing engine, designed to provide a generic shader interface which will scale by GPU and scale to future NVidia architectures.

Seems like the core CUDA of OptiX is shader and ray software scheduling via uber-kernel and persistent threads. Where each ray has a shader defined state packet, and shaders can spawn rays.

There are two very important programming constructs likely being used here (arg, missed what I think was a HPG talk on this before Siggraph):

1. CONTINUATION BASED PROGRAMMING. Looks like the shader defined ray state packet gets passed between sub-kernels. Likely ray creation places ray on queue with packet.

2. PERSISTENT THREADS WITH CO-OPERATIVE MULTI-TASKING? Looks like shaders can have local data and spawn rays. So likely at some set points while shader is running (ie after it needs to wait on a ray result), local state is returned to continuation queue, and uber-kernel thread scheduler pulls another warp collection of coherent jobs to run? I need to think this through in detail and perhaps try it before I will know for sure.

Something Awesome Is Brewing

Look at what constructs are core to OptiX, and think about how OptiX is designed for both current and future GPUs! Everyone is in the race towards redefining graphics in terms of a "general purpose parallel machine" and NVidia is showing an elegant, easy to program model which scales forward via OptiX.

A little future hardware acceleration to efficiently support thread continuations (and dynamic grouping for efficient parallel execution), and you've got something truly wonderful.

That might not be it, "it" could be something totally different, but whatever "it" is, "it" is going to be awesome, and I'm overflowing with ideas on how to build things with "it".

Via Siggraph you get a glimpse into the minds of some of those involved in shaping the future of graphics. There is something magical about the insane competition going on between Intel, AMD, and NVidia. Something which I believe is going to move us into the age of fully programmable graphics in the not so distant future!

3D Display

First I don't see 3D as a gimic, in fact I just about only goto the movie theater for 3D movies, everything else is Blue-ray or DVD for me, so I'm very biased. I'd like the technology to take off.

Had a chance to try out NVidia's 3D glasses, and was impressed. The high refresh rate was very smooth. Unfortunately (for me, but great for their target audience) the example they had was an interactive three screen volume raycast of a scanned human body, so actual display update was 10 fps or so. Didn't get a chance to see it under high motion, or high depth range (between ultra foreground and horizon). Apparently you render both frames then do the swap, so one gets 60Hz per eye of physical scan out, with rendered frames at or under 60Hz.

Other vendor had a real 3D display, no glasses required (perhaps something over the display to insure interleaved vertical lines only hit the left or right eye). Another vendor had a 2 monitor solution, one upright, the other table top, and a half angle physical sheet to combine both polarized views.

OpenGL BOF

Missed it, got busy after the GRAMPS talk, did walk in at the end and got rewarded with a great 3.2 API Quick Reference Card.

Didn't realize it at the time, since I don't really know people yet, but likely gave Barthold at NVidia trouble about the divergent DX11 and GL paths (how DX11 mixes compute and graphics, and how since GL and CL are separate, that it is going to be a problem to port apps, and things like tessellation which likely could benefit from compute features, seems like a problem).

What I got in return was some very insightful responses, and good humor even given my very direct comments! Thanks Barthold, or whoever at Khronos/NVidia, for putting up with me!

One Day Down

Finished off Wednesday at Siggraph, and realized I'm a lot more scatter brained in person than I thought I would be when meeting many people I've never talked to before in person. Must be the programmers curse, great at language through code, traded off for bad at language through speech?

Off to Siggraph

2009-08-04T14:47:00.003-05:00

Off to Siggraph 2009 tomorrow, be in around Wednesday at noon.

Any of you run into this ugly dude below and want to talk about graphics, please do!

OpenGL 3.2

2009-08-04T09:42:00.000-05:00

Link to GL and GLSL specs.

The Khronos Group continues their steady stride on OpenGL improvements! Really appreciate the hard work, dedication, and the ability to download spec with changes highlighted!

- 64-bit integer support
- geometry shader
- BGRA support
- Draw*BaseVertex(), provides index[i]+base_vertex indexing
- support for binding buffer objects to index targets
- texture fetch from multi-sample textures
- texture fetch from integer buffer texture
- texture fetch from integer 2D rect texture
- seamless cube map filtering
- SampleMaski() for multi-sample coverage mask
- new simple FramebufferTexture() call
- fence and wait on fence support

Aras's Compact Normal Storage

2009-08-04T08:14:00.002-05:00

Aras Pranckevičius : Compact Normal Storage - Impressive method of compressing view space normals into 2 bytes using a sphere-map transform! Page compares with other methods and includes GPU assembly output, timing, and error screen shots.

EAA Airshow

2009-08-03T09:14:00.002-05:00

Had a great time at the annual EAA Airventure in Oshkosh Wisconsin this weekend. Each year pyro junkies get to experience the instant wave of heat and delayed boom of the "wall of flame". However, the best thing at the show each year IMO is the talk given by Burt Rutan, in which you get a glimpse into the mind the most important air and space visionary of our time!

As Seen on Twitter and More

2009-08-03T07:32:00.003-05:00

@mike_acton: A quick sketch on why qsort is not a concurrent algorithm. Feedback welcome, as always! link

I'm a big fan of Mike's presentations, the above covering parallel merge sort. I'd also highly suggest, Typical C++ BS, also posted on Twitter (in fact it was after seeing this only on twitter that I decided to get an account).

@bkaradzic: Watching winning 4kb intro from Evoke09... Lunaquatic by BluFlame

Blade Runner Meets CryEngine 2 - Great indirect advertising for Cry Engine 2.

GPU Address Space Mapping

2009-08-02T15:43:00.002-05:00

This post addresses some of the issues with texture and 2D/3D array address mappings when mixing compute and fixed function as well as dealing with cache lines and the higher granularity of global memory access of GPUs.

For those developers moving on to compute (CUDA, OpenCL, DX11 CS4/CS5), this is going to be an important thing to have a full understanding of!

I would suggest reading Jake Cannell's Winning my own little battle against Cuda Post for a great example of the problems with mixing fixed function texturing and compute.

Review

CACHE LINES

I'm using the term cache line to refer to both the cached and un-cached case. In the un-cached case, it would be the optimal size of a global memory operation. Cache line size for example, is often 64-bytes for current PCs, nearly sure 64-bytes for Larrabee, and at least for later model NVidia GPUs, global memory segment size is a mix of 32/64/128-bytes depending on coalescing (not sure what actual texture cache line size is).

LINEAR MAPPING

2D array elements (or image pixels) are addressed in linear order. A 64-byte cache line of 32-bit elements provides the following,

0123456789ABCDEF

Often array width is a power of two, and then element address for this 32-bit/element array can be computed using the following equation (3 OPS : 2 SHIFTs, 1 ADD),

address = (x << 2) + (y << constant)

Of which,

constant = 2 + log_base_2(width)

SPACE FILLING CURVE AND/OR TILED MAPPING

Elements are addressed using some kind of space filling curve. This space filling curve in should result in a cache line representing a tile of elements in the array. The tiles themselves can be addressed in linear order (easier for software to do the address translation). Or perhaps the tiles are also addressed in a space filling curve (which might provide benefits in the case of hardware predictive cache line pre-fetching). Or perhaps the entire array is accessed using a consistent space filling curve (no set tiles), in the case where the hardware provides a full bit-interleave instruction (Larrabee provides 2 such instructions).

The space filling curve might be Z-order (or full bit-interleaved), which would be a very expensive transform to do in software without special instructions (following the 16 elements per cache line example),

0145 2367 89CD ABEF

Or perhaps the tiles are wide Z-order (multi-bit-interleaved), which is much easier to compute in software when the ISA is lacking a bit-interleave instruction,

0123 4567 89AB CDEF

Still address translation can be expensive (11 OPS : 4 ANDs, 4 SHIFTs, 3 ADDs) even for just linear tile address with multi-bit-interleave for the tiles and power of two sized array width (note only need 3 shifts if elements are one byte, or if one can make use of a ISA which provides a "free" shift in the instruction addressing, like x86),

hiX = x & mask0 hiY = y & mask1 loX = x & mask2 loY = y & mask3 address = hiX << shift0 + hiY << shift1 + loX << shift2 + loY << shift3

The Problems

2D/3D LOCALITY AND JOB DISTRIBUTION

GPUs group pixels and compute jobs into 2D tiles (or blocks) for efficient computation (to maintain good 2D data locality). DX11 provides 3D blocks for CS5. Job or block distribution across cores in the GPU is NDA material, but it is possible to reverse engineer this via careful programming (in some cases). There are all sorts of complexities here, such as a cluster of cores sharing a texture cache.

Linear addressing of 2D data often leads to cache line utilization being low, and cache line thrashing.

COMPUTE ACCESS TO TEXTURES LIMITED COMPARED TO FF/ROP

To my knowledge, texture access in compute is limited to linear addressed textures. This either may or might NOT change for DX11. And if it does, it might not change for all vendors. In theory it should be possible for Larrabee, all depends on what the drivers do!

ROP units provide coalescing, might be in the form of write-combine caching, and/or in perhaps in the form of queuing ROP requests. Either way fixed function render targets are rather impressively fast, even for fine granularity scatter (try on a middle-end NVidia DX10 card to point scatter via drawing points in GL vs global memory scatter in CUDA). Fixed function ROP also provides write to MSAA render targets.

Compute global accesses are coarse granularity accesses, and likely might stay un-cached on some GPUs.

The point here is that compute access to render targets might NOT be as fast as the fixed function rendering path. Clearly Larrabee could be an exception here.

MIXING CS AND PS

If a CS pass builds something for a PS pass, then the PS pass might have to sample from a linear addressed texture. If texture cache access or memory bandwidth becomes a bottleneck, and only IF, there might be a few possible work-arounds.

If the texture is not filtered then one can work in a linear texture and do address translation manually. Even if one doesn't know the actual GPU texture address to texel mapping, one could test different software efficient mappings to see if a good ALU to TEX trade-off can be made.

Often thread groups (or blocks) in compute passes will all be accessing a shared "tile" of data anyway, so tile address mapping in CS might not be as bad if one computes a linear tile address, and then uses constants for the in-tile pixel ordering.

In the filtered case, CS producing data for filtered TEX lookups, the graphics and compute APIs don't currently provide a good solution (unless that solution is hidden behind the driver and hardware).

BYTES PER PIXEL VS BYTES PER CACHE LINE

In the physical cached case, as the number of bytes per pixel approaches bytes per cache line, linear vs tiled address order "should" become less important. For example with 64-byte cache lines,

4... 8... 16 32 --- bytes per pixel ---- ---- -- -- 0123 0123 01 01 --- possible software cache line tile mapping 4567 4567 23 89AB CDEF

A G-buffer or some other fat "logical texture", might have a logical 16-bytes/pixel, but this is divided perhaps across 4 textures with 4-bytes/pixel physically, resulting in the problem case again.

DISCLOSURE OF CACHE DESIGN

In the CPU space, cache details are mostly fully public knowledge. Often the OS even provides a way for software to query cache line details of all the cache levels. Some algorithms actually adapt themselves to cache line size (cache-aware), and there is growing research into cache-oblivious algorithms which naturally work well regardless of cache details.

In the GPU space, cache details (cache as in TEX) are to my knowledge all NDA material, and also likely a bit part of the secret sauce which enables some vendor advantage. GPUs caches likey have various important differences from CPU caches, such as in N-way where N is larger because DX/GL provide access to a much larger number of simultaneous arrays (textures) than would be typical in CPU space, or perhaps different pre-fetching logic (if there is even any automatic pre-fetching at all).

Note a smart programmer can reverse engineer cache design using careful programming, however often those who don't have NDA access also don't have a full set of graphics cards representing the market to test.

In terms of cache as program managed local store, NVidia's CUDA docs provide great documentation on most of the important issues regarding global memory access on NVidia's architecture. Exceptions such as how memory requests are parallalized to the memory controllers can be reversed engineered via software, but would again require having a large number of cards to test (cards have different numbers of MCs).

ATI also likely provides great docs, but I haven't looked into detail on their latest R600/R700 memory architecture yet.

DISCLOSURE OF TEXTURE ADDRESS MAPPING

All NDA material, sometimes possible to reverse engineer. Details such as tile size, space filling curve, etc, all hidden, and likely for good reason. Many programmers would cry themselves to sleep at night attempting to deal with all these details on all GPUs.

Possible Future Solutions To This Mess

The largest problem here is that vendors would all have to agree on a good solution to make it a "good" solution. Vendors opening up their texture address mapping would result in a mess of lots of vendor or card specific code in engine.

Larrabee's 2 bit-interleave instructions and NO ROP design, is an interesting solution to the problem. Yet I don't know anything about Larrabee's texture address mapping to know if this actually is a solution to the problem of CS stages writing to filterable cache friendly textures.

I've got another idea in mind for a general solution, but it requires a change in hardware, which likely is NOT going to be adopted via all vendors unless some central controlling entity like Microsoft decided to force it! So never going to happen!

Arch Linux

2009-07-29T23:36:00.002-05:00

Arch Linux Website | x86-64 Package List

At home for OpenGL and CUDA development I now have Arch Linux on my new AMD64 machine. Lots of thanks to Judd Vinet and Aaron Griffin for the best Linux distro out there!

pacman -S awesome

MinWM

2009-07-29T23:07:00.002-05:00

Not sure why I didn't do this years ago, I open sourced my tiny X window manager on google code.

Mimimal Window Manager [MinWM]

Project Hosted on Google Code

Simple and tiny (10KB) single screen tiling X Window Manager. No messy configuration files. Provides the following keyboard interface to manage windows,

ALT+ESC .......... Close window. ALT+TAB .......... Cycle through window list (just like Windows). ALT+` ............ Cycle window shape between full, and left/right halves.

And that is it!

MinWM requires that the user starts a terminal, prior to running MinWM, which can be used to start other applications.

History

Moving windows around the screen is a waste of time. I got sick of traditional X Window Managers and decided to write my own. 100% of the time I either wanted one window full screen, or two windows split screen, and a very simple way to manage windows without the mouse. From this MinWM was born.

I have been using it for maybe a decade, and still the current tiny tile based X Window Managers are too complicated and get it all wrong. So MinWM is now open source so anyone can enjoy working fast!

Compile From Source

Grab minwm.c from google code. Might have to modify your include and library path, but something like this should work,

gcc minwm.c -Os -o minwm -I/usr/X11/include -L/usr/X11/lib -lX11

Install

Compile yourself a binary. Setup your .xinitrc file to startup a terminal and then minwm,

xterm -rv -ls +sb -sl 4096 & $HOME/minwm

Then run xinit to start X.

Morphological AA Part III - GPU Ideas

2009-07-29T08:31:00.003-05:00

Filtering Via Fractional Pixel Offset

Seems like a filtering pass could be done using one bilinear fetch per pixel using a computed fractional pixel offset to blend between the neighborhood of 4 pixels. This should be enough to remove aliasing. One could look at the simple case of a horizontal span,

        XXXX           .,xXX -- example horizontal span
    XXXXXXXX       .,xXXXXXX
XXXXXXXXXXXX to ,xXXXXXXXXXX
XXXXXXXXXXXX    XXXXXXXXXXXX

In this case (horizontal) the fractional pixel offset would be mostly a vertical adjustment. This offset would always be on the side of the pixel center with the highest intensity gradient. As drawn in ascii below, offset would be highest in the pixels CC, would taper off to low offset by the Zs, and would reverse direction in area CC (left side would point down, right side would point up). The CC area is where the edge crosses the pixel center.

   Z  CC  Z  
   |  ||  |   
       XXXXXXX -- talking about this span
XXXXXXXXXXXXXX
XXXXXXXXXXXXXX

Temporal Amortizing Searching?

After re-thinking this a little: not worth it!

The expensive part is finding the "CC" points on the rows and columns of pixels in the image, and then the "Z" to "Z" length, which is used to compute the sub pixel offset to filter with.

The first part, finding "CC" points, is NOT going to have much temporal coherence other than being close to edges in the image. The span length ("Z" to "Z" length) however should have some good temporal coherence, but this span length can be easily computed from the "CC" point data.

Making This Fast on the GPU

This is just a bunch of notes for some later point when I have both some free time to try stuff and a desire to use this...

I have a feeling that using the offset method one could fully separate this into vertical and horizontal passes (if one wanted to do this, which I wouldn't). Meaning fractional pixel offset in X direction is dependent on Y spans only, and Y direction is dependent on X spans only.

Problem with the GPU is that one needs tens of thousands of threads to be fast. So the trivial case of one scan line (or vertical line) per thread isn't going to work. The algorithm will have to be tiled.

Searching for "CC" points in parallel might be possible with the following process.

1. First map the image into X and Y intensity gradients (fold this into the first reduction pass), then reduce by highest gradient, keeping position of highest gradient during the reduction. With CUDA/CS/CL one might be able to do the entire reduction in one kernel pass (unlike with traditional GPGPU methods).

2. Then run a full screen pass parallelized in tiles, use the reduction tree to search for a set of important "CC" points per the vertical and horizontal spans in the tile (note the search will have to extend outside the area of the tile to support AA of near vertical and horizontal edges). Use the diff in closest "CC" coordinates between lines to compute the "Z" to "Z" length per span. From this information, the fractional pixel offset could be computed...

Morphological AA Part II

2009-07-28T16:19:00.006-05:00

This is a follow up to Justin's question on the Morphological Antialiasing Post and more crazy AA ideas.

What I think is most interesting about these types of post filtering techniques, is that it likely might help with things like specular, cubemap, and alpha test aliasing which can be tough to solve using other methods.

Filtering via Fractional Pixel Offset

EDIT. Nope that idea needed a lot of work... so I killed it!
EDIT. Actually, idea is sound, see third post!

Temporal Feedback

I did a prior test on a polygon based renderer in which I only used motion vector based framebuffer reprojection feedback and a little bit of smooth camera breathing motion even when still (as you often get in a FPS) to remove aliasing. Results were promising, but I never did have the opportunity to use it anywhere (yet).

Note, any method which is amortized temporally will require some form of data feedback, which will at a minimum be another RT, and thus be costly like a full screen post pass. Another important thing to note is that under motion the eye won't see over-filtering or noise as much, as long as it converges to a good enough looking sharp but not aliased edge. One could take advantage of noise or excess filtering under motion when temporal feedback likely won't help as much.

Crazy Depth Pre-Pass for AA

Jake Cannell's Deferred Rendering Without MSAA but With MSAA Z Prepass Idea Post talks about doing a 4x MSAA depth only pass, followed by careful (to keep accelerated coarse Z cull working) Z restore at non-MSAA resolution for shading, followed by bilateral up-sample based on MSAA depth for antialiasing. This would be crazy expensive (Z pre-pass with deferred, and bilateral up-sample), but with DX11 CS maybe not as bad.

Could use depth information to compute and X,Y sub-pixel offset instead of the full bilateral up-sample.

Twitter: Database Lag or Account Broken?

2009-07-28T10:59:00.006-05:00

Decided to start using Twitter: @TimothFarrar (Timothy sans y)

Except the interface seems rather broken right now (at least for me)?

Sorry if a few of you got multiple follow notifications!

Follow links act as if they work. Got emails back, yet nothing shows in my following list. Tweets from those I followed showed in my home page until today and everything vanished. Nice. Help says "database replication lag"... also Firefox likes to spin forever on twitter.com. I smell multidimensional lag.

Morphological Antialiasing

2009-07-28T01:15:00.003-05:00

Linked off Real-Time Rendering, Alexander Reshetov's Morphological Antialiasing Paper is quite awesome. Unoptimized approach is about 150 cycles/pixel in a CPU implementation. For real-time performance, I bet you could amortize the cost of something like this over multiple frames!

EDIT: Yeah, after a second look, amortizing this over multiple frames would be tough! See second post...

Final DX11 Hardware Predictions

2009-07-27T09:02:00.003-05:00

My last few round(s) of hardware predictions (or better guesses) was quite wild, this one is perhaps more conservative! BTW, I don't ask or receive any info about yet to be released hardware but what I gather from public sources which anyone has access to, otherwise I wouldn't be writing any of this.

A little on the progression of NVidia hardware. Naming might be a little off because I don't follow NVidia's wild naming changes all too well...

G80 - base GPU arch
G92 - global atomics
GT200 - shared memory atomics, global coalescing, double precision
G210 - added DX10.1 support
G300 - ?

Each line of GPUs sees a steady progression of improvements on the base arch set in the G80 line. Many of these improvements are things required for DX11. Almost like each generation of hardware was a chance to test and later improve upon the design of individual components required for DX11. NVidia was early on CS, while ATI was early on framebuffer and tessellation.

IMO, DX11 = DX10.1 + NVidia's CUDA + ATI's Tessellation

Seems like nothing revolutionary will be required this DX round, perhaps both primary vendors will be running with a similar design as for DX10. I would NOT be surprised if some DX11 features are "emulated" in software instead of having dedicated hardware. For example, append/consume, or maybe part of the tessellation pipeline pipes through memory instead of having direct GPU routing, or even interleaved MPMD via software emulated dynamic branching (link a bunch of pipeline stage kernels into one binary)?

Also not surprised that ATI is ready with DX11 in the R800 series for desktop GPUs, but perhaps DX11 notebook GPUs not until next year. R900 is likely to be the real interesting arch from ATI for the DX11 generation.

Likely DX11 Games Will be Designed for DX9/DX10 Hardware First

Lets rewind to DX10 generation history. It seems as if the DX10 generation was largely a generation of faster DX9 cards in the eyes of the average user. We can place the blame on many factors: such as a Vista upgrade requirement, consoles based on mostly DX9 generation hardware, developers not being in a situation conducive to pushing DX10 feature set, etc.

Not much of this has changed. DX11 is way beyond consoles, and next generation of consoles is well off. Many PC gamers still using XP (like 50% on Valve Survey). Hardcore PC devs are leaving the PC market behind (Id, Epic, Crytec all now have a console focus). When DX10 and DX11 support is used only to improve a core engine which must run on DX9+ consoles, this is crippling for tech.

Given this, I'd bet that DX11 supporting engines will be mostly designed for DX11 CS4 for some time. Which means that many games will likely run just fine on current generation DX10 hardware under DX11 drivers...

What about NVidia and DX11?

My current guess is that the 300 series will be all about increased ALU performance and power efficiency.

Not sure if we see a arch revision or a arch derived from G80 evolution. Could be no core arch changes from a programming perspective, simply better DX9 and DX10 performance all around, with a tiny minor amount of DX11 stuff added (maybe FF tessellation unit). Looks like at least some rumors agree with this speculation. Speculation that NVidia adopted a G80 derived arch for DX11 launch, and the next major revolution arch happens post DX11 launch GPUs (perhaps like ATI getting R800 early, with a later DX11 arch in R900)?

What about CUDA 3.0?

If the above speculation is true, something must be new for 3.0 and 300 series hardware. Maybe now the hardware can run different kernels simultaneously instead of in series? Maybe CUDA 3.0 is this plus something like DX11's better dynamic linking via better JIT and shader patching?

Or perhaps 300 series is a completely new arch?

It is all a very intriguing mystery!

Pixel Junk Shooter Video with Ice

2009-07-23T13:22:00.002-05:00

Lots of really nice attention to detail in this game: shooting water out of pools from under the surface, ice melting and cooling your player, dunking to quench overheating.

Gather Centric Parallel Programming

2009-07-23T08:43:00.002-05:00

The tool box of ideas to efficiently tackle complex problems on single threaded serial machines is rather well explored. IMO one of the core themes is scatter centric design. Many linear time algorithms depend on efficient scatter. One common use for scatter is binning (linear pass through N objects, output object pointer or index into 1 or more of M bins). A step up from binning is full sorting at which point when N is 1000 (or larger), radix sort is used. Radix sort of 32-bit integers on PCs is often fastest with one parallel histogram generation pass (binning), followed by 3 scatter passes. Binning or sorting (both of which depend on scatter) is crucial for many forms of optimization. For example factoring out branching by sorting objects into categories and processing all similar cases together.

Given that general purpose APIs for GPU parallel programming are relatively new, IMO there is very much left to discover (or re-discover) in terms of methods and tools to solve complex problems on highly parallel GPUs. IMO one of the core themes in GPU programming is gather centric design, which contrasts greatly with serial optimization.

Why Gather Centric?

(1.) LIMITED ATOMIC OPERATION FUNCTIONALITY. Atomic operations are designed for unordered usage because fine grain ordered atomic operations simply do NOT scale to massively parallel machines. Scattering data via global atomics is limited to a small subset of single ALU operations, and these operations are relatively expensive. So global communication via global atomic operations is limited in usefulness. Instead data must be gathered to an independent in-order thread for complex logic.

(2.) READ ONLY CACHES. Excluding Larrabee, current GPUs effectively have only read-only caches. Write-able caches used for scatter centric communication does NOT scale because of cache contention. The most important feature caches provide on GPUs is more bandwidth efficient fine granularity gather (at the cost of high latency). Global vector gather is expensive because performance is a function of the number of cache lines or memory segments required to service the gather.

(3.) SCATTER IS EXPENSIVE. Global scatter is not bandwidth efficient as it is a function of the number of cache lines or memory segments touched. In the case of un-cached global masked writes (non-full writes), hardware would have to both read and write back all touched lines or segments.

Gather Centric Rules for SIMD GPUs

(1.) Algorithm split into independent stages.
(2.) Coarse gather from vector sized and aligned global reads.
(3.) Fine granularity gather from texture reads.
(4.) Reuse read data in a program managed local store.
(5.) All threads do the same amount of work.

Likely this means that a tree structure is accessed via TEX reads but object data is carefully setup in vector sized packets (for efficient global memory access). Parallel methods where all threads of a vector process the same data are fast on all platforms. Parallel methods where all threads of a vector gather from different local objects require an extra transpose pass on Larrabee but on CUDA only require that objects be loaded with different bank offsets (or that threads access data from a different bank given objects loaded on the same bank offset).

Gather Centric Toolbox

Note, some serial programming techniques live on inside the thread group (or CUDA block). More serial program techniques live on within a Larrabee thread because of the larger write-able cache.

What to do with no scatter? Scatter via gather.

Scatter via gather means emulating scatter by having the scatter destination search for the result of the scatter. This REQUIRES data in a prior pass to be efficiently searchable. Which means one of two things,

(1.) Either do a hierarchical parallel reduction in a prior pass to generate some kind of acceleration structure to reduce search cost to log time. If the data is static then this prior pass was done offline.

(2.) Or have the data kept and updated at run-time in a hierarchical tree structure such that no extra parallel reduction is required to generate the acceleration structure for searching.

Gather centric optimization has more to do with the following,

(1.) Limiting the amount of gather searching by "good for the GPU" acceleration structures. Often this means having fixed N-ary trees where N packs well with the SIMD vector size of the machine.

(2.) Keeping thread groups doing the searching grouped by good data locality, and maintaining a high amount of local store data reuse (to avoid getting bandwidth bound).

(3.) Keeping all threads doing a fixed (instead of variable) amount of work. To me this screams "amortize sorting and/or searching over multiple passes".

The "Holy Grail" of GPU Ray Tracing

This post wouldn't be complete without some talk about GPU ray tracing, given that ray tracing is a perfect parallel example of gather centric design.

GPU ray tracing breaks rule (5) which states that all threads do the same amount of work. GPUs lack ability to efficiently do fine granularity regrouping of divergent threads. However a recent CUDA paper found 2x speed up from bringing down the cost of long running threads from thread block granularity to warp granularity. This is the current state of the art.

The problems of dynamic objects and performance have yet not been solved for ray tracing (or ray casting). This fact alone has made the topic a worthy background obsession of mine.

Temporal Coherence

I'm convinced that a solution can be found in use of temporal (in combination with spacial) coherence to lower search cost. I've managed getting part way to a solution using a scatter centric design in the form of what I call stochastic visibility. However if I want something which will scale on GPUs and something which I can port to OpenCL to run on any GPU, I need to re-formulate a gather centric solution to the problem of real-time tree-traversal of dynamic data.

A solution has been eluding me for some time!

Last night / this morning, a new possible solution became evident thanks to exploring a combination of unrelated topics (including thinking about doing a hierarchical cellular automata to increase the "speed of sound" in a fluid computation), writing out the above rules again of gather centric design for SIMD GPUs, and rethinking how to solve reprojection without scatter. The solution involves maintaining a GPU side scene tree structure of which I use the previous frame's tree traversal to accelerate the search for the solution to the next frame's visibility...

More on this later, if it works!

My Ideal Parallel Machine

2009-07-20T10:39:00.002-05:00

Clearly I am NOT a hardware engineer, but I am seeing a very consistent pattern of what I want in massively parallel hardware from what I attempt to create in software, and from this I've derived my ideal parallel machine below.

The machine is a very dense 3D computing/memory matrix of which a 2D slice looks somewhat like this,

(CNM)(CNM)(CNM)(CNM)(CNM)...(CNM) (CNM)(CNM)(CNM)(CNM)(CNM)...(CNM) (CNM)(CNM)(CNM)(CNM)(CNM)...(CNM) (CNM)(CNM)(CNM)(CNM)(CNM)...(CNM) (CNM)(CNM)(CNM)(CNM)(CNM)...(CNM) ................................. .............SCALES_IN_X_Y_Z..... ................................. (CNM)(CNM)(CNM)(CNM)(CNM)...(CNM)

- Each (CNM) is a compute/network/memory node in the matrix.
- Each node is connected to neighboring nodes in the matrix.
- All memory is distributed among the nodes.
- There is no external memory.
- There is no hardware cache coherency between nodes.

The network (N) part of each node would route data from local memory (M) in the background of computation (C). Total interconnect bandwidth would be tremendous, do to the highly parallel neighbor to neighbor routing. Network latency would be mostly a function of the distance between the sending and receiving nodes.

Crucial Change from Current Practice

I am just starting to really understand how I would program for this kind of machine in a general purpose way, and have algorithms scale to a variable number of compute nodes.

The crucial change is that memory is no longer external to the CPU, and we abandon the "load, compute, store" model of programming at larger scales.

Parallel computation on current hardware is either distributed to fixed locations or variable locations based on core load balancing, and data is either distributed among local and distant caches, or distant banks of memory. Effectively computation becomes fine grain loop of loading data from distant locations, doing work, and then storing the results back to other distant locations. This does NOT scale well at all.

The problem is that the physical data flow in the machine does NOT match the needs of the program, and computation does NOT follow the required data flow.

Micro Scale Example

When looking at the current model of CPU atomic operations, parallel computation interfacing with shared data is distributed to many nodes, requiring an expensive serialization of computation because the atomic logic plays out finely interleaved across nodes in which there is a long communication latency between nodes.

For example, with CPU nodes fetching or writing to a shared queue, the queue pointer ends up on a shared cache line, and each add to the pointer happens on a different node requiring a synchronous transfer of the full cache line over long distances before an operation can finish.

Data and computation is NOT flowing through the machine in sync with the needs of the algorithm. Data and computational locality should be in parallel with each other!

I believe GPUs, excluding Larrabee, have traditionally had atomic operations on globally shared memory computed at the point of communication synchronization near the memory controller (MC) or ROP, with no need to transfer "cache lines" back and forth around the machine for the ALU operations which compose the atomic operation.

In this case data is flowing through the machine in sync with the needs of the algorithm. Computation on data which needs to be joined physically happens at the point where data is routed to a common location.

The DX11 CS, OpenCL, and CUDA parallel computing models (on non-Larrabee hardware) logically terminate here, and are vastly limited by a small microcosm of transfer of "compute locality" expressed in the set of supported atomic operations. Any other transfer of "compute locality" requires complex routing to/from distant memory and node-to-node synchronization. Clearly in the case of Larrabee, this routing might be of lower cost (or lower latency) because of the coherent caches, but still suffers from the same problem, just less so!

Macro Scale Example

A more macro example would be the non-general GPU fixed function ROP and framebuffer logic. Data is transferred on chip, results pass from compute units to distributed serialization points (ROP units) for a given tile of the framebuffer. Work distribution is to my knowledge a fixed layout to provide good load balancing and to better parallelize framebuffer memory requests. Note that a complex set of logic happens with respect to Z, stencil, and blending in the ROP units themselves. Compute locality transfers from ALU units to ROP units as does the data.

This is a macro example because the entire set of ROP blend operations is conditionally dependent on fine granularity Z and stencil checks. This sort of thing cannot be done efficiently in software via the small subset of unordered atomic operations provided by DX11 CS, OpenCL, or CUDA.

An attempt to emulate this ROP unit in software on non-Larrabee GPUs is a disaster. The problem results from high overhead in all the many possible ways to do the required node to node communication to provide framebuffer functionality. At some point node to node communication must go through global memory, or data must be fetched from global memory (or TEX cached) many times (gather/searching) to avoid the node communication (from scatter, etc). Either way, performance is screwed compared to the fixed function hardware. Note, Larrabee should be better at software ROP emulation because of lower overhead (lower communication latency) do to caches, but looks to me as if it digresses into worst case behavior as the granularity of writes to the framebuffer become small (either the binning costs grow huge, or the synchronization on Z buffer check/write grow huge, choose your poison).

The Big Picture

The programming model on my suggested compute/memory matrix becomes one of direct data flow. Data flows through the machine based on the needs of program synchronization and communication, sync points become physical nodes where data is transferred to be jointly processed (or serialized), the program's next step of computation (the joint processing) happens at the destination node, instead of the compute node where the data originated (as with traditional programming models).

In the ROP example, compute nodes would processes fragment shader kernels, and then output a network packet of fragment shader intermediate results. This network packet gets background routed to the node(s) assigned to the corresponding ROP tile(s) hit by the output samples. These ROP nodes do any programmable blending and can even continue doing more programmable shading (which might include transferring new packets between other nodes for computation). The ROP tile nodes can route back coarse Z/stencil data packets to raster nodes to early cull.

For load balancing, nodes themselves could easily service many different destination compute domains at the same time. Also note with this programming model, data need not be transferred to and from fixed long term set locations of memory, data structures can instead continuously be dynamically moved around in the distributed memory.

Undoubtedly more posts on this later when I get more time...

Windows SDK, Windows from eSATA, ETC

2009-07-19T12:22:00.004-05:00

This post is an open request for comments from anyone who wants to share advice regarding the topics below.

Visual Studio Express 2008 or Windows SDK

For personal development I really need just one thing, a compiler/linker binary which runs from the command line. Clearly I also need the system headers, but I absolutely DO NOT need an IDE or debugger, those just slow me down.

Anyone have experience with the free Visual Studio Express 2008?

I didn't even know it existed until today, and it seems easier to just to get Express than attempting to maintain a manual install of MinGW/MSYS for GCC.

What about just installing the Windows SDK without Express?

Looks like I really only need the SDK and NOT Express. If this is the case, damn, thanks Microsoft! Love to get any feedback anyone has as to this. For example, are there any limitations as to the Windows SDK compilers? Like hidden missing support for assemply or something crippling like that? What about with Windows 7 Beta? Are Windows 7 GPU drivers good or buggy?

Bootable Thumb/USB or eSATA

The dead machine is helping me rethink my development environment. I'm now looking at the option of building bootable external drives for development, to make it easy to switch between operating systems and physical machines (I used to switch internal drives). I made sure to get eSATA on the new motherboard. Too bad the iMac requires boot-camp.

XP from USB : Cedrick Collomb Blog | Ngine.de

Anyone run XP or Vista from eSATA?

Switching Linux Distro

My first Linux experience I think was Slackware, and post that I did my own from scratch because I wanted pthreads before it was available and later wanted everything to run from a ZIP drive. I've also played with NetBSD, OpenBSD, and FreeBSD. Settled with Crux for a long while (KISS), but went to Gentoo on the last machine for easy 64-bit support (big mistake). Now Slackware is just getting 64-bit support (in current branch only), and FreeBSD needs some kernel changes to get 64-bit NVidia drivers, so those are out for me.

Turns out that Arch Linux is a very Crux like (KISS) distro with good 64-bit support, so I'm going to move to it on the new machine.

Weak Computer Took One for the Team

2009-07-18T18:15:00.002-05:00

Insanely Twisted Shadow Planet looks really cool.

Had a weekend of freedom to really work (wife is out of town), and the damn PC Linux dev machine died today. After testing 3 different GPUs and 3 different power supplies I've ruled out everything but one or more of motherboard, memory, or CPU. So instead my bother is driving up to visit and I'm catching up on PS3 games. It is amazing how much waiting on cheep parts costs in opportunity cost.

I'd really like to just immediately transition development to my older iMac (with it's 8800M GTS GPU), but STILL waiting on September for Snow Leopard for OpenGL 3 support (I'm not a ADC Premier or Select Member so don't have access to the beta, assuming GL3 was in the beta, which I don't really know). Going back to pre-GL3 isn't really an option.

Atom Update

Tested out having 4K unique scene tree roots with heavy overlap and sparse l-systems. Visibility worked with everything moving. Drawing is mostly constant regardless of how many root nodes I have (as long as I stay within a low percent of root nodes to total nodes). Will have no problem with the 64K dynamic object target. Also pushed the GTX275 to 8M nodes (4x my high-end GPU target), but couldn't hit 60Hz. BTW, I'm now targeting 60Hz on low end, 30 just isn't smooth enough for me.

I have a new better way to store the node tree in a modified 360 fisheye projection mapping alone instead of having to reproject between an octahedron (scene tree storage) and fisheye mapping (drawn for player). The new mapping is a spherical 360 degree fisheye with resulting xy coordinates modified by sin() to push the projection into half way between a rectangle and sphere. Screenshot would have been included if computer was alive.

This was an accidental discovery, I didn't expect it to work well. It has the side effect of a more rectangular (normal) view in the center, and actually performs better with the stochastic visibility. The view expansion near the end of projection (the edge of the projection is from one "point" behind the viewer) actually helps leave room for more of the tree structure without collisions, and the view "compression" under fast forward motion as objects fly behind the viewer requires less frames to fully hole fill.

The result of all this is that the scene tree itself it stored directly in a very thick "frame buffer", and it works do to having some kind of 360 degree view mapping.

If this sounds alien, or like I'm in crazy land, it should. I spend more than 40 hours a week programming in reality (at Humanhead), so going off the deep end is quite refreshing indeed!

Console Art Detailed

2009-07-18T18:04:00.002-05:00

Neat to see how others build things, awesome character artist at work shared these links from Ben Sprout Artworks : Prefabs in KZ2 and Environments in MP3.

Alaska : Juneau

2009-07-18T16:27:00.002-05:00

We ended our trip in Juneau to catch our flight back to Chicago (which went indirectly through Anchorage). Juneau is one of the very common ports for Alaska cruise ships so be warned. I'd suggest skipping the Mount Roberts Tramway and instead hiking up/down the deadly path (very dangerous in the rain) which leads to/from the top (it is only 1800 feet, and the local kids run up the path).

Best local food we found in a very short amount of time was the Sandpiper Cafe. Other places of interest included the Alaska State Museum (where they had a great Science on a Sphere exhibit with hundreds of different visualizations of planetary data), the Alaskan Brewery Company tour, and the salmon hatchery sporting a really nice set of aquariums to keep people visiting outside of spawning season,

This is it From Alaska Folks

After this I will stick to the usual crazy programming stuff for a while. By the way, nearly all the photos thanks to the lovely lady below,

Hope you all enjoyed a glimpse of Alaska!

One Last Departing Gift : Alaska Stream Crossing Epic Fail!

An attempt to do something really stupid for the camera turned out horribly right, we laughed non-stop for about an hour after this...

Where is my Raytraced Physics Toy?

2009-07-17T14:34:00.003-05:00

Someone really needs to join GPU physics with GPU raytracing into an interactive physics toy. The scene has shown the way, interesting 11k Intro by FRequency,

Another 4k by FRequency,

Another 1k by FRequency,

DX10/GL3 GPGPU : Stream Out with Scatter

2009-07-16T16:06:00.000-05:00

In regards to DX10/OpenGL3 level hardware and point scatter, moving PS work to the VS stage was a huge win. In the case where there isn't enough VS work to amortize the setup cost, perhaps using stream out (or GL transform feedback) would enable adding a secondary task to be done in the VS stage. So VS would stream out a set of outputs which are not passed along to the PS stage (PS stage gets separate VS outputs).

This might be perfect for the case where an object needs to self update (do that via stream out), and then scatter its index into another data structure (via point drawing).

Wow.

2009-07-14T11:39:00.003-05:00

Damn, Naughty Dog wizards have set the graphics bar again! Particles colored to match scene tint, particles getting shadowed, perhaps sub-surface scattering, huge numbers of triangles, and more...

Hybrid Cached Raycast and Reproject Raster

2009-07-14T10:52:00.004-05:00

I've been suggesting this idea and permutations of this idea partly because I'd like to see someone try it, and I'm busy "traversing" a different solution to a similar problem (just geared towards dynamic rather than more static geometry). Undoubtedly there are those who have thought of this, just I haven't seen a paper on it yet.

The rough idea is to amortize the ray traversal cost over many frames and to make use of both spacial and temporal coherence. This shares a lot of similarities to video compression, in regards to reprojection being like motion prediction in B/P-frames, and the tile cache being like bits of I-frames.

The idea uses a traversal tile cache. A tile being a warp or vector (SIMD) friendly packet of the result of raycasts into an object in the scene. For example for CUDA, a tile might be a 8x4 grid of raycasts. An object might be the entire static world, or a character, or a dynamic object. Tiles are associated with an object ID so that the tile can be correctly reprojected as the object and view moves.

Tracing into a skinned character is a separate problem for which I don't have a good direct solution. However, given a solution to this problem, one could use the per tile object ID to "re-skin" the tile during reprojection. Actually one would want to store the tile in the base pose coordinate space, and simply apply the bones to the fragments in the tile to do the reprojection.

Tiles in the cache become less valid as objects move or the view changes. Each frame the cache gets updated to fill areas of the view(s) (note, there might be multiple views such as shadowmaps, etc) which have sparse coverage. Tile cache update rate would scale to match the capacity of the GPU.

Drawing the frame would be done via a composite pass which re-projects, lights, and blends all the cached tiles (and fills in any holes). Part of lighting could very well be done in tile space, and at a frequency lower than the screen refresh for low frequency lighting, and perhaps in image space at the screen refresh rate for high frequency lighting like specular and hard shadows.

Ultimately the success and failure of this type of system depends on if the composite/reprojection cost is lower than the full traversal.

Compared to What I'm Doing

With classic ray traversal the structure is geared towards answering the question "if a ray intersects the object, what does it hit?". With my structure the traversal is geared towards answering the question, "given a lower quality version of the object in this 3D region, what does the next higher quality level of detail look like". In my case the traversal could be of any tree, even if the tree self intersects. In the case of a skinned object, classic ray traversal is a complex problem, in my case it is trivial.

I go directly down the tree without divergence, because my traversal doesn't go by ray intersection. However I depend on a reprojection scatter pass (which can be redone as a coarse binned scatter, local gather, or whatever form is fastest on the hardware). This reprojection scatter pass is used to prune out traversals of the tree which are not adding to visible nodes.

My scene is that of a forest of l-systems (64K roots, 2M nodes). Each traversal step is amazingly expensive (decompress rule, apply rule) compared to a triangle or voxel raycast or raytrace, and in my case the composite/reprojection step isn't as costly as traversal, so the hybrid traverse/raster works.

Not sure if the hybrid cached trace/raster approach will work well for classic raycasting, but I'd bet one of you out there will be answering this question in the next year or so!

Amusement Machine Video

2009-07-13T10:51:00.003-05:00

The Amusement Machine guy posted an interesting video (a month ago), real-time gpu raycasting on ATI 4850.

Alaska : Gustavus

2009-07-13T08:12:00.004-05:00

Gustavus, with its somewhere over 300 residents (you know a town has its priorities right when the airport is larger than the town itself), is a launching point for all things done in Glacier Bay. We stayed at Aimee's Guest House and would highly recommend Aimee to anyone interested in staying in Gustavus,

The most dangerous animal in Alaska (2nd to people of course) viewed from the deck of the guest house during breakfast,

Dinner was mostly a combination of biking to Pep's Packing for some smoked salmon and the local grocery. The best dinner in town was the Gustavus Inn where you have to reserve your seat before noon (so they know how much fish to get), food is served family style, and the highlight of dinner is a mix of local fish/crab and greens/fruit from the owner's extensive garden.

Eight mile bike ride from the guest house is Glacier Bay National Park which only really sports 3 hiking trails from the lodge (and only one sized for a day trip). Aimee's Guest House provides vintage 1-speeds for the journey,

And on route is path to a military plane crash site (sans the awesome radial motors),

On the other direction is the road into town, from that, on route to the pier is the Mount Fairweather Golf Course,

Video of Particles via L-System and Stochastic Visibility

2009-07-10T10:13:00.003-05:00

I talked about this a while back and never did post a video, and I had to post something cool before SigGraph 2009, so below are screen shots and a quick video of an old stochastic visibility with particle generation effect test. The OpenGL2 (NVidia Specific) GLSL demo was built a long time ago for a 8600 GTS, and runs way too fast on the GTX275 (even bottle-necked by triangle setup in point scatter). Unfortunately all I had time to do was add a very slow serial frame capture and write MJPEG, which when turned on slows the program down too much to make a good video (it is hard to think in slow motion). The conversion from MJPEG to MPEG4 isn't ideal either (it is a huge file) at 60fps 720p.

There are a lot of things I never bothered fixing in the demo, such as the boundary between black background and the fractal (you will see a blurry 16x16 block pattern outline late in the video), or the pops when a bad floating point number gets into the particle buffer. This only shows the result of the tree expansion limited to one level per frame (fill artifacts from edges of occluders). Also hole filling was broken. So it isn't really representative of what I'm planning on for Atom, but it was a huge milestone in learning GPGPU techniques!

Video

Link to MPEG4 Video - 720p 60fps (115 MB)

Other Screen Shots

What is it?

I'd qualify it as bad coder art used to test GPGPU tree data structures (everything is computed on the GPU, the CPU just sends in view position and view direction). The full fractal structure is huge, something like 2^24 pixels wide before precision problems sink in?

The scene tree is an 8-arry tree with one simple l-system rule for the 8 children based on the parent (position/scale/quaternion). It isn't limited to just one rule, I'm just lazy, can lookup any rule from a texture. The engine maintains the scene tree on the GPU (256K coarse nodes, 2M fine nodes). The projection is a 360 fisheye, and the particle effect runs in projected fisheye space based on motion vectors and parent tree position relative to child tree position. The visibility is computed in an octahedron space (different mapping from the view), so the edges of the fisheye projection loose quality fast. Color is instanced via a direct visualization of eye relative world position,

vec3 rgb = vec3(0.0); rgb += vec3(0.7,0.625,0.5) * vec3(pow(abs(sin(sqrt(abs(y/4096.0))*2.0)),16.0)); rgb += vec3(0.3,0.4,0.5) * vec3(pow(abs(sin(sqrt(abs(x/4096.0))*32.0)),4.0)); rgb += vec3(0.5,0.4,0.3) * vec3(pow(abs(sin(z/4096.0*256.0)),4.0)); return rgb * rgb;

Future

I'm still fascinated with the problem of solving visibility for fully dynamic geometry requiring very expensive ray traversals (if ray cast or traced), but without ray casting (or tracing) and instead using what I've termed as "stochastic visibility" and blogged about. Effectively keeping a consistent tree structure of the scene as required by visibility (and shading), then only expanding and contracting the tree structure enough each frame to service rendering to the quality level required for human frame to frame perception. Stochastic visibility collides a point per scene tree node in an view space mapping to both help compute tree updates and solve for visibility. It is stochastic because the points collided are randomly positioned inside the bounding volume of the node. Collisions help prune out nodes which yield non-visible geometry. Collisions also directly solve the tree node memory allocation problem (there is no memory allocation).

Vacation has provided some new ideas to improve upon my old non-CUDA/OpenCL "stochastic visibility" which I will be trying soon,

(1.) IMPROVED TEMPORAL CONSISTENCY VS HISTORY BUFFER. Collisions would sometimes prune out a part of the required scene tree branch, and keeping persistent nodes was way too costly to do in OpenGL3 with GPGPU methods on a 8600 level graphics card. One trivial solution to this problem would be to run a pass which checks the source data for collisions and does a reduction of the source nodes keeping the highest priority node which had a collision. If I had a set of reducing resolution history buffers, I could likely insure good temporal consistency (solving the random pruning problem).

(2.) ADD PARENT LINKS. My scene tree nodes didn't have links to the parents. Was a problem for animated l-systems and particle fluid dynamics, because I wanted a physical constraint which pulled nodes back to the static position in parent (as defined as a possibly animated l-system rule). Turns out that the history buffer pass enables me to correct parent links in child nodes, even though parents move to different memory locations each frame. If this works, this GPU data structure truly becomes awesome because it solves the 1M node memory allocation problem per frame with fully dynamic trees, automatic "defragmentation" of memory, automatic regrouping for good data/branch locality and cache performance, and now nodes maintaining parent links even though all memory locations change per frame.

(3.) MULTI-LEVEL TREE UPDATE PER FRAME. With LOD transparent blend-in and a triangle based scene, I've found in previous results that I can add/prune nodes at only a LOD level per frame and effectively service visibility if I have a conservative amount of overlap. However when the screen tree goes down to the pixel or near pixel/level, the tree must be able to expand by more than one level per frame to fill visibility gaps in dynamic geometry. I'm planning on a new method where I first re-project the current node set, then do a hierarchical image space reduction to choose the highest priority nodes for varying level multi-level tree update per frame.

Cool Blog Posts

2009-07-09T03:25:00.003-05:00

I'm starting to really like the Followers tool on the blog,

Sanders' Blog : Deferred Virtual Texture Shading Post Series - Sander van Rossen has a really neat idea to combine virtual texturing with texture space shading and do shading in the physical pages of the texture cache itself. Dynamic geometry would require updating normals (and coordinates) of the texels in the physical texture page cache. But lighting calculations could be cached (or not done every frame) and transparent rendering is simple. The challenge is efficiently rendering lights into physical texture page cache. I think this would be a highly worthy problem to solve in DX11 with CS5. With unique texturing, for each texture page in the page cache, have an associated bounding geometry to intersect with lights. Bin lights GPU side per texture page in the CS. Do highly coherent shading in the page cache with the lights in the page bin. This could be done by running enough threads to fill the hardware and just extracting work from the bins to load balance all the shading.

EnterTheSingularity - Jake Cannell has a bunch of cool ideas, such as this, "Taking the spatial-temporal coherence optimization idea to its logically conclusion results in a very interesting radical way to approach a renderer. Really what you want is something like a real-time video compression loop, but without the encoding step. When you need new fresh image macrotiles, you invoke the rasterizer, sometimes rendering them at reduced spatial resolution depending on the amount of motion blur, and then these are fed into a forward reprojection system which composites most of the image, along with some holes and errors introduced by new/unpredicted moving objects, lighting changes, etc. But that would be an entire new line of research, and is more suited to a ray tracer or really tight tree-based micro-rasterizer that can effeciently handle micro draw call invocations."

GPU Ray Traversal Efficiency

2009-07-09T01:21:00.000-05:00

Understanding the Efficiency of Ray Traversal on GPUs
by Timo Aila and Samuli Laine of NVidia Research

I really like this HPG2009 paper which explores the primary factors effecting performance of ray traversal on NVidia GPUs. They compared actual results with results from a simulator which had zero cost memory performance. So actuals vs upper bound on best (unobtainable) GPU performance of only ALU cycle counts.

They found work distribution yielding lower SIMD efficiency was the primary bottleneck for ray traversal and NOT the memory system. Quote from the paper, "We have also shown that, contrary to conventional wisdom, ray tracing performance of GTX285 is not significantly hampered by the lack of cache hierarchy. In fact, we can also expect good scaling to more complex scenes as a result of not relying on caches."

Upper End Result in the Paper on a GTX285

- 140M rays/sec for primary rays (coherent)
- 40M rays/sec for secondary randomly shuffled GI rays (non-coherent)

They found around (and often better than) a 2x performance increase via two primary optimizations. With these changes they managed between 77% to 91% of the theoretical maximum GPU performance assuming infinite memory speed.

(1.) PERSISTENT THREADS. Instead of launching a thread per pixel, launch enough threads to fill the GPU once, and have the threads pull work from a global pool. The problem is that ray traversal has varying execution time. The advantage of threads grabbing their own work is that long running warps only stall the warp instead of stalling the entire block (another block would have to wait until the longest running warp finished before being issued). Also note paper pulls work from the global pool into small local (shared memory) pools to reduce pressure on a single global atomic counter.

(2.) PER-RAY TRAVERSAL. Typical CPU SIMD tracers do packet traversal partly because of lack of vector gather and because CPUs don't automatically handle branch divergence. The paper found that per-ray traversal is faster than packet traversal on the GTX285 in all cases, and for diffuse primary rays about 2x faster (even though per-ray traversal decreases memory coherency)!

Implementation

- 192 threads/block
- 21-25 registers/thread
- node data fetched via 1D texture fetch (cached)
- triangle data fetched via direct global memory access (non-cached)
- per-ray stacks via global memory access (non-cached)
- binary BVH traversal with maximum leaf size of 8
- tree built using greedy surface-area heuristic
- 64 bytes per the two BVH nodes (one full global memory segment access)
- pair of BVH nodes always fetched and tested together
- Woop's unit triangle intersection test
- 48 bytes per triangle (1 to 2 global memory segments per access)
- rays assigned in Morton order (good node texture coherency on primary rays)

I'm making some assumptions here (which I didn't see in the paper directly but should be obvious?), that the pairs of BVH nodes and 48 byte triangle leaf data (referred to below as a "data block") are fetched using 1-2 global memory transaction into shared memory.

(1.) Threads write address of data block into shared memory.
(2.) All threads share in loading the data block for one thread at a time.
(3.) Data blocks are loaded with increasing bank offsets.

This insures very fast memory transactions both globally (perfect coalescing on GT200), and no bank conflicts locally afterward when threads move from coordinated loading to individual parallel computation.

Notice stack data is also thread-local storage which is non-cached global memory accesses. I'm going to assume again that they insured good global memory segment coalescing (GT200) via keeping per thread stacks spaced out by a stride of 16. This way global memory segments required for stack update would be a function of divergence in stack size. So stacks of similar size would get serviced in a single global memory transaction (aside, this is exactly the same thing required on Larrabee to do fast parallel stacks since gather performance is a function of cache lines accessed).

Dedicated Warp-Wide Scan or Prefix Sum Instruction

This enables lower software overhead for fine thread granularity work regrouping on the GPU. IMO NVidia needs this (or something like Larrabee's 2 related instructions), covered it in prior blog posts, have hopes that this might be in NVidia's GT300 arch. The paper covers this in the section on "Improving the SIMD efficiency further".

Paper found that current software overhead for replacing terminated rays didn't yield a consistent performance update. However simulator suggested a max 20% improvement with the addition of two warp-wide hardware instructions (prefix sum and population count) assuming no memory bottleneck.

Paper also has some very rough speculation towards going to fully dynamic "Work queues" using the simulated new warp-wide instructions showing a max potential of 40% and 80% performance improvements in two cases assuming no memory bottleneck.

SIMD Width : 32-wide vs 16-wide vs 1-wide (Scaler)

Given impossibly infinitely fast memory, and identical computational power, simulator results showed 16-wide SIMD to be 6-19% faster, and scaler to be 30-144% faster than 32-wide SIMD.

Alaska : Glacier Bay by Air

2009-07-08T16:11:00.003-05:00

A 75 minute flight-seeing tour of Glacier Bay Area is the best way to see the park. We waited for a clear day before doing the tour and got rewarded by the ability to fly over the Brady Ice Field. Flight-seeing provides a way to see many glaciers up close and personal, or choose specific areas of the park to look at, such as the peaks of mountains.

Or the glaciers themselves.

Here are some of Margarie Glacier,

Unfortunately a majority of the shots taken didn't turn out well (mixture of new camera with auto-exposure not adapting to the snow and overexposing everything, and rather would see stuff instead of take photos). The experience was hugely better than what you will ever see from a set of photos or even videos.

Giga Voxels at Siggraph 2009

2009-07-08T12:10:00.001-05:00

Beyond Triangles : GigaVoxels Effects In Video Games
Friday, 7 August | 3:45 PM | Room 260-262

Looks like a really neat talk,

Mixing with existing triangle scenes.
More efficient CUDA implementation of the pipeline.
Faster octree traversal.
Fully GPU side brick cache management.
Efficient soft shadows and depth-of-field approximation.
Instancing of millions of octree-based volumetric objects.
Image space method supporting transparency with overlapping instances.

Alaska : Glacier Bay by Boat

2009-07-06T11:28:00.003-05:00

The marine wilderness of Glacier Bay includes tidewater glaciers, snow-capped mountain ranges, ocean coastlines, deep fjords, and freshwater rivers and lakes. This diverse land and seascape hosts a mosaic of plant communities and a variety of marine and terrestrial wildlife and presents many opportunities for adventure. Here is the Glacier Bay Area viewed from google,

View Larger Map

The park provides an 8 hour boat tour of the area which visits many of the glaciers in the bay. The tour starts from the park dock at Bartlet Cove (which can be seen by zooming in the map above East of Gustavus).

Kathryn and I decided to take the tour after seeing photos of the tour where people were getting on and off the boat. Turns out this was misleading, those photos were of picking up and dropping off kayakers and not off boat excursions as part of the tour as I expected. Eight hours is a long time even on the large boat, and I spend part of the tour (most of the return) sleeping. But in their defense, they did provide a great lunch, and the views in route were spectacular.

Stellar's Sea Lions (above), Brown Bear (aka Grizzly below), and other animals were seen on the tour. At one point we watched a Brown Bear traverse a patch of mostly vertical rock in route to the beach were we ended up picking up the first batch of kayakers. Apparently a few other kayakers were picked up before from the same spot with a bear taking a nap on the beach.

Glacier views are provided half way in the tour along with lunch. Below is Margarie Glacier plus detail shots. The gray silt when dispersed yields the turquoise water color which proliferates in the Glacier Bay area. We did see some calving at Margarie, but didn't get any good shots of it.

The Lamplugh Glacier currently sports an internal waterfall,

A far better way to see the bay is by air, and that post will have to wait until Kathryn gets the photos finished...

Alaska : Mendenhall Glacier

2009-07-06T10:15:00.004-05:00

This 12 mile long and 1.5 mile wide glacier is located in Mendenhall Valley, about 12 miles from Juneau.

Trails in the valley lead around the lake created by the receding ice mass. This lake is filled with very blue icebergs. In the glacial areas, ice takes on this blue color due to air bubbles trapped inside the ice getting squeezed out in the glacier.

Around the lake there are 2 primary trails leading towards the glacier, the shorter easy hike ends at a huge waterfall (picture below). A secondary trail loop off this waterfall hike was closed (probably due to the trail being destroyed by a landslide).

On the other end of the lake starting at the camp area is another longer 7 mile trail which leads to a view over the glacier. We would have liked to do the trail, but we had to catch our flight back to Chicago in a few hours, so we opted to cheat and take a helicopter tour which landed on the glacier itself.

Helicopter Tour

Looking back towards the visitors center. The trail which follows the shore leading to the huge waterfall (which is out of view) is on the left side of this image.

Detail of calved icebergs at the glacier's terminus,

Glacial detail viewed from the chopper was fantastic,

Looking towards the landing spot,

Similar view after touching down on the glacial surface,

And yes, glacial melt water is quite refreshing,

These melt water streams carve through the top of the glacier surface,

Views can be quite stunning,

But eventually you have to return to the helicopter holding your hat on, to return the ice boot covers they provided you. For some reason I felt compelled to duck my head down, perhaps as a subconscious reaction that I could self terminate by jumping upward. Also note how they position someone to help, so stupid tourists don't deviate into the rear tail rotor.

The helicopter tour was the most expensive excursion we did. Something to do with turbine motor vs piston motor and amortizing the cost of the flying machine over the number of tourists... It was totally a last minute decision made along the lines of "we haven't done this before, and we probably are not going back to get a second chance to land on a glacier, might have enough time to do this before flying back to Chicago". Alaska's coastal weather is unpredictable, after a morning of near ground level clouds, and seeing the overcast cloud ceiling lifting, we took off for the airport looking for a building with lots of helicopters...

Google Map View of the Area

View Larger Map

Alaska : Eagles

2009-07-06T09:55:00.002-05:00

Gustavus has a lot of Bald Eagles which hang out around the pier at low tide. One evening, a bunch of them were fighting for what was likely fisherman scraps,

One eagle managed to steal some skin, and flew away to an isolated post on the pier to consume it. Food flying into the mouth, not out...

The eagle was close enough to get a good video without the camera's sorry excuse of digital zoom for lack of telephoto,

Multicore IC for Realtime Ray Tracing

2009-07-05T12:39:00.004-05:00

800TFLOPS Multicore IC for Realtime Ray Tracing describes a monster multi-core board dedicated to real-time ray tracing. Quick stats pulled from the diagram,

Board Target

- 1000W.
- 792 TFLOPs (only)!
- 73 cores on chip x 9 chips.
- 1920 x 1080 output resolution.
- Images traced using 35 channel color.
- Chips at 750MHz on 45nm process.
- 130 million gates/chip at 17mm footprint.

Totals Across 9 Chips on Board

- 864 KB L2 I$.
- 7056 KB L2 D$.
- 4608 KB L1 I$.
- 9216 KB L1 D$.
- 954 KB registers.
- 928 KB L0 memory.
- 2592 GB/s bandwidth.

Each Chip

- 8 clusters of 9 cores.
- 1 64-bit RISC CPU.
- 96 KB L2 I$.
- 784 KB L2 D$.
- 512 KB L1 I$ (total over 8 clusters).
- 1024 KB L1 D$ (total over 8 clusters).
- 106 KB registers (total over 8 clusters).
- 103 KB L0 memory (total over 8 clusters).
- 288 GB/s bandwidth (total over 8 clusters).

Each Cluster

- 64 KB L1 I$.
- 128 KB L1 D$.
- 13.25 KB registers.
- 12.9 KB L0 memory.
- 36 GB/s via 3 interconnect buses each 12 GB/s and 128-bit.
- 9 cores sharing registers (?).

CORE 0,
- Ray generation.
- Reserved register file.
- Register file : 8 bytes x 16 registers x 4 banks (512 B).
- 256 B L0 memory.

CORE 1,
- Pixel Traverse.
- 2-way SIMD.
- Register file : (4x2) bytes x 16 registers x 4 banks (512 B).
- ? B CAM (content accessible memory) L0 memory.

CORE 2,
- Bounding box processing.
- 2-way SIMD x 3.
- (4x2) bytes x 16 registers x 8 banks (1 KB)
- 4 KB L0 memory.

CORE 3,
- Intersection processing.
- 8-way SIMD x 24.
- (4x8) bytes x 16 registers x 16 banks (8 KB)
- 5.4 KB L0 memory.

CORE 4,
- Other ray-tracing processing.
- 4 bytes x 16 registers x 16 banks (256 B)
- 256 B L0 memory.

CORES 5, 6, AND 7,
- "Light Processing", write processing.
- 4 bytes x 16 registers x 16 banks (1 KB)
- 1 KB L0 memory.

CORE 8,
- Not shown in the diagram?

Odd Transparent Rendering Ideas

2009-07-05T12:17:00.002-05:00

Transparent content tends to require a trade-off between overdraw and sharpness. The ideal solution in my mind is to able to have sharp features (high resolution) with low amounts of overdraw (more opaque) intermixed with smooth features (low resolution) and high amounts of overdraw (more transparent).

With 100% order-independent rendering, such as only additive transparency, one possible solution would be as follows,

(1.) Have a set of accumulation buffers at reducing resolutions.
(2.) Distribute draws to buffers based on sharpness and overdraw.
(3.) Combine all buffers.

Combine buffers requires repeated passes of up-sample and merge in order of the smallest buffer to the largest buffer. Note it might be a good idea to filter better in the up-sample passes to avoid the typical mip maximize artifacts. Also note these combine passes are going to be bandwidth bound, so its a good idea to pair with some ALU bound work in the same kernel if possible!

TANGENT ON BLEND MODES

Typical blend modes used in order-independent rendering are ADD and MUL. ADD to lighten, and MUL to darken. Often using both requires 2 passes, and when mixed with multiple resolutions, artifacts are found such as background MUL particles (say smoke) drawing over foreground ADD particles (say fire).

Another issue is gamma-correct (linear space) blending, are transparent surfaces combined in linear space or after the linear to gamma transform?

In the case of gamma space (non-linear) blending, I've never liked ADD blending at all for additive transparency. Perhaps this is because of years of photography where I go to complex lengths to avoid over-exposure. HDR in games is often (IMO) ridiculously overdone (really, should 25-50% of the frame be white). A visually better option (IMO) is Photoshop's SCREEN blend mode. SCREEN blending results in a soft approach to white, with blending limited at white. From a photography perspective, SCREEN works more like the way film responds to over-exposure. SCREEN is an inverted multiply,

d = 1 - ( (1-a) * (1-b) )

The inversion required by SCREEN can be factored out to doing MUL blending with an inverted framebuffer when attempting to port this over to DX or GL or GCM.

In the case of linear blending, at least in Photoshop, I've found MUL for darkening unacceptable with both 8-bit and 16-bit modes. To test this in Photoshop, one must actually build a linear color profile (which is quite easy). The problem is the loss of precision (banding, etc) when the MUL results in colors near black (even in 16-bit mode). Perhaps with 16-bit float framebuffer outside of Photoshop this isn't a problem.

BILLBOARDS AND SPHERICAL PARTICLES

Particle billboards suffer from the problem of hard intersection with solid geometry, and in the case of sorted billboards, pop when sort order changes.

Both of these problems have similar solutions, dynamically adjusting transparency across the particle to avoid the artifact. Spherical (or Z buffer aware) particles easily solves the first problem of intersecting with solid geometry. The second problem of pop on order change is more complex.

TRANSFORMING ORDER DEPENDENT TO ORDER INDEPENDENT?

This is a wild completely untested idea I haven't had time to try. Note the application of this would be for a non-triangle renderer which doesn't have standard opaque triangle surfaces. Application of this to a typical game engine is likely not worth it (requires two alphas, one for weight and one for transparency, or other complicated things)!

The idea is to solve the following problems in one go,

(1.) Removal of the order change pop on billboards.
(2.) Ability to easily work in multiple resolutions for performance.
(3.) Ability to do blend, lighten and darken.
(4.) Ability to do all this in one "pass".

The idea is similar in function to front first blending in which RGB accumulates color and A accumulates coverage. Except now RGB accumulates weighted color, and A accumulates weight. Actual RGB color is weighted RGB divided by weight A (assuming a FP16 buffer).

However, drawing would be order-independent, billboards rendered by computing a weight factor of visibility across the billboard's surface (or perhaps for smaller billboards, constant weight factor across the billboard's surface). This visibility factor would take into account both the billboard's transparency and expected occlusion by other billboards and surfaces. RGB would be billboard color times weight. Also for lighten and darken, weight stored in A could be adjusted.

Clearly this is adding another problem to solve (visibility estimation for weight factor). However I believe the weight factor could easily be computed at a lower effective frame rate and at a much lower resolution, possibly leading to better performance overall.

Alaska : Whale Watching

2009-07-03T19:57:00.012-05:00

Hopefully flickr won't mind the deep linking...

Alaska was incredible, not the typical vacation location, and IMO perfect for the independent traveler. My wife Kathryn and I spent over a week in the Glacier Bay area, and over a few blog posts I'm going to share a little of the experience.

The TAZ

I'd highly suggest the Icy Straits afternoon whale watching excursion of Tod Seben's Cross Sound Express, also known as The TAZ in Gustavus. Easy to see how the Humpbacks were nearly hunted to extinction. They are simple to find if one knows where they feed, grab a set of binoculars and head towards the sprays. Being up-wind of the spray provides the opportunity to experience "whale's breath". It is amazing that something so foul smelling could travel so far without dispersing (more).

Watching the Humpbacks is something else indeed, especially watching a 40 foot animal jump out of the water. Besides breaching, they spy-hop, tail-lob, slap, and wave. At one point in the afternoon, a pod of whales turned and closed the distance to the boat while we were stationary and watching (regulations specify a minimum distance to approach the whales, however they can choose to approach the boat on their own). Size becomes apparent when seeing something the size of the boat swim by.

Well worth it. At one point, Tod let Kathryn pilot the TAZ,

Bubble Net Feeding Viewed via Air Taxi

Towns and cities in the Glacier Bay area are connected via air and water only. Roads don't exist except locally. Our taxi pilot mentioned that he had seen a pod of Humpbacks from the air on route, and said he would attempt to find them again as we took the air taxi from Gustavus to Juneau. By some act of random chance, he did indeed find the pod, and they were bubble net feeding, which according to our pilot must have been rare to see in the air because he had never see that before. We got a first hand look via a bunch of descending tight banked loops.

Below a shot of the pod after feeding. BTW, that wasn't telephoto, all we had was a small point and shoot, which apparently has some issues with white point computation given light angle to the water...

iPhone 3GS GLBench

2009-07-01T23:55:00.003-05:00

The GLBenchmark site provides a way to compare the GL performance of all the iPhone models including the 3GS! Results are very interesting indeed showing 3GS with 10x the triangle performance of the prior models, and 3x the texture fill rate.

Yes 18 hour days but no bugs.

2009-06-30T18:13:00.015-05:00

Back today from Alaska via overnight flight, will have a blog post on the trip once the photos are sorted through (and BTW I think Wisconsin has worse mosquitos). Playing catchup with everything missed over the past 9 days.

What looks (edit) NOT (edit, thanks Cyril) to be some DX11 realtime GI screenshots, one with indirect specular highlights. Found from the AMD DX11 Whitepaper Article, but with no link to the actual white paper, no new news as to DX11 (give us something real like limits on Append/Consume buffers mapped per shader). I'm also still disapointed at the CS4.x spec in how it doesn't provide support for features in current shipping NVidia DX10 cards (referring to G80 support for shared memory write access and later model atomic operations). Perhaps went to ATI's favor, or will ultimately help NVidia force DX11 upgrades?

Cuda 2.3 Beta is Available for Registered Developers. According to the forum post, provides float2half and half2float which I believe matches the DX11 functionality. Also some interesting unanswered questions regarding __synchronous_start(int s) and __synchronous_end() which looked to be found in the intermediate nvcc output.

Interesting stuff in the Shack's John Carmack Interview, "Now I am very excited about what I can do from a hardware and graphics standpoint with the 3GS. With vertex fragment shaders and OpenGL 2.0, I'm pretty convinced that I can actually run the MegaTexture id Tech 5 content creation pipeline on there. And I'm not sure what game I want to do that with yet, but the combination of seeing people download 700mb files of Myst on there, and the new capabilities, I could do some mind-blowingly cool stuff on there."

Battle Field Heroes is live.

18 Hour Days, Lots of Bugs, but No Computer?

2009-06-19T16:04:00.007-05:00

My wife and I are going up to Alaska for next week. Packing light, not going to bring the pro camera this time (sticking to a real vacation), between the two of us, just two carry-on backpacks. Going to be awesome.

Requesting Siggraph 2009 Advice

2009-06-17T23:55:00.002-05:00

I'm looking for any advice as to what to see and do this Siggraph 2009, since I've never been to a Siggraph before, and I've got to book everything this week (going on vacation next week before the early deadline ends). Also looking to meet up with any of you who read this blog and want to talk about tech and stuff! Please either email my gmail address (farrarfocus) or post comments below.

Currently planning on flying in 10am Wednesday and doing two day passes for the following,

Wed 5th: 1:45pm-5:30pm - Efficient Substitutes for Subdivision Surfaces
Thr 6th: 8:30am-12:15pm - Beyond Programmable Shading I
Thr 6th: 1:45pm-5:30pm - Beyond Programmable Shading II

Stream Compaction for Deferred Shading

2009-06-16T10:58:00.002-05:00

Stream Compaction for Deferred Shading
Jared Hoberock, Victor Lu, Yuntao Jia, John C. Hart
Proc. High Performance Graphics, Aug. 2009

Paper addresses the issue of SIMD shader divergence when deferred shading on Geforce 8/9 series hardware using the following tests,

UNORDERED IMPLICIT - All N shaders in one program, use possible divergent branching (at pixel level) to shade pixel. In the case of a small number of shaders (or low run-time shaders) the overhead of stream compaction or sort was higher than divergent branching.

UNORDERED EXPLICIT - All N shaders in separate programs, run one full pass for each shader (N passes). Was always slower than the UNORDERED IMPLICIT path, even though max register usage (controls max warps running in parallel for latency hiding) would be smaller for explicit shaders. Overhead of fetching shader type just to mask threads was too high.

COMPACTION IMPLICIT - Overhead proportional to the number of required shaders. Use stream compaction to group pixels by shader and use possible divergent branching (at beginning and end of groups) to shade. They run N passes of a masked parallel scan (one per shader) to create offsets for stream compaction. Then one shading pass using possibly divergent branching (at the N shader boundaries in the compacted stream) to shade. The shading pass fetches pixel offset from the scan results, gathers the pixel data, shades, and scatters the results back.

COMPACTION EXPLICIT - Same as above, but run a pass per shader when shading. This case looks to always be faster than COMPACTION IMPLICIT because of better register usage when shading (better latency hiding).

SORT IMPLICIT - Overhead proportional to the log of the number of shaders. Instead of running N scans, radix sort by shader for the first pass. Shading is the same as COMPACTION IMPLICIT except using the radix sorted offsets instead the scan offsets.

SORT EXPLICIT - Save as above, but run a pass per shader when shading. This case provided the best results overall when overhead of running the sort was smaller than the cost of divergent branching in the UNORDERED IMPLICIT case.

LOCAL IMPLICIT - Do COMPACTION IMPLICIT in small tiles and only use shared memory. They were only able to support a working size of 8 warps (256 pixels). I think they ment 8 warps total per block and one block per multiprocessor, which would not be enought to hide memory latency, and would stall the multiprocessor on sync threads. Given the small tile size, shading costs alone were always higher than the UNORDERED IMPLICIT case.

About the Results

The paper results are in milliseconds on 2 types of 9800 cards. They don't really provide enough information (no number of samples shaded). My guess working backwards from 9800 GTX+ specs and average instructions per shader numbers in the paper, is about ~0.2G (G=billion) shader invocations in the Car scene (for a 512x512 image estimated from the PDF, also assuming ALU bound, which probably isn't the case) for 337ms of shading time (note sort overhead was an extra 244ms). I believe a 8800 GTX can do a 1M (million) entry parallel scan in under 1ms according to a previous paper I read. So 244ms overhead on a 9800 GTX+ likely collaborates my estimation that they are working in numbers greater than 50M at a time?

There is much more to talk about with regards to this paper, deferred shading, and compute/DX11, but too busy now...

Braid for the Mac

2009-06-16T10:15:00.002-05:00

Old news which I never got, Braid is Available for the Mac. Woohoo (and yes I'm sans PC or 360 at home).

Factoring Out the Job Scheduler

2009-06-11T16:03:00.002-05:00

Last night I got another idea on how I might be able to both simplify my job scheduler (in fact remove it) and better handle a dynamic CPU load. All while keeping a very simple design with minimal sync points. Be warned, I have NOT had time to try this yet, it should work in theory...

Threading

(1.) The program entry starts N threads then goes to sleep forever.
(2.) N threads to match the N hardware threads of the machine.
(3.) N threads are locked down via thread affinity.
(4.) Each of these N threads runs the exact same function.
(5.) Each thread starts with function(X, N).
(6.) Where X is the thread index ranging {0 to N-1}.

The basis of parallelism here is running N instances of a program. However since the N instances can have completely divergent branching, the program can use X to mix and match task and data parallel computation anywhere in the program. Note all threads share the same memory.

Synchronization

For efficient computation only very coarse grain synchronization is supported between the N threads. Coarse as in you can count the number of sync points per screen refresh on your fingers. Examples below show lines for each thread and letters as follows,

- -> thread doing work x -> thread blocking E -> event S -> sync point, thread blocks until all threads finished E

During normal program operation sync points don't ever block because enough work is provided in between by a good program architect. Note in this model threads have a lot of slop and don't ever actually fully sync up in time unless something goes really wrong.

--------E-------S------- -----------E------S----- ------E--------S-------- -------E------------S---

When something goes wrong, threads do block.

--------E-------S------- ---------------E--S----- ------E-----Sxxx-------- ---E-----Sxxxxxx--------

The program would be composed of a sequence of E and S points per frame.

-----E1----E2S1-----S2--

In the above example, after S1, data generated prior to E1 is free to be read-only shared between all threads. Likewise after S2 data generated prior to E2 (ie between E1 and E2) is now safe to share as read-only.

Data outside sync points is considered "unordered" (just as UAVs are in DX11) and can be shared only with DX11 or CUDA like unordered atomic operations. Data which would otherwise require fine-grain synchronization to share, is simply re-computed in each thread.

Dynamic Load Balancing

This is where things get interesting. The goal case is where threads reach E at the same time (E is in the same point in code on all threads). There is enough slop in the system so that the E point in time can be a bit different between threads. To load balance one would want to dynamically adjust work distribution such that at the next E point, that the threads would be closer together in timing.

To enable dynamic balancing, all threads have a secondary work distribution factor Y (the primary work distribution factor is the thread index X). This Y factor provides a segment {start,end} along the range of {0 to 1}. This Y segment can be used by the thread to choose the ranges of data parallel work to operate on.

(1.) At E, each thread stores CPU clock cycle time.
(2.) This requires CPU clock cycle counter bases to be re-synced periodically.
(3.) Also at E, each thread records diff in cycle time to the previous E.
(4.) At S, threads read all the E times for all the threads.
(5.) Note the E times are read-only.
(6.) Each thread uses the E times to compute a new future Y segment.
(7.) This new future Y segment stored to be used at the next sync point.

Example Forms of Task and Data Parallel Programming

To have one thread handle an OpenGL interface, the program would just dynamic branch on X==0 (X being the thread index). Note, a GL context is usually tied to one thread. With DX11 there is also the concept of the primary thread, secondary threads build deferred commands. Note it is ok to lock down a task to just one thread because the dynamic load balancing would redistribute the data parallel work via the Y segment (from above).

To process M objects, each thread would compute an exclusive set of the M objects using the Y segment (something like Y.start * M to Y.end * M, with fixup to insure no overlap).

Advantages

With this system, there isn't a worklet system or a scheduler, there is just the program running multiple instances of itself. Synchronization is flexible, and rarely used (very parallel friendly). Read-only data would likely be well shared on a shared L1 or L2 or L3 because the same program is running on all threads and the dynamic load balancing works to bring together the phase (the time) of the threads.

iPhone 3GS

2009-06-11T11:20:00.001-05:00

AnandTech: The iPhone 3GS Hardware Exposed & Analyzed

The new iPhone 3GS looks to be awesome, according to AnandTech,

- PowerVR SGX GPU
- OpenGL ES 2 (shaders)
- likely a 3.5x to 7x triangle throughput increase
- ARM Cortex A8
- 600 MHz
- dual issue instructions
- 32KB L1 I$, 32KB L1 D$, 256KB L2
- Neon
- likely a 2x processor performance increase

Yeah really cool! As long as VBOs and other SGX GPU hardware features work in iPhone OS 3.0, developers are bound to get really happy.

The Last Guardian Images

2009-06-10T15:11:00.003-05:00

The Shack has images of The Last Guardian. Looks very much like SotC to me in muted color palette choice.

Some Assembly Required

2009-06-09T10:40:00.002-05:00

This mornings googling turned up a really great blog, Some Assembly Required, written by Elan Ruskin at Valve. Great topics + disassembly + timings = awesome!

DX11 Binning

2009-06-05T09:03:00.004-05:00

BTW, I over simplified my DX11 binning example in a previous post. In reality there are some important considerations for speed. For example on Larrabee since scatter time is a function of the number of cache lines, it might NOT be wise to bin into queues arranged in memory like this,

int Queue4Pixel0[16], int Queue4Pixel1[16], int Queue4Pixel2[16], ... offset = (pixel << 4) + bin

Because only 1 queue would be on a cache line, so the scatter at best would be 16 cycles. Read back of the data later also might not be ideal in this format. Other option is this which has better cache locality at the expense of more complex offset logic,

int QueueBin0ForPixels0to16[16], int QueueBin1ForPixels0to16[16], ... offset = (pixel & 15) + (((pixel & (~15)) + bin) << 4);

So for each grouping of 16 pixels, each queue bin is on the same cache line. Which in the completely coherent case would take a clock cycle to scatter to, and would degrade in performance as a function of how many different queue lengths there are in the groupings of 16 pixels.

The same above logic goes for global scatter in CUDA (NVidia) since performance is a function of the number of segments touched. Similar logic could work for shared memory (or shared register) scatter but there are other options in CUDA since performance is a function of bank conflicts (and not "cache lines" or "segments").

Clearly this is assuming via ATI and NVidia DX11 hardware that texture UAVs are built from CUDA/OpenCL/DX11CS global memory accesses. This assumption could be dead wrong, and if another functional unit is involved (like UAV enabled ROPs) then I would need to re-explore this...

CS vs PS

Oops, another important consideration I left out is the assumption that the binning is happening from the CS stage (so block arrangement insures vector alignment and coalescing).

PS stage provides some complex problems in that pixel grouping in the SIMD units isn't likely always coalesce/vector friendly with the address calculations above. Full SIMD width number of pixels (full course raster tile) is likely to get packed in a very vector cache line friendly way on Larrabee, but you'd have to know if the output addressing was Morton order (bit interleaved or not) and adjust the offset calculation. Grouping of 2x2 pixel quads (or just pixels) into SIMD vector groups for small triangles probably breaks down the binning into something vector scatter unfriendly on Larrabee (extra different cache lines). Banked memory could lessen this limitation (so who knows what NVidia's GT300 performance will be).

Thread Scheduling Part 3

2009-06-05T01:44:00.002-05:00

My Personal Views on Solutions to This Mess

*** Personal views for the independent developer. Clearly things are different with the constraints a huge legacy code bases, etc...

First a little background, back when I started programming, is was practical to write your own operating systems and drivers (I did this in my teens). I would boot DOS, then switch into protected mode with the DOS interrupt table and low memory intact so if necessary I could switch into virtual x86 mode and use DOS for things I didn't feel like programming my self (like access to the FAT files system on the main harddrive). Everything else, sound, video, etc was done accessing the hardware directly. Actually things where much EASIER back then (IMO). Since then, the industry had added layer upon layer to abstract out the hardware...

WHAT DO I MEAN BY EASIER? - Back in the early days you got control over everything. Computers were single processor, no virtual memory, no threading. Applications (at least the stuff I did) would work with non-blocking asynchronous IO requests and coarse polling with a small set of interrupt service routines to handle hardware interrupts. It was rather easy to poll with one bitmask read to check for completion of any IO request, then cooperatively "multi-task" in application. Things were actually more interactive and responsive back then as well. Polling was cheep (read from memory, no kernel call), task switch was cheap because it was cooperative, no state to save and restore. Programs were tiny, systems were tiny. Check out Contiki for a modern example!

I hold no illusions that somehow we could go back to the way things were before, however there are a huge number of lessons from back then which should be applied now.

What I am Doing Circa 2009

First I have one thread per hyperthread in the system (or per processor in a non-hyperthreaded system). These threads are locked down via setting thread affinity. These threads are my "compute threads" which don't ever do blocking IO. Whenever possible I use low-level as possible non-blocking IO interfaces (this covers networking and raw file IO on most systems).

For all the other cases, meaning when the operating system or API doesn't provide a 100% non-blocking interface, I resort to blocking IO threads to which I non-locking queue requests to. Note because scheduling granularity is 1ms or so at best, queues get groups of requests at one time (to lower atomic operation overhead and overhead of kernel call to wake a IO thread which was blocking on an empty queue). Blocking IO threads run at a higher priority so they will be insured to preempt a "compute thread" whenever they get marked as runnable. Blocking IO threads are designed to have very short runtime and thus a minor effect on the "compute thread" which it prempted.

My job scheduler is quite different from what others do (it is modeled directly on how GPUs process draw calls and what is now the OpenCL model). I break my program down into an array/list of small jobs (like draw calls). I handle handle different sections of the program by turning on/off jobs in this list.

Each job has an index to a prior dependent job (which controls when the job can be scheduled). I keep track of job runtime (free profiling) and use this to pre-schedule all jobs into one queue per "compute thread" a frame in advance. So "compute threads" simply pull jobs in order off their queue and spin if in the unlikely case the job's dependency hasn't finished execution.

It is up to the program designer (ie me) to insure that the program pipeline doesn't ever need to stall and thus has enough non-dependent jobs to handle the variability of the preemption by the blocking IO threads.

Note that with this system I don't ever need any other thread-2-thread synchronization other than atomic operations. All that "work" and "mess" is factored out into proper program "pipeline" design!

To handle scalability to machines of varying number of processors, I have two types of jobs: (BATCH) This job is attempted to be scheduled on all "compute threads" at the same time. Each job entry point gets the "compute thread" index and the number of "compute threads". It is then up to the job to carve up the data and process its group. Batches are for data parallel work. (TASK) This job gets scheduled on only one processor. Tasks are for task parallel work.

For fast development cycles, I have all my code in a single dynamic (or shared on unix platforms) library, so I can re-compile, unload the old library, and load the new library to instantly test code changes without needing to restart the program. This is my re-attachable code model. BTW, I don't use the debugger ever at home, I don't need to!

Thread Scheduling Part 2

2009-06-04T23:31:00.007-05:00

Part 2 is all about the IO side of Thread Scheduling, CPU IO in a parallel world.

Review, Blocking vs Non-Blocking

BLOCKING - Thread which issues the IO call is put to sleep until the operation is "finished". At some point in the IO call, a kernel call is made (to service the IO request) which results in a run-level task switch to kernel mode. In kernel mode, the kernel queues the IO request, sleeps the calling thread and does a task switch to another runnable thread. At some point later after the kernel services the IO request, the original thread (which issued the IO) is again scheduled to run (latency can be a function of IO latency plus how often the blocking thread gets rescheduled to run).

LIKELY BLOCKING, POSSIBLY NON-BLOCKING - In the case of reads, sometimes the data is cached (perhaps disk cache) and thus the IO call can be serviced right away. This can also be the case for writes when the kernel call copies the write request to a queue (or to the page cache) and immediately returns (note when queue is full or page isn't in page cache, kernel call can become a blocking call).

MOST LIKELY NON-BLOCKING, POSSIBLY BLOCKING - This is the troubling case where an interface is almost completely non-blocking, but has some caveats. Caveats in that the interface doesn't return some error code indicating that the call would block. Often this is the case when buffer or queue space is full. The v-sync section in the previous post is a good example of this!

NON-BLOCKING - Correct fully non-blocking API, if a call would block, the call returns with an error code. Non-blocking APIs might provide a polling function to return the status of pending operations. Non-blocking APIs might also provide a callback interface. In this case a user supplied function is called on pending IO. Callback interface can be lower latency at the expense of loosing control over job synchronization.

Disk Access Re-ordering To Reduce Seek Time

All operating systems which have good file IO performance on rotating disc based storage do this (else performance would suck worse than it does now). The operating system queues up the disk accesses required to service the file IO (virtual memory IO, etc) from all processes and threads, and orders the accesses to reduce seek time.

This should be kept in mind when thinking about application file IO. It is often smart to queue up file IO requests from many files at once, instead of just following the read->process->read->process model. Likewise, having just one asych thread servicing file requests using a blocking interface might not yield the best throughput depending on access patterns.

Read Ahead

The operating system provides a read-ahead mechanism which attempts to predict which pages will be needed in the future (based on past access patterns) and effectively adds future requests to current requests when reading from the device. In some cases (Linux) it can get various patterns like sequential + stride, etc.

Why Your Desktop User Experience Is Often a Waiting Game

(1.) Same thread processing the application "GUI" event queue does file IO. Epic Fail!

(2.) Application is creating a huge data structure, it writes into memory which was just allocated, but hasn't yet been mapped to physical pages by the OS. System is under heavy load, and not enough non-dirty pages are available to discard. Operating system needs to write out dirty pages to disc to provide pages for the application to write to. Application blocks until disc IO is finished on the dirty pages. Epic Fail!

(3.) Application suffers from code bloat. App could even be tiny but suffer from a bloated hierarchy of dynamic link libraries. Non-accessed code likely isn't resident (or if loaded, was paged out to provide space for something else). Example, user clicks on menu and opens up a seldom used dialog, paged-out code is accessed, and App blocks on file IO. Epic Fail!

(4.) Application access a huge number of small files, all with long path names (think bloated file structure with huge number of files per directory). Seek nightmare for two reasons, first just accessing the (possibly uncached) directory structure to find the files, and second accessing the (likely uncached) files. Epic Fail!

Solving the first case is just good application design. Solving the third case is expert application design. One has to insure all GUI or user-input processing is in a small core amount of common code (and data) which will always be resident (often accessed likely equals always resident). Possibly blocking functions (ie ones which might access code which is seldom called, ie which is paged out) needs to be done in an asynchronous thread. This way the app is always responsive!

Point here is anything which might result (directly or indirectly) in file IO must be on a separate thread from whatever thread or thread(s) interact with the user.

Flavors of File IO

CACHED STREAM FILE IO - Think the standard C library interface: open(), close(), read(), write(), etc. Reads and writes end up as a mix of blocking and non-blocking kernel calls. Nothing new here. I avoid this stuff.

MEMORY MAPPED FILE IO - The case where the operating system provides file IO directly using the virtual memory system. All the mapping kernel call has to do is to setup the page tables to invalid and then setup the os pages mapping to map to the file. At this point it can return (however it might prefetch based on mapping hints). When the app then attempts to first access memory in the mapped range, a page fault will happen, the kernel will then queue the disk IO, and sleep the thread (assuming the access isn't beyond the end of the file, also assuming the disk page isn't cached in memory). Depending on access patterns, the OS might prefetch N extra pages on fault (read-ahead). Nothing says that the entire file is loaded.

A scheduling pitfall of memory mapped IO is possible lengthy page faults even after the first access. So if using memory mapped IO in a blocking service thread, that thread would have to manually touch (read a integer from) all the pages to insure they are all actually loaded (skipping pages isn't an option because the read-ahead might detect that an then only service the skipped pages). This way when the async service thread marks the IO as complete, the program can assume a compute thread will not block. Problem with a blocking async thread touching pages to insure they are loaded is that each read will be a full (wasted) cache line miss out to main memory (if each miss is 1000 cycles, on a 2 GHz machine with 4KB pages, 1 ms of miss time is only good for at best ~7.8 MB of touched blocks, costing you 1/16 of your 60Hz frame). So totally NOT practical!

Pages modified are automatically marked by the CPU so on unmap the OS can only write out modified pages for free. Also in theory, memory mapped IO could provide the operating system a way to avoid an extra copy on read and write from the file system. This is assuming the OS does direct file DMA to/from the pages. Avoiding the copy would be avoiding the cache pollution and stalls caused by huge CPU to CPU copies. Avoiding the extra copy sitting around in the disk cache also saves memory.

Memory mapped file IO is available on Windows/Unix/Mac/Etc, but not on consoles, and other embedded systems!

RAW DIRECT ASYNCHRONOUS FILE IO - Given the problems with cached or memory mapped file IO, for databases and other critical high performance applications, operating systems often provide an interface for raw file IO. Interfaces for raw file IO are also commonly non-blocking and asynchronous (for obvious reasons) and likely batch request friendly.

Windows Asynchronous Disk I/O

Quick refresher, I'm not personally interested in things like accessing a huge number of random files. IMO that is just broken by design. What I do want is background IO with simple polling for completion (no heavy weight kernel call), and no duplication of memory (ie copy of the file block in the file cache). So I'm not going to cover things like I/O completion ports...

Microsoft has a great article on this which should be read first. Windows AIO has one key design flaw in that the operating system can choose to go synchronous for an operation without giving the application a choice. This IMO is an Epic Fail case for any asynchronous interface.

Cases of going synchronous are: NTFS compression, NTFS encryption, extending a file's length, and large number of cached file IO requests. FILE_FLAG_NO_BUFFERING (raw direct un-cached IO) best insures an actual asynchronous call. However from their example (of no-buffering async calls), only 500 requests got queued in 0.22 seconds. Which shows a 0.45ms average kernel call time. Yikes, that is either certainly not asynchronous, or just spin locking, or the thread is getting heavily preempted!!!

Windows does provide fail cases with ERROR_INVALID_USER_BUFFER or ERROR_NOT_ENOUGH_MEMORY whenever there are too many outstanding asynchronous I/O requests.

Windows also does provide a HasOverlappedIoCompleted() macro to check for AIO completion without using a kernel call.

MacOSX Asynchronous Disk I/O

POSIX asynchronous IO wasn't supported until MacOSX 10.4. POSIX AIO provides a true non-blocking interface, if a request cannot be enqueued, then functions fail. POSIX AIO requires a aio_error() "system call" to check on status of an async operation. I believe that POSIX AIO on OSX is a wrapper for kqueue() so polling does indeed require a system call (yuck). According to some internet sources a process is limited to 16 AIO requests at a time on OSX. Probably better to just use kqueue() directly on OSX.

Linux Asynchronous Disk I/O

Included in 2.6 kernel (patch for 2.4). Works on files or block devices opened with O_DIRECT or raw devices. Doesn't work on sockets or pipes. One option is to directly interface with the io_*() system calls through the libaio userspace wrapper library. Basically,

#include < libaio.h > // setup context aio_context_t ctx; bzero(&ctx, sizeof(aio_context_t)); io_setup(max_events, &ctx); // submit events struct iocb tab[events]; io_prep_pread(tab+0, fd, buf, bytes, offset); io_prep_pwrite(tab+1, fd, buf, bytes, offset); ... io_submit(ctx, events, tab); // poll on completion using syscall // will want to wrap this in something which disables signals struct io_event tab2[events]; num_events_or_error = io_getevents(ctx, 0, events, tab2, 0); if(tab2[0].obj == tab+0) // then that event is finished // tab2[0].res2 == 0 then ok, else error // tab2[0].res == amount read/written

Links: Linux System Calls, Neat Interactive Map of Linux Kernel

Other option would be to use libposix-aio to get the POSIX AIO interface (which uses io syscalls). The link above is to the source which is a good reference to the syscalls (for example they wrap io_getevents() in something which masks interrupts, stuff like that likely learned the hard way through trial and error!).

BTW, Linux too fails on the most basic level of simply directly updating the iocb with the completion status so a kernel call isn't needed to poll. Arg.

Non-Blocking Network Interface

I personally don't use TCP, and only do UDP. Via UDP, the problem of the "connection" interface to the operating system goes away, and a single socket can gather all incoming packets from all clients and/or send packets to all clients (recvfrom()/sendto()). A quick ioctl(), or ioctlsocket() (with Windows WSA emulation of BSD sockets), can be used to set non-blocking operation.

Audio

Going to leave this for another blog post...

More DX11 Thoughts

2009-06-04T21:56:00.002-05:00

Was going through Bill Bilodeau's Presentation off the repiloque blog, and realized I had missed some important info from slide 20 saying that the "3 times faster and 1/100th the size" is comparing ATI DX9 (actually DX10, with vertex texture fetch) tessellation (requiring at least one pass to export per edge tessellation factors, and a second to tessellate) to running the high poly mesh through the standard DX9 pipeline, also that "DX11 Tessellator algorithms can usually be done in one pass". Guessing this "usually" is suggesting that hardware might still optionally spill intermediate results to main memory under a few rare conditions?

Given that tessellation is a data amplification process and that pixel shading is (for the most part) also a data amplification process, it isn't too hard to understand that the hardware could likely keep intermediate data all on chip. GS is the wild card there, and I don't plan on using GS at all anyway.

UAV Limits

- Pixel Shader limited to 8 RTVs+UAVs total
- Compute Shader limited to 8 UAVs

So the limit of the union of Render Targets and Unordered Resource Views is the same as the limit of Render Targets... says to me that at least one DX11 vendor will be using the same interface for both RTVs and UAVs. Also UAV texture access is still unordered, I'm thinking any deep buffer magic (k-buffer, etc) is going to require atomic operations...

Speaking of a combined RTV+UAV ROP unit. So does the new DX11 path support better fine granularity (32-bit) scattering (assuming the scatter has good data locality)? Nick Thibieroz's Shader Model 5.0 and Compute Shader slides describe "Order-Independent Transparency" and "Data Binning Operations" as applications of UAVs. Is this really going to be fast enough make binning pixels in a single pass useful? With Larrabee the answer is YES (I think), what about ATI and NVidia? Currently with UAVs emulated via CUDA on a DX10 GT200, the answer is a clear NO.

Binning literally means something like this (in C like code, not HLSL),

bin = atomic_add(queue_head + pixel_offset, 1); if(bin < max_bin) { queue[pixel_offset * max_bin + bin] = value; }

Note the first line is a fully coalesced vector friendly GLOBAL atomic add (this is an unordered view). In theory this could be fast (coalesced) but latency unfriendly because of the needed return value. The second line scatter requires a fast 32-bit scatter (value is a 32-bit pixel) to a "cache" or ROP write-combining buffer to be bandwidth friendly. If binning of 32-bit values with good locality gets fast, then an amazing amount of algorithms become possible on DX11. I've been waiting for this since the CUDA PTX .surf space has been documented but never implemented in hardware...

Append / Consume

No limits posted yet that I can find (I haven't tried/installed the DX11 software reference), but ATI presentations seem to show only Raw or Structured Buffers can be mapped as Append or Consume in a shader but seems like NOT as both Append and Consume at the same time? I'm also guessing Append output or Consume input is to/from main memory since it seems most likely that the Append to Consume boundary is at the draw call. I'm also thinking the limits on Append/Consume buffers mapped in a shader will be small.