CUDA：在内核执行时从主机修改映射内存

Question

我想在内核执行时从主机修改一块映射内存，然后从内核中读取这个值。

我正在尝试通过以下方式执行此操作。我有一个这样的内核：

__global__ void kernel(int* d_ptr)
{
    *d_ptr = 1;
    while( *d_ptr);
}

d_ptr 映射到主机可访问的一块内存。

我也有一个类似这样的宿主线程

void run( void* input )
{
    int* h_ptr = (int*)input;
    while( kernel_running)
        *h_ptr = 0;
}

因此，主机线程反复向内核重复读取的位置写入 0，直到它看到 0。理论上，内核应该在读取主机线程写入的值后立即停止。问题是内核从不读取这个 0，所以它永远不会终止。

奇怪的是，如果我像这样在内核中添加打印语句

__global__ void kernel(int* d_ptr)
{
    *d_ptr = 1;
    while( *d_ptr) printf("%d\n", *d_ptr);
}

然后它确实读取 0 并终止。我不知道发生了什么。 cuda programming guide 没有提供很多关于并发访问映射内存的信息，而且也很难找到解决这个问题的问题。任何指针？我正在使用 Windows 并且正在编译

nvcc -g -arch=sm_20 -lineinfo

整个代码如下所示：

bool kernel_running = 0;

__global__ void kernel(int* d_ptr)
{
    *d_ptr = 1;
    while( *d_ptr) printf("%d\n", *d_ptr);
}

void run( void* input )
{
    int* h_ptr = (int*)input;
    while( kernel_running)
    {
        *h_ptr = 0;
    }
}

int main()
{
    // HOST AND DEVICE POINTERS
    int* h_ptr = 0;
    int* d_ptr = 0;

    // INITIALIZE POINTERS
    assert( cudaHostAlloc(&h_ptr, sizeof(int), cudaHostAllocMapped) == cudaSuccess);
    assert( cudaHostGetDevicePointer(&d_ptr, h_ptr, 0) == cudaSuccess);

    // RUN KERNEL
    kernel_running = 1;
    _beginthread( run, 0, h_ptr);
    kernel<<<1,1>>>(d_ptr);
    assert( cudaDeviceSynchronize() == cudaSuccess);
    kernel_running = 0;
}

Answer 1

建议：

• 在致电cudaHostAlloc之前添加：

  cudaSetDeviceFlags(cudaDeviceMapHost);
  

  documentation 建议这样做。

• 将您的标志变量声明为volatile：

  __global__ void kernel(volatile int* d_ptr)
  

以下代码无需使用线程即可为我工作：

#include <stdio.h>

#define cudaCheckErrors(msg) \
    do { \
        cudaError_t __err = cudaGetLastError(); \
        if (__err != cudaSuccess) { \
            fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \
                msg, cudaGetErrorString(__err), \
                __FILE__, __LINE__); \
            fprintf(stderr, "*** FAILED - ABORTING\n"); \
            exit(1); \
        } \
    } while (0)

__global__ void mykernel(volatile int *data){

  while (*data) {};
  printf("finished\n");
}

int main(){

  int *d_data, *h_data;
  cudaSetDeviceFlags(cudaDeviceMapHost);
  cudaCheckErrors("cudaSetDeviceFlags error");
  cudaHostAlloc((void **)&h_data, sizeof(int), cudaHostAllocMapped);
  cudaCheckErrors("cudaHostAlloc error");
  cudaHostGetDevicePointer(&d_data, h_data, 0);
  cudaCheckErrors("cudaHostGetDevicePointer error");
  *h_data = 1;
  printf("kernel starting\n");
  mykernel<<<1,1>>>(d_data);
  cudaCheckErrors("kernel fail");
  getchar();
  *h_data = 0;
  cudaDeviceSynchronize();
  cudaCheckErrors("kernel fail 2");
  return 0;
}