How to assert in kernel code?

What is the equivalent technique of an assertion in CUDA kernel code?

There does not seem to be an assert for CUDA kernel code. I want a way to catch programmer mistakes easily in kernel code. A mechanism where I can set conditions that need to be true and the kernel should bail out when the condition is false with an error message.

You can do assert as usual, for example

#include <stdio.h>

#define  assert( X ) if ( !(X) ) \

    printf("tid %d: %s, %d\n", threadIdx.x, __FILE__, __LINE__);\

    return ;

__global__ void foo(int *data )

{

    int tid = threadIdx.x ;

assert ( tid >= 5 ) ;

}

int main(int argc, char *argv[])

{

foo<<<1,256 >>>(NULL);

cudaThreadExit();

}

result is

tid 0: main.cu, 12

tid 1: main.cu, 12

tid 2: main.cu, 12

tid 3: main.cu, 12

tid 4: main.cu, 12

However when assertion fails, not all threads terminate. If you have some synchronization issue,

then program may hang.

Thanks so much, this is perfect! External Image