What happens when a kernel launch in unsuccessful ??

  • There are 2 kinds of errors in computer programs:

      1. Compile time errors:

          • Errors discovered during compilation of a computer program
          • E.g.: syntax errors, undefined variable name, etc.

      2. Run time errors:

          • Errors that occur when the program is executed (run)
          • E.g.: array index out of bound, using a null pointer, etc.

A kernel launch operation is a run time operation

What happens when a kernel launch in unsuccessful ??

Consider the following CUDA program with an error launching the kernel function: 

#include <stdio.h>
#include <unistd.h>

__global__ void hello( )
{
    printf("threadIdx.x=%d\n", threadIdx.x);
}

int main()
{
   hello<<< 1, 1025 >>>( );  // Error: #threads >= 1024 !!! 

   printf("I am the CPU: Hello World ! \n");

   cudaDeviceSynchronize();
}

It compiles, but you do not see the outputs from the CUDA threads !!!
DEMO: /home/cs355001/demo/CUDA/1-intro/hello-error.cu

Error handling in kernel launching operation
 

Background information:

  • If a CUDA kernel function completes normally (= no error) then:

      • The kernel function will set the system error variable to the value cudaSuccess

        (cudaSuccess is a symbolic constant)

  • The CDUA library function cudaGetLastError( ) returns:

      • The value of the system error variable set by the last completed CUDA kernel function        

 

In next slide, I will show an example on how to use the cudaGetLastError( ) function.

Example on how to use the cudaGetLastError( ) function

Consider the CUDA program with an error launching the kernel function: 

#include <stdio.h>
#include <unistd.h>

__global__ void hello( )
{
    printf("threadIdx.x=%d\n", threadIdx.x);
}

int main()
{
   hello<<< 1, 1025 >>>( );  // Error: #threads >= 1024 !!! 

   printf("I am the CPU: Hello World ! \n");

   cudaDeviceSynchronize();
}

We will add a call to cudaGetLastError() after the launch code to check error status....

Example on how to use the cudaGetLastError( ) function

You should (must ?) always check for error after you launch a grid: 

__global__ void hello( )
{
    printf("threadIdx.x=%d\n", threadIdx.x);
}

int main()
{
   hello<<< 1, 1025 >>>( ); // Error: #threads >= 1024 !!! 

   cudaError_t err = cudaGetLastError(); // Get last error code 
   if ( err != cudaSuccess )                                    
   {                                                            
      printf("CUDA Error: %s\n", cudaGetErrorString(err));           
      exit(-1);                                                 
   }                                                            

   printf("I am the CPU: Hello World ! \n");
   cudaDeviceSynchronize();
}

DEMO: /home/cs355001/demo/CUDA/1-intro/hello-error2.cu