Optix, recursion and templates

Gaming and Visualization Technologies Visualization OptiX
1

Hello everyone,

I have encountered some issues which I’ve reduced to the following:

Take the SDK’s optixMeshViewer example code, more precisely the cuda file triangle_mesh.cu and replace the mesh_intersect program with:

class Data
{
};

template<bool v>
void __device__ recurse(int const level, int const primIdx, Data const& x)
{
    if(level >0)
    {
        Data y;

        recurse<true>(level - 1, primIdx, y);
    }
    else
    {
        meshIntersect<false>( primIdx );
    }
}

RT_PROGRAM void mesh_intersect( int primIdx )
{
    int const level = 1;
    Data const x;
    recurse<true>(level, primIdx, x );
}

Then at runtime:

OptiX Error: 'Unknown error (Details: Function "RTresult _rtProgramCreateFromPTXString(RTcontext, const char*, const char*, RTprogram_api**)" caught exception: Compile Error: Only intersection programs can call rtPotentialIntersection.
================================================================================
Backtrace:
	(0) () +0x711547
	(1) () +0x70dd31
	(2) () +0x70dddb
	(3) () +0x6fd7bf
	(4) () +0x369e5c
	(5) () +0x36a995
	(6) () +0x1feacb
	(7) () +0x200350
	(8) () +0x4342aa
	(9) () +0x17ded1
	(10) rtProgramCreateFromPTXString() +0x3c3
	(11) () +0x65a61
	(12) loadMesh(std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, OptiXMesh&, optix::Matrix<4u, 4u> const&) +0x901
	(13) loadMesh(std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&) +0x17f
	(14) main() +0x372
	(15) __libc_start_main() +0xe7
	(16) _start() +0x2a

================================================================================
)'

However, the modified optixMeshViewer runs fine if I either set level = 0 in mesh_intersect, remove the template in recurse or remove the passing of Data. Is there something I’m doing wrong?

Thanks in advance for your help,

Jean-Marie

2

The compile error seems accurate:

Compile Error: Only intersection programs can call rtPotentialIntersection.

The intersection program here is mesh_intersect. Special OptiX functions like rtPotentialIntersection or rtReportIntersection must be called directly in mesh_intersect, not in a helper function like recurse() or meshIntersect().

However, if nvcc inlines the helper functions when it produces PTX, then as far as OptiX is concerned this is legal; OptiX would just see the one function for mesh_intersect. Thus the recommendation in the programming guide to use the “inline” or even “forceinline” decorators on all helper functions. So you could try adding one of those to “recurse”, but with the caveat that it could generate a lot of code.

3

Thank you, that explains a lot.