2

考虑以下程序：

#include <iostream>
#include <array>
#include <unistd.h>

using clock_value_t = long long;

__device__ void gpu_sleep(clock_value_t sleep_cycles)
{
    clock_value_t start = clock64();
    clock_value_t cycles_elapsed;
    do { cycles_elapsed = clock64() - start; }
    while (cycles_elapsed < sleep_cycles);
}

__global__ void dummy(clock_value_t duration_in_cycles)
{
    gpu_sleep(duration_in_cycles);
}

int main()
{
    const clock_value_t duration_in_clocks = 1e7;
    const size_t buffer_size = 2e7;
    constexpr const auto num_streams = 8;

    std::array<char*, num_streams> host_ptrs;
    std::array<char*, num_streams> device_ptrs;
    std::array<cudaStream_t, num_streams> streams;
    for (auto i=0; i<num_streams; i++) {
        cudaMallocHost(&host_ptrs[i], buffer_size);
        cudaMalloc(&device_ptrs[i], buffer_size);
        cudaStreamCreateWithFlags(&streams[i], cudaStreamNonBlocking);
    }
    cudaDeviceSynchronize();
    for (auto i=0; i<num_streams; i++) {
        cudaMemcpyAsync(device_ptrs[i], host_ptrs[i], buffer_size, cudaMemcpyDefault, streams[i]);
        dummy<<<128, 128, 0, streams[i]>>>(duration_in_clocks);
        cudaMemcpyAsync(host_ptrs[i], device_ptrs[i], buffer_size, cudaMemcpyDefault, streams[i]);
    }
    usleep(50000);
    for (auto i=0; i<num_streams; i++) { cudaStreamSynchronize(streams[i]); }
    for (auto i=0; i<num_streams; i++) {
        cudaFreeHost(host_ptrs[i]);
        cudaFree(device_ptrs[i]);
    }
}

我在 GTX Titan X 上运行它，CUDA 8.0.61，Fedora 25，驱动程序 375.66。我看到的时间线是这样的：

在此处输入图像描述

这张图有几点不对：

  • 据我所知，一次只能进行一次 HtoD 传输。
  • 所有的内存传输都应该花费基本相同的时间——它们具有相同的数据量；PCIe 总线对传输速率的影响如此之大，没有什么其他有趣的事情发生了。
  • 一些 DtoH 条就像它们被拉长一样，直到另一个流上发生某些事情。
  • 有一个巨大的差距，似乎没有计算机，也没有真正的 I/O。即使所有先前完成的内核的 D​​toH 都占据了这个空白，那仍然会留下非常大量的时间。这实际上看起来像是一个调度问题，而不是一个分析错误。

那么，我应该如何解读这个时间线呢？问题出在哪里？（希望不是程序员...）

我应该提到，使用较少的流（例如 2），时间线在相同的 SW+HW 上看起来非常好：

在此处输入图像描述

4

0 回答 0