$ cat t815.cu
#include <stdio.h>

#define DELAY_VAL 5000000000ULL

__global__ void child(){

  unsigned long long start = clock64();
  while (clock64()< start+DELAY_VAL);
}

__global__ void parent(){

  child<<<1,1>>>();
}

int main(int argc, char* argv[]){

  cudaStream_t st1, st2;
  cudaStreamCreate(&st1);
  cudaStreamCreate(&st2);
  parent<<<1,1,0,st1>>>();
  if (argc > 1){
    printf("running double kernel\n");
    parent<<<1,1,0,st2>>>();
    }
  cudaDeviceSynchronize();
}

$ nvcc -arch=sm_35 -rdc=true -lcudadevrt t815.cu -o t815
$ time ./t815
3.65user 1.88system 0:05.65elapsed 97%CPU (0avgtext+0avgdata 82192maxresident)k
0inputs+0outputs (0major+2812minor)pagefaults 0swaps
$ time ./t815 double
running double kernel
3.68user 1.83system 0:05.64elapsed 97%CPU (0avgtext+0avgdata 82200maxresident)k
0inputs+0outputs (0major+2814minor)pagefaults 0swaps
$ time cuda-memcheck ./t815
========= CUDA-MEMCHECK
========= ERROR SUMMARY: 0 errors
3.16user 2.25system 0:05.68elapsed 95%CPU (0avgtext+0avgdata 87040maxresident)k
0inputs+0outputs (0major+4573minor)pagefaults 0swaps
$ time cuda-memcheck ./t815 double
========= CUDA-MEMCHECK
running double kernel
========= ERROR SUMMARY: 0 errors
7.27user 3.04system 0:10.46elapsed 98%CPU (0avgtext+0avgdata 87116maxresident)k
0inputs+0outputs (0major+4594minor)pagefaults 0swaps
$