Skip to content

Instantly share code, notes, and snippets.

@brycelelbach
Last active May 27, 2020 21:44
Show Gist options
  • Save brycelelbach/111ff48d6de40d2fe3d924c665032af7 to your computer and use it in GitHub Desktop.
Save brycelelbach/111ff48d6de40d2fe3d924c665032af7 to your computer and use it in GitHub Desktop.
// 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();
}
}
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment