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

Question

我希望能够测量CUDA中内核启动的开销。

我知道有各种参数会影响这个开销，我对以下内容感兴趣：

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

我主要是为了测量使用在CUDA 6.0中引入的管理内存的优势。我会根据评论更新代码。谢谢！

- pranith

---

您可能希望查看 N. Wilt 的 CUDA 手册第 6.1.1 节。 - Vitality

《CUDA Handbook》并发目录中的源代码有测量内核启动开销的应用程序。我还没有测试托管内存的影响，但在GTC听说驱动程序会将所有托管内存复制回主机内存进行同步。因此，当您分配更多的托管内存时，同步内核启动开销应该会增加。https://github.com/ArchaeaSoftware/cudahandbook/blob/master/concurrency/nullKernelSync.cu - ArchaeaSoftware

Answer 1

如何测量CUDA中内核启动的开销在N. Wilt的《CUDA手册》书籍的6.1.1节中有所涉及。基本思想是启动一个空内核。以下是示例代码片段：

#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.00245毫秒。

如果你想检查这个时间与启动线程数量的依赖关系，只需要改变内核启动配置<<<*,*>>>。结果表明，随着启动的线程数量增加，计时时间并没有显著变化，这与书中的说法一致，即大部分时间都花费在驱动程序中。

- Vitality

---

是的，这回答了线程数量方面的开销。我会尝试实现不同大小的内存。谢谢！ - pranith

Answer 2

也许你会对弗吉尼亚大学的这些测试结果感兴趣：

内存传输开销：http://www.cs.virginia.edu/~mwb7w/cuda_support/memory_transfer_overhead.html 内核启动开销：http://www.cs.virginia.edu/~mwb7w/cuda_support/kernel_overhead.html 它们的测量方式与JackOLantern提出的方法类似。

- Alejandro Silvestri