参考代码：
#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，是同步的。
