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();