memory - 为什么 CUDA 内存复制速度会这样，有一些恒定的驱动程序开销？

Question

在我的旧 GeForce 8800GT 上使用 CUDA 内存时，我总是会遇到奇怪的 0.04 毫秒开销。我需要将 ~1-2K 传输到设备的常量内存中，处理其中的数据并从设备中仅获取一个浮点值。

我有一个使用 GPU 计算的典型代码:

//allocate all the needed memory: pinned, device global
for(int i = 0; i < 1000; i++)
{
    //Do some heavy cpu logic (~0.005 ms long)        
    cudaMemcpyToSymbolAsync(const_dev_mem, pinned_host_mem, mem_size, 0, cudaMemcpyHostToDevice);
    my_kernel<<<128, 128>>>(output);
    //several other calls of different kernels
    cudaMemcpy((void*)&host_output, output, sizeof(FLOAT_T), cudaMemcpyDeviceToHost);
    // Do some logic with returned value 
}

我决定使用以下代码测量 GPU 内存的工作速度(注释了所有内核调用，添加了 cudaDeviceSynchronize 调用):

//allocate all the needed memory: pinned, device global
for(int i = 0; i < 1000; i++)
{
    //Do some heavy cpu logic (~0.001 ms long)        
    cudaMemcpyToSymbolAsync(const_dev_mem, pinned_host_mem, mem_size, 0, cudaMemcpyHostToDevice);
    cudaMemcpyAsync((void*)&host_output, output, sizeof(FLOAT_T), cudaMemcpyDeviceToHost);
    cudaDeviceSynchronize();
    // Do some logic with returned value 
}

我测量了周期的执行时间，结果约为 0.05 秒(因此，每次迭代 0.05 毫秒)。奇怪的是，当我尝试做更多内存工作(添加额外的 cudaMemcpyToSymbolAsync 和 cudaMemcpyAsync 调用)时，每次调用都会获得额外的 <0.01 毫秒时间。与这个人的研究相对应:http://www.cs.virginia.edu/~mwb7w/cuda_support/memory_transfer_overhead.html

他还获得了每次将 1K block 传输到 GPU 所需的时间为 0.01 毫秒。那么这 0.04 毫秒 (0.05 - 0.01) 的开销从何而来呢？有任何想法吗？也许我应该在较新的卡上尝试此代码？

在我看来，在 cudaDeviceSynchronize 和 CPU 代码之后，我的 GeForce 会进入某种省电模式或类似的模式。

Answer 1

我建议您增加正在实现的线程数量

    //Use malloc() to allocate memory on CPU. 
    //Change mem_size to the total memory to be tranferred to GPU.        
    cudaMemcpyToSymbolAsync(const_dev_mem, pinned_host_mem, mem_size, 0, cudaMemcpyHostToDevice);
    dim3 dimBlock(128,2);
    dim3 dimGrid(64000,1);
    my_kernel<<<dimGrid, dimBlock>>>(output);
    //several other calls of different kernels
    //change size field to 1000*sizeof(FLOAT_T)
    cudaMemcpy((void*)&host_output, output, sizeof(FLOAT_T), cudaMemcpyDeviceToHost);
    // Do some logic with returned value 

如果代码崩溃(由于更多线程或更多 GPU 内存)，请使用循环。但是，减少它们。