CUDA多设备基于页式内存的异步执行实现方案咨询

1. 单线程+CUDA流+设备上下文切换 #

无需依赖OpenMP多线程，单线程内通过显式切换设备上下文，为每个设备创建独立流并绑定异步操作，即可实现多设备并行。核心逻辑是：主机依次为每个设备提交拷贝、核函数任务，切换设备上下文后，当前设备的异步操作会在后台执行，主机无需阻塞等待，直接处理下一个设备的任务。

示例代码框架：

int device_count;
cudaGetDeviceCount(&device_count);

// 为每个设备创建独立流
cudaStream_t streams[device_count];
for (int i = 0; i < device_count; i++) {
    cudaSetDevice(i);
    cudaStreamCreate(&streams[i]);
}

// 批量提交多设备任务
for (int i = 0; i < device_count; i++) {
    cudaSetDevice(i);
    // 主机到设备异步拷贝（页式内存）
    cudaMemcpyAsync(d_data[i], h_data[i], data_size, cudaMemcpyHostToDevice, streams[i]);
    // 绑定到对应流的核函数启动
    compute_kernel<<<grid_dim, block_dim, 0, streams[i]>>>(d_data[i], d_result[i]);
    // 设备到主机异步拷贝（页式内存）
    cudaMemcpyAsync(h_result[i], d_result[i], result_size, cudaMemcpyDeviceToHost, streams[i]);
}

// 等待所有设备任务完成并清理资源
for (int i = 0; i < device_count; i++) {
    cudaSetDevice(i);
    cudaStreamSynchronize(streams[i]);
    cudaStreamDestroy(streams[i]);
}

关键：每次操作前必须切换到目标设备上下文，确保流和操作绑定到对应硬件，实现真正的并行调度。

2. 统一内存（Unified Memory）简化实现 #

若设备支持计算能力3.0及以上，可使用CUDA统一内存替代页式内存，无需手动调用拷贝接口，由系统自动管理主机-设备数据迁移，大幅简化代码。数据量较大时，Pascal及以后架构的细粒度迁移机制能保证较好性能。

示例代码框架：

// 分配跨设备共享的统一内存
float *u_data, *u_result;
cudaMallocManaged(&u_data, data_size * device_count);
cudaMallocManaged(&u_result, result_size * device_count);

// 为每个设备分配任务
for (int i = 0; i < device_count; i++) {
    cudaSetDevice(i);
    compute_kernel<<<grid_dim, block_dim>>>(u_data + i*data_size, u_result + i*result_size);
}

// 等待所有设备任务完成
cudaDeviceSynchronize();

// 后续直接访问u_result即可获取结果，无需手动拷贝
cudaFree(u_data);
cudaFree(u_result);

3. 事件依赖控制（可选） #

如果需要跨设备的任务同步（如某设备任务依赖另一设备的输出），可使用CUDA事件实现精细调度：

cudaEvent_t device0_done;
cudaSetDevice(0);
cudaEventCreate(&device0_done);
cudaEventRecord(device0_done, streams[0]);

// 设备1的任务等待设备0完成
cudaSetDevice(1);
cudaStreamWaitEvent(streams[1], device0_done, 0);
// 提交设备1的依赖任务
compute_kernel<<<grid_dim, block_dim, 0, streams[1]>>>(...);

方案对比 #

• 单线程流调度：无多线程同步开销，符合CUDA原生调度模型，代码轻量可控。
• 统一内存：代码最简洁，自动处理数据迁移，适合对性能要求不极端、追求开发效率的场景。

内容的提问来源于stack exchange，提问作者Samuel