CUDA: Modify mapped memory from host while kernel is executing

Question

I want to modify a piece of mapped memory from the host WHILE the kernel is executing, and then read this value from the kernel.

I'm trying to do this the following way. I have a kernel like this:

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

d_ptr is mapped to a piece of memory accessible from the host.

I also have a host thread that looks like this

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

So, the host thread is repeatedly writing a 0 to a place that the kernel is repeatedly reading from until it sees a 0. Theoretically, the kernel should stop right after it reads the value written by the host thread. The problem is that the kernel never reads this 0, so it never terminates.

The weird part is that if I add a print statement in the kernel like this

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

Then it does read the 0 and terminates. I have no clue on what's happening. The cuda programming guide does not provide much information about concurrent access to mapped memory and it's been pretty hard to find a question that addresses this either. Any pointers? I'm using Windows and I'm compiling with

nvcc -g -arch=sm_20 -lineinfo

The whole code looks like this:

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

recommendations:

• add this before your call to cudaHostAlloc:

  cudaSetDeviceFlags(cudaDeviceMapHost);
  

  the documentation suggests this.

• declare your flag variable as volatile:

  __global__ void kernel(volatile int* d_ptr)
  

The following code works for me without the use of threads:

#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;
}