首页 文章

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

提问于
浏览 次
4

当我尝试重叠数据传输和内核执行时看起来卡正在按顺序执行所有内存传输，无论我使用什么流 .

那么，如果我发出以下内容：

  • stream 1： MemcpyA_HtoD_1; Kernel_1; MemcpyA_DtoH_1

  • stream 2： MemcpyA_HtoD_2; Kernel_2; MemcpyA_DtoH_2

The MemcpyA_HtoD_2 will wait till the MemcpyA_DtoH_1 is completed . 因此没有实现重叠 . 无论我使用什么配置的流，Memcpy操作总是按顺序发布 . 因此，实现重叠的唯一方法包括缓冲输出或延迟输出传输直到下一次迭代 .

我使用CUDA 5.5，Windows 7 x64和GTX Titan . 固定所有cpu内存，使用异步版本完成data_transfers .

请参阅以下具有以下行为的屏幕：

发出，host_to_device - > kernel - > device_to_host（正常行为）并且无法获得重叠 .

non overlapping

发出host_to_device - >内核（避免内核后的device_to_host）重叠...因为无论我尝试什么流配置，所有内存副本都按顺序执行 .

overlapping

UPDATE

如果有人有兴趣重现这个问题，我编写了一个合成程序，显示出这种不良行为 . 它是使用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();
}
cuda

1 回答

  • 1

    TL; DR：问题是由Nsight Monitor中的WDDM TDR延迟选项引起的！设置为false时，会出现问题 . 相反，如果将TDR延迟值设置为非常高的数字，并将“enabled”选项设置为true，则问题就会消失 .

    请阅读下面的其他（较旧）步骤，直到我找到上述解决方案，以及其他一些可能的原因 .

    我刚才能够部分解决这个问题！我认为它特定于windows和aero . 请尝试这些步骤并发布结果以帮助其他人！我在GTX 650和GT 640上试过了 .

    在你做任何事情之前，考虑使用板载gpu（作为显示器）和离散的gpu（用于计算），因为nvidia驱动程序已经验证了windows的问题！当你使用板载gpu时，说驱动程序没有满载，所以很多bug都被规避了 . 此外，工作时保持系统响应能力！

    • 确保您的并发问题与旧驱动程序（包括BIOS）等其他问题无关 .

    • 转到计算机>属性

    • 选择左侧的高级系统设置

    • 转到“高级”选项卡

    • 在性能上单击设置

    • 在“视觉效果”选项卡中，选择"adjust for best performance"项目符号 .

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

    或者，您可以：

    • 右键单击桌面，选择个性化

    • 从基本主题中选择一个没有空气动力学的主题 .

    这也将如上所述，但启用了更多的可视化选项 . 对于我的两个设备，这个设置也有效，所以我保留了它 .

    当您尝试这些解决方案时，请回到此处并发布您的发现！

    对我来说，它解决了大多数情况下的问题（我已经制作了一个平铺的dgemm），但是请注意，我仍然无法正常运行“simpleStreams”并实现并发...

    更新：新窗口安装完全解决了问题！以前的步骤改进了某些情况下的行为，但只有全新安装解决了所有问题！

    我会尝试找到一种解决这个问题的不那么激进的方法，也许只需恢复注册表即可 .

    回复于

相关问题