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

Question

在 CUDA C Best Practices Guide Version 5.0, Section 6.1.2 中，写道:

In contrast with cudaMemcpy(), the asynchronous transfer version requires pinned host memory (see Pinned Memory), and it contains an additional argument, a stream 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()调用而不是在每个内存传输调用中。

这绝对是预期的行为。