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

Advertisements

5 thoughts on “CUDA: Assertion in Kernel Code

  1. the return made my nvidia compiler generate some falty code. it changed my shared memory although the assert didn’t include any access/check on it!!!
    my version
    #define CudaAssert( expression ) \
    if ( !(expression)) { \
    printf( “Assert failed %d:%d\n”, blockIdx.x, threadIdx.x ); \
    }

    Like

    1. Shimi: The conventional way an assert is used (in non-parallel environments or languages) is to alert the programmer to a pre or post condition in his code that is not compliant at run-time. I think the above mechanism achieves that, since the programmer comes to know that something is wrong with a certain condition he is testing. Maybe he cannot see all the threads with this error, but even one failure is enough help IMO.

      Like

Leave a Reply

Fill in your details below or click an icon to log in:

WordPress.com Logo

You are commenting using your WordPress.com account. Log Out /  Change )

Google photo

You are commenting using your Google account. Log Out /  Change )

Twitter picture

You are commenting using your Twitter account. Log Out /  Change )

Facebook photo

You are commenting using your Facebook account. Log Out /  Change )

Connecting to %s

This site uses Akismet to reduce spam. Learn how your comment data is processed.