NVirt
This past weekend, I attended the I3D conference for the first time. The conference was great. There was a ton of really interesting work, and the keynote and capstone talks gave me a really fresh perspective on some of the aspects of computer interaction (graphical and otherwise) that we, as graphics researchers, often neglect or take for granted. The talk by Austin Robinson, however, contained a completely unexpected "reveal". Apparently, Nvidia has developed a raytracing/raycasting API (NVirt).
Nvidia seems to envision NVirt as a low level API, providing something akin to an OpenGL or Direct 3D for raycasting. In reality, however, it seems to be a bit more high level. Austin's slides (when they become available) have some great images demonstrating how the pipeline works, and where the hooks are introduced for the user to customize the tracing procedure. The geometry is loaded into the system through a scene graph esque interface. In addition to the basic information (position and direction), the user can customized the ray structures they want to trace to carry along any extra information they deem necessary. Once the geometry and rays are defined, the user writes a number of kernels that define hooks for various events which occur during ray traversal. First, there is a hook for ray generation, where one can generate rays using what appears to be an arbitrary kernel. Thus, it is possible to simply implement a pinhole camera generator, or to pass in a buffer of rays, or many things in between. There are many other hooks for which one can customized the kernels, though I do not remember them all offhand. There is one that is called for any intersection of a ray, one that is called for the first intersection of a ray, one that allows the selective traversal of children, and another that is called when a ray fails to intersect any geometry.
Perhaps some of the most interesting things that NVirt brings to the table are those we will never be able to see. Apparently, they guys and gals at Nvidia have been doing some compiler level voodoo to present the NVirt user with capabilities that we don't currently see in CUDA. For example, it appears that one will be able to call the ray tracing routine from within a user written kernel. In CUDA, this would be akin to the ability to call a __global__ function from within a CUDA kernel; something which is currently verboten. Hopefully, Nvidia will educate us on some of this compiler technology, as it seems very useful in enabling more general kernel behavior even beyond the bounds of ray tracing.
Finally, Austin made a few comments on the "level" of the API. In particular, the reasons they chose to hide certain details from the user. The idea seems to be that the user should have as much control over ray intersections as possible to allow a wide variety of raytracing-esque algorithms. For example, you should be able to use this technology in implementing a photon map. However, enough detail should be hidden so that the interface can remain consistent while implementation changes. Clearly, G300 class cards will be different from G200 cards, which are different from G80 cards. Nvidia may discover optimizations that have a tremendous effect on performance on current of future generation hardware. As long as these optimizations don't change the interface, they should be implementable. If the user was given complete control over every detail, then things like alteration of the traversal structure could not be done. This means that people implemented in the more "fundamental" parallel ray tracing research (like developing spatial subdivision and ray traversal structures) will still have to roll their own GPU ray tracing solutions. However, for the large majority of us, NVirt seems like a very interesting tool. I hope, and expect, that the research community will use it to produce some very interesting results, and that those in industry will take commercial eye candy to new heights.
