localacc@scrooge:~/src/cuda_samples/kerneltest$ cat simple1.cu
#include <stdio.h>

__global__ void kernel2 ()
{
	int iIdx = threadIdx.x ;
	printf ("hi from id : %d\n", iIdx) ;
}

__global__ void kernel1 ()
{
	kernel2 <<<1,2>>> () ;
}

int some()
{
	kernel1 <<<2,2>>> () ;
	cudaDeviceSynchronize () ;

	printf( "Hello, World! \n");
	return 0;
}

int main (void)
{
        printf ("Hi there !!\n") ;

        some () ;

        return 0 ;
}
localacc@scrooge:~/src/cuda_samples/kerneltest$ nvcc -arch=sm_35 -rdc=true simple1.cu -o simple1 -lcudadevrt
localacc@scrooge:~/src/cuda_samples/kerneltest$ ./simple1
Hi there !!
hi from id : 0
hi from id : 1
hi from id : 0
hi from id : 1
hi from id : 0
hi from id : 1
hi from id : 0
hi from id : 1
Hello, World!
localacc@scrooge:~/src/cuda_samples/kerneltest$