About getting OptiX closer to double precision

[catcondo]

Using the following approach, it should be possible to report entirely double precision results from custom written __intersection__() OptiX programs. Basically, pack those doubles into ints and pass them out through the optixReportIntersection() method, as in this simple example. Sure, there will be a performance hit compared to single precision, but then those of us involved in applications requiring double precision have already crossed that bridge willingly.

union DblAsInt2 {
 double a;
 int2 b;
};

extern "C" __global__ void __intersection__blort()
{
  double u, v, hitT;

  //...code to calculate some intersection in double precision

  DblAsInt2 uConverted {uv.x};
  DblAsInt2 vConverted {uv.y};
    
  optixReportIntersection(hitT, 1,
                            uConverted._int2.x, uConverted._int2.y,
                            vConverted._int2.x, vConverted._int2.y);

}

This can get things really close to true double precision ray tracing, the only slight issue being that the optixGetWorldRayOrigin() and optixGetWorldRayDirection() available to an __intersection__ routine are only single precision.

So my question is, would the developers consider supporting some new optixTrace() method/interface that has the ability to provide double3 ray origin and direction values? Things like triangular facet ray tracing, BVH traversal, curves, transformations, etc can remain in single precision. I would just like a way for an __intersection__ program to acquire double precision ray info so that it can generate double precision results.

Thanks.

[droettger]

optixReportIntersection arguments are float hitT, unsigned int hitKind, and zero to eight 32bit unsigned int attribute registers.

I would change your code snippet to this:

union DoubleAsUint2 {
 double d;
 uint2  ui2; // attribute registers are unsigned int
};

extern "C" __global__ void __intersection__blort()
{
  double hitT, u, v; // You could also define these as DoubleAsUint2 directly and only use the `.d` view of that inside your intersection routine.

  // ... code to calculate some intersection in double precision

  // Split 64bit double into 2x32bit unsigned int.
  DoubleAsUint2 T;
  DoubleAsUint2 U;
  DoubleAsUint2 V;

   T.d = hitT;
   U.d = u;
   V.d = v;
    
  optixReportIntersection(float(hitT), 1,
                          T.ui2.x, T.ui2.y,
                          U.ui2.x, U.ui2.y,
                          V.ui2.x, V.ui2.y); // Needs OptixPipelineCompileOptions::numAttributeValues = 6;
}

This way you also return the double precision hit intersection you calculated.

Also note that the intersection program works in object coordinates.
Using optixGetWorldRayOrigin() and optixGetWorldRayDirection() there is potentially slower than optixGetObjectRayOrigin() and optixGetObjectRayDirection().
See https://raytracing-docs.nvidia.com/optix8/guide/index.html#device_side_functions#ray-information

So my question is, would the developers consider supporting some new optixTrace() method/interface that has the ability to provide double3 ray origin and direction values?

That shouldn’t be needed since you could track your double precision ray data yourself already via the per ray payload.

The intersection program has access to the per ray payload with the optixGetPayload functions.
Means you could route your double ray origin and direction through to the intersection program as well.

Everything else would remain in float, so if you also stored your geometric primitives in double precision, you would need to make sure that your custom primitives’ floating point AABB input to optixAccelBuild has been rounded outwards to the next bigger float extents to surely cover your double coordinates.
Then the BVH traversal done against floating point rays would hit the floating point AABBs and you could do all calculation in double precision inside your own device programs.

Some care would need to be taken if you need double precision hit results because you need to span huge distances. Then this outward AABB rounding would need to be tuned for the maximum deviation of the double vs. floating point ray directions.

[catcondo]

Thank you for the quick, kind, and thorough explanation. I had not thought of providing a double precision origin/direction through payload values in the optixTrace() call. Previously I had assumed incorrectly that payloads were only set by hit/miss programs and eventually passed back to optixTrace(). I took a closer look at the OptiX documentation, and sure enough, it clearly states:

Payload values are passed to and returned from optixTrace, and follow a copy-in/copy-out semantic.

[droettger]

You can also use two per ray payload (PRD) registers to encode a 64-bit pointer to a local structure you declared inside the ray generation program and access more than the 32 payload registers available today. (It was limited to 8 in earlier OptiX 7 versions.)

While that has the benefit of allowing bigger PRD structs, it comes with a performance penalty over using registers directly due to the local memory accesses and potentially higher stack size requirements.
So if you can fit your PRD into the available registers, benchmark that.

I’m using that split and merge pointer method throughout my examples like this:
https://github.com/NVIDIA/OptiX_Apps/blob/master/apps/MDL_renderer/shaders/per_ray_data.h#L100
https://github.com/NVIDIA/OptiX_Apps/blob/master/apps/MDL_renderer/shaders/raygeneration.cu#L123
https://github.com/NVIDIA/OptiX_Apps/blob/master/apps/MDL_renderer/shaders/hit.cu#L78
Mind that this pointer never changes so I don’t ever need to touch the PRD values after the split into two unsigned ints.