2021 update

Makefile

@@ -51,6 +51,7 @@ ALL_DEPENDS := $(DEPENDS) $(K_DEPENDS) $(T_DEPENDS) $(L_DEPENDS)
 MAKEFILES	:= Makefile
 
 CC			:= gcc
+#CC			:= clang-mp-11
 
 CFLAGS 		+= -O3 -march=native -Wall -Wno-unused-function
 CFLAGS		+= -I./include -I./traces/include
@@ -59,9 +60,9 @@ LDLIBS		+= -lm
 ifeq ($(ARCH),DARWIN)
 LDLIBS		+= -framework OpenGL
 else
-CFLAGS		+= -rdynamic
+CFLAGS		+= -pthread -rdynamic
 LDFLAGS		+= -export-dynamic
-LDLIBS		+= -lGL -lpthread -ldl
+LDLIBS		+= -lGL -ldl
 endif
 
 # Vectorization

README.md

@@ -13,3 +13,5 @@ Most of the parameters can be specified as command line arguments, when running
 * interactive mode / performance mode
 * monitoring mode
 * and much more!
+
+Please read the [Getting Started Guide](https://gforgeron.gitlab.io/easypap/doc/Getting_Started.pdf) before you start!

data/rle/greyship.rle

+#C small extensible c/4 greyship: Hartmut Holzwart, 24 March 2006
+x = 95, y = 133, rule = B3/S23
+23bo16bo$22bobo14bobo$22bobo14bobo$23bo16bo$23boo15boo$22b3o14b3o$$21b
+oboo13boboo$20boobo13boobo$22bobo14bobo$22bobo14bobo$21bo16bo$22bo16bo
+$20boo15boo$15b4oboo10b4oboo$15bobb3o11bobb3o$20bo16bo$15b3ob3o10b3ob
+3o$20boo15boo$15bobb3o11bobb3o$15b4obbo10b4obbo$20bobbo13bobbo$20bobo
+14bobo$23bo16bo$16boobo3bo9boobo3bo$15bobboo12bobboo$15boo7bo7boo7bo$
+16bobob4o9bobob4o$$16bo16bo27boobo$17bo16bo25bob3o$15b3o14b3o25boboo$
+16boo15boo24bobo$59bobo12boobo$15bobbob3o9bobbob3o19boo12bob3o$15b4obo
+11b4obo24b3o8boboo$15bo5boo9bo5boo22b3o7bobo$16b7o10b7o24bo7bobo$16bob
+o14bobo26boo8boo$17bo16bo27bobbo9b3o$16boboo13boboo25bo12b3o$17b3o14b
+3o25bobbo11bo$62b4o9boo$17boo15boo26boo11bobbo$17boo3bo11boo3bo21boo
+12bo$23bo16bo19b3o12bobbo$20booboo12booboo18bobo12b4o$21boboo13boboo
+19boo12boo$10bobo8bobo14bobo19boo12boo$5boobbobbo8bobo14bobo22bo9b3o$
+4b3obo3bo47bo12bobo$5boobbo13boo15boo21bo10boo$12bobbobo5bo4b3obobo5bo
+4b3obobo7b3o11boo$11b5o3bo4bo3b5o3bo4bo3b5o3bo5boo15bo$13bo6bobbo13bo
+bbo13bo18bo$10bobb6obo3b12obo3b12obo3bo17bo$5bobobobo47boobo9b3o$boobb
+ob4ob47obobo9boo$ooboobo53bob3obbo$bo5b53o3b3obo3bo$oo59bo10boobo13bo
+bboo$bo5b65obobo13bobb3o$ooboobo66bob3obbo8bo4bo$boobbob66o3b3obo3b5o$
+5bobo66bo10boboo$8b77o$$8b77o$5bobo66bo10boboo$boobbob66o3b3obo3b5o$oo
+boobo66bob3obbo8bo4bo$bo5b65obobo13bobb3o$oo59bo10boobo13bobboo$bo5b
+53o3b3obo3bo$ooboobo53bob3obbo$boobbob4ob47obobo9boo$5bobobobo47boobo
+9b3o$10bobb6obo3b12obo3b12obo3bo17bo$13bo6bobbo13bobbo13bo18bo$11b5o3b
+o4bo3b5o3bo4bo3b5o3bo5boo15bo$12bobbobo5bo4b3obobo5bo4b3obobo7b3o11boo
+$5boobbo13boo15boo21bo10boo$4b3obo3bo47bo12bobo$5boobbobbo8bobo14bobo
+22bo9b3o$10bobo8bobo14bobo19boo12boo$21boboo13boboo19boo12boo$20booboo
+12booboo18bobo12b4o$23bo16bo19b3o12bobbo$17boo3bo11boo3bo21boo12bo$17b
+oo15boo26boo11bobbo$62b4o9boo$17b3o14b3o25bobbo11bo$16boboo13boboo25bo
+12b3o$17bo16bo27bobbo9b3o$16bobo14bobo26boo8boo$16b7o10b7o24bo7bobo$
+15bo5boo9bo5boo22b3o7bobo$15b4obo11b4obo24b3o8boboo$15bobbob3o9bobbob
+3o19boo12bob3o$59bobo12boobo$16boo15boo24bobo$15b3o14b3o25boboo$17bo
+16bo25bob3o$16bo16bo27boobo$$16bobob4o9bobob4o$15boo7bo7boo7bo$15bobb
+oo12bobboo$16boobo3bo9boobo3bo$23bo16bo$20bobo14bobo$20bobbo13bobbo$
+15b4obbo10b4obbo$15bobb3o11bobb3o$20boo15boo$15b3ob3o10b3ob3o$20bo16bo
+$15bobb3o11bobb3o$15b4oboo10b4oboo$20boo15boo$22bo16bo$21bo16bo$22bobo
+14bobo$22bobo14bobo$20boobo13boobo$21boboo13boboo$$22b3o14b3o$23boo15b
+oo$23bo16bo$22bobo14bobo$22bobo14bobo$23bo16bo!

data/templates/kernel_template.c

+
+#include "easypap.h"
+
+#include <omp.h>
+
+// If defined, the initialization hook function is called quite early in the
+// initialization process, after the size (DIM variable) of images is known.
+// This function can typically spawn a team of threads, or allocated additionnal
+// OpenCL buffers.
+// A function named <kernel>_init_<variant> is search first. If not found, a
+// function <kernel>_init is searched in turn.
+void <template>_init (void)
+{
+  PRINT_DEBUG ('u', "Image size is %dx%d\n", DIM, DIM);
+  PRINT_DEBUG ('u', "Tile size is %dx%d\n", TILE_W, TILE_H);
+  PRINT_DEBUG ('u', "Press <SPACE> to pause/unpause, <ESC> to quit.\n");
+}
+
+// The image is a two-dimension array of size of DIM x DIM. Each pixel is of
+// type 'unsigned' and store the color information following a RGBA layout (4
+// bytes). Pixel at line 'l' and column 'c' in the current image can be accessed
+// using cur_img (l, c).
+
+static unsigned compute_color (int i, int j)
+{
+  return cur_img (i, j);
+}
+
+// The kernel returns 0, or the iteration step at which computation has
+// completed (e.g. stabilized).
+
+///////////////////////////// Simple sequential version (seq)
+// Suggested cmdline(s):
+// ./run --size 1024 --kernel <template> --variant seq
+// or
+// ./run -s 1024 -k <template> -v seq
+//
+unsigned <template>_compute_seq (unsigned nb_iter)
+{
+  for (unsigned it = 1; it <= nb_iter; it++) {
+
+    for (int i = 0; i < DIM; i++)
+      for (int j = 0; j < DIM; j++)
+        next_img (i, j) = compute_color (i, j);
+
+    swap_images ();
+  }
+
+  return 0;
+}
+
+
+///////////////////////////// Tiled sequential version (tiled)
+// Suggested cmdline(s):
+// ./run -k <template> -v tiled -g 16 -m
+// or
+// ./run -k <template> -v tiled -ts 64 -m
+//
+static void do_tile (int x, int y, int width, int height, int who)
+{
+  // Calling monitoring_{start|end}_tile before/after actual computation allows
+  // to monitor the execution in real time (--monitoring) and/or to generate an
+  // execution trace (--trace).
+  // monitoring_start_tile only needs the cpu number
+  monitoring_start_tile (who);
+
+  for (int i = y; i < y + height; i++)
+    for (int j = x; j < x + width; j++)
+      next_img (i, j) = compute_color (i, j);
+
+  // In addition to the cpu number, monitoring_end_tile also needs the tile
+  // coordinates
+  monitoring_end_tile (x, y, width, height, who);
+}
+
+unsigned <template>_compute_tiled (unsigned nb_iter)
+{
+  for (unsigned it = 1; it <= nb_iter; it++) {
+
+    for (int y = 0; y < DIM; y += TILE_H)
+      for (int x = 0; x < DIM; x += TILE_W)
+        do_tile (x, y, TILE_W, TILE_H, 0 /* CPU id */);
+
+    swap_images ();
+  }
+
+  return 0;
+}

include/api_funcs.h

+#ifndef API_FUNCS_IS_DEF
+#define API_FUNCS_IS_DEF
+
+typedef enum {
+    VEC_TYPE_CHAR,
+    VEC_TYPE_INT,
+    VEC_TYPE_FLOAT,
+    VEC_TYPE_DOUBLE
+} vec_type_t;
+
+typedef enum {
+    DIR_HORIZONTAL,
+    DIR_VERTICAL
+} direction_t;
+
+unsigned easypap_requested_number_of_threads (void);
+unsigned easypap_number_of_cores (void);
+unsigned easypap_number_of_gpus (void);
+int easypap_mpi_rank (void);
+int easypap_mpi_size (void);
+void easypap_check_mpi (void);
+void easypap_check_vectorization (vec_type_t vec_type, direction_t dir);
+int easypap_proc_is_master (void);
+
+
+#endif
\ No newline at end of file

include/arch_flags.h

@@ -5,13 +5,21 @@
 
 #if __AVX2__ == 1
 
-#define VEC_SIZE 8
-#define AVX2     1
+#define VEC_SIZE_CHAR   32
+#define VEC_SIZE_INT     8
+#define VEC_SIZE_FLOAT   8
+#define VEC_SIZE_DOUBLE  4
+
+#define AVX2 1
 
 #elif __SSE__ == 1
 
-#define VEC_SIZE 4
-#define SSE      1
+#define VEC_SIZE_CHAR   16
+#define VEC_SIZE_INT     4
+#define VEC_SIZE_FLOAT   4
+#define VEC_SIZE_DOUBLE  2
+
+#define SSE 1
 
 #endif
 

include/constants.h

@@ -3,6 +3,9 @@
 #define CONSTANTS_IS_DEF
 
 #define DEFAULT_DIM 1024
+#define DEFAULT_CPU_TILE_SIZE 32
+#define DEFAULT_GPU_TILE_SIZE 16
+
 #define DEFAULT_KERNEL "none"
 #define DEFAULT_VARIANT "seq"
 #define DEFAULT_OCL_VARIANT "ocl"

include/cpustat.h

@@ -15,7 +15,7 @@ void cpustat_init (int x, int y);
 void cpustat_reset (long now);
 void cpustat_start_work (long now, int who);
 long cpustat_finish_work (long now, int who);
-void cpustat_start_idle (long now, int who);
+void cpustat_deduct_idle (long duration, int who);
 void cpustat_freeze (long now);
 void cpustat_display_stats (void);
 void cpustat_clean (void);

include/debug.h

@@ -17,12 +17,13 @@
 //      'u' -- user
 //      'M' -- MPI
 
+#include "global.h"
+#include "api_funcs.h"
+
 #include <stdarg.h>
 #include <stdio.h>
 #include <stdlib.h>
 
-#include "global.h"
-
 void debug_init (char *flags);
 int debug_enabled (char flag);
 

include/easypap.h

@@ -3,6 +3,9 @@
 
 
 #include "global.h"
+#include "api_funcs.h"
+#include "img_data.h"
+#include "hooks.h"
 #include "arch_flags.h"
 #include "debug.h"
 #include "error.h"
@@ -16,4 +19,5 @@
 #include <mpi.h>
 #endif
 
+
 #endif

include/global.h

@@ -2,8 +2,21 @@
 #ifndef GLOBAL_IS_DEF
 #define GLOBAL_IS_DEF
 
-#include "hooks.h"
-#include "img_data.h"
+// Images are DIM * DIM arrays of pixels
+// Tiles have a size of CPU_TILE_H * CPU_TILE_W
+// An image contains CPU_NBTILES_Y * CPU_NBTILES_X
+
+extern unsigned DIM;
+
+extern unsigned TILE_W;
+extern unsigned TILE_H;
+extern unsigned NB_TILES_X;
+extern unsigned NB_TILES_Y;
+
+extern unsigned GPU_TILE_W;
+extern unsigned GPU_TILE_H;
+extern unsigned GPU_SIZE_X;
+extern unsigned GPU_SIZE_Y;
 
 extern unsigned do_display;
 extern unsigned vsync;
@@ -17,14 +30,6 @@ extern char *draw_param;
 extern unsigned opencl_used;
 extern unsigned easypap_mpirun;
 
-
 extern char *kernel_name, *variant_name;
 
-unsigned easypap_requested_number_of_threads (void);
-unsigned easypap_number_of_cores (void);
-int easypap_mpi_rank (void);
-int easypap_mpi_size (void);
-void easypap_check_mpi (void);
-int easypap_proc_is_master (void);
-
 #endif

include/graphics.h

@@ -4,11 +4,9 @@
 
 #ifdef ENABLE_SDL
 
-#include <SDL.h>
-
 #include "global.h"
 
-extern unsigned WIN_WIDTH, WIN_HEIGHT;
+#include <SDL.h>
 
 void graphics_init (void);
 void graphics_alloc_images (void);

include/hooks.h

@@ -5,15 +5,16 @@ typedef void (*void_func_t) (void);
 typedef unsigned (*int_func_t) (unsigned);
 typedef void (*draw_func_t) (char *);
 
-extern void_func_t the_first_touch;
+extern draw_func_t the_config;
 extern void_func_t the_init;
+extern void_func_t the_first_touch;
 extern draw_func_t the_draw;
 extern void_func_t the_finalize;
 extern int_func_t the_compute;
 extern void_func_t the_refresh_img;
 
 void *hooks_find_symbol (char *symbol);
-void hooks_establish_bindings (void);
+void hooks_establish_bindings (int silent);
 
 // Call function ${kernel}_draw_${suffix}, or default_func if symbol not found
 void hooks_draw_helper (char *suffix, void_func_t default_func);

include/img_data.h

 #ifndef IMG_DATA_IS_DEF
 #define IMG_DATA_IS_DEF
 
-#include <stdint.h>
+#include "global.h"
 
-extern unsigned DIM, GRAIN, TILE_SIZE;
+#include <stdint.h>
 
 extern uint32_t *restrict image, *restrict alt_image;
 
@@ -54,4 +54,8 @@ static inline uint32_t rgba (int r, int g, int b, int a)
   return (r << 24) | (g << 16) | (b << 8) | a;
 }
 
