cudaMemcpyAsync一定是异步的吗？

参考代码：

#include <iostream>
#include <device_launch_parameters.h>
#include <cuda_runtime.h>
#include <string.h>
#include <assert.h>

#define N 1048576

__global__ void MyKernel(int *p0, int n)
{
    int i = threadIdx.x + blockIdx.x * blockDim.x;
    if (i < n)
    {
        p0[i] = 1;
    }
}

#define CUDA_CALL(x) {const cudaError_t a = (x); if (a != cudaSuccess) { printf("\nCUDA \
Error: %s (err_num¼%d) \n", cudaGetErrorString(a), a); cudaDeviceReset(); assert(0);} }

// #define USE_MALLOC_HOST

int main(void)
{

    int* p0;
    long nbytes = sizeof(int) * N;

#ifdef USE_MALLOC_HOST
    CUDA_CALL(cudaMallocHost((void**)&p0, nbytes));
#else
    p0 = new int[N];
#endif
    memset(p0, 0, nbytes);

    int* devp0;
    // GPU端分配内存
    CUDA_CALL(cudaMalloc((void**)&devp0, nbytes));

    dim3 dimBlock(1024);
    dim3 dimGrid(N/1024);

    // 执行kernel
    cudaStream_t stm;
    CUDA_CALL(cudaStreamCreate(&stm));

    MyKernel<<<dimGrid, dimBlock, 0, stm>>>(devp0, N);
    // 将在GPU端计算好的结果拷贝回CPU端
    CUDA_CALL(cudaMemcpyAsync(p0, devp0, nbytes, cudaMemcpyDeviceToHost, stm));

    // 理论上，cudaMemcpyAsync是异步的，下面的断言会失败，但实际上。。。
    assert(p0[0] == 1);
    assert(p0[N/3] == 1);
    assert(p0[N/3*2] == 1);
    assert(p0[N-1] == 1);

    cudaStreamDestroy(stm);
    cudaFree(devp0);
#ifdef USE_MALLOC_HOST
    cudaFreeHost((void*)p0);
#else
    delete[] p0;
#endif
}

理论上，cudaMemcpyAsync是异步的，所以assert会失败。
但上面的代码运行却不会出错：

image.png

如果把上面代码中的USE_MALLOC_HOST宏打开，编译运行：

image.png

这次失败了。
所以看起来，当通过cudaMemcpyAsync从GPU拷贝数据到CPU时，如果目的地的host memory是new/malloc出来的，调用是同步的；如果host memory是cudaMallocHost出来的，则才真是异步的。
下面在nsys-ui中看看两种情况下的区别（要先把代码中的assert注释掉，不能让进程dump掉）。

host memory为new/malloc的情况

image.png

从图中可以看出，cudaMemcpyAsync返回的时候，HW中的Memcpy已经结束了。
host memory为cudaMallocHost的情况

image.png

从图中可以看出，基本上cudaMemcpyAsync返回后，HW中的Memcpy才开始工作。

官方文档说明见：CUDA Runtime API :: CUDA Toolkit Documentation (nvidia.com)
“For transfers from device memory to pageable host memory, the function will return only once the copy has completed.”这句话似乎也说明，从GPU拷贝数据到CPU上new/malloc出来的memory，是同步的。

cudaMemcpyAsync一定是异步的吗？

推荐阅读更多精彩内容