12

在 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没有页面锁定内存的情况下执行？我在这里错过了什么吗？

4

1 回答 1

18

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()调用时阻塞，而不是在每个内存传输调用时阻塞。

这是绝对预期的行为。

于 2012-12-30T19:50:33.187 回答