+unsigned hsv_to_rgb (float h, float s, float v);
+unsigned heat_to_rgb (float h); // 0.0 = cold, 1.0 = hot
+unsigned heat_to_3gauss_rgb (double v); // 0.0 = cold, 1.0 = hot
+
 #endif

include/monitoring.h

@@ -49,7 +49,19 @@
     if (do_gmonitor | do_trace) {                                              \
       long t = what_time_is_it ();                                             \
       gmonitor_end_tile (t, (c), (x), (y), (w), (h));                          \
-      trace_record_end_tile (t, (c), (x), (y), (w), (h));                      \
+      trace_record_end_tile (t, (c), (x), (y), (w), (h), TASK_TYPE_COMPUTE);   \
+    }                                                                          \
+  } while (0)
+
+#define monitoring_gpu_tile(x, y, w, h, c, s, e, tr)                           \
+  do {                                                                         \
+    if (do_gmonitor | do_trace) {                                              \
+      if ((tr) == TASK_TYPE_COMPUTE)                                           \
+        gmonitor_start_tile ((s), (c));                                        \
+      trace_record_start_tile ((s), (c));                                      \
+      if ((tr) == TASK_TYPE_COMPUTE)                                           \
+        gmonitor_end_tile ((e), (c), (x), (y), (w), (h));                      \
+      trace_record_end_tile ((e), (c), (x), (y), (w), (h), (tr));              \
     }                                                                          \
   } while (0)
 
@@ -87,11 +99,19 @@
     }                                                                          \
   } while (0)
 
-#define monitoring_end_tile(x, y, w, h, c)                                     \
+#define monitoring_end_tile(x, y, w, h, c, tr)                                 \
   do {                                                                         \
     if (do_trace) {                                                            \
       long t = what_time_is_it ();                                             \
-      trace_record_end_tile (t, (c), (x), (y), (w), (h));                      \
+      trace_record_end_tile (t, (c), (x), (y), (w), (h), TASK_TYPE_COMPUTE);   \
+    }                                                                          \
+  } while (0)
+
+#define monitoring_gpu_tile(x, y, w, h, c, s, e, tr)                           \
+  do {                                                                         \
+    if (do_trace) {                                                            \
+      trace_record_start_tile ((s), (c));                                      \
+      trace_record_end_tile ((e), (c), (x), (y), (w), (h), (tr));              \
     }                                                                          \
   } while (0)
 
@@ -103,6 +123,7 @@
 #define monitoring_end_iteration() (void)0
 #define monitoring_start_tile(c) (void)0
 #define monitoring_end_tile(x, y, w, h, c) (void)0
+#define monitoring_gpu_tile(x, y, w, h, c, s, e, tr) (void)0
 
 #endif
 

include/ocl.h

