linux - 为什么 Nvidia Visual Profile 在纯同步代码的时间线中显示重叠的数据传输？

Question

Nsight Visual Profile 生成的时间线看起来很奇怪。我没有编写任何传输重叠代码，但您可以看到 MemCpy 和 Compute 内核之间的重叠。

这让我无法调试真正重叠的代码。

我使用 CUDA 5.0、Tesla M2090、Centos 6.3、2x CPU Xeon E5-2609

有人遇到过类似的问题吗？它只发生在某些 Linux 发行版上吗？如何解决？

这是代码。

#include <cuda.h>
#include <curand.h>
#include <cublas_v2.h>
#include <thrust/device_vector.h>
#include <thrust/host_vector.h>
#include <thrust/device_ptr.h>

int main()
{
    cublasHandle_t hd;
    curandGenerator_t rng;
    cublasCreate(&hd);
    curandCreateGenerator(&rng, CURAND_RNG_PSEUDO_MTGP32);

    const size_t m = 5000, n = 1000;
    const double alpha = 1.0;
    const double beta = 0.0;

    thrust::host_vector<double> h(n * m, 0.1);
    thrust::device_vector<double> a(m * n, 0.1);
    thrust::device_vector<double> b(n * m, 0.1);
    thrust::device_vector<double> c(m * m, 0.1);
    cudaDeviceSynchronize();

    for (int i = 0; i < 10; i++)
    {
        curandGenerateUniformDouble(rng,
                thrust::raw_pointer_cast(&a[0]), a.size());
        cudaDeviceSynchronize();

        thrust::copy(h.begin(), h.end(), b.begin());
        cudaDeviceSynchronize();

        cublasDgemm(hd, CUBLAS_OP_N, CUBLAS_OP_N,
                m, m, n, &alpha,
                thrust::raw_pointer_cast(&a[0]), m,
                thrust::raw_pointer_cast(&b[0]), n,
                &beta,
                thrust::raw_pointer_cast(&c[0]), m);
        cudaDeviceSynchronize();
    }

    curandDestroyGenerator(rng);
    cublasDestroy(hd);

    return 0;
}

这是捕获的配置文件时间线。

Answer 1

Compute Capability 2.* (Fermi) 设备能够实现内核级并发以及内核和复制并发。为了跟踪并发内核，内核开始和结束时间戳被收集在一个单独的时钟域中，而不是内存复制时间戳。该工具负责关联这些不同的时钟。在您的屏幕截图中，我相信存在不同的比例因数(相关性差)，因为您可以看到每个内存副本不是以恒定值偏移，而是以缩放偏移量偏移。

如果您在 nvprof 中使用选项 --concurrent-kernels off 我认为问题将会消失。当并发内核被禁用时，内存复制和内核时序使用相同的源时钟作为时间戳。

Compute Capability 3.* (Kepler) 和 5.* (Maxwell) 具有不同的计时计算内核机制。对于这些设备，可以在工具中看到与内核的结束时间戳和内存副本或内核的开始重叠。工作不重叠。在工具中有一个设计决策，是在有重叠的可能性(通常 <500ns)还是将其作为依赖工作之间的恒定开销引入之间。这些工具决定以可能在序列化工作上显示非常小的重叠为代价来避免引入开销。