CUDA 映射内存：设备 -> 主机写入在主机上不可见

Question

我试图做的是修改一个驻留在映射内存中的变量，这会导致主程序退出。

但不是这个，主程序一直在while (var == 0) ; 行上旋转。我不知道如何刷新新值，以便它在主机端也可见。

顺便说一句。该变量在任何地方都声明为volatile，我尝试使用__threadfence_system() 函数但没有成功。

主机->设备方向效果很好。

系统：Windows 7 x64，驱动程序 358.50，GTX 560

这是我无法工作的一段代码：

static void handleCUDAError(cudaError_t err, const char *file, int line)
{
    if (err != cudaSuccess) {
        printf("%s in %s at line %d\n", cudaGetErrorString(err), file, line);
        exit(EXIT_FAILURE);
    }
}

#define CUDA_ERROR_CHECK(err) (handleCUDAError(err, __FILE__, __LINE__ ))

__global__ void echoKernel(volatile int* semaphore)
{
    *semaphore = 1;

    __threadfence_system();
}

int main()
{
    CUDA_ERROR_CHECK(cudaSetDevice(0));
    CUDA_ERROR_CHECK(cudaSetDeviceFlags(cudaDeviceMapHost));

    volatile int var = 0;
    volatile int *devptr;

    CUDA_ERROR_CHECK(cudaHostRegister((int*)&var, sizeof (int), cudaHostRegisterMapped));
    CUDA_ERROR_CHECK(cudaHostGetDevicePointer(&devptr, (int*)&var, 0));

    echoKernel <<< 1, 1 >>> (devptr);

    while (var == 0) ;

    CUDA_ERROR_CHECK(cudaDeviceSynchronize());

    CUDA_ERROR_CHECK(cudaHostUnregister((int*)&var));
    CUDA_ERROR_CHECK(cudaDeviceReset());

    return 0;
}

Answer 1

当我在 linux 上运行您的代码时，它按原样运行，没有问题。

但是在 Windows 上，WDDM 命令批处理存在问题。实际上，在您进入挂起的 while 循环之前，您的内核不会启动并且不会启动。

WDDM 命令队列是最终将发送到 GPU 设备的命令队列。各种事件将导致此队列“刷新”，并将内容作为“批量”命令传递给 GPU。

各种 cuda 运行时 API 调用可能会有效地强制“刷新”命令队列，例如 cudaDeviceSynchronize() 或 cudaMemcpy()。然而，在内核启动之后，您在进入 while 循环之前不会发出任何运行时 API 调用。结果，在这种情况下，内核调用似乎“卡在”队列中并且从未“刷新”。

您可以通过多种方式解决此问题，例如在内核启动后记录一个事件，然后查询该事件的状态。这将产生刷新队列的效果，从而启动内核。

这是对我有用的代码修改示例：

#include <stdio.h>
static void handleCUDAError(cudaError_t err, const char *file, int line)
{
    if (err != cudaSuccess) {
        printf("%s in %s at line %d\n", cudaGetErrorString(err), file, line);
        exit(EXIT_FAILURE);
    }
}

#define CUDA_ERROR_CHECK(err) (handleCUDAError(err, __FILE__, __LINE__ ))

__global__ void echoKernel(volatile int* semaphore)
{
    *semaphore = 1;

    __threadfence_system();
}

int main()
{
    CUDA_ERROR_CHECK(cudaSetDevice(0));
    CUDA_ERROR_CHECK(cudaSetDeviceFlags(cudaDeviceMapHost));

    volatile int var = 0;
    volatile int *devptr;

    CUDA_ERROR_CHECK(cudaHostRegister((int*)&var, sizeof(int), cudaHostRegisterMapped));
    CUDA_ERROR_CHECK(cudaHostGetDevicePointer(&devptr, (int*)&var, 0));

    cudaEvent_t my_event;
    CUDA_ERROR_CHECK(cudaEventCreate(&my_event));

    echoKernel << < 1, 1 >> > (devptr);
    CUDA_ERROR_CHECK(cudaEventRecord(my_event));
    cudaEventQuery(my_event);

    while (var == 0);

    CUDA_ERROR_CHECK(cudaDeviceSynchronize());

    CUDA_ERROR_CHECK(cudaHostUnregister((int*)&var));
    CUDA_ERROR_CHECK(cudaDeviceReset());

    return 0;
}

在 CUDA 7.5、驱动程序 358.50、Win7 x64 发布项目、GTX460M 上测试。

请注意，我们不会将 cudaEventQuery 调用包装在标准错误检查器中，因为它的预期行为是在事件尚未完成时返回非零状态。