Welcome to Fractal Forums

Fractal Software => Programming => Topic started by: ker2x on March 26, 2012, 07:48:26 PM




Title: NVidia Optix 2.5 is out
Post by: ker2x on March 26, 2012, 07:48:26 PM
http://www.nvidia.com/content/newsletters/web/nv-newsletter/optix-email-web.html

Highlights include:
  • Out-of-Core Memory Paging — scene sizes can now exceed physical memory on professional GPUs (Quadro or Tesla, host RAM permitting).
  • Unlimited Textures — More than 127 textures can now be used (127 texture limitation still applies to textures used in graphics interop).
  • GPU BVH Builder — "Lbvh" GPU algorithm can build acceleration structures far faster than previously possible, and far faster than on the CPU.
  • Visual Studio 2010 support.


Title: Re: NVidia Optix 2.5 is out
Post by: cbuchner1 on September 28, 2012, 03:25:49 PM
Currently trying the Optix 3.0.0 beta...

What I find annoying is that their precompiled executables assume sm_20 (Geforce 2xx models) as minimum requirement. That corresponds to CUDA compute capability 2.0 and higher. I will try to recompile the samples to see if I can tweak them to run on older GPUs, like the one in my laptop.

UPDATE: no success, going back to Optix 2.6.0 (the latest stable release) worked for me.

A quick update: I was able to get two Optix samples that are advertised to work only with sm_20 to run on sm_11. The samples in question are mcmc_sampler, and the zoneplate sample. The problem is lack of an atomicAdd() for float values on this hardware. But the following function (if added to helpers.h) can replace atomicAdd() - it will just be somewhat slower. Just be sure to replace occurences of atomicAdd in some .cu files with MyAtomicAdd. In CmakeLists.txt, now edit options to require sm_11 instead of sm_2.0.

Code:
OPTIONS -arch sm_11

Code:
static __device__ __inline__ void MyAtomicAdd(float *address, float value)
{
   int oldval, newval, readback;
 
   oldval = __float_as_int(*address);
   newval = __float_as_int(__int_as_float(oldval) + value);
   while ((readback=atomicCAS((int *)address, oldval, newval)) != oldval)
     {
      oldval = readback;
      newval = __float_as_int(__int_as_float(oldval) + value);
     }
}