BookGPU/Chapters/chapter17/code/collem_kernels.cl

   1 #include "collem_structures.h"
   2
   3 kernel void reduce(
   4     global CollemWorld *world,
   5     global int *populations,
   6     global int *patchesOwners,
   7     global int *plotsPopulations)
   8 {
   9     const int i = get_global_id(0);
  10     const int j = get_global_id(1);
  11
  12     // Retrieve the owner parcel
  13     const int plotId = CELL(world, patchesOwners, i, j);
  14     if (plotId == -1) return;
  15
  16     // Add cell population to the owner parcel population
  17     atomic_add(plotsPopulations + plotId, CELL(world, populations, i, j));
  18 }
  19
  20 kernel void diffuse(
  21     global CollemWorld *world,
  22     global int *populations,
  23     global int *overflows,
  24     global int *newPopulations)
  25 {
  26     const int i = get_global_id(0);
  27     const int j = get_global_id(1);
  28
  29     // Compute the population to diffuse for this cell
  30     OVERFLOW(world, overflows, i, j) = CELL(world, populations, i, j) * world->diffusionRate;
  31
  32     // _syncthreads() in CUDA
  33     barrier(CLK_GLOBAL_MEM_FENCE);
  34
  35     // Retrieve neighbors surplus and add them to the current cell population
  36     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);
  37
  38     CELL(world, newPopulations, i, j) = CELL(world, populations, i, j) + surplus / 8.0 - OVERFLOW(world, overflows, i, j);
  39 }