Моя программа представляет собой конвейер, который содержит несколько ядер и memcpys. Каждая задача будет проходить через один и тот же конвейер с разными входными данными. Хост-код сначала выбирает канал, инкапсуляцию памяти блокнота и объекты CUDA при обработке задачи. А после последнего этапа я запишу событие, а затем перейду к обработке следующего задания.
Логика основного конвейера заключается в следующем. Проблема заключается в том, что операции в разных потоках не перекрываются. Я приложил график обработки 10 задач. Вы можете видеть, что ни одна операция в потоках не перекрывается. Для каждого ядра имеется 256 потоков в блоке и 5 блоков в сетке.
Все буферы, используемые для memcpy, закреплены, я уверен, что я удовлетворяю этим требованиям для перекрывающегося выполнения ядра и передачи данных. Может кто-нибудь помочь мне выяснить причину? Спасибо.
Информация об окружающей среде
GPU: Tesla K40m (GK110)
Макс. Деформации / SM: 64
Максимальное количество нитей блоков / SM: 16
Макс. Потоков / 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) );
}
