Device assertions

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).

Does a thread being aborted cause the entire kernel to be aborted? Or do the rest continue merrily on their way?

The thread will exit and the others will continue.

EDIT: If you want it to abort the entire kernel, replace asm(“exit;”); with asm(“trap;”);

Or you abandon the PTX inlining and simply use __trap().