10 files changed, 315 insertions, 0 deletions
diff --git a/cuda-global/.DS_Store b/cuda-global/.DS_Store
new file mode 100644
index 0000000..96423a5
--- /dev/null
+++ b/cuda-global/.DS_Store
diff --git a/cuda-global/Makefile b/cuda-global/Makefile
new file mode 100644
index 0000000..fd802dc
--- /dev/null
+++ b/cuda-global/Makefile
@@ -0,0 +1,20 @@
+.DEFAULT_GOAL := all
+
+INCLUDES = -I include/
+
+game.o: include/game.cuh src/game.cu
+	nvcc -c src/game.cu $(INCLUDES) -o build/game.o
+
+file.o: game.o include/file.cuh src/file.cu
+	nvcc -c src/file.cu $(INCLUDES) -o build/file.o
+
+create_grid.o: file.o game.o include/create_grid.cuh src/create_grid.cu
+	nvcc -c src/create_grid.cu $(INCLUDES) -o build/create_grid.o
+
+gol: game.o file.o create_grid.o
+	nvcc $(INCLUDES) -o gol build/game.o build/file.o build/create_grid.o src/main.cu
+
+clean:
+	$(RM) build/* gol output/*
+
+all: gol
diff --git a/cuda-global/gol b/cuda-global/gol
new file mode 100755
index 0000000..3c94c50
--- /dev/null
+++ b/cuda-global/gol
diff --git a/cuda-global/include/create_grid.cuh b/cuda-global/include/create_grid.cuh
new file mode 100644
index 0000000..3402468
--- /dev/null
+++ b/cuda-global/include/create_grid.cuh
@@ -0,0 +1,12 @@
+#include <stdio.h>
+#include <stdlib.h>
+#include "file.cuh"
+#include "game.cuh"
+
+#ifndef CREATE_GRID_H
+#define CREATE_GRID_H
+
+void print_grid(struct GAME* game);
+void create_grid(int argc, char** argv);
+
+#endif // CREATE_GRID_H
diff --git a/cuda-global/include/file.cuh b/cuda-global/include/file.cuh
new file mode 100644
index 0000000..594eb4b
--- /dev/null
+++ b/cuda-global/include/file.cuh
@@ -0,0 +1,11 @@
+#include "game.cuh"
+#include <stdlib.h>
+#include <stdio.h>
+
+#ifndef FILE_H
+#define FILE_H
+
+void read_in(char* filename, struct GAME* game); 
+void write_out(char* filename, struct GAME* game); 
+
+#endif //FILE_H
diff --git a/cuda-global/include/game.cuh b/cuda-global/include/game.cuh
new file mode 100644
index 0000000..873a4cf
--- /dev/null
+++ b/cuda-global/include/game.cuh
@@ -0,0 +1,21 @@
+#include <stdlib.h>
+#include <stdio.h>
+#include <string.h>
+#include <omp.h>
+
+#ifndef GAME_H
+#define GAME_H
+
+
+struct GAME {
+  unsigned char* grid;
+  int padding;
+  int width;
+  int height;
+};
+
+__device__ int neighbors(struct GAME game, int x, int y);
+__global__ void next(struct GAME game, unsigned char* newGrid);
+void randomize(struct GAME* game);
+
+#endif // GAME_H
diff --git a/cuda-global/src/create_grid.cu b/cuda-global/src/create_grid.cu
new file mode 100644
index 0000000..fa208f2
--- /dev/null
+++ b/cuda-global/src/create_grid.cu
@@ -0,0 +1,38 @@
+#include "create_grid.cuh"
+
+void print_grid(struct GAME* game) {
+  printf("\n===GRID===\n");
+  for (int y = 0; y < game->height; y++) {
+    for (int x = 0; x < game->width; x++) {
+      printf("%i ", game->grid[y*(game->width+game->padding*2) + x]);
+    }
+    printf("\n");
+  }
+}
+
+void create_grid(int argc, char** argv) {
+  char* filename;
+  struct GAME game;
+  game.padding = 0;
+  if (argc == 5) {
+    game.width = atoi(argv[2]);
+    game.height = atoi(argv[3]);
+    filename = argv[4];
+  } else {
+    printf("Usage: ./gol create-grid <width> <height> <filename>\n");
+    exit(1);
+  }
+
+  int size = (game.width+game.padding*2) * (game.height+game.padding*2);
+  unsigned char* grid = (unsigned char*)malloc(sizeof(unsigned char) * size);
+  for (int y = 0; y < game.height; y++) {
+    printf("Row %i: ", y);
+    for (int x = 0; x < game.width; x++) {
+      char temp;
+      scanf("%i%c", (unsigned int*)&game.grid[y*(game.width+game.padding*2) + x],&temp);
+    }
+  }
+  game.grid = grid;
+  write_out(filename, &game);
+  print_grid(&game);
+}
diff --git a/cuda-global/src/file.cu b/cuda-global/src/file.cu
new file mode 100644
index 0000000..b1df5f9
--- /dev/null
+++ b/cuda-global/src/file.cu
@@ -0,0 +1,17 @@
+#include "file.cuh"
+
+void read_in(char* filename, struct GAME* game) {
+  FILE* file = fopen(filename, "rb");
+  for (int i = game->padding; i < game->height+game->padding; i++) {
+    fread(&game->grid[i*(game->width + 2*game->padding) + game->padding], sizeof(unsigned char), game->width, file);
+  }
+  fclose(file);
+}
+
+void write_out(char* filename, struct GAME* game) {
+  FILE* file = fopen(filename, "w+");
+  for (int i = game->padding; i < game->height+game->padding; i++) {
+    fwrite(&game->grid[i*(game->width + 2*game->padding) + game->padding], sizeof(unsigned char), game->width, file);
+  }
+  fclose(file);
+}
diff --git a/cuda-global/src/game.cu b/cuda-global/src/game.cu
new file mode 100644
index 0000000..9021916
--- /dev/null
+++ b/cuda-global/src/game.cu
@@ -0,0 +1,46 @@
+#include "game.cuh"
+
+__device__ int neighbors(struct GAME game, int x, int y) {
+  int n = 0;
+
+  for (int dy = -1; dy <= 1; dy++) {
+    for (int dx = -1; dx <= 1; dx++) {
+      if (!(dx == 0 && dy == 0) && (x+dx) > 0 && (y+dy) > 0 && (x+dx) < game.width+(game.padding*2) && (y+dy) < game.height+(game.padding*2)) {
+        if (game.grid[(y+dy) * (game.width+game.padding*2) + (x+dx)]) {
+          n++;
+        }
+      }
+    }
+  }
+  return n;
+}
+
+__global__ void next(struct GAME game, unsigned char* newGrid) {
+  int idy = blockDim.y * blockIdx.y + threadIdx.y;
+  int idx = blockDim.x * blockIdx.x + threadIdx.x;
+
+  if (idy <= game.height+game.padding*2 && idx <= game.width+game.padding*2) {
+    int my_neighbors = neighbors(game, idx, idy);
+    int my_coord = idy * (game.width+game.padding*2) + idx;
+    newGrid[my_coord] = 0; // It's possible that there are artifacts from the last iteration
+    if (game.grid[my_coord]) {
+      if (my_neighbors < 2 || my_neighbors > 3) {
+        newGrid[my_coord] = 0;
+      } else  {
+        newGrid[my_coord] = 1;
+      }
+    } else {
+      if (my_neighbors == 3) {
+        newGrid[my_coord] = 1;
+      }
+    }
+  }
+}
+
+void randomize(struct GAME* game) {
+  for (int y = game->padding; y < game->height+game->padding; y++) {
+    for (int x = game->padding; x < game->width+game->padding; x++) {
+      game->grid[y*(game->width+game->padding*2) + x] = (unsigned char) rand() & 1;
+    }
+  }
+}
diff --git a/cuda-global/src/main.cu b/cuda-global/src/main.cu
new file mode 100644
index 0000000..2b11fe1
--- /dev/null
+++ b/cuda-global/src/main.cu
@@ -0,0 +1,150 @@
+#include <stdlib.h>
+#include <time.h>
+#include <stdio.h>
+#include <string.h>
+#include <cstring>
+
+#include "file.cuh"
+#include "game.cuh"
+#include "create_grid.cuh"
+
+
+/*
+  Rules for life:
+  Any live cell with fewer than two live neighbors dies (underpopulation).
+  Any live cell with two or three live neighbors continues to live.
+  Any live cell with more than three live neighbors dies (overpopulation).
+  Any dead cell with exactly three live neighbors becomes a live cell (reproduction).
+ */
+#define BLOCK 32
+#define PADDING 10
+//#define VERBOSE 1
+#define SEED 100
+
+// gpuErrchk source: https://stackoverflow.com/questions/14038589/what-is-the-canonical-way-to-check-for-errors-using-the-cuda-runtime-api
+#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
+inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort =
+true) {
+  if (code != cudaSuccess) {
+    fprintf(stderr, "GPUassert: %s %s %d\n", cudaGetErrorString(code), file,
+            line);
+    if (abort)
+      exit(code);
+  }
+}
+
+void simulate(int argc, char** argv) {
+  srand(SEED);
+  clock_t totalStart = clock();
+  char* filename;
+  struct GAME game;
+  game.padding = PADDING;
+  int iterations, log_each_step, block_size;
+  if (argc == 8) {
+    filename = argv[2];
+    game.width = atoi(argv[3]);
+    game.height = atoi(argv[4]);
+    iterations = atoi(argv[5]);
+    log_each_step = atoi(argv[6]);
+    block_size = atoi(argv[7]);
+  } else {
+    printf("Usage: ./gol simulate <filename | random> <width> <height> <iterations> <log-each-step?1:0> <block-size>\n");
+    filename = "random";
+    game.height = 10;
+    game.width = 10;
+    iterations = 5;
+    log_each_step = 0;
+  }
+
+  // Allocate space for current grid (1 byte per tile)
+  int size = (game.height+(2*game.padding)) * (game.width+(2*game.padding)) * sizeof(unsigned char);
+  game.grid = (unsigned char*)malloc(size);
+  memset(game.grid, 0, size);
+
+  if (strcmp(filename, "random") == 0) {
+    randomize(&game);
+  } else {
+    read_in(filename, &game);
+  }  
+
+  char iteration_file[1024];
+
+  unsigned char* grid_d;
+  unsigned char* newGrid_d;
+  gpuErrchk(cudaMalloc(&grid_d, size));
+  gpuErrchk(cudaMemcpy(grid_d, game.grid, size, cudaMemcpyHostToDevice));
+  gpuErrchk(cudaMalloc(&newGrid_d, size));
+
+  unsigned char* grid_h = (unsigned char*)malloc(size);
+  unsigned char* temp;
+
+  game.grid = grid_d;
+
+  int grid_num = (int)ceil((game.width+(2*game.padding))/(float)block_size);
+  dim3 dim_grid(grid_num, grid_num, 1);
+  dim3 dim_block(block_size, block_size, 1);
+
+  cudaEvent_t startLife, stopLife;
+  cudaEventCreate(&startLife);
+  cudaEventCreate(&stopLife);
+  double timeComputingLife = 0;
+  float localTime = 0;
+
+  for (int i = 0; i <= iterations; i++) {
+    if (i > 0) {
+      cudaEventRecord(startLife);
+      next<<<dim_grid, dim_block>>>(game, newGrid_d);
+      cudaEventRecord(stopLife);
+      cudaEventSynchronize(stopLife);
+      cudaEventElapsedTime(&localTime, startLife, stopLife);
+      timeComputingLife += localTime/1000;
+
+      temp = game.grid;
+      game.grid = newGrid_d;
+      newGrid_d = temp;
+    }
+    if (log_each_step) {
+      gpuErrchk(cudaMemcpy(grid_h, game.grid, size, cudaMemcpyDeviceToHost));
+      #ifdef VERBOSE
+        printf("\n===Iteration %i===\n", i);
+        for (int y = game.padding; y < game.height+game.padding; y++) {
+          for (int x = game.padding; x < game.width+game.padding; x++) {
+            printf("%s ", grid_h[y*(game.width+2*game.padding) + x] ? "X" : " ");
+          }
+          printf("\n");
+        }
+        printf("===End iteration %i===\n", i);
+      #endif
+      sprintf(iteration_file, "output/iteration-%07d.bin", i);
+      temp = game.grid;
+      game.grid = grid_h;
+      write_out(iteration_file, &game);
+      game.grid = temp;
+    }
+  }
+
+  clock_t totalEnd = clock();
+  printf("\n===Timing===\nTime computing life: %f\nClock time: %f\n", timeComputingLife, ((double)totalEnd - (double)totalStart)/CLOCKS_PER_SEC);
+
+  cudaFree(&newGrid_d);
+  cudaFree(&grid_d);
+  cudaFree(&game.grid);
+  free(grid_h);
+}
+
+int main(int argc, char** argv) {
+  if (argc >= 2) {
+    if (strcmp(argv[1], "simulate") == 0) {
+      simulate(argc, argv);
+    } else if (strcmp(argv[1], "create-grid") == 0) {
+      create_grid(argc, argv);
+    } else {
+      printf("Unknown input: %s\n", argv[1]);
+      exit(1);
+    }
+  } else {
+    printf("Usage: ./gol <simulate | create-grid>\n");
+    exit(1);
+  }
+  return 0;
+}