BookGPU/Chapters/chapter17/code/collem_kernel_diffuse.cl

   1 #include "collem_structures.h"
   2
   3 kernel void diffuse(
   4     global CollemWorld *world,
   5     global int *populations,
   6     global int *overflows,
   7     global int *newPopulations)
   8 {
   9     const int i = get_global_id(0);
  10     const int j = get_global_id(1);
  11
  12     // Compute the population to diffuse for this cell
  13     OVERFLOW(world, overflows, i, j) = CELL(world, populations, i, j) * world->diffusionRate;
  14
  15     // _syncthreads() in CUDA
  16     barrier(CLK_GLOBAL_MEM_FENCE);
  17
  18     // Retrieve neighbors surplus and add them to the current cell population
  19     int surplus = OVERFLOW(world, overflows, i + 1, j) + OVERFLOW(world, overflows, i - 1, j) + OVERFLOW(world, overflows, i, j - 1) + OVERFLOW(world, overflows, i, j + 1);
  20
  21     CELL(world, newPopulations, i, j) = CELL(world, populations, i, j) + surplus / 8.0 - OVERFLOW(world, overflows, i, j);
  22 }