arrive_wait_stencil_solve.cpp

// Every thread must wait for every other thread to read its
// neighbors and finish its computation before it can start storing.
__host__ __device__
void solve(double* u, size_t n, size_t steps) {
  for (size_t t = 0; t < steps; ++t) {
    double un = stencil(u[n - 1], u[n], u[n + 1]);
    __syncthreads();
    u[n] = un;
    __syncthreads();
  }
} 

// Every thread must wait for every other thread to read its
// neighbors before it can start computing.
__host__ __device__
void solve(double* u, size_t n, size_t steps) {
  for (size_t t = 0; t < steps; ++t) {
    double left = u[n – 1], right = u[n + 1];
    __syncthreads();
    u[n] = stencil(left, u[n], right);
    __syncthreads();
  }
} 

__host__ __device__
void solve(double* u, size_t n, size_t steps, barrier& b) {
  for (size_t t = 0; t < steps; ++t) {
    double left = u[n – 1], right = u[n + 1];
    auto tok = b.arrive();
    double un = stencil(left, u[n], right);
    b.wait(tok);
    u[n - 1] = un;
    b.arrive_and_wait();
  }
}