如何确保两个CUDA流中的kernel同时提交到GPU上运行？

你的硬件环境（GTX1070帕斯卡架构、CUDA 11.3）本身支持多Kernel并行执行，以下是可落地的优化方案，以及你当前实现的问题：

现有实现的不足 #

你当前用双CPU线程加条件变量同步的方案完全没有必要，反而会引入额外开销：

• CUDA Kernel的启动语法<<<>>>本身是异步调用，提交后立刻返回CPU，不需要多线程来实现“同时提交”
• 双CPU线程的同步开销、线程调度延迟，反而会拉长两个Kernel提交的时间差，效果不如单线程连续提交
• CPU层面的线程同步只能保证两个线程几乎同时调用启动接口，无法控制CUDA驱动层的调度逻辑，不能保证Kernel在GPU端同时启动

推荐优化方案 #

方案1：单线程多流提交（最简单，性能优于现有实现） #

直接在单线程内连续向两个不同的非阻塞流提交Kernel即可，CUDA驱动会自动识别两个流无依赖，尽可能调度两个Kernel并行执行，代码大幅简化：

#include <cuda_runtime.h>

__global__ void kernelZero(){/* 你的逻辑 */}
__global__ void kernelOne(){/* 你的逻辑 */}

int main(void) {
    cudaStream_t streamZero, streamOne;
    // 创建非阻塞流（默认参数即为非阻塞，支持并发）
    cudaStreamCreate(&streamZero);
    cudaStreamCreate(&streamOne); 

    // 单线程连续提交两个Kernel，提交时间差远小于多线程同步方案
    kernelZero<<<blocksPerGridZero, threadsPerBlockZero, 0, streamZero>>>();
    kernelOne<<<blocksPerGridOne, threadsPerBlockZero, 0, streamOne>>>();

    // 等待两个流执行完成
    cudaStreamSynchronize(streamZero);
    cudaStreamSynchronize(streamOne);

    cudaStreamDestroy(streamZero);  
    cudaStreamDestroy(streamOne);  
    return 0;
}

方案2：使用CUDA Graph实现最高概率的并行启动（满足“完全同时下发执行”需求） #

CUDA 11.3完全支持CUDA Graph特性，你可以预先把两个流的执行逻辑捕获为静态执行图，一次性提交执行，驱动会提前知道两个流的任务无依赖，会最大程度保证两个Kernel同时调度到GPU执行，启动延迟远低于动态提交的方案，示例代码如下：

#include <cuda_runtime.h>

__global__ void kernelZero(){/* 你的逻辑 */}
__global__ void kernelOne(){/* 你的逻辑 */}

int main(void) {
    cudaStream_t streamZero, streamOne;
    cudaStreamCreate(&streamZero);
    cudaStreamCreate(&streamOne); 

    cudaGraph_t graph;
    cudaGraphExec_t graphExec;

    // 开始捕获流的执行逻辑
    cudaStreamBeginCapture(streamZero, cudaStreamCaptureModeGlobal);
    kernelZero<<<blocksPerGridZero, threadsPerBlockZero, 0, streamZero>>>();
    // 同时捕获另一个流的Kernel
    kernelOne<<<blocksPerGridOne, threadsPerBlockZero, 0, streamOne>>>();
    cudaStreamEndCapture(streamZero, &graph);

    // 实例化执行图
    cudaGraphInstantiate(&graphExec, graph, NULL, NULL, 0);

    // 一次性提交整个执行图，两个Kernel会被驱动尽量并行调度
    cudaGraphLaunch(graphExec, streamZero);
    cudaStreamSynchronize(streamZero);

    // 资源释放
    cudaGraphExecDestroy(graphExec);
    cudaGraphDestroy(graph);
    cudaStreamDestroy(streamZero);  
    cudaStreamDestroy(streamOne);  
    return 0;
}

必要的注意事项 #

要实现两个Kernel实际并行执行，还需要满足硬件资源要求：

• 两个Kernel占用的SM总数、寄存器总数、共享内存总数不超过GTX1070的硬件上限（GTX1070共15个SM，每个SM有65536个32位寄存器、96KB共享内存），如果单个Kernel已经占满所有硬件资源，另一个Kernel只能排队执行
• 不要开启CUDA_LAUNCH_BLOCKING环境变量，该变量会强制所有Kernel同步执行，无法并发
• 不要在核函数内部使用printf、全局同步等会导致执行停顿的逻辑，否则会影响并行效果

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