为什么对 thrust::inclusive_scan 的调用比后续调用慢得多？

Question

当我多次调用thrust::inclusive_scan时，为什么第一次比后续调用慢很多？

这是代码

float ttime;
for(int i=0;i<5;i++){
     cudaEvent_t start,stop; 
     cudaEventCreate(&start); 
     cudaEventCreate(&stop);
     cudaEventRecord(start,0);

     thrust::device_ptr<int > din(device_input);
     thrust::device_ptr<int > dout(device_output);
     thrust::inclusive_scan(din,din+N,dout);

     cudaEventRecord(stop,0); 
     cudaEventSynchronize(stop); 
     cudaEventElapsedTime(&ttime,start,stop);
     printf("cost %fms\n",ttime);
}

我在GTX1080上运行，结果是

cost 39.180702ms
cost 0.200704ms
cost 0.201728ms
cost 0.202752ms
cost 0.197632ms

谁能帮忙解释一下？

Answer 1

Thrust 是使用 CUDA 运行时 API 构建的，API 使用惰性 context initialisation。

没有记录确切的初始化顺序，并且有经验证据表明它随着时间的推移而改变。但是，上下文设置似乎是在上完成的。

第一次调用缓慢很可能与程序中包含推力代码的模块的加载和初始化有关。您可以通过分析您的代码并查看分析执行时间与第一次调用的挂钟时间来验证这一点。

Answer 2

为 Talonmies 的有效答案添加几句话：

在我的 this question 中，有一些鸡尾酒餐巾纸计算 CUDA 初始化需要多少时间。
我还建议，要将 libthrust 的加载+初始化时间与运行时 API 初始化开销分开，请执行以下三个阶段：
1. 执行一些不写入任何输出的虚拟内核 - 两次 (T_1, T2_)
2. 使用（几乎）没有数据的 libthrust 调用（但确实会启动内核）- 两次（T_3、T_4）
3. 现在为您的真实通话计时

粗略地说，(T_1 - T_2) 是 CUDA 负载和初始时间，(T_3 - T_4) 是推力负载和初始时间。

查看分析时间线很有帮助； CUDA 将 "shove" 大部分初始化到您的一个 API 调用中 - 但不一定是第一个。

为什么对 thrust::inclusive_scan 的调用比后续调用慢得多？

Why the call to thrust::inclusive_scan is much slower than subsequent calls?

cuda

thrust