Почему я не получаю перекрытие ввода-вывода с этим кодом?

Question

Следующая программа:

#include <iostream>
#include <array>

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 = 5e7;
    constexpr const auto num_streams = 2;

    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]);
    }
    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]);
    }
}

должно приводить к перекрытию операций ввода-вывода и вычислений между работой в первом и втором потоках: когда хост-устройство первого потока заканчивается, ядро ​​первого потока может запуститься, но может и передача потока-хост-устройства второго потока. Вместо этого я получаю следующий график без наложения:

Я думаю, что я покрыл свои основы, чтобы гарантировать перекрытие. Потоки являются неблокирующими (и на самом деле постановка работы на работу завершается задолго до того, как это сделает первый HtoD); память хоста закреплена... так что мне не хватает, чтобы увидеть перекрытие?

Использование CUDA 8.0.61 в GNU/Linux Mint 18.2 с NVIDIA GTX 650 Ti Boost. Но драйвер v384.59.

Answer 1

Хорошо, это должно быть что-то с моей моделью GPU, потому что с Fedora 25 и GTX Titan X я получаю: