inline host function - unexpected behavior strange things happen every day :)

Hello :smile:

First of all my apology for the typo in the toppic name - it is all about device function not host as written by mistake. Sooory :angel:

I just find out that my device function defined without “inline” convention will do some unexpected strange things. :nuke: I thought that every device function is inlined by default and explicitly inlining will have no effect so i removed the keyword …

Take look at this little code

inline __device__  float4 GetTex3D16(float3 pos)

{

   return tex3D(tex16, pos.x, pos.y, pos.z);

}

Yes, very simple but still problematic… it produces (without inline) many artefacts in my program. Basicaly returned texture’s values are not correct …

I dont know what to do, because i found this bug when trying to update to template version (see below) but even with inline this function will still make problems and i dont accept that :biggrin:

template <int bpc>

inline __device__  float4 GetTex3D(float3 pos)

{

   if (bpc > 8) return tex3D(tex16, pos.x, pos.y, pos.z);

   else return tex3D(tex8, pos.x, pos.y, pos.z);

}

And one last thing - when i looked into cubin file i found out that without inline, one of my kernel used additional 28 bytes of local memory! LOL :angel: And decuded asm is also quite different … hmmm

Anybody have any idea what is happeing ?

Thnkx in advance :biggrin:

PS.: Emulation mode is doing correctly :D