@@ -2,6 +2,7 @@
 #define OCL_IS_DEF
 
 #define CL_USE_DEPRECATED_OPENCL_1_2_APIS
+#define CL_TARGET_OPENCL_VERSION 120
 
 #ifdef __APPLE__
 #include <OpenCL/opencl.h>
@@ -17,15 +18,18 @@
 #endif
 
 #include "error.h"
+#include "monitoring.h"
 
-void ocl_init (int show_config_and_quit);
+void ocl_init (int show_config, int silent);
+void ocl_build_program (int list_variants);
 void ocl_alloc_buffers (void);
 void ocl_map_textures (GLuint texid);
-void ocl_send_image (unsigned *image);
-void ocl_retrieve_image (unsigned *image);
+void ocl_send_data (void);
+void ocl_retrieve_data (void);
 unsigned ocl_invoke_kernel_generic (unsigned nb_iter);
 void ocl_wait (void);
 void ocl_update_texture (void);
+unsigned easypap_number_of_gpus (void);
 size_t ocl_get_max_workgroup_size (void);
 
 #define check(err, format, ...)                                                \
@@ -34,10 +38,17 @@ size_t ocl_get_max_workgroup_size (void);
       exit_with_error (format " [OCL err %d]", ##__VA_ARGS__, err);            \
   } while (0)
 
