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.