Compile error: PHINode should have one entry for each predecessor of its parent basic block

I’m working on an OptiX project and on cuda file compilation I get this error:

PHINode should have one entry for each predecessor of its parent basic block!
1>    %__cuda_local_var_533998_11_non_const_bsdfDirPdfW.5 = phi float [ %__cuda_local_var_533998_11_non_const_bsdfDirPdfW.2617, %569 ], [ %581, %579 ], [ %__cuda_local_var_533998_11_non_const_bsdfDirPdfW.2617, %569 ], !dbg !535

Specs:
Windows 8.1 x64, Cuda 6, Optix 3.6.0. Driver 334.88, Visual Studio 2012
The file was compiled as ptx, compiler flags -gencode=arch=compute_20,code=“sm_20,compute_20”
Same code compiles using Cuda 5.5, Optix 3.5.1 and works.

I have identified that it happens due to an inlined function the function updateMisTermsOnScatter() in the sample below. It was there first, the issue appeared when I changed how bsdfFactor on line 3 is computed, e.g. using a function call (not inlined). If I remove inline directive from updateMisTermsOnScatter() it compiles using Cuda 6 as well.

RT_PROGRAM void closestHitLight()
{
    // some omitted code that doesn't affect
    float bsdfDirPdfW;
    float cosThetaOut;
    float3 bsdfFactor = lightVertex.bsdf.vcmSampleF(
                                &subpathPrd.direction, bsdfSample, &bsdfDirPdfW);
    // some omitted code that doesn't affect
    updateMisTermsOnScatter(subpathPrd, cosThetaOut, bsdfDirPdfW, bsdfRevPdfW, 
                            misVcWeightFactor, misVmWeightFactor, &vertexPickPdf);
}

// signature
__inline__ __device__ void updateMisTermsOnScatter(
    SubpathPRD & aPathPrd, const float & aCosThetaOut, const float & aBsdfDirPdfW,
    const float & aBsdfRevPdfW, const float & aMisVcWeightFactor, const float & aMisVmWeightFactor,
    const float const * aVertexPickPdf = NULL)
 { // omitted code }

Could anyone explain meaning of the error? Possible causes? Best practices to avoid something like that?
From searching online it seems it LLVM stuff, but I don’t pretty much anything about that.

I tried to reproduce this with a small sample in another project, but couldn’t. I don’t really have time currently to disassemble code piece by piece to pinpoint what exactly is the key problem in this case, so I guess I’ll use Cuda 5.5 for now.

I have not seen an error message like this before, but it seems to indicate some sort of internal compiler error, basically a constraint violation or an assertion failure.

I would suggest filing a bug, using the bug reporting form linked from the registered developer website. Please attach a self-contained repro code (small code is better, but large code can be dealt with as logn as you rpovide precise build instructions). The compiler team may have a suggestion for a work-around once the issue has been root-caused. If you are restricted from sharing your actual source code, there may be other ways to repor, like providing one of the files that contain the intermediate representation. The compiler team can advise what is possible in this regard.

Thank you for your help.

Moving to the OptiX forum.

Please try if using forceinline instead of inline helps. The latter is just a hint to the compiler but OptiX requires all functions to be inlined. Only callable programs are real calls.
I’ve seen CUDA compilers stopping to inline code with increasing number of function parameters.

I’m always using a “#define RT_FUNCTION forceinline device” for any function in OptiX programs to match the look of the code to the RT_PROGRAM and RT_CALLABLE_PROGRAM ones.

@njuffa as I mentioned in the post I don’t have time currently to make small self-contained repro project, I tried to make one, but I couldn’t reproduce it there. I could do that in 2 weeks. If you wish to look at this bug sooner, I could push code to github (I’m extending an open source renderer).

@Detlef forceinline didn’t help.

It’s weird. Today it doesn’t compile also when inline is removed from updateMisTermsOnScatter() which worked yesterday. Not sure how it did yesterday.

What’s “not real” about function calls in Optix?
If Cuda compiler doesn’t inline some function, the program JIT compilation will fail at runtime?

I have one ~100 line function that isn’t inlined and it does work. I’d think that Cuda compiler won’t try to inline that.

If you look into your compiled PTX code you give to OptiX and that already contains a “call” instruction, I would expect the context->compile() step to fail. Only if you use callable programs OptiX will generate call instructions in the kernel.

OptiX takes the input PTX apart and puts it back together. The resulting PTX sent to the CUDA driver can be rather different. This paper explains that in chapter 6: http://graphics.cs.williams.edu/papers/OptiXSIGGRAPH10/Parker10OptiX.pdf

Actually I got carried away by the OptiX connection. If the compiler error already occurs at the nvcc compilation step and CUDA 5.5 works and CUDA 6.0 doesn’t, getting the reproducing *.cu file would be nice to investigate that.