-extern unsigned SIZE, TILE, TILEX, TILEY;
+// Kernels get executed by GPU_SIZE_Y * GPU_SIZE_X threads
+// Tiles have a size of GPU_TILE_H * GPU_TILE_W
+extern unsigned GPU_SIZE_X, GPU_SIZE_Y, GPU_TILE_W, GPU_TILE_H, GPU_TILE_W;
+
 extern cl_context context;
 extern cl_kernel compute_kernel;
 extern cl_command_queue queue;
 extern cl_mem cur_buffer, next_buffer;
+extern long _calibration_delta;
+
+long ocl_monitor (cl_event evt, int x, int y, int width, int height,
+                  task_type_t task_type);
 
 #endif

kernel/c/blur.c

@@ -5,7 +5,7 @@
 
 ///////////////////////////// Sequential version (tiled)
 // Suggested cmdline(s):
-// ./run -l images/1024.png -k blur -v seq
+// ./run -l images/1024.png -k blur -v seq -si
 //
 static void do_tile_reg (int x, int y, int width, int height)
 {
@@ -51,7 +51,7 @@ unsigned blur_compute_seq (unsigned nb_iter)
 
 ///////////////////////////// Tiled sequential version (tiled)
 // Suggested cmdline(s):
-// ./run -l images/1024.png -k blur -v tiled -g 16 -m
+// ./run -l images/1024.png -k blur -v tiled -ts 32 -m si
 //
 static inline void do_tile (int x, int y, int width, int height, int who)
 {
@@ -66,9 +66,9 @@ unsigned blur_compute_tiled (unsigned nb_iter)
 {
   for (unsigned it = 1; it <= nb_iter; it++) {
 
-    for (int y = 0; y < DIM; y += TILE_SIZE)
-      for (int x = 0; x < DIM; x += TILE_SIZE)
-        do_tile (x, y, TILE_SIZE, TILE_SIZE, 0);
+    for (int y = 0; y < DIM; y += TILE_H)
+      for (int x = 0; x < DIM; x += TILE_W)
+        do_tile (x, y, TILE_W, TILE_H, 0);
 
     swap_images ();
   }

kernel/c/invert.c

@@ -3,9 +3,11 @@
 
 #include <omp.h>
 
+#define INV_MASK  ((unsigned)0xFFFFFF00)
+
 static inline unsigned compute_color (int i, int j)
 {
-  return (unsigned)0xFFFFFF00 ^ cur_img (i, j);
+  return INV_MASK ^ cur_img (i, j);
 }
 
 ///////////////////////////// Simple sequential version (seq)
@@ -32,7 +34,6 @@ static void do_tile_reg (int x, int y, int width, int height)
       cur_img (i, j) = compute_color (i, j);
 }
 
-
 static void do_tile (int x, int y, int width, int height, int who)
 {
   monitoring_start_tile (who);
@@ -42,18 +43,17 @@ static void do_tile (int x, int y, int width, int height, int who)
   monitoring_end_tile (x, y, width, height, who);
 }
 
-
 ///////////////////////////// Tiled sequential version (tiled)
 // Suggested cmdline(s):
-// ./run -l images/shibuya.png -k invert -v tiled -g 32 -i 100 -n
+// ./run -l images/shibuya.png -k invert -v tiled -ts 32 -i 100 -n
 //
 unsigned invert_compute_tiled (unsigned nb_iter)
 {
   for (unsigned it = 1; it <= nb_iter; it++) {
 
-    for (int y = 0; y < DIM; y += TILE_SIZE)
-      for (int x = 0; x < DIM; x += TILE_SIZE)
-        do_tile (x, y, TILE_SIZE, TILE_SIZE, 0 /* CPU id */);
+    for (int y = 0; y < DIM; y += TILE_H)
+      for (int x = 0; x < DIM; x += TILE_W)
+        do_tile (x, y, TILE_W, TILE_H, 0 /* CPU id */);
 
   }
 

kernel/c/life.c

@@ -26,7 +26,6 @@ static inline cell_t *table_cell (cell_t *restrict i, int y, int x)
 
 void life_init (void)
 {
-
   // life_init may be (indirectly) called several times so we check if data were
   // already allocated
   if (_table == NULL) {
@@ -70,8 +69,8 @@ static inline void swap_tables (void)
 
 static int compute_new_state (int y, int x)
 {
-  unsigned n  = 0;
-  unsigned me = cur_table (y, x) != 0;
+  unsigned n      = 0;
+  unsigned me     = cur_table (y, x) != 0;
   unsigned change = 0;
 
   if (x > 0 && x < DIM - 1 && y > 0 && y < DIM - 1) {
@@ -146,9 +145,9 @@ unsigned life_compute_tiled (unsigned nb_iter)
   for (unsigned it = 1; it <= nb_iter; it++) {
     unsigned change = 0;
 
-    for (int y = 0; y < DIM; y += TILE_SIZE)
-      for (int x = 0; x < DIM; x += TILE_SIZE)
-        change |= do_tile (x, y, TILE_SIZE, TILE_SIZE, 0);
+    for (int y = 0; y < DIM; y += TILE_H)
+      for (int x = 0; x < DIM; x += TILE_W)
+        change |= do_tile (x, y, TILE_W, TILE_H, 0);
 
     swap_tables ();
 
@@ -161,18 +160,9 @@ unsigned life_compute_tiled (unsigned nb_iter)
   return res;
 }
 
-
 ///////////////////////////// Initial configs
 
-void life_draw_stable (void);
 void life_draw_guns (void);
-void life_draw_random (void);
-void life_draw_clown (void);
-void life_draw_diehard (void);
-void life_draw_bugs (void);
-void life_draw_otca_off (void);
-void life_draw_otca_on (void);
-void life_draw_meta3x3 (void);
 
 static inline void set_cell (int y, int x)
 {
@@ -200,7 +190,7 @@ static void inline life_rle_generate (char *filename, int x, int y, int width,
 
 void life_draw (char *param)
 {
-  if (access (param, R_OK) != -1) {
+  if (param && (access (param, R_OK) != -1)) {
     // The parameter is a filename, so we guess it's a RLE-encoded file
     life_rle_parse (param, 1, 1, RLE_ORIENTATION_NORMAL);
   } else
@@ -232,7 +222,7 @@ static void at_the_four_corners (char *filename, int distance)
                   RLE_ORIENTATION_HINVERT | RLE_ORIENTATION_VINVERT);
 }
 
-// Suggested cmdline: ./run -k life -s 2176 -a otca_off -ts 64 -r 10
+// Suggested cmdline: ./run -k life -s 2176 -a otca_off -ts 64 -r 10 -si
 void life_draw_otca_off (void)
 {
   if (DIM < 2176)
@@ -241,7 +231,7 @@ void life_draw_otca_off (void)
   otca_autoswitch ("data/rle/otca-off.rle", 1, 1);
 }
 
-// Suggested cmdline: ./run -k life -s 2176 -a otca_on -ts 64 -r 10
+// Suggested cmdline: ./run -k life -s 2176 -a otca_on -ts 64 -r 10 -si
 void life_draw_otca_on (void)
 {
   if (DIM < 2176)
@@ -250,7 +240,7 @@ void life_draw_otca_on (void)
   otca_autoswitch ("data/rle/otca-on.rle", 1, 1);
 }
 
-// Suggested cmdline: ./run -k life -s 6208 -a meta3x3 -ts 64 -r 50
+// Suggested cmdline: ./run -k life -s 6208 -a meta3x3 -ts 64 -r 50 -si
 void life_draw_meta3x3 (void)
 {
   if (DIM < 6208)
@@ -265,9 +255,27 @@ void life_draw_meta3x3 (void)
 // Suggested cmdline: ./run -k life -a bugs -ts 64
 void life_draw_bugs (void)
 {
-  for (int y = 0; y < DIM / 2; y += 32) {
-    life_rle_parse ("data/rle/tagalong.rle", y + 1, y + 8, RLE_ORIENTATION_NORMAL);
-    life_rle_parse ("data/rle/tagalong.rle", y + 1, (DIM - 32 - y) + 8, RLE_ORIENTATION_NORMAL);
+  for (int y = 16; y < DIM / 2; y += 32) {
+    life_rle_parse ("data/rle/tagalong.rle", y + 1, y + 8,
+                    RLE_ORIENTATION_NORMAL);
+    life_rle_parse ("data/rle/tagalong.rle", y + 1, (DIM - 32 - y) + 8,
+                    RLE_ORIENTATION_NORMAL);
+  }
+}
+
+// Suggested cmdline: ./run -k life -v omp -a ship -s 512 -m -ts 16
+void life_draw_ship (void)
+{