I have been hoping for a device equivalent of std::assert for quite a while now, and finally have something that actually works that I thought I would share:
// Preprocessor macros
#ifdef device_assert
#undef device_assert
#endif
#define device_assert(x) _assert(x, #x, __FILE__, __LINE__)
__noinline__ __device__ void _assert(bool condition, const char* expression,
const char* filename, int line)
{
if(!condition)
{
printf("%s:%i - assertion '%s' failed!\n", filename, line, expression);
asm("exit;");
}
}
The trick involves inlining the PTX ‘exit’ instruction, which will immediately abort the
calling thread.
I don’t think that this is kosher PTX since it creates a device function that never returns,
but it does work on current hardware (Fermi).