cuda - 如何测量 CUDA 中内核启动的开销

Question

我想测量 CUDA 中内核启动的开销。

据我了解，有多种参数会影响此开销。我对以下内容感兴趣:

• 创建的线程数
• 正在复制的数据大小

我这样做主要是为了衡量使用 CUDA 6.0 中引入的托管内存的优势。我将用我开发的代码和评论来更新这个问题。谢谢!

Answer 1

N. Wilt 编写的《CUDA 手册》一书的第 6.1.1 节介绍了如何测量 CUDA 中内核启动的开销。基本思想是启动一个空内核。这是示例代码片段

#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;
}

在我的笔记本电脑 GeForce GT540M 卡上，内核启动开销为 0.00245ms .

如果想从启动的线程数来查看本次的依赖关系，那么只需更改内核启动配置 <<<*,*>>> 。看来时间并没有随着启动的线程数量而显着变化，这与书中的说法一致，即大部分时间都花在驱动程序上。