4

CUDA 内核是否可以在没有任何主机端调用（例如 of ）的情况下将写入同步到设备映射内存cudaDeviceSynchronize？当我运行以下程序时，内核似乎并没有在终止之前等待对设备映射内存的写入完成，因为在内核启动后立即检查页面锁定的主机内存不会显示对内存的任何修改（除非插入延迟或cudaDeviceSynchronize取消注释调用）：

#include <stdio.h>
#include <cuda.h>

__global__ void func(int *a, int N) {
    int idx = threadIdx.x;

    if (idx < N) {
        a[idx] *= -1;
        __threadfence_system();
    }
}

int main(void) {
    int *a, *a_gpu;
    const int N = 8;
    size_t size = N*sizeof(int);

    cudaSetDeviceFlags(cudaDeviceMapHost);
    cudaHostAlloc((void **) &a, size, cudaHostAllocMapped);
    cudaHostGetDevicePointer((void **) &a_gpu, (void *) a, 0);

    for (int i = 0; i < N; i++) {
        a[i] = i;
    }
    for (int i = 0; i < N; i++) {
        printf("%i ", a[i]);
    }
    printf("\n");

    func<<<1, N>>>(a_gpu, N);
    // cudaDeviceSynchronize();

    for (int i = 0; i < N; i++) {
        printf("%i ", a[i]);
    }
    printf("\n");

    cudaFreeHost(a);
}

我在 Linux 上使用 CUDA 4.2.9 为 sm_20 编译上述内容，并在 Fermi GPU (S2050) 上运行它。

4

1 回答 1

4

在任何内核活动发生之前，内核启动将立即返回到主机代码。内核执行以这种方式与主机执行异步并且不会阻塞主机执行。因此，您必须稍等片刻或使用屏障（如 cudaDeviceSynchronize()）来查看内核的结果也就不足为奇了。

如此处所述：

为了方便主机和设备之间的并发执行，一些函数调用是异步的：在设备完成请求的任务之前，控制权返回给主机线程。这些都是：

  • 内核启动；
  • 两个地址之间的内存复制到同一个设备内存；
  • 从主机到设备的内存复制 64 KB 或更少的内存块；
  • 由以 Async 为后缀的函数执行的内存拷贝；
  • 内存设置函数调用。

当然，这都是故意的，以便您可以同时使用 GPU 和 CPU。如果您不想要这种行为，您已经发现的一个简单解决方案是插入一个屏障。如果您的内核正在生成您将立即复制回主机的数据，那么您不需要单独的屏障。内核之后的 cudaMemcpy 调用将等到内核完成后再开始它的复制操作。

我想回答你的问题，你希望内核启动是同步的，甚至不需要使用屏障（你为什么要这样做？添加 cudaDeviceSynchronize() 调用有问题吗？）可以这样做：

“程序员可以通过将 CUDA_LAUNCH_BLOCKING 环境变量设置为 1 来全局禁用系统上运行的所有 CUDA 应用程序的异步内核启动。此功能仅用于调试目的，绝不能用作使生产软件可靠运行的一种方式。”

如果你想要这种synchronous行为，最好只使用障碍（或依赖于另一个后续的 cuda 调用，如 cudaMemcpy）。如果您使用上述方法并依赖它，那么一旦其他人尝试在未设置环境变量的情况下运行它，您的代码就会中断。所以这真的不是一个好主意。

于 2012-12-05T01:40:06.293 回答