cuda 异步函数

为了提高cuda效率使用异步函数是一个很常规的选择，但是异步函数并没有我自己想象的这么智能。

它要你要异步传输的数据在主机端(host)不能被更改，即异步函数只是指示了一个传输的位置指针，并没有对这个数据进行缓存，到真正需要的时候，才会去主机内存中去寻找这个值。所以在做异步的时候要保证异步传输的主机端不能在异步完成(或者说拷贝完成)之前被修改。要不然这个数据就成了修改后的了。

实验代码如下

#include <stdio.h>
#include <cuda_runtime.h>

__global__ void async_kernel(int *x)
{
    printf("%d\n",x[0]);
}

int main(int argc,char *argv[])
{
    int *x,*xd;

    cudaStream_t stm;
    cudaStreamCreate(&stm);

    cudaMallocHost(&x,sizeof(int));
    cudaMalloc(&xd,sizeof(int));

    printf("Start\n");

    x[0]=1;
    cudaMemcpyAsync(xd,x,sizeof(int),cudaMemcpyHostToDevice,stm);
    x[0]=10;

    async_kernel<<<1,1,0,stm>>>(xd);

    cudaDeviceSynchronize();

    printf("End\n");
    fflush(stdout);
}

执行效果如下

当然别忘了同步函数 cudaDeviceSynchronize();