diff --git a/dumps/dump-life-seq-dim-2176-iter-1000.png b/dumps/dump-life-seq-dim-2176-iter-1000.png new file mode 100644 index 0000000..1fe2167 Binary files /dev/null and b/dumps/dump-life-seq-dim-2176-iter-1000.png differ diff --git a/dumps/dump-life-seq-dim-2176-iter-500.png b/dumps/dump-life-seq-dim-2176-iter-500.png new file mode 100644 index 0000000..490aec5 Binary files /dev/null and b/dumps/dump-life-seq-dim-2176-iter-500.png differ diff --git a/dumps/dump-life-seq-dim-2176-iter-5000.png b/dumps/dump-life-seq-dim-2176-iter-5000.png new file mode 100644 index 0000000..61b3ba8 Binary files /dev/null and b/dumps/dump-life-seq-dim-2176-iter-5000.png differ diff --git a/dumps/dump-life-seq-dim-2176-iter-50000.png b/dumps/dump-life-seq-dim-2176-iter-50000.png new file mode 100644 index 0000000..96ef163 Binary files /dev/null and b/dumps/dump-life-seq-dim-2176-iter-50000.png differ diff --git a/kernel/c/life.c b/kernel/c/life.c index eefabcd..92d6e9d 100644 --- a/kernel/c/life.c +++ b/kernel/c/life.c @@ -57,6 +57,19 @@ void life_refresh_img (void) cur_img (i, j) = cur_table (i, j) * color; } +// Only called when --dump or --thumbnails is used +void life_refresh_img_ocl(void) +{ + cl_int err; + + err = clEnqueueReadBuffer(queue, cur_buffer, CL_TRUE, 0, + sizeof(unsigned) * DIM * DIM, _table, 0, NULL, + NULL); + check(err, "Failed to read buffer from GPU"); + + life_refresh_img(); +} + static inline void swap_tables (void) { cell_t *tmp = _table; @@ -132,6 +145,38 @@ unsigned life_compute_tiled (unsigned nb_iter) return res; } +// omp tiled version +unsigned life_compute_omp_tiled (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) + { + temp = do_tile (x, y, TILE_W, TILE_H, omp_get_thread_num()); + + #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); diff --git a/kernel/ocl/life.cl b/kernel/ocl/life.cl new file mode 100644 index 0000000..c9aac40 --- /dev/null +++ b/kernel/ocl/life.cl @@ -0,0 +1,14 @@ +__kernel void life_ocl (__global unsigned *in, __global unsigned *out) +{ + +} + +// DO NOT MODIFY: this kernel updates the OpenGL texture buffer +// This is a life-specific version (generic version is defined in common.cl) +__kernel void life_update_texture (__global unsigned *cur, __write_only image2d_t tex) +{ + 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)); +} \ No newline at end of file diff --git a/validate_life b/validate_life new file mode 100755 index 0000000..4df35bf --- /dev/null +++ b/validate_life @@ -0,0 +1,31 @@ +#!/usr/bin/env bash + +if [ -z $1 ] +then + exit 1 +fi + +ITERATIONS="500 1000 5000" +SIZE=2176 +VERSION="octa_off" +LIFE_CONFIG=$1 + +for iter in $ITERATIONS +do + ./run -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" + + diff $GENERATED $TRUTH + equal=$? + + rm $GENERATED + + if [ $equal -ne 0 ] + then + exit 1 + fi +done + +exit 0