📅 2011-Mar-02 ⬩ ✍️ Ashwin Nanjappa ⬩ 🏷️ assertions, cuda ⬩ 📚 Archive
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>
void fooKernel( const int* vals )
__global__
{// ...
0 ) && "Input data not valid!" );
CudaAssert( ( vals[ threadIdx.x ] < // ...
}
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