cuda - 将设备内存访问与主机线程同步

Question

CUDA 内核是否可以在没有任何主机端调用(例如，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) 上运行它。

Answer 1

在任何内核事件发生之前，内核启动将立即返回主机代码。通过这种方式，内核执行与主机执行异步，并且不会阻止主机执行。因此，您必须稍等一下，或者使用屏障(如 cudaDeviceSynchronize())来查看内核的结果，这并不奇怪。

如上所述here :

In order to facilitate concurrent execution between host and device, some function calls are asynchronous: Control is returned to the host thread before the device has completed the requested task. These are:

• Kernel launches;
• Memory copies between two addresses to the same device memory;
• Memory copies from host to device of a memory block of 64 KB or less;
• Memory copies performed by functions that are suffixed with Async;
• Memory set function calls.

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

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

"Programmers can globally disable asynchronous kernel launches for all CUDA applications running on a system by setting the CUDA_LAUNCH_BLOCKING environment variable to 1. This feature is provided for debugging purposes only and should never be used as a way to make production software run reliably. "

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