I’ve run into an odd bug after a driver update. I have an inlined device function which does shadow checks, and the function is called inside a for loop over the lights in the scene. With the 527.27 driver, the value of a variable in the function is correct for the first light, but when the function runs for the second light, the variable’s value reverts to the previous light’s value partway through the function. I cannot share the actual code, but I’m attaching a pared down psuedocode version of the function that hopefully illustrates what I’m seeing.
I can work around the issue with some extra copying and volatile variables. Are there any known optimizer issues that might cause something like this?
Current test configuration, with variable reversion problem:
Windows 10 21H2
dual Quadro RTX 4000
527.27 driver
Visual Studio 2017
CUDA 10.0
OptiX 6.5
Configurations without the problem:
Same machine as above with 516.59 driver. Various other Windows/Linux machines with 510 series and previous drivers and Maxwell/Pascal/Turing/Ampere GPUs (some Quadro, some GeForce).
__device__ __forceinline__ void lighting(/*etc*/)
{
for ( int currentLight = 0; currentLight < lightCount; ++currentLight ) {
const float shadowScale = shadowCheck( hitPoint, lights[ currentLight ], normal );
//etc
}
}
__device__ __forceinline__ float shadowCheck(
const optix::float3& hitPoint,
const Light& light,
const optix::float3& normal, /*etc*/ )
{
optix::float3 lightDirection;
bool useFiniteRay;
if ( light.type == X ) {
lightDirection = optix::normalize(/*etc*/);
useFiniteRay = true;
} else if (light.type == Y ) {
lightDirection = /*etc*/;
useFiniteRay = false;
} else {/*etc*/}
// These prints are OK for both lights.
rtPrintf( "lightDirection %f, %f, %f\n", lightDirection.x, lightDirection.y, lightDirection.z );
rtPrintf( "useFiniteRay %i\n", useFiniteRay );
if ( /*lightDirection vs normal and various flag checks*/) {
RayShadowPayload payload;
//etc;
do {
// This print is wrong when the function is called for light 1. It has the
// lightDirection from the light 0 call.
rtPrintf( "lightDirection %f, %f, %f\n", lightDirection.x, lightDirection.y, lightDirection.z );
// This print is OK until I add a workaround for the lightDirection issue (copying and restoring
// lightDirection from a volatile struct). Then useFiniteRay starts showing the light 0 value when
// running for light 1. Making useFiniteRay volatile works around that issue.
rtPrintf( "useFiniteRay %i\n", useFiniteRay );
const float maxRayLength = useFiniteRay ? /*etc*/ : RT_DEFAULT_MAX;
const optix::Ray shadowRay = optix::make_Ray(lightDirection, maxRayLength, /*etc*/);
rtTrace(/*etc*/);
} while (/*various flags*/);
}
return payload.value;
}