为什么两个 CUDA 流中的操作不重叠？

Question

我的程序是一个流水线，其中包含多个内核和memcpys。每个任务将通过具有不同输入数据的相同管道。宿主代码在处理任务时会首先选择一个Channel，它是对暂存器内存和CUDA对象的封装。在最后一个阶段之后，我会记录一个事件然后去处理下一个任务。
主要流水线逻辑如下。问题是不同流中的操作不重叠。我附上处理10个任务的时间表。您可以看到流中的 none 操作是重叠的。对于每个内核，一个块中有 256 个线程，一个网格中有 5 个块。用于 memcpy 的所有缓冲区都已固定，我确信我已经满足那些 requirements 用于重叠内核执行和数据传输的要求。有人可以帮我找出原因吗？谢谢

环境信息
GPU：Tesla K40m (GK110)
最大值 Warps/SM: 64
最大线程数 Blocks/SM：16
最大值 Threads/SM: 2048
CUDA 版本：8.0

    void execute_task_pipeline(int stage, MyTask *task, Channel *channel) {
    assert(channel->taken);
    assert(!task->finish());

    GPUParam *para = &channel->para;

    assert(para->col_num > 0);
    assert(para->row_num > 0);

    // copy vid_list to device
    CUDA_ASSERT( cudaMemcpyAsync(para->vid_list_d, task->vid_list.data(),
                sizeof(uint) * para->row_num, cudaMemcpyHostToDevice, channel->stream) );

    k_get_slot_id_list<<<WK_GET_BLOCKS(para->row_num),
        WK_CUDA_NUM_THREADS, 0, channel->stream>>>(
                vertices_d,
                para->vid_list_d,
                para->slot_id_list_d,
                config.num_buckets,
                para->row_num);

    k_get_edge_list<<<WK_GET_BLOCKS(para->row_num),
        WK_CUDA_NUM_THREADS, 0, channel->stream>>>(
                vertices_d,
                para->slot_id_list_d,
                para->edge_size_list_d,
                para->offset_list_d,
                para->row_num);

    k_calc_prefix_sum(para, channel->stream);

    k_update_result_table_k2u<<<WK_GET_BLOCKS(para->row_num),
        WK_CUDA_NUM_THREADS, 0, channel->stream>>>(
            edges_d,
            para->vid_list_d,
            para->updated_result_table_d,
            para->prefix_sum_list_d,
            para->offset_list_d,
            para->col_num,
            para->row_num);

    para->col_num += 1;
    // copy result back to host
    CUDA_ASSERT( cudaMemcpyAsync(&(channel->num_new_rows), para->prefix_sum_list_d + para->row_num - 1,
            sizeof(uint), cudaMemcpyDeviceToHost, channel->stream) );
    // copy result to host memory
    CUDA_ASSERT( cudaMemcpyAsync(channel->h_buf, para->updated_result_table_d,
                channel->num_new_rows * (para->col_num + 1), cudaMemcpyDeviceToHost, channel->stream) );

    // insert a finish event in the end of pipeline
    CUDA_ASSERT( cudaEventRecord(channel->fin_event, channel->stream) );
}

Answer 1

您是否试图重叠 82 微秒期间的治疗？

由于您已经分析了您的应用程序，所以线索可以在两次内核执行之间的橙色大框中（在您的图像中是不可读的）。

如果这是一个同步删除它。

如果这是一条类似 cudaLaunch_KernelName 的轨迹，请尝试使您的处理更大（更多数据或更多计算），因为您向 GPU 发送命令所花费的时间多于执行它所花费的时间，因此您无法对这些不同的流进行并行计算。

为什么两个 CUDA 流中的操作不重叠？

Why operations in two CUDA Streams are not overlapping?

cuda

nvprof

cuda-streams

nvvp