使用可分页内存进行异步内存复制的效果？

Question

在CUDA C最佳实践指南5.0版,第6.1.2节中,写道:

与cudaMemcpy()相比,异步传输版本需要固定主机内存(请参阅固定内存),它还包含一个附加参数,即流ID.

这意味着cudaMemcpyAsync如果我使用简单的内存,函数应该会失败.

但事实并非如此.

出于测试目的,我尝试了以下程序:

核心:

__global__ void kernel_increment(float* src, float* dst, int n)
{
    int tid = blockIdx.x * blockDim.x + threadIdx.x;

    if(tid<n)   
        dst[tid] = src[tid] + 1.0f;
}

主要:

int main()
{
    float *hPtr1, *hPtr2, *dPtr1, *dPtr2;

    const int n = 1000;

    size_t bytes = n * sizeof(float);

    cudaStream_t str1, str2;

    hPtr1 = new float[n];
    hPtr2 = new float[n];

    for(int i=0; i<n; i++)
        hPtr1[i] = static_cast<float>(i);

    cudaMalloc<float>(&dPtr1,bytes);
    cudaMalloc<float>(&dPtr2,bytes);

    dim3 block(16);
    dim3 grid((n + block.x - 1)/block.x);

    cudaStreamCreate(&str1);
    cudaStreamCreate(&str2);

    cudaMemcpyAsync(dPtr1,hPtr1,bytes,cudaMemcpyHostToDevice,str1);
    kernel_increment<<<grid,block,0,str2>>>(dPtr1,dPtr2,n);
    cudaMemcpyAsync(hPtr2,dPtr2,bytes,cudaMemcpyDeviceToHost,str1);

    printf("Status: %s\n",cudaGetErrorString(cudaGetLastError()));

    cudaDeviceSynchronize();

    printf("Status: %s\n",cudaGetErrorString(cudaGetLastError()));

    cudaStreamDestroy(str1);
    cudaStreamDestroy(str2);

    cudaFree(dPtr1);
    cudaFree(dPtr2);

    for(int i=0; i<n; i++)
        std::cout<<hPtr2[i]<<std::endl;

    delete[] hPtr1;
    delete[] hPtr2;

    return 0;
}

该程序给出了正确的输出.数组成功递增.

如何在cudaMemcpyAsync没有页面锁定内存的情况下执 我在这里错过了什么吗？

Answer 1

cudaMemcpyAsync从根本上说是一个异步版本cudaMemcpy.这意味着它在发出复制调用时不会阻塞调用主机线程.这是电话的基本行为.

可选的,如果呼叫发射到非默认流,如果主机内存是一个固定配置,该装置具有自由DMA复制引擎,可以发生在复制操作,而GPU同时进行其他操作:无论是内核执行或另一个副本(在具有两个DMA复制引擎的GPU的情况下).如果不满足任何这些条件,则GPU上的操作在功能上与标准cudaMemcpy呼叫相同,即.它在GPU上串行操作,不会同时发生复制内核执行或同时多次复制.唯一的区别是操作不会阻塞调用主机线程.

在示例代码中,主机源和目标内存未固定.因此,内存传输不能与内核执行重叠(即,它们在GPU上串行操作).这些调用在主机上仍然是异步的.所以你拥有的功能相当于:

cudaMemcpy(dPtr1,hPtr1,bytes,cudaMemcpyHostToDevice);
kernel_increment<<<grid,block>>>(dPtr1,dPtr2,n);
cudaMemcpy(hPtr2,dPtr2,bytes,cudaMemcpyDeviceToHost);

除了所有调用在主机上都是异步的,因此主机线程在cudaDeviceSynchronize()调用时阻塞,而不是在每次内存传输调用时阻塞.

这是绝对预期的行为.