Add GPU version of life and start lazy parallel

kernel/c/life.c

@@ -13,6 +13,7 @@ static unsigned color = 0xFFFF00FF; // Living cells have the yellow color
 typedef unsigned cell_t;
 
 static cell_t *restrict _table = NULL, *restrict _alternate_table = NULL;
+static unsigned* _changed = NULL;
 
 static inline cell_t *table_cell (cell_t *restrict i, int y, int x)
 {
@@ -30,23 +31,33 @@ void life_init (void)
   // already allocated
   if (_table == NULL) {
     const unsigned size = DIM * DIM * sizeof (cell_t);
+    const changed_size = TILE_H * TILE_W * sizeof(unsigned);
 
-    PRINT_DEBUG ('u', "Memory footprint = 2 x %d bytes\n", size);
+    PRINT_DEBUG ('u', "Memory footprint = 2 x %d bytes (classic) + %d bytes (lazy)\n", size, changed_size);
 
     _table = mmap (NULL, size, PROT_READ | PROT_WRITE,
                    MAP_PRIVATE | MAP_ANONYMOUS, -1, 0);
 
     _alternate_table = mmap (NULL, size, PROT_READ | PROT_WRITE,
                              MAP_PRIVATE | MAP_ANONYMOUS, -1, 0);
+
+
+    _changed = mmap (NULL, changed_size, PROT_READ | PROT_WRITE,
+                             MAP_PRIVATE | MAP_ANONYMOUS, -1, 0);
+
+    for (unsigned i = 0; i < TILE_H * TILE_W; ++i)
+      _changed[i] = 1;
   }
 }
 
 void life_finalize (void)
 {
   const unsigned size = DIM * DIM * sizeof (cell_t);
+  const changed_size = TILE_H * TILE_W * sizeof(unsigned);
 
   munmap (_table, size);
   munmap (_alternate_table, size);
+  munmap (_changed, changed_size);
 }
 
 // This function is called whenever the graphical window needs to be refreshed
@@ -57,6 +68,35 @@ void life_refresh_img (void)
       cur_img (i, j) = cur_table (i, j) * color;
 }
 
+// unsigned life_invoke_ocl (unsigned nb_iter)
+// {
+//   size_t global[2] = {GPU_SIZE_X, GPU_SIZE_Y};
+//   size_t local[2]  = {GPU_TILE_W, GPU_TILE_H};
+//   cl_int err;
+
+//   monitoring_start_tile (easypap_gpu_lane (TASK_TYPE_COMPUTE));
+
+//   for (unsigned it = 1; it <= nb_iter; it++) {
+
+//     // Set kernel arguments
+//     //
+//     err = 0;//   for (unsigned it = 1; it <= nb_iter; it++) {
+//     err |= clSetKernelArg (compute_kernel, 0, sizeof (cl_mem), &cur_buffer);
+//     err |= clSetKernelArg (compute_kernel, 1, sizeof (cl_mem), &next_buffer);
+//     check (err, "Failed to set kernel arguments");
+
+//     err = clEnqueueNDRangeKernel (queue, compute_kernel, 2, NULL, global, local,
+//                                   0, NULL, NULL);
+//     check (err, "Failed to execute kernel");
+//   }
+
+//   clFinish (queue);
+
+//   monitoring_end_tile (0, 0, DIM, DIM, easypap_gpu_lane (TASK_TYPE_COMPUTE));
+
+//   return 0;
+// }
+
 // Only called when --dump or --thumbnails is used
 void life_refresh_img_ocl(void)
 {
@@ -177,6 +217,46 @@ unsigned life_compute_omp_tiled (unsigned nb_iter)
   return res;
 }
 
+int tiles_around_changed()
+{
+
+}
+
+// omp tiled version
+unsigned life_compute_omp_tiled_lazy (unsigned nb_iter)
+{
+  unsigned res = 0, change = 0, temp = 0;
+
+  for (unsigned it = 1; it <= nb_iter; it++) {
+    change = 0;
+    temp = 0;
+
+    #pragma omp parallel for schedule(runtime) collapse(2) shared(change) private(temp)
+    for (int y = 0; y < DIM; y += TILE_H)
+    {
+      for (int x = 0; x < DIM; x += TILE_W)
+      {
+        if (tiles_around_changed())
+          temp = do_tile (x, y, TILE_W, TILE_H, omp_get_thread_num());
+
+        _changed[y + (x / TILE_W)] = temp;
+
+        #pragma omp critical
+        change |= temp;
+      }
+    }
+
+    swap_tables ();
+
+    if (!change) { // we stop if all cells are stable
+      res = it;
+      break;
+    }
+  }
+
+  return res;
+}
+
 ///////////////////////////// Initial configs
 
 void life_draw_guns (void);

kernel/ocl/life.cl

@@ -1,6 +1,22 @@
+#include "kernel/ocl/common.cl"
+
 __kernel void life_ocl (__global unsigned *in, __global unsigned *out)
 {
+  int x = get_global_id (0);
+  int y = get_global_id (1);
+
+  if (y > 0 && y < DIM - 1 && x > 0 && x < DIM - 1) {
+    unsigned n  = 0;
+    unsigned me = in[y * DIM + x];
+
+    for (int yloc = y - 1; yloc < y + 2; yloc++)
+      for (int xloc = x - 1; xloc < x + 2; xloc++)
+        n += in[yloc * DIM + xloc];
+
+    n = (n == 3 + me) | (n == 3);
 
+    out[y * DIM + x] = n;
+  }
 }
 
 // DO NOT MODIFY: this kernel updates the OpenGL texture buffer
@@ -10,5 +26,5 @@ __kernel void life_update_texture (__global unsigned *cur, __write_only image2d_
   int y = get_global_id (1);
   int x = get_global_id (0);
 
-  // write_imagef (tex, (int2)(x, y), color_scatter (cur [y * DIM + x] * 0xFFFF00FF));
+  write_imagef (tex, (int2)(x, y), color_scatter (cur [y * DIM + x] * 0xFFFF00FF));
 }

validate_life

@@ -9,10 +9,16 @@ ITERATIONS="500 1000 5000"
 SIZE=2176
 VERSION="octa_off"
 LIFE_CONFIG=$1
+OCL_FLAG=""
+
+if [[ $LIFE_CONFIG =~ "ocl" ]]
+then
+    OCL_FLAG="-o"
+fi
 
 for iter in $ITERATIONS
 do
-    ./run -k life -v $LIFE_CONFIG -s $SIZE -a $VERSION -du -n -i $iter
+    ./run $OCL_FLAG -k life -v $LIFE_CONFIG -s $SIZE -a $VERSION -du -n -i $iter
 
     TRUTH="dumps/dump-life-seq-dim-$SIZE-iter-$iter.png"
     GENERATED="dump-life-$LIFE_CONFIG-dim-$SIZE-iter-$iter.png"