CUDA主机-设备同步疑问：函数内调用核函数是否自动同步？

首先直接给你明确结论：

• 不管是直接在main里调用核函数，还是通过自定义的initializedevvar()函数调用，主机都会立即恢复控制，完全不会等待核函数在GPU上执行完成。
• 想要让"gpu call"的打印操作在核函数执行完毕后才进行，不可能通过把核函数放到自定义函数里调用的方式规避同步操作，必须显式使用CUDA的同步API。

为什么会这样？ #

CUDA核函数的启动本质是异步非阻塞的：当你写下initialize<<<1,1>>>(...)这条语句时，主机端只是把核函数的启动指令提交到了默认流（Stream 0）的任务队列里，然后就立刻转头去执行下一行代码了——根本不会等GPU把核函数的逻辑跑完。这个异步特性是CUDA的核心机制之一，和核函数是在哪个函数里被调用的没有任何关系，哪怕你把核函数调用包在十层自定义函数里，主机依然会直接返回，不会等待GPU任务。

所以在你的代码里，两种调用方式之后的std::cout<<"gpu call"<<endl;都会在核函数还没执行（或者刚启动）的时候就被执行，完全达不到你想要的“核函数跑完再打印”的效果。

那怎么实现核函数执行完再打印？ #

必须显式添加同步操作，常用的方式有两种：

1. 全局设备同步：调用cudaDeviceSynchronize()，这个函数会阻塞当前主机线程，直到GPU上所有未完成的任务都执行完毕。
2. 流级同步：调用cudaStreamSynchronize(cudaStreamDefault)，只等待默认流里的任务完成（比全局同步更高效，不会阻塞其他流的任务）。

修改你的代码示例的话，应该是这样：

void main(){
    bool *state1, *state2, *state3, *state4;
    cudaMalloc(&state1, sizeof(bool));
    cudaMalloc(&state2, sizeof(bool));
    cudaMalloc(&state3, sizeof(bool));
    cudaMalloc(&state4, sizeof(bool));
    
    // 直接调用核函数后同步
    initialize<<<1,1>>>(state1, state2, state3, state4);
    cudaStreamSynchronize(cudaStreamDefault); // 等待核函数执行完成
    std::cout<<"gpu call"<<endl;
    
    ... ...
    
    // 自定义函数调用后同步
    auto dummy = initializedevvar(state1, state2, state3, state4);
    cudaStreamSynchronize(cudaStreamDefault);
    std::cout<<"gpu call"<<endl;
    
    cudaFree(state1);
    cudaFree(state2);
    cudaFree(state3);
    cudaFree(state4);
}

另外提一句：默认流在某些场景下会触发隐式同步（比如主机调用cudaMemcpy这类数据传输API时，会等待默认流的任务完成），但单纯的核函数调用后直接执行CPU端的打印操作，不会触发这种隐式同步，所以必须手动加同步。

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