/**
   * pre-OpenCL 1.1 compatibility macros
  */
  
  #if __OPENCL_VERSION__ <= CL_VERSION_1_0
      #define atomic_add(p, v) atom_add(p, v)
      #define atomic_inc(p) atom_inc(p)
  #endif
  
  /**
   * DATA STRUCTURES
   */
   
  // Microbial colony
  typedef struct MM {
      float   x;
      float   y;
      int     carbon;
      int     dormancy;
  } MM;
  
  // Carbon deposit
  typedef struct OM {
      float   x;
      float   y;
      int     carbon;
      int     lock;
  } OM;
  
  // Environment
  typedef struct MiorWorld {
      int     nbMM;
      int     nbOM;
      int     RA;
      float   RR;
      float   GR;
      float   K;
      int     width;
      int     minSize;
      int     CO2;
      int     lock;
  } MiorWorld;
  
  /**
   * HELPERS
   */
  
  #define CAPACITY(store) store[0]
  #define NTH_LIST(store, index) (store + 1 + index * (CAPACITY(store) + 1))
  #define LIST_SIZE(list) list[0]
  #define LIST_GET(list, index) list[1 + index]
  #define LIST_SET(list, index, val) list[1 + index] = val
  #define LIST_APPEND(list, val) LIST_SET(list, atomic_inc(list), val)
  
  /**
   * KERNELS
   */
  
  /*
   * Dimensions expectations
   *
   * Global : nbMM, nbOM
   */
  kernel void topology(
      global MM mmList[],
      global OM omList[],
      global MiorWorld * world,
      global int mmCSR[],
      global int omCSR[],
      global int topo[])
  {
      const int iMM = get_global_id(0);
      const int iOM = get_global_id(1);
      
      const float dist = hypot(
          omList[iOM].x - mmList[iMM].x,
          omList[iOM].y - mmList[iMM].y
      );
      
      // If the two agents can interact
      if (dist <= world->RA) {
          /*global int * list = NTH_LIST(omCSR, iMM);
          LIST_SET(list, atomic_inc(list), 0);
          
          list = NTH_LIST(mmCSR, iOM);
          LIST_SET(list, atomic_inc(list), 0);
          */
          LIST_APPEND(NTH_LIST(omCSR, iMM), iOM);
          LIST_APPEND(NTH_LIST(mmCSR, iOM), iMM);
          
  #ifndef NOTOPO
          topo[iMM * world->nbOM + iOM] = 0;
  #endif
  	}
  }
  
  /**
   * Dimensions expectations
   *
   * Global : nbOM
   */
  kernel void scatter(
      global OM omList[],
      global MiorWorld * world,
      global int mmCSR[],
      global int parts[])
  {
      // Retrieve current OM index
      const int iOM = get_global_id(0);
      
      // Retrieve list and number of accessible MM (CSR)
      global int * mmIndexes = NTH_LIST(mmCSR, iOM);
      const int nbIndexes = LIST_SIZE(mmIndexes);
      
      // Compute the carbon part to allocate to each MM
      const int carbonPart = (world->K * omList[iOM].carbon) / nbIndexes;
      
      // Allocate parts
      for (int i = 0; i < nbIndexes; i++) {
          const int iMM = LIST_GET(mmIndexes, i);       
          parts[iMM * world->nbOM + iOM] = carbonPart;
      }
      
      omList[iOM].carbon -= nbIndexes * carbonPart;
  }
  
  /**
   * Dimensions expectations
   *
   * Global : nbMM
   */
  kernel void live(
      global MM mmList[],
      global MiorWorld * world,
      global int omCSR[],
      global int parts[])
  {
      // Retrieve current MM
      const int iMM = get_global_id(0);
      global MM * currentMM = mmList + iMM;
      
  #ifdef COPY_PARTS
      int localParts[NBOM];
      
      for (int i = 0; i < NBOM; i++) {
          localParts[i] = parts[iMM * world->nbOM + i];
      }
  #endif
      
      // Compute needs
      const int breathNeed = currentMM->carbon * world->RR;
      const int growthNeed = currentMM->carbon * world->GR;
      
      // Retrieve list and number of accessible OM (CSR)
      global int * omIndexes = NTH_LIST(omCSR, iMM);
      const int nbIndexes = LIST_SIZE(omIndexes);
      
      // BREATHING CHECK
      
      int i = 0;
      int remainingNeed = breathNeed;
      
      while (remainingNeed > 0 && i < nbIndexes) {
          const int iOM = LIST_GET(omIndexes, i);
  #ifdef COPY_PARTS
          remainingNeed -= localParts[iOM];
  #else
          remainingNeed -= parts[iMM * world->nbOM + iOM];
  #endif
          i++;
      }
      
      // DORMANCY CHECK
      
      if (remainingNeed > 0) {
          currentMM->dormancy = 1;
          return;
      }
      
      // ACTUAL BREATHING
      
      i = 0;
      remainingNeed = breathNeed;
      currentMM->dormancy = 0;
      
      while (i < nbIndexes) {
          const int iOM = LIST_GET(omIndexes, i);
  #ifdef COPY_PARTS
          const int offer = localParts[iOM];
  #else
          const int offer = parts[iMM * world->nbOM + iOM];
  #endif
          const int consum = min(offer, remainingNeed);
          
          remainingNeed -= consum;
          parts[iMM * world->nbOM + iOM] = offer - consum;
  #ifdef COPY_PARTS
          localParts[iOM] = offer - consum;
  #endif
          if (remainingNeed == 0) {
              break; // don't increment i !
          }
          
          i++;
      }
      
      // ENVIRONMENT UPDATE
      
      atomic_add(&(world->CO2), breathNeed);
      
      // GROWTH
      
      remainingNeed = growthNeed;
      
      while (remainingNeed > 0 && i < nbIndexes) {
          const int iOM = LIST_GET(omIndexes, i);
  
  #ifdef COPY_PARTS
          const int offer = localParts[iOM];
  #else
          const int offer = parts[iMM * world->nbOM + iOM];
  #endif
          const int consum = min(offer, remainingNeed);
          
          remainingNeed -= consum;
          parts[iMM * world->nbOM + iOM] = offer - consum;
  
  #ifdef COPY_PARTS
          localParts[iOM] = offer - consum;
  #endif
          i++;
      }
          
      // MM CARBON UPDATE
      
      currentMM->carbon += (growthNeed - remainingNeed);
      
      /*for (int i = 0; i < 310; i++) {
          parts[iMM * world->nbOM + i] = localParts[i];
      }*/
  }
      
      
  /**
   * Dimensions expectations
   *
   * Global : nbOM
   */
  kernel void gather(
      global OM omList[],
      global MiorWorld * world,
      global int mmCSR[],
      global int parts[])
  {
      // Retrieve current OM index
      const int iOM = get_global_id(0);
      
      // Retrieve list and number of accessible MM (CSR)
      global int * mmIndexes = NTH_LIST(mmCSR, iOM);
      const int nbIndexes = LIST_SIZE(mmIndexes);
      
      // Update current OM carbon value
      int carbon = omList[iOM].carbon;
      
      for (int i = 0; i < nbIndexes; i++) {
          const int iMM = LIST_GET(mmIndexes, i);
          carbon += parts[iMM * world->nbOM + iOM];
      }
      
      omList[iOM].carbon = carbon;
  }
  
  /**
   * Dimensions expectations
   *
   * Global : max(nbOM, nbMM)
   */
  kernel void autolive(
      global MM mmList[],
      global OM omList[],
      global MiorWorld * world,
      global int mmCSR[],
      global int omCSR[],
      global int parts[])
  {
      const int i = get_global_id(0);
      int CO2total = -1;
      
      while (CO2total == -1 || CO2total != world->CO2) {
          if (i < world->nbOM) {
              scatter(omList, world, mmCSR, parts);
          }
          
          barrier(CLK_GLOBAL_MEM_FENCE);
          
          if (i < world->nbMM) {
              live(mmList, world, omCSR, parts);
          }
          
          if (i < world->nbOM) {
              gather(omList, world, mmCSR, parts);
          }
          
          barrier(CLK_GLOBAL_MEM_FENCE);
          
          //if (CO2total == world->CO2) break;
          CO2total = world->CO2;
      }
  
  }