How to measure overhead of a kernel launch in CUDA

Question

I want to measure the overhead of a kernel launch in CUDA.

I understand that there are various parameters which affect this overhead. I am interested in the following:

• number of threads created
• size of data being copied

I am doing this mainly to measure the advantage of using managed memory which has been introduced in CUDA 6.0. I will update this question with the code I develop and from the comments. Thanks!

Answer 1

How to measure the overhead of a kernel launch in CUDA is dealt with in Section 6.1.1 of the "CUDA Handbook" book by N. Wilt. The basic idea is to launch an empty kernel. Here is a sample code snippet

#include <stdio.h>

__global__ void EmptyKernel() { }

int main() {

    const int N = 100000;

    float time, cumulative_time = 0.f;
    cudaEvent_t start, stop;
    cudaEventCreate(&start);
    cudaEventCreate(&stop);

    for (int i=0; i<N; i++) { 

        cudaEventRecord(start, 0);
        EmptyKernel<<<1,1>>>(); 
        cudaEventRecord(stop, 0);
        cudaEventSynchronize(stop);
        cudaEventElapsedTime(&time, start, stop);
        cumulative_time = cumulative_time + time;

    }

    printf("Kernel launch overhead time:  %3.5f ms \n", cumulative_time / N);
    return 0;
}

On my laptop GeForce GT540M card, the kernel launch overhead is 0.00245ms.

If you want to check the dependence of this time from the number of threads launched, then just change the kernel launch configuration <<<*,*>>>. It appears that the timing does not significantly change with the number of threads launched, which is consistent with the statement of the book that most of that time is spent in the driver.