/** * 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; } }