BookGPU/Chapters/chapter17/code/mior_kernels.cl

   1 kernel void simulate(
   2     global         MM        allMMList[],
   3     global         OM        allOMList[],
   4     global         World     allWorlds[],
   5     constant       int       allMMOffsets[],
   6     constant       int       allOMOffsets[],
   7     global const   int       allMMCSR[],
   8     global const   int       allOMCSR[],
   9     global         int       allParts[])
  10 {
  11     const int iSim = get_group_id(0);
  12     const int i    = get_local_id(0);
  13
  14     global World *     world     = allWorlds + iSim;
  15     global MM *        mmList    = allMMList + iSim * NB_MM;
  16     global OM *        omList    = allOMList + iSim * NB_OM;
  17     constant int *     mmOffsets = allMMOffsets + iSim * NB_MM;
  18     constant int *     omOffsets = allOMOffsets + iSim * NB_OM;
  19     global const int * mmCSR     = allMMCSR;
  20     global const int * omCSR     = allOMCSR;
  21     global int *       parts     = allParts  + iSim * PARTS_SIZE;
  22
  23     if (i < world->nbOM) {
  24         scatter(omList, omOffsets, omCSR, parts, world);
  25     }
  26
  27     barrier(CLK_GLOBAL_MEM_FENCE);
  28
  29     if (i < world->nbMM) {
  30         live(mmList, mmOffsets, mmCSR, parts, world);
  31     }
  32
  33     barrier(CLK_GLOBAL_MEM_FENCE);
  34
  35     if (i < world->nbOM) {
  36         gather(omList, omOffsets, omCSR, parts, world);
  37     }
  38
  39     barrier(CLK_GLOBAL_MEM_FENCE);
  40 }