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. External Image 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