问映射内存:设备->主机写入在主机上不可见
EN

Stack Overflow用户

提问于 2015-10-31 10:58:21

回答 1查看 371关注 0票数 0

我想要做的是修改一个驻留在映射内存中的变量，该变量将导致主程序退出。

但是，主程序没有这样做，而是继续在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;
}

cuda

单节点MySQL

低成本高体验，解决您的基础业务数据需求

回答 1

Stack Overflow用户

回答已采纳

发布于 2015-10-31 15:36:26

当我在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;
}