cuda - 在 GTX Titan 卡中重叠内核执行和数据传输的最佳策略是什么？

Question

当我尝试重叠数据传输和内核执行时，无论我使用什么流，卡似乎都按顺序执行所有内存传输。

所以，如果我发出以下命令:

• 流 1:MemcpyA_HtoD_1；内核_1； MemcpyA_DtoH_1
• 流 2:MemcpyA_HtoD_2；内核_2； MemcpyA_DtoH_2

MemcpyA_HtoD_2 将等待 MemcpyA_DtoH_1 完成。所以没有实现重叠。无论我使用什么流配置，Memcpy 操作总是按顺序发出。因此，实现重叠的唯一方法是缓冲输出或将输出传输延迟到下一次迭代。

我使用 CUDA 5.5、Windows 7 x64 和 GTX Titan。所有 CPU 内存均已固定，并且 data_transfers 均使用异步版本完成。

查看以下屏幕的行为:

发出，host_to_device -> kernel -> device_to_host(正常行为)并且无法重叠。

发出 host_to_device -> 内核(避免在内核之后使用 device_to_host)会重叠...因为无论我尝试什么流配置，所有内存副本都会按顺序执行。

更新

如果有人有兴趣重现这个问题，我已经编写了一个综合程序来显示这种不良行为。它是使用 CUDA 5.5 的完整 VS2010 解决方案

VS2010 Streams Not Working link

有人可以在 Linux 上执行这个来测试重叠吗？

#include "cuda_runtime.h"
#include "device_launch_parameters.h"


#include <stdio.h>
#define N 1024*1024

__global__ void someKernel(int *d_in, int *d_out) {
    for (int i = threadIdx.x; i < threadIdx.x + 1024; i++) {
        d_out[i] = d_in[i];
    }
}

int main () {
    int *h_bufferIn[100];
    int *h_bufferOut[100];
    int *d_bufferIn[100];
    int *d_bufferOut[100];

    //allocate some memory
    for (int i = 0; i < 100; i++) {
        cudaMallocHost(&h_bufferIn[i],N*sizeof(int));
        cudaMallocHost(&h_bufferOut[i],N*sizeof(int));
        cudaMalloc(&d_bufferIn[i], N*sizeof(int));
        cudaMalloc(&d_bufferOut[i], N*sizeof(int));
    }

    //create cuda streams
    cudaStream_t st[2];
    cudaStreamCreate(&st[0]);
    cudaStreamCreate(&st[1]);

    //trying to overlap computation and memcpys
    for (int i = 0; i < 100; i+=2) {
        cudaMemcpyAsync(d_bufferIn[i], h_bufferIn[i], N*sizeof(int), cudaMemcpyHostToDevice, st[i%2]);
        someKernel<<<1,256, 0, st[i%2]>>>(d_bufferIn[i], d_bufferOut[i]);
        cudaMemcpyAsync(h_bufferOut[i], d_bufferOut[i], N*sizeof(int), cudaMemcpyDeviceToHost, st[i%2]);
        cudaStreamQuery(0);

        cudaMemcpyAsync(d_bufferIn[i+1], h_bufferIn[i+1], N*sizeof(int), cudaMemcpyHostToDevice, st[(i+1)%2]);
        someKernel<<<1,256, 0, st[(i+1)%2]>>>(d_bufferIn[i+1], d_bufferOut[i+1]);
        cudaMemcpyAsync(h_bufferOut[i+1], d_bufferOut[i+1], N*sizeof(int), cudaMemcpyDeviceToHost, st[(i+1)%2]);
        cudaStreamQuery(0);
    }
    cudaDeviceSynchronize();
}

Answer 1

TL;DR: The issue is caused by the WDDM TDR delay option in Nsight Monitor! When set to false, the issue appears. Instead, if you set the TDR delay value to a very high number, and the "enabled" option to true, the issue goes away.

阅读下面的内容，了解我找到上述解决方案之前所遵循的其他(较旧的)步骤，以及其他一些可能的原因。

我最近才部分解决了这个问题!我认为这是特定于 windows 和 aero 的。请尝试这些步骤并发布您的结果以帮助其他人!我已经在 GTX 650 和 GT 640 上尝试过。

Before you do anything, consider using both onboard gpu(as display) and the discrete gpu (for computations), because there are verified issues with the nvidia driver for windows! When you use onboard gpu, said drivers don't get fully loaded, so many bugs are evaded. Also, system responsiveness is maintained while working!

1. 确保您的并发问题与旧驱动程序(包括 BIOS)等其他问题无关。
2. 转到计算机>属性
3. 选择左侧的高级系统设置
4. 转到“高级”选项卡
5. 效果点击设置
6. 在“视觉效果”选项卡中，选择“调整以获得最佳性能”项目符号。

这将禁用空气动力学和几乎所有视觉效果。如果此配置有效，您可以尝试逐一启用视觉效果框，直到找到导致问题的精确框!

Alternatively, you can:

1. 右键单击桌面，选择个性化
2. 从基本主题中选择一个没有 Aero 的主题。

这也将像上面一样工作，但启用了更多视觉选项。对于我的两台设备，这个设置也有效，所以我保留了它。

Please, when you try these solutions, come back here and post your findings!

对我来说，它解决了大多数情况下的问题(我制作的平铺dgemm)，但请注意我仍然无法正确运行“simpleStreams”并且实现并发...

UPDATE: The problem is fully solved with a new windows installation!! The previous steps improved the behavior for some cases, but ONLY a fresh install solved ALL the problems!

我会尝试找到一种不太激进的方法来解决这个问题，也许只恢复注册表就足够了。