CUDA: Assertion in Kernel Code

Assertions are very useful to catch programmer mistakes. Sadly, there is no mechanism to trigger an actual assert() in CUDA kernel code. Lung Sheng Chien showed me a simple way to create a assert for CUDA kernel code using macros:

// Macro definition
#define CudaAssert( X ) if ( !(X) ) { printf( "Thread %d:%d failed assert at %s:%d!", blockIdx.x, threadIdx.x, __FILE__, __LINE__ ); return; }

// Usage

#include <cstdio>

__global__ void fooKernel( const int* vals )
{
    // ...
    CudaAssert( ( vals[ threadIdx.x ] < 0 ) && "Input data not valid!" );
    // ...
}

Note:

• To be able to use printf() in kernel code, the compiler must be compiling for a GPU architecture of sm_20 or later.

• A cudaThreadSynchronize() call is needed immediately after the kernel call to flush the buffer that printf() writes to.

Tried with: CUDA 3.2 and Visual Studio 2008