Using assert within kernel invocation

Question

Is there convenient way for using asserts within the kernels invocation on device mode?

Answer 1

CUDA now has a native assert function. Use assert(...). If its argument is zero, it will stop kernel execution and return an error. (or trigger a breakpoint if in CUDA debugging.)

Make sure to include "assert.h". Also, this requires compute capability 2.x or higher, and is not supported on MacOS. For more details see CUDA C Programming Guide, Section B.16.

The programming guide also includes this example:

#include <assert.h>
__global__ void testAssert(void)
{
   int is_one = 1;
   int should_be_one = 0;
   // This will have no effect
   assert(is_one);
   // This will halt kernel execution
   assert(should_be_one);
}
int main(int argc, char* argv[])
{
   testAssert<<<1,1>>>();
   cudaDeviceSynchronize();
   return 0;
}

Answer 2

#define MYASSERT(condition) \
  if (!(condition)) { return; }

MYASSERT(condition);

if you need something fancier you can use cuPrintf() which is available from the CUDA site for registered developers.