为什么两个 CUDA 流中的操作不重叠?
Why operations in two CUDA Streams are not overlapping?
我的程序是一个流水线,其中包含多个内核和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) );
}
您是否试图重叠 82 微秒期间的治疗?
由于您已经分析了您的应用程序,所以线索可以在两次内核执行之间的橙色大框中(在您的图像中是不可读的)。
如果这是一个同步删除它。
如果这是一条类似 cudaLaunch_KernelName 的轨迹,请尝试使您的处理更大(更多数据或更多计算),因为您向 GPU 发送命令所花费的时间多于执行它所花费的时间,因此您无法对这些不同的流进行并行计算。
我的程序是一个流水线,其中包含多个内核和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) );
}
您是否试图重叠 82 微秒期间的治疗?
由于您已经分析了您的应用程序,所以线索可以在两次内核执行之间的橙色大框中(在您的图像中是不可读的)。
如果这是一个同步删除它。
如果这是一条类似 cudaLaunch_KernelName 的轨迹,请尝试使您的处理更大(更多数据或更多计算),因为您向 GPU 发送命令所花费的时间多于执行它所花费的时间,因此您无法对这些不同的流进行并行计算。