// ================================================================================================
// Tim Backus
// CIS 450 - High Performance Computing
// 3D Game of Life - CUDA Version
// ================================================================================================

#define GOL_IO_FILENAME "gol3DOutput.dat"
#define GOL_CUDA_THREADS_SIZE 8

#include <stdio.h>
#include <string.h>
#include <ctype.h>
#include <time.h>
#include <stdlib.h>

// ------------------------------------------------------------------------------------------------
// CUDA kernel (Gather & Map) - Adds up the number of neighbors for a cell in a 3x3x3 cube and
//   sets each cell to alive or dead depending on its number of neighbors and the rules for this 
//   current game.
// ------------------------------------------------------------------------------------------------
__global__
void lifeItrKernel(const char* const d_in, char* d_out, const unsigned int xsize,
                   const unsigned int ysize, const unsigned int zsize, const unsigned int alow,
                   const unsigned int ahigh) {
  
  extern __shared__ char shMem[];
               
  // Calculate block and thread IDs
  const int threadPosX = blockIdx.x * blockDim.x + threadIdx.x;
  const int threadPosY = blockIdx.y * blockDim.y + threadIdx.y;
  const int threadPosZ = blockIdx.z * blockDim.z + threadIdx.z;
  const unsigned int stepX = ysize * zsize;
  const unsigned int arrayPos = threadPosX * stepX + threadPosY * zsize + threadPosZ;
  const unsigned int threadID = threadIdx.x * blockDim.y * blockDim.z + 
                                threadIdx.y * blockDim.z + threadIdx.z;

  // Ensure thread bounds
  if(threadPosX >= xsize) return;
  if(threadPosY >= ysize) return;
  if(threadPosZ >= zsize) return;
  
  // Copy global into shared memory
  shMem[threadID] = d_in[arrayPos];
  __syncthreads();
  
  // Begin adding neighbors
  char sum = 0;
  
  // X-Axis neighbors
  int xc, xcoord;
  for(xc = threadPosX - 1; xc <= threadPosX + 1; xc++) {
    
    // Wrap X-Axis
    xcoord = xc;
    if(xc < 0) xcoord = xsize;
    else if(xc >= xsize) xcoord = 0;
    
    // Y-Axis neighbors
    int yc, ycoord;
    for(yc = threadPosY - 1; yc <= threadPosY + 1; yc++) {
      
      // Wrap Y-Axis
      ycoord = yc;
      if(yc < 0) ycoord = ysize;
      else if(yc >= ysize) ycoord = 0;
      
      // Z-Axis neighbors
      int zc, zcoord;
      for(zc = threadPosZ - 1; zc <= threadPosZ + 1; zc++) {
        
        // Wrap Z-Axis
        zcoord = zc;
        if(zc < 0) zcoord = zsize;
        else if(zc >= zsize) zcoord = 0;
        
        // Don't count the cell itself
        if(threadPosX != xcoord || threadPosY != ycoord || threadPosZ != zcoord) {
          
          // Use shared memory instead of global memory if the current coord is in the thread block
          if((xcoord >= blockDim.x * blockIdx.x && xcoord < (blockDim.x) * blockIdx.x + 1) &&
             (ycoord >= blockDim.y * blockIdx.y && ycoord < (blockDim.y) * blockIdx.y + 1) &&
             (zcoord >= blockDim.z * blockIdx.z && zcoord < (blockDim.z) * blockIdx.z + 1)) {
            sum += shMem[(xcoord % blockDim.x) * blockDim.y * blockDim.z + (ycoord % blockDim.y) * 
                          blockDim.z + (zcoord % blockDim.z)];
          } else {
            sum += d_in[xcoord * stepX + ycoord * zsize + zcoord];
          }
        }
      }
    }
  }
  
  // Set the cell's dead or alive status based on its neighbor count
  if (sum < alow || sum > ahigh) {
    d_out[arrayPos] = 0;
  } else if (sum >= alow && sum <= ahigh) {
    d_out[arrayPos] = 1;
  }
  
}


// ------------------------------------------------------------------------------------------------
// CUDA kernel (Gather) - Adds up the number of neighbors for a cell in a 3x3x3 cube.
// ------------------------------------------------------------------------------------------------
__global__
void sumNeighborsKernel(const char* const d_in, char* d_out, const unsigned int xsize,
                        const unsigned int ysize, const unsigned int zsize) {
  
  
  // Calculate block and thread IDs
  const int threadPosX = blockIdx.x * blockDim.x + threadIdx.x;
  const int threadPosY = blockIdx.y * blockDim.y + threadIdx.y;
  const int threadPosZ = blockIdx.z * blockDim.z + threadIdx.z;
  const unsigned int stepX = ysize * zsize;
  const unsigned int arrayPos = threadPosX * stepX + threadPosY * zsize + threadPosZ;
  
  // printf("TID=%d,%d,%d\n", threadIdx.x, threadIdx.y, threadIdx.z);
  // printf("TPOS=%d,%d,%d\n", threadPosX, threadPosY, threadPosZ);
  // printf("APOS=%d\n", arrayPos);
  
  // Ensure thread bounds
  if(threadPosX > xsize - 1) return;
  if(threadPosY > ysize - 1) return;
  if(threadPosZ > zsize - 1) return;
  
  char sum = 0;
  
  // X-Axis neighbors
  int xc, xcoord;
  for(xc = threadPosX - 1; xc <= threadPosX + 1; xc++) {
    
    // Wrap X-Axis
    xcoord = xc;
    if(xc < 0) xcoord = xsize;
    else if(xc >= xsize) xcoord = 0;
    
    // Y-Axis neighbors
    int yc, ycoord;
    for(yc = threadPosY - 1; yc <= threadPosY + 1; yc++) {
      
      // Wrap Y-Axis
      ycoord = yc;
      if(yc < 0) ycoord = ysize;
      else if(yc >= ysize) ycoord = 0;
      
      // Z-Axis neighbors
      int zc, zcoord;
      for(zc = threadPosZ - 1; zc <= threadPosZ + 1; zc++) {
        
        // Wrap Z-Axis
        zcoord = zc;
        if(zc < 0) zcoord = zsize;
        else if(zc >= zsize) zcoord = 0;
        
        // Don't count the cell itself
        if(threadPosX != xcoord || threadPosY != ycoord || threadPosZ != zcoord) {
          sum += d_in[xcoord * stepX + ycoord * zsize + zcoord];
        }
        
      }
    }
  }
  
  d_out[arrayPos] = sum;
}

// ------------------------------------------------------------------------------------------------
// CUDA kernel (Map) - Sets each cell to alive or dead depending on its number of neighbors and
//   the rules for this current game.
// ------------------------------------------------------------------------------------------------
__global__
void setAliveDeadKernel(const char* const d_nei, char* d_out, const unsigned int xs, 
                        const unsigned int ys, const unsigned int zs, const unsigned int alow, 
                        const unsigned int ahigh) {
  
  // Calculate block and thread IDs
  const int threadPosX = blockIdx.x * blockDim.x + threadIdx.x;
  const int threadPosY = blockIdx.y * blockDim.y + threadIdx.y;
  const int threadPosZ = blockIdx.z * blockDim.z + threadIdx.z;
  const int stepX = ys * zs;
  const int arrayPos = threadPosX * stepX + threadPosY * zs + threadPosZ;

  // Ensure thread bounds
  if(threadPosX > xs - 1) return;
  if(threadPosY > ys - 1) return;
  if(threadPosZ > zs - 1) return;
  
  // Set the cell alive or dead as according to the rules
  if (d_nei[arrayPos] < alow || d_nei[arrayPos] > ahigh) {
    d_out[arrayPos] = 0;
  } else if (d_nei[arrayPos] >= alow && d_nei[arrayPos] <= ahigh) {
    d_out[arrayPos] = 1;
  }
}

// ------------------------------------------------------------------------------------------------
// Returns the 1D position of a simulated 3D array
// ------------------------------------------------------------------------------------------------
int getArrIndex(const unsigned int xp, const unsigned int yp, const unsigned int zp,
                const unsigned int ys, const unsigned int zs) {
  return xp * ys * zs + yp * zs + zp;
}

// ------------------------------------------------------------------------------------------------
// Prints a 3D array.
// ------------------------------------------------------------------------------------------------
void print3DArray(char* arr, unsigned const int x, unsigned const int y, unsigned const int z) {
  int i;
  for(i = 0; i < x; ++i) {
    printf("Dimension %d:\n", i);
    int j;
    for(j = 0; j < y; ++j) {
      int k;
      for(k = 0; k < z; ++k) {
        printf("%d ", (char)arr[getArrIndex(i, j, k, y, z)]);
      }
      printf("\n");
    }
    printf("\n");
  }
}

// ------------------------------------------------------------------------------------------------
// Writes cells to alive or dead, randomly.
// ------------------------------------------------------------------------------------------------
void randomizeGrid(char* grid, unsigned const int size, unsigned const int chance) {
 
  srand(time(NULL));
 
  int i;
  for(i = 0; i < size; i++) {
    grid[i] = (char)((rand() % 100 <= chance) ? 1 : 0);
  }
  
}

// ------------------------------------------------------------------------------------------------
// Initializes the game data file.
// Line 1: <iteration count> <x-size> <y-size> <z-size>
// Line 2: Blank
// ------------------------------------------------------------------------------------------------
void initGameFile(const unsigned int itrs, const unsigned int x, const unsigned int y,
                  const unsigned int z) {
  FILE *fp;
  fp = fopen(GOL_IO_FILENAME, "w+");
  fprintf(fp, "%d %d %d %d\n\n", itrs, x, y, z);
  fclose(fp);
}

// ------------------------------------------------------------------------------------------------
// Writes a game to a file for visualization within Java.
// For every iteration, a block of text is created of the format:
//   "<x-coord>:<z-coords for y=0>, <z-coords for y=1>, ..."
//   Z-coords are represented by a 0 or 1 for each z-coordinate
// Example: Game with 5 iterations, x=3, y=7, z=4
//   5 3 7 4
//   
//   0:0000,0000,0000,0000,0000,0000,0000
//   1:0000,0000,0000,0100,0000,0000,0001
//   2:0000,0000,0010,0100,0001,0011,0000
//   0:0000,0000,0000,0000,0000,0000,0000
//   1:0000,0000,0000,0100,0000,0000,0001
//   2:0000,0000,0010,0100,0001,0011,0000
//   0:0000,0000,0000,0000,0000,0000,0000
//   1:0000,0000,0000,0100,0000,0000,0001
//   2:0000,0000,0010,0100,0001,0011,0000
//   0:0000,0000,0000,0000,0000,0000,0000
//   1:0000,0000,0000,0100,0000,0000,0001
//   2:0000,0000,0010,0100,0001,0011,0000
//   0:0000,0000,0000,0000,0000,0000,0000
//   1:0000,0000,0000,0100,0000,0000,0001
//   2:0000,0000,0010,0100,0001,0011,0000
// 
// ------------------------------------------------------------------------------------------------
void writeGameStep(char* arr, unsigned const int x, unsigned const int y, unsigned const int z) {
  FILE *fp;
  fp = fopen(GOL_IO_FILENAME, "a");
  
  int i;
  for(i = 0; i < x; i++) {
    fprintf(fp, "%d:", i);
    int j;
    for(j = 0; j < y; j++) {
      
      if(j > 0) {
        fprintf(fp, ",");
      }
      // Print Z-Dim values
      int k;
      for(k = 0; k < z; k++) {
        fprintf(fp, "%d", arr[getArrIndex(i, j, k, y, z)]);
      }
    }
    fprintf(fp, "\n");
  }
  fclose(fp);
}

// ------------------------------------------------------------------------------------------------
// Runs the Game of Life.
// ------------------------------------------------------------------------------------------------
void runLife(const unsigned int iterations, unsigned int xsize, const unsigned int ysize, 
             const unsigned int zsize, const unsigned int initc, const unsigned int alow,
             const unsigned int ahigh, const unsigned int printArr, const unsigned int writeOut) {
  
  // Memory values
  const unsigned int arrSize = xsize * ysize * zsize;
  const unsigned int arrMem = arrSize * sizeof(char);
  
  // GPU grid dimensions
  const int gx = ceil((double) xsize / GOL_CUDA_THREADS_SIZE);
  const int gy = ceil((double) ysize / GOL_CUDA_THREADS_SIZE);
  const int gz = ceil((double) zsize / GOL_CUDA_THREADS_SIZE);
  printf("Grid dimension: %d,%d,%d\n", gx, gy, gz);
  dim3 gridDim(gx, gy, gz);
  
  // GPU thread dimensions
  const int tx = GOL_CUDA_THREADS_SIZE;
  const int ty = GOL_CUDA_THREADS_SIZE;
  const int tz = GOL_CUDA_THREADS_SIZE;
  printf("Block dimension: %d,%d,%d\n", tx, ty, tz);
  dim3 blockDim(tx, ty, tz);
  
  // Initialize game space
  char *h_in = (char *) malloc(arrMem);
  printf("Randomizing initial game (could take a while)...\n");
  randomizeGrid(h_in, arrSize, initc);
  
  // Print the initial array if enabled
  if(printArr) {
    printf("Initial grid:\n");
    print3DArray(h_in, xsize, ysize, zsize);
  }
  
  // Initialize the output file if enabled
  if(writeOut) {
    initGameFile(iterations, xsize, ysize, zsize);
    writeGameStep(h_in, xsize, ysize, zsize);
  }
  
  // Pointers for GPU game data
  char *d_in;
  char *d_out;
  
  // Allocate input array on GPU
  printf("Allocating %d bytes of memory on the GPU...\n", 
         (int)(xsize * ysize * zsize * sizeof(char)));
  cudaMalloc(&d_in, arrMem);
  
  // Allocate output array on GPU
  cudaMalloc(&d_out, arrMem);
  
  // Copy the host data to the GPU
  cudaMemcpy(d_in, h_in, arrMem, cudaMemcpyHostToDevice);
  
  // Do Game of Life iterations
  int itrNum;
  for(itrNum = 0; itrNum < iterations; itrNum++) {
    
    printf("Iteration %d ", itrNum);
    
    // Run the kernel to simulate an iteration of 3D life
    clock_t start = clock();
    lifeItrKernel<<<gridDim, blockDim, (tx * ty * tz * sizeof(char))>>>(d_in, d_out, xsize, ysize, zsize, alow, ahigh);
    cudaError_t cerr = cudaDeviceSynchronize();
    if(cerr != cudaSuccess) {
      printf("Kernel lifeItr failed with error \"%s\".\n", cudaGetErrorString(cerr));
    }
    clock_t end = clock();
    
    // Copy the memory back to the input array for the next iteration
    cudaMemcpy(d_in, d_out, arrMem, cudaMemcpyDeviceToDevice);
    
    printf("took %d ticks.\n", (end - start));
    
    // Print and write out if enabled
    if(printArr || writeOut) {
      cudaMemcpy(h_in, d_out, arrMem, cudaMemcpyDeviceToHost);
      if(printArr) {
        print3DArray(h_in, xsize, ysize, zsize);
      }
      if(writeOut) {
        writeGameStep(h_in, xsize, ysize, zsize);
      }
    }
  }
  
  // Free memory
  cudaFree(d_in);
  cudaFree(d_out);
  free(h_in);
}

// ------------------------------------------------------------------------------------------------
// Prints the usage message if a bad number of runtime arguments are passed.
// ------------------------------------------------------------------------------------------------
void printUsage() {
  printf("Arguments (separated by spaces):\n");
  printf("  MAX_ITERATIONS\n  SIZE_X\n  SIZE_Y\n  SIZE_Z\n  INITIAL_ALIVE_CHANCE (1-100)\n");
  printf("  ALIVE_THRESHOLD_LOW (inclusive)\n  ALIVE_THRESHOLD_HIGH (inclusive)\n");
  printf("  PRINT_ARRAY? (0=no, 1=yes)\n  WRITE_TO_FILE? (0=no, 1=yes)\n");
}

// ------------------------------------------------------------------------------------------------
// Main Method
// ------------------------------------------------------------------------------------------------
int main(int argc, char *argv[]) {
  
  // Ensure proper runtime argument count
  if(argc != 10) {
    printUsage();
    return EXIT_SUCCESS;
  }
  
  // Parse iteration count
  unsigned const int iterations = atoi(argv[1]);
  
  // Parse X-Size
  unsigned const int sizeX = atoi(argv[2]);
  
  // Parse Y-Size
  unsigned const int sizeY = atoi(argv[3]);
  
  // Parse Z-Size
  unsigned const int sizeZ = atoi(argv[4]);
  
  // Parse initial alive chance
  unsigned const int initChance = atoi(argv[5]);
  
  // Parse alive low threshold (inclusive)
  unsigned const int aliveLow = atoi(argv[6]);
  
  // Parse alive high threshold (inclusive)
  unsigned const int aliveHigh = atoi(argv[7]);
  
  // Parse whether or not to print the array
  unsigned const int printArray = atoi(argv[8]);
  
  // Parse whether or not to output to disk
  unsigned const int writeOut = atoi(argv[9]);
  
  // Print game information to the console
  printf("Starting %d iteration Game of Life (CUDA) with sizes x=%d, y=%d, z=%d\n", iterations,
         sizeX, sizeY, sizeZ);
  printf("  initial alive chance=%d, neighbors for alive=%d to %d\n", initChance, 
         aliveLow, aliveHigh);
  if(writeOut) {
    printf("  File output enabled.\n");
  }
  runLife(iterations, sizeX, sizeY, sizeZ, initChance, aliveLow, aliveHigh, printArray, writeOut);
  
  return EXIT_SUCCESS;
}
	.file	"tmpxft_002538cc_00000000-6_Cuda3DLife.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2037:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2037:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z11getArrIndexjjjjj
	.type	_Z11getArrIndexjjjjj, @function
_Z11getArrIndexjjjjj:
.LFB2027:
	.cfi_startproc
	endbr64
	imull	%ecx, %edi
	addl	%esi, %edi
	imull	%r8d, %edi
	leal	(%rdi,%rdx), %eax
	ret
	.cfi_endproc
.LFE2027:
	.size	_Z11getArrIndexjjjjj, .-_Z11getArrIndexjjjjj
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"Dimension %d:\n"
.LC1:
	.string	"%d "
.LC2:
	.string	"\n"
	.text
	.globl	_Z12print3DArrayPcjjj
	.type	_Z12print3DArrayPcjjj, @function
_Z12print3DArrayPcjjj:
.LFB2028:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movl	%edx, %eax
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	imull	%ecx, %eax
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	xorl	%r12d, %r12d
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	xorl	%ebp, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movl	%ecx, %ebx
	subq	$56, %rsp
	.cfi_def_cfa_offset 112
	movl	%eax, 40(%rsp)
	movl	%esi, %eax
	movq	%rdi, 32(%rsp)
	movl	%edx, 12(%rsp)
	movq	%rax, 24(%rsp)
.L4:
	movq	24(%rsp), %rax
	movl	%ebp, %edx
	cmpq	%rax, %rbp
	je	.L12
	leaq	.LC0(%rip), %rsi
	xorl	%eax, %eax
	leal	(%rbx,%r12), %r15d
	movl	%r12d, %r14d
	movl	$2, %edi
	xorl	%r13d, %r13d
	call	__printf_chk@PLT
	movl	12(%rsp), %eax
	movq	%rax, 16(%rsp)
.L5:
	movq	16(%rsp), %rax
	cmpq	%rax, %r13
	je	.L13
	movl	%r14d, %ecx
.L7:
	cmpl	%r15d, %ecx
	je	.L14
	movq	32(%rsp), %rdi
	movslq	%ecx, %rax
	leaq	.LC1(%rip), %rsi
	movl	%ecx, 44(%rsp)
	movsbl	(%rdi,%rax), %edx
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	44(%rsp), %ecx
	incl	%ecx
	jmp	.L7
.L14:
	leaq	.LC2(%rip), %rsi
	xorl	%eax, %eax
	incq	%r13
	addl	%ebx, %r14d
	movl	$2, %edi
	addl	%ebx, %r15d
	call	__printf_chk@PLT
	jmp	.L5
.L13:
	leaq	.LC2(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	incq	%rbp
	call	__printf_chk@PLT
	movl	40(%rsp), %eax
	addl	%eax, %r12d
	jmp	.L4
.L12:
	addq	$56, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2028:
	.size	_Z12print3DArrayPcjjj, .-_Z12print3DArrayPcjjj
	.globl	_Z13randomizeGridPcjj
	.type	_Z13randomizeGridPcjj, @function
_Z13randomizeGridPcjj:
.LFB2029:
	.cfi_startproc
	endbr64
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	movl	$100, %r13d
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	movl	%edx, %r12d
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	movl	%esi, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	movq	%rdi, %rbx
	xorl	%edi, %edi
	pushq	%rcx
	.cfi_def_cfa_offset 48
	addq	%rbx, %rbp
	call	time@PLT
	movl	%eax, %edi
	call	srand@PLT
.L16:
	cmpq	%rbp, %rbx
	je	.L19
	call	rand@PLT
	cltd
	idivl	%r13d
	cmpl	%edx, %r12d
	setnb	(%rbx)
	incq	%rbx
	jmp	.L16
.L19:
	popq	%rax
	.cfi_def_cfa_offset 40
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%rbp
	.cfi_def_cfa_offset 24
	popq	%r12
	.cfi_def_cfa_offset 16
	popq	%r13
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2029:
	.size	_Z13randomizeGridPcjj, .-_Z13randomizeGridPcjj
	.section	.rodata.str1.1
.LC3:
	.string	"w+"
.LC4:
	.string	"gol3DOutput.dat"
.LC5:
	.string	"%d %d %d %d\n\n"
	.text
	.globl	_Z12initGameFilejjjj
	.type	_Z12initGameFilejjjj, @function
_Z12initGameFilejjjj:
.LFB2030:
	.cfi_startproc
	endbr64
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	movl	%ecx, %r12d
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	movl	%edi, %ebp
	leaq	.LC4(%rip), %rdi
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	subq	$16, %rsp
	.cfi_def_cfa_offset 48
	movl	%esi, 12(%rsp)
	leaq	.LC3(%rip), %rsi
	movl	%edx, 8(%rsp)
	call	fopen@PLT
	movl	%ebp, %ecx
	movl	$2, %esi
	pushq	%rdx
	.cfi_def_cfa_offset 56
	movq	%rax, %rbx
	movq	%rax, %rdi
	leaq	.LC5(%rip), %rdx
	pushq	%r12
	.cfi_def_cfa_offset 64
	movl	24(%rsp), %r9d
	xorl	%eax, %eax
	movl	28(%rsp), %r8d
	call	__fprintf_chk@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 32
	movq	%rbx, %rdi
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	jmp	fclose@PLT
	.cfi_endproc
.LFE2030:
	.size	_Z12initGameFilejjjj, .-_Z12initGameFilejjjj
	.section	.rodata.str1.1
.LC6:
	.string	"a"
.LC7:
	.string	"%d:"
.LC8:
	.string	","
.LC9:
	.string	"%d"
	.text
	.globl	_Z13writeGameStepPcjjj
	.type	_Z13writeGameStepPcjjj, @function
_Z13writeGameStepPcjjj:
.LFB2031:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	xorl	%r12d, %r12d
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movl	%esi, %ebp
	leaq	.LC6(%rip), %rsi
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movl	%ecx, %ebx
	subq	$56, %rsp
	.cfi_def_cfa_offset 112
	movq	%rdi, 32(%rsp)
	leaq	.LC4(%rip), %rdi
	movl	%edx, 12(%rsp)
	call	fopen@PLT
	movq	%rax, (%rsp)
	movl	12(%rsp), %eax
	imull	%ebx, %eax
	movl	%eax, 40(%rsp)
	movl	%ebp, %eax
	xorl	%ebp, %ebp
	movq	%rax, 24(%rsp)
.L23:
	movq	24(%rsp), %rax
	movl	%ebp, %ecx
	cmpq	%rax, %rbp
	je	.L34
	movq	(%rsp), %rdi
	leaq	.LC7(%rip), %rdx
	xorl	%eax, %eax
	movl	%r12d, %r14d
	movl	$2, %esi
	leal	(%rbx,%r12), %r15d
	xorl	%r13d, %r13d
	call	__fprintf_chk@PLT
	movl	12(%rsp), %eax
	movq	%rax, 16(%rsp)
.L24:
	movq	16(%rsp), %rax
	cmpq	%rax, %r13
	je	.L35
	testq	%r13, %r13
	je	.L25
	movq	(%rsp), %rdi
	leaq	.LC8(%rip), %rdx
	movl	$2, %esi
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
.L25:
	movl	%r14d, %r8d
.L26:
	cmpl	%r15d, %r8d
	je	.L36
	movq	32(%rsp), %rdx
	movslq	%r8d, %rax
	movq	(%rsp), %rdi
	movl	$2, %esi
	movl	%r8d, 44(%rsp)
	movsbl	(%rdx,%rax), %ecx
	leaq	.LC9(%rip), %rdx
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
	movl	44(%rsp), %r8d
	incl	%r8d
	jmp	.L26
.L36:
	incq	%r13
	addl	%ebx, %r14d
	addl	%ebx, %r15d
	jmp	.L24
.L35:
	movq	(%rsp), %rdi
	leaq	.LC2(%rip), %rdx
	xorl	%eax, %eax
	incq	%rbp
	movl	$2, %esi
	call	__fprintf_chk@PLT
	movl	40(%rsp), %eax
	addl	%eax, %r12d
	jmp	.L23
.L34:
	movq	(%rsp), %rdi
	addq	$56, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	jmp	fclose@PLT
	.cfi_endproc
.LFE2031:
	.size	_Z13writeGameStepPcjjj, .-_Z13writeGameStepPcjjj
	.section	.rodata.str1.1
.LC10:
	.string	"Arguments (separated by spaces):\n"
.LC11:
	.string	"  MAX_ITERATIONS\n  SIZE_X\n  SIZE_Y\n  SIZE_Z\n  INITIAL_ALIVE_CHANCE (1-100)\n"
.LC12:
	.string	"  ALIVE_THRESHOLD_LOW (inclusive)\n  ALIVE_THRESHOLD_HIGH (inclusive)\n"
.LC13:
	.string	"  PRINT_ARRAY? (0=no, 1=yes)\n  WRITE_TO_FILE? (0=no, 1=yes)\n"
	.text
	.globl	_Z10printUsagev
	.type	_Z10printUsagev, @function
_Z10printUsagev:
.LFB2033:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	leaq	.LC10(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	leaq	.LC11(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	leaq	.LC12(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	leaq	.LC13(%rip), %rsi
	xorl	%eax, %eax
	popq	%rdx
	.cfi_def_cfa_offset 8
	movl	$2, %edi
	jmp	__printf_chk@PLT
	.cfi_endproc
.LFE2033:
	.size	_Z10printUsagev, .-_Z10printUsagev
	.globl	_Z41__device_stub__Z13lifeItrKernelPKcPcjjjjjPKcPcjjjjj
	.type	_Z41__device_stub__Z13lifeItrKernelPKcPcjjjjjPKcPcjjjjj, @function
_Z41__device_stub__Z13lifeItrKernelPKcPcjjjjjPKcPcjjjjj:
.LFB2059:
	.cfi_startproc
	endbr64
	subq	$168, %rsp
	.cfi_def_cfa_offset 176
	movq	%rdi, 24(%rsp)
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movl	%edx, 12(%rsp)
	leaq	32(%rsp), %rdx
	movl	%ecx, 8(%rsp)
	leaq	40(%rsp), %rcx
	movl	%r8d, 4(%rsp)
	movl	%r9d, (%rsp)
	movq	%fs:40, %rax
	movq	%rax, 152(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 96(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 128(%rsp)
	movq	%rsp, %rax
	movq	%rax, 136(%rsp)
	leaq	176(%rsp), %rax
	movq	%rax, 144(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L39
	pushq	40(%rsp)
	.cfi_def_cfa_offset 184
	leaq	_Z13lifeItrKernelPKcPcjjjjj(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 192
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	112(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 184
	popq	%rdx
	.cfi_def_cfa_offset 176
.L39:
	movq	152(%rsp), %rax
	subq	%fs:40, %rax
	je	.L41
	call	__stack_chk_fail@PLT
.L41:
	addq	$168, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2059:
	.size	_Z41__device_stub__Z13lifeItrKernelPKcPcjjjjjPKcPcjjjjj, .-_Z41__device_stub__Z13lifeItrKernelPKcPcjjjjjPKcPcjjjjj
	.globl	_Z13lifeItrKernelPKcPcjjjjj
	.type	_Z13lifeItrKernelPKcPcjjjjj, @function
_Z13lifeItrKernelPKcPcjjjjj:
.LFB2060:
	.cfi_startproc
	endbr64
	jmp	_Z41__device_stub__Z13lifeItrKernelPKcPcjjjjjPKcPcjjjjj
	.cfi_endproc
.LFE2060:
	.size	_Z13lifeItrKernelPKcPcjjjjj, .-_Z13lifeItrKernelPKcPcjjjjj
	.section	.rodata.str1.1
.LC15:
	.string	"Grid dimension: %d,%d,%d\n"
.LC16:
	.string	"Block dimension: %d,%d,%d\n"
.LC17:
	.string	"Randomizing initial game (could take a while)...\n"
.LC18:
	.string	"Initial grid:\n"
.LC19:
	.string	"Allocating %d bytes of memory on the GPU...\n"
.LC20:
	.string	"Iteration %d "
.LC21:
	.string	"Kernel lifeItr failed with error \"%s\".\n"
.LC22:
	.string	"took %d ticks.\n"
	.text
	.globl	_Z7runLifejjjjjjjjj
	.type	_Z7runLifejjjjjjjjj, @function
_Z7runLifejjjjjjjjj:
.LFB2032:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movl	%edx, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	imull	%ecx, %ebp
	movl	%edi, %ebx
	subq	$104, %rsp
	.cfi_def_cfa_offset 160
	movl	160(%rsp), %eax
	movl	%r9d, 40(%rsp)
	movl	%esi, 16(%rsp)
	movl	168(%rsp), %r15d
	movl	%edx, 20(%rsp)
	movl	176(%rsp), %r12d
	movl	%ecx, 32(%rsp)
	movl	%r8d, 36(%rsp)
	movl	%eax, 44(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 88(%rsp)
	xorl	%eax, %eax
	movl	%esi, %eax
	cvtsi2sdq	%rax, %xmm0
	mulsd	.LC14(%rip), %xmm0
	imull	%eax, %ebp
	call	ceil@PLT
	movl	20(%rsp), %eax
	cvttsd2sil	%xmm0, %r13d
	cvtsi2sdq	%rax, %xmm0
	mulsd	.LC14(%rip), %xmm0
	call	ceil@PLT
	movl	32(%rsp), %eax
	cvttsd2sil	%xmm0, %r14d
	cvtsi2sdq	%rax, %xmm0
	mulsd	.LC14(%rip), %xmm0
	call	ceil@PLT
	movl	%r14d, %ecx
	movl	%r13d, %edx
	movl	$2, %edi
	cvttsd2sil	%xmm0, %r8d
	leaq	.LC15(%rip), %rsi
	xorl	%eax, %eax
	movl	%r8d, 8(%rsp)
	call	__printf_chk@PLT
	movl	8(%rsp), %r8d
	xorl	%eax, %eax
	movl	$8, %edx
	movl	$8, %ecx
	movl	$2, %edi
	movl	%r13d, 64(%rsp)
	leaq	.LC16(%rip), %rsi
	movl	%r8d, 72(%rsp)
	movl	$8, %r8d
	movl	%r14d, 68(%rsp)
	call	__printf_chk@PLT
	movl	%ebp, %edi
	movl	$8, 84(%rsp)
	movabsq	$34359738376, %rax
	movq	%rax, 76(%rsp)
	movl	%ebp, %eax
	movq	%rax, 24(%rsp)
	call	malloc@PLT
	leaq	.LC17(%rip), %rsi
	movl	$2, %edi
	movq	%rax, 8(%rsp)
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	36(%rsp), %edx
	movq	8(%rsp), %rdi
	movl	%ebp, %esi
	call	_Z13randomizeGridPcjj
	testl	%r15d, %r15d
	je	.L51
	leaq	.LC18(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	32(%rsp), %ecx
	movl	20(%rsp), %edx
	movl	16(%rsp), %esi
	movq	8(%rsp), %rdi
	call	_Z12print3DArrayPcjjj
.L51:
	testl	%r12d, %r12d
	je	.L52
	movl	32(%rsp), %ecx
	movl	20(%rsp), %edx
	movl	%ebx, %edi
	movl	16(%rsp), %esi
	call	_Z12initGameFilejjjj
	movl	32(%rsp), %ecx
	movl	20(%rsp), %edx
	movl	16(%rsp), %esi
	movq	8(%rsp), %rdi
	call	_Z13writeGameStepPcjjj
.L52:
	movl	%ebp, %edx
	leaq	.LC19(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	24(%rsp), %rsi
	leaq	48(%rsp), %rdi
	xorl	%r13d, %r13d
	call	cudaMalloc@PLT
	movq	24(%rsp), %rsi
	leaq	56(%rsp), %rdi
	call	cudaMalloc@PLT
	movq	24(%rsp), %rdx
	movq	8(%rsp), %rsi
	movl	$1, %ecx
	movq	48(%rsp), %rdi
	call	cudaMemcpy@PLT
.L53:
	movl	%r13d, %edx
	cmpq	%rbx, %r13
	je	.L79
	leaq	.LC20(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	call	clock@PLT
	movl	84(%rsp), %ecx
	movl	72(%rsp), %esi
	xorl	%r9d, %r9d
	movq	76(%rsp), %rdx
	movq	64(%rsp), %rdi
	movl	$512, %r8d
	movq	%rax, %r14
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L54
	pushq	%rax
	.cfi_def_cfa_offset 168
	movl	52(%rsp), %eax
	pushq	%rax
	.cfi_def_cfa_offset 176
	movl	56(%rsp), %r9d
	movl	48(%rsp), %r8d
	movl	36(%rsp), %ecx
	movl	32(%rsp), %edx
	movq	72(%rsp), %rsi
	movq	64(%rsp), %rdi
	call	_Z41__device_stub__Z13lifeItrKernelPKcPcjjjjjPKcPcjjjjj
	popq	%rdx
	.cfi_def_cfa_offset 168
	popq	%rcx
	.cfi_def_cfa_offset 160
.L54:
	call	cudaDeviceSynchronize@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L55
	call	cudaGetErrorString@PLT
	leaq	.LC21(%rip), %rsi
	movl	$2, %edi
	movq	%rax, %rdx
	xorl	%eax, %eax
	call	__printf_chk@PLT
.L55:
	call	clock@PLT
	movq	24(%rsp), %rdx
	movq	56(%rsp), %rsi
	movl	$3, %ecx
	movq	48(%rsp), %rdi
	movq	%rax, %rbp
	call	cudaMemcpy@PLT
	movq	%rbp, %rdx
	xorl	%eax, %eax
	movl	$2, %edi
	subq	%r14, %rdx
	leaq	.LC22(%rip), %rsi
	call	__printf_chk@PLT
	movl	%r15d, %eax
	orl	%r12d, %eax
	je	.L57
	movq	24(%rsp), %rdx
	movq	56(%rsp), %rsi
	movl	$2, %ecx
	movq	8(%rsp), %rdi
	call	cudaMemcpy@PLT
	testl	%r15d, %r15d
	je	.L58
	movl	32(%rsp), %ecx
	movl	20(%rsp), %edx
	movl	16(%rsp), %esi
	movq	8(%rsp), %rdi
	call	_Z12print3DArrayPcjjj
.L58:
	testl	%r12d, %r12d
	je	.L57
	movl	32(%rsp), %ecx
	movl	20(%rsp), %edx
	movl	16(%rsp), %esi
	movq	8(%rsp), %rdi
	call	_Z13writeGameStepPcjjj
.L57:
	incq	%r13
	jmp	.L53
.L79:
	movq	48(%rsp), %rdi
	call	cudaFree@PLT
	movq	56(%rsp), %rdi
	call	cudaFree@PLT
	movq	88(%rsp), %rax
	subq	%fs:40, %rax
	je	.L61
	call	__stack_chk_fail@PLT
.L61:
	movq	8(%rsp), %rdi
	addq	$104, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	jmp	free@PLT
	.cfi_endproc
.LFE2032:
	.size	_Z7runLifejjjjjjjjj, .-_Z7runLifejjjjjjjjj
	.section	.rodata.str1.1
.LC23:
	.string	"Starting %d iteration Game of Life (CUDA) with sizes x=%d, y=%d, z=%d\n"
.LC24:
	.string	"  initial alive chance=%d, neighbors for alive=%d to %d\n"
.LC25:
	.string	"  File output enabled.\n"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2034:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	subq	$24, %rsp
	.cfi_def_cfa_offset 80
	cmpl	$10, %edi
	je	.L81
	call	_Z10printUsagev
	jmp	.L82
.L81:
	movq	8(%rsi), %rdi
	movq	%rsi, %r14
	call	atoi@PLT
	movq	16(%r14), %rdi
	movl	%eax, 4(%rsp)
	call	atoi@PLT
	movq	24(%r14), %rdi
	movl	%eax, 8(%rsp)
	call	atoi@PLT
	movq	32(%r14), %rdi
	movl	%eax, 12(%rsp)
	call	atoi@PLT
	movq	40(%r14), %rdi
	movl	%eax, %ebx
	call	atoi@PLT
	movq	48(%r14), %rdi
	movl	%eax, %ebp
	call	atoi@PLT
	movq	56(%r14), %rdi
	movl	%eax, %r12d
	call	atoi@PLT
	movq	64(%r14), %rdi
	movl	%eax, %r13d
	call	atoi@PLT
	movq	72(%r14), %rdi
	movl	%eax, %r15d
	call	atoi@PLT
	movl	8(%rsp), %ecx
	movl	4(%rsp), %edx
	movl	%ebx, %r9d
	movl	12(%rsp), %r8d
	movl	%eax, %r14d
	movl	$2, %edi
	xorl	%eax, %eax
	leaq	.LC23(%rip), %rsi
	call	__printf_chk@PLT
	xorl	%eax, %eax
	movl	%r13d, %r8d
	movl	%r12d, %ecx
	movl	%ebp, %edx
	leaq	.LC24(%rip), %rsi
	movl	$2, %edi
	call	__printf_chk@PLT
	testl	%r14d, %r14d
	je	.L83
	leaq	.LC25(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
.L83:
	pushq	%rax
	.cfi_def_cfa_offset 88
	movl	%r12d, %r9d
	movl	%ebp, %r8d
	movl	%ebx, %ecx
	pushq	%r14
	.cfi_def_cfa_offset 96
	pushq	%r15
	.cfi_def_cfa_offset 104
	pushq	%r13
	.cfi_def_cfa_offset 112
	movl	44(%rsp), %edx
	movl	40(%rsp), %esi
	movl	36(%rsp), %edi
	call	_Z7runLifejjjjjjjjj
	addq	$32, %rsp
	.cfi_def_cfa_offset 80
.L82:
	addq	$24, %rsp
	.cfi_def_cfa_offset 56
	xorl	%eax, %eax
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2034:
	.size	main, .-main
	.text
	.globl	_Z44__device_stub__Z18sumNeighborsKernelPKcPcjjjPKcPcjjj
	.type	_Z44__device_stub__Z18sumNeighborsKernelPKcPcjjjPKcPcjjj, @function
_Z44__device_stub__Z18sumNeighborsKernelPKcPcjjjPKcPcjjj:
.LFB2061:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movq	%rdi, 24(%rsp)
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movl	%edx, 12(%rsp)
	leaq	32(%rsp), %rdx
	movl	%ecx, 8(%rsp)
	leaq	40(%rsp), %rcx
	movl	%r8d, 4(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 96(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 128(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L88
	pushq	40(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z18sumNeighborsKernelPKcPcjjj(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 176
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	112(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 168
	popq	%rdx
	.cfi_def_cfa_offset 160
.L88:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L90
	call	__stack_chk_fail@PLT
.L90:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2061:
	.size	_Z44__device_stub__Z18sumNeighborsKernelPKcPcjjjPKcPcjjj, .-_Z44__device_stub__Z18sumNeighborsKernelPKcPcjjjPKcPcjjj
	.globl	_Z18sumNeighborsKernelPKcPcjjj
	.type	_Z18sumNeighborsKernelPKcPcjjj, @function
_Z18sumNeighborsKernelPKcPcjjj:
.LFB2062:
	.cfi_startproc
	endbr64
	jmp	_Z44__device_stub__Z18sumNeighborsKernelPKcPcjjjPKcPcjjj
	.cfi_endproc
.LFE2062:
	.size	_Z18sumNeighborsKernelPKcPcjjj, .-_Z18sumNeighborsKernelPKcPcjjj
	.globl	_Z46__device_stub__Z18setAliveDeadKernelPKcPcjjjjjPKcPcjjjjj
	.type	_Z46__device_stub__Z18setAliveDeadKernelPKcPcjjjjjPKcPcjjjjj, @function
_Z46__device_stub__Z18setAliveDeadKernelPKcPcjjjjjPKcPcjjjjj:
.LFB2063:
	.cfi_startproc
	endbr64
	subq	$168, %rsp
	.cfi_def_cfa_offset 176
	movq	%rdi, 24(%rsp)
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movl	%edx, 12(%rsp)
	leaq	32(%rsp), %rdx
	movl	%ecx, 8(%rsp)
	leaq	40(%rsp), %rcx
	movl	%r8d, 4(%rsp)
	movl	%r9d, (%rsp)
	movq	%fs:40, %rax
	movq	%rax, 152(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 96(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 128(%rsp)
	movq	%rsp, %rax
	movq	%rax, 136(%rsp)
	leaq	176(%rsp), %rax
	movq	%rax, 144(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L93
	pushq	40(%rsp)
	.cfi_def_cfa_offset 184
	leaq	_Z18setAliveDeadKernelPKcPcjjjjj(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 192
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	112(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 184
	popq	%rdx
	.cfi_def_cfa_offset 176
.L93:
	movq	152(%rsp), %rax
	subq	%fs:40, %rax
	je	.L95
	call	__stack_chk_fail@PLT
.L95:
	addq	$168, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2063:
	.size	_Z46__device_stub__Z18setAliveDeadKernelPKcPcjjjjjPKcPcjjjjj, .-_Z46__device_stub__Z18setAliveDeadKernelPKcPcjjjjjPKcPcjjjjj
	.globl	_Z18setAliveDeadKernelPKcPcjjjjj
	.type	_Z18setAliveDeadKernelPKcPcjjjjj, @function
_Z18setAliveDeadKernelPKcPcjjjjj:
.LFB2064:
	.cfi_startproc
	endbr64
	jmp	_Z46__device_stub__Z18setAliveDeadKernelPKcPcjjjjjPKcPcjjjjj
	.cfi_endproc
.LFE2064:
	.size	_Z18setAliveDeadKernelPKcPcjjjjj, .-_Z18setAliveDeadKernelPKcPcjjjjj
	.section	.rodata.str1.1
.LC26:
	.string	"_Z18setAliveDeadKernelPKcPcjjjjj"
.LC27:
	.string	"_Z18sumNeighborsKernelPKcPcjjj"
.LC28:
	.string	"_Z13lifeItrKernelPKcPcjjjjj"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2066:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	pushq	$0
	.cfi_def_cfa_offset 24
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC26(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z18setAliveDeadKernelPKcPcjjjjj(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC27(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z18sumNeighborsKernelPKcPcjjj(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC28(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z13lifeItrKernelPKcPcjjjjj(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2066:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.nvFatBinSegment,"aw"
	.align 8
	.type	_ZL15__fatDeviceText, @object
	.size	_ZL15__fatDeviceText, 24
_ZL15__fatDeviceText:
	.long	1180844977
	.long	1
	.quad	fatbinData
	.quad	0
	.local	_ZL20__cudaFatCubinHandle
	.comm	_ZL20__cudaFatCubinHandle,8,8
	.section	.rodata.cst8,"aM",@progbits,8
	.align 8
.LC14:
	.long	0
	.long	1069547520
	.ident	"GCC: (Ubuntu 13.3.0-6ubuntu2~24.04) 13.3.0"
	.section	.note.GNU-stack,"",@progbits
	.section	.note.gnu.property,"a"
	.align 8
	.long	1f - 0f
	.long	4f - 1f
	.long	5
0:
	.string	"GNU"
1:
	.align 8
	.long	0xc0000002
	.long	3f - 2f
2:
	.long	0x3
3:
	.align 8
4:

// ================================================================================================
// Tim Backus
// CIS 450 - High Performance Computing
// 3D Game of Life - CUDA Version
// ================================================================================================

#define GOL_IO_FILENAME "gol3DOutput.dat"
#define GOL_CUDA_THREADS_SIZE 8


#include <hip/hip_runtime.h>
#include <stdio.h>
#include <string.h>
#include <ctype.h>
#include <time.h>
#include <stdlib.h>

// ------------------------------------------------------------------------------------------------
// CUDA kernel (Gather & Map) - Adds up the number of neighbors for a cell in a 3x3x3 cube and
//   sets each cell to alive or dead depending on its number of neighbors and the rules for this 
//   current game.
// ------------------------------------------------------------------------------------------------
__global__
void lifeItrKernel(const char* const d_in, char* d_out, const unsigned int xsize,
                   const unsigned int ysize, const unsigned int zsize, const unsigned int alow,
                   const unsigned int ahigh) {
  
  extern __shared__ char shMem[];
               
  // Calculate block and thread IDs
  const int threadPosX = blockIdx.x * blockDim.x + threadIdx.x;
  const int threadPosY = blockIdx.y * blockDim.y + threadIdx.y;
  const int threadPosZ = blockIdx.z * blockDim.z + threadIdx.z;
  const unsigned int stepX = ysize * zsize;
  const unsigned int arrayPos = threadPosX * stepX + threadPosY * zsize + threadPosZ;
  const unsigned int threadID = threadIdx.x * blockDim.y * blockDim.z + 
                                threadIdx.y * blockDim.z + threadIdx.z;

  // Ensure thread bounds
  if(threadPosX >= xsize) return;
  if(threadPosY >= ysize) return;
  if(threadPosZ >= zsize) return;
  
  // Copy global into shared memory
  shMem[threadID] = d_in[arrayPos];
  __syncthreads();
  
  // Begin adding neighbors
  char sum = 0;
  
  // X-Axis neighbors
  int xc, xcoord;
  for(xc = threadPosX - 1; xc <= threadPosX + 1; xc++) {
    
    // Wrap X-Axis
    xcoord = xc;
    if(xc < 0) xcoord = xsize;
    else if(xc >= xsize) xcoord = 0;
    
    // Y-Axis neighbors
    int yc, ycoord;
    for(yc = threadPosY - 1; yc <= threadPosY + 1; yc++) {
      
      // Wrap Y-Axis
      ycoord = yc;
      if(yc < 0) ycoord = ysize;
      else if(yc >= ysize) ycoord = 0;
      
      // Z-Axis neighbors
      int zc, zcoord;
      for(zc = threadPosZ - 1; zc <= threadPosZ + 1; zc++) {
        
        // Wrap Z-Axis
        zcoord = zc;
        if(zc < 0) zcoord = zsize;
        else if(zc >= zsize) zcoord = 0;
        
        // Don't count the cell itself
        if(threadPosX != xcoord || threadPosY != ycoord || threadPosZ != zcoord) {
          
          // Use shared memory instead of global memory if the current coord is in the thread block
          if((xcoord >= blockDim.x * blockIdx.x && xcoord < (blockDim.x) * blockIdx.x + 1) &&
             (ycoord >= blockDim.y * blockIdx.y && ycoord < (blockDim.y) * blockIdx.y + 1) &&
             (zcoord >= blockDim.z * blockIdx.z && zcoord < (blockDim.z) * blockIdx.z + 1)) {
            sum += shMem[(xcoord % blockDim.x) * blockDim.y * blockDim.z + (ycoord % blockDim.y) * 
                          blockDim.z + (zcoord % blockDim.z)];
          } else {
            sum += d_in[xcoord * stepX + ycoord * zsize + zcoord];
          }
        }
      }
    }
  }
  
  // Set the cell's dead or alive status based on its neighbor count
  if (sum < alow || sum > ahigh) {
    d_out[arrayPos] = 0;
  } else if (sum >= alow && sum <= ahigh) {
    d_out[arrayPos] = 1;
  }
  
}


// ------------------------------------------------------------------------------------------------
// CUDA kernel (Gather) - Adds up the number of neighbors for a cell in a 3x3x3 cube.
// ------------------------------------------------------------------------------------------------
__global__
void sumNeighborsKernel(const char* const d_in, char* d_out, const unsigned int xsize,
                        const unsigned int ysize, const unsigned int zsize) {
  
  
  // Calculate block and thread IDs
  const int threadPosX = blockIdx.x * blockDim.x + threadIdx.x;
  const int threadPosY = blockIdx.y * blockDim.y + threadIdx.y;
  const int threadPosZ = blockIdx.z * blockDim.z + threadIdx.z;
  const unsigned int stepX = ysize * zsize;
  const unsigned int arrayPos = threadPosX * stepX + threadPosY * zsize + threadPosZ;
  
  // printf("TID=%d,%d,%d\n", threadIdx.x, threadIdx.y, threadIdx.z);
  // printf("TPOS=%d,%d,%d\n", threadPosX, threadPosY, threadPosZ);
  // printf("APOS=%d\n", arrayPos);
  
  // Ensure thread bounds
  if(threadPosX > xsize - 1) return;
  if(threadPosY > ysize - 1) return;
  if(threadPosZ > zsize - 1) return;
  
  char sum = 0;
  
  // X-Axis neighbors
  int xc, xcoord;
  for(xc = threadPosX - 1; xc <= threadPosX + 1; xc++) {
    
    // Wrap X-Axis
    xcoord = xc;
    if(xc < 0) xcoord = xsize;
    else if(xc >= xsize) xcoord = 0;
    
    // Y-Axis neighbors
    int yc, ycoord;
    for(yc = threadPosY - 1; yc <= threadPosY + 1; yc++) {
      
      // Wrap Y-Axis
      ycoord = yc;
      if(yc < 0) ycoord = ysize;
      else if(yc >= ysize) ycoord = 0;
      
      // Z-Axis neighbors
      int zc, zcoord;
      for(zc = threadPosZ - 1; zc <= threadPosZ + 1; zc++) {
        
        // Wrap Z-Axis
        zcoord = zc;
        if(zc < 0) zcoord = zsize;
        else if(zc >= zsize) zcoord = 0;
        
        // Don't count the cell itself
        if(threadPosX != xcoord || threadPosY != ycoord || threadPosZ != zcoord) {
          sum += d_in[xcoord * stepX + ycoord * zsize + zcoord];
        }
        
      }
    }
  }
  
  d_out[arrayPos] = sum;
}

// ------------------------------------------------------------------------------------------------
// CUDA kernel (Map) - Sets each cell to alive or dead depending on its number of neighbors and
//   the rules for this current game.
// ------------------------------------------------------------------------------------------------
__global__
void setAliveDeadKernel(const char* const d_nei, char* d_out, const unsigned int xs, 
                        const unsigned int ys, const unsigned int zs, const unsigned int alow, 
                        const unsigned int ahigh) {
  
  // Calculate block and thread IDs
  const int threadPosX = blockIdx.x * blockDim.x + threadIdx.x;
  const int threadPosY = blockIdx.y * blockDim.y + threadIdx.y;
  const int threadPosZ = blockIdx.z * blockDim.z + threadIdx.z;
  const int stepX = ys * zs;
  const int arrayPos = threadPosX * stepX + threadPosY * zs + threadPosZ;

  // Ensure thread bounds
  if(threadPosX > xs - 1) return;
  if(threadPosY > ys - 1) return;
  if(threadPosZ > zs - 1) return;
  
  // Set the cell alive or dead as according to the rules
  if (d_nei[arrayPos] < alow || d_nei[arrayPos] > ahigh) {
    d_out[arrayPos] = 0;
  } else if (d_nei[arrayPos] >= alow && d_nei[arrayPos] <= ahigh) {
    d_out[arrayPos] = 1;
  }
}

// ------------------------------------------------------------------------------------------------
// Returns the 1D position of a simulated 3D array
// ------------------------------------------------------------------------------------------------
int getArrIndex(const unsigned int xp, const unsigned int yp, const unsigned int zp,
                const unsigned int ys, const unsigned int zs) {
  return xp * ys * zs + yp * zs + zp;
}

// ------------------------------------------------------------------------------------------------
// Prints a 3D array.
// ------------------------------------------------------------------------------------------------
void print3DArray(char* arr, unsigned const int x, unsigned const int y, unsigned const int z) {
  int i;
  for(i = 0; i < x; ++i) {
    printf("Dimension %d:\n", i);
    int j;
    for(j = 0; j < y; ++j) {
      int k;
      for(k = 0; k < z; ++k) {
        printf("%d ", (char)arr[getArrIndex(i, j, k, y, z)]);
      }
      printf("\n");
    }
    printf("\n");
  }
}

// ------------------------------------------------------------------------------------------------
// Writes cells to alive or dead, randomly.
// ------------------------------------------------------------------------------------------------
void randomizeGrid(char* grid, unsigned const int size, unsigned const int chance) {
 
  srand(time(NULL));
 
  int i;
  for(i = 0; i < size; i++) {
    grid[i] = (char)((rand() % 100 <= chance) ? 1 : 0);
  }
  
}

// ------------------------------------------------------------------------------------------------
// Initializes the game data file.
// Line 1: <iteration count> <x-size> <y-size> <z-size>
// Line 2: Blank
// ------------------------------------------------------------------------------------------------
void initGameFile(const unsigned int itrs, const unsigned int x, const unsigned int y,
                  const unsigned int z) {
  FILE *fp;
  fp = fopen(GOL_IO_FILENAME, "w+");
  fprintf(fp, "%d %d %d %d\n\n", itrs, x, y, z);
  fclose(fp);
}

// ------------------------------------------------------------------------------------------------
// Writes a game to a file for visualization within Java.
// For every iteration, a block of text is created of the format:
//   "<x-coord>:<z-coords for y=0>, <z-coords for y=1>, ..."
//   Z-coords are represented by a 0 or 1 for each z-coordinate
// Example: Game with 5 iterations, x=3, y=7, z=4
//   5 3 7 4
//   
//   0:0000,0000,0000,0000,0000,0000,0000
//   1:0000,0000,0000,0100,0000,0000,0001
//   2:0000,0000,0010,0100,0001,0011,0000
//   0:0000,0000,0000,0000,0000,0000,0000
//   1:0000,0000,0000,0100,0000,0000,0001
//   2:0000,0000,0010,0100,0001,0011,0000
//   0:0000,0000,0000,0000,0000,0000,0000
//   1:0000,0000,0000,0100,0000,0000,0001
//   2:0000,0000,0010,0100,0001,0011,0000
//   0:0000,0000,0000,0000,0000,0000,0000
//   1:0000,0000,0000,0100,0000,0000,0001
//   2:0000,0000,0010,0100,0001,0011,0000
//   0:0000,0000,0000,0000,0000,0000,0000
//   1:0000,0000,0000,0100,0000,0000,0001
//   2:0000,0000,0010,0100,0001,0011,0000
// 
// ------------------------------------------------------------------------------------------------
void writeGameStep(char* arr, unsigned const int x, unsigned const int y, unsigned const int z) {
  FILE *fp;
  fp = fopen(GOL_IO_FILENAME, "a");
  
  int i;
  for(i = 0; i < x; i++) {
    fprintf(fp, "%d:", i);
    int j;
    for(j = 0; j < y; j++) {
      
      if(j > 0) {
        fprintf(fp, ",");
      }
      // Print Z-Dim values
      int k;
      for(k = 0; k < z; k++) {
        fprintf(fp, "%d", arr[getArrIndex(i, j, k, y, z)]);
      }
    }
    fprintf(fp, "\n");
  }
  fclose(fp);
}

// ------------------------------------------------------------------------------------------------
// Runs the Game of Life.
// ------------------------------------------------------------------------------------------------
void runLife(const unsigned int iterations, unsigned int xsize, const unsigned int ysize, 
             const unsigned int zsize, const unsigned int initc, const unsigned int alow,
             const unsigned int ahigh, const unsigned int printArr, const unsigned int writeOut) {
  
  // Memory values
  const unsigned int arrSize = xsize * ysize * zsize;
  const unsigned int arrMem = arrSize * sizeof(char);
  
  // GPU grid dimensions
  const int gx = ceil((double) xsize / GOL_CUDA_THREADS_SIZE);
  const int gy = ceil((double) ysize / GOL_CUDA_THREADS_SIZE);
  const int gz = ceil((double) zsize / GOL_CUDA_THREADS_SIZE);
  printf("Grid dimension: %d,%d,%d\n", gx, gy, gz);
  dim3 gridDim(gx, gy, gz);
  
  // GPU thread dimensions
  const int tx = GOL_CUDA_THREADS_SIZE;
  const int ty = GOL_CUDA_THREADS_SIZE;
  const int tz = GOL_CUDA_THREADS_SIZE;
  printf("Block dimension: %d,%d,%d\n", tx, ty, tz);
  dim3 blockDim(tx, ty, tz);
  
  // Initialize game space
  char *h_in = (char *) malloc(arrMem);
  printf("Randomizing initial game (could take a while)...\n");
  randomizeGrid(h_in, arrSize, initc);
  
  // Print the initial array if enabled
  if(printArr) {
    printf("Initial grid:\n");
    print3DArray(h_in, xsize, ysize, zsize);
  }
  
  // Initialize the output file if enabled
  if(writeOut) {
    initGameFile(iterations, xsize, ysize, zsize);
    writeGameStep(h_in, xsize, ysize, zsize);
  }
  
  // Pointers for GPU game data
  char *d_in;
  char *d_out;
  
  // Allocate input array on GPU
  printf("Allocating %d bytes of memory on the GPU...\n", 
         (int)(xsize * ysize * zsize * sizeof(char)));
  hipMalloc(&d_in, arrMem);
  
  // Allocate output array on GPU
  hipMalloc(&d_out, arrMem);
  
  // Copy the host data to the GPU
  hipMemcpy(d_in, h_in, arrMem, hipMemcpyHostToDevice);
  
  // Do Game of Life iterations
  int itrNum;
  for(itrNum = 0; itrNum < iterations; itrNum++) {
    
    printf("Iteration %d ", itrNum);
    
    // Run the kernel to simulate an iteration of 3D life
    clock_t start = clock();
    lifeItrKernel<<<gridDim, blockDim, (tx * ty * tz * sizeof(char))>>>(d_in, d_out, xsize, ysize, zsize, alow, ahigh);
    hipError_t cerr = hipDeviceSynchronize();
    if(cerr != hipSuccess) {
      printf("Kernel lifeItr failed with error \"%s\".\n", hipGetErrorString(cerr));
    }
    clock_t end = clock();
    
    // Copy the memory back to the input array for the next iteration
    hipMemcpy(d_in, d_out, arrMem, hipMemcpyDeviceToDevice);
    
    printf("took %d ticks.\n", (end - start));
    
    // Print and write out if enabled
    if(printArr || writeOut) {
      hipMemcpy(h_in, d_out, arrMem, hipMemcpyDeviceToHost);
      if(printArr) {
        print3DArray(h_in, xsize, ysize, zsize);
      }
      if(writeOut) {
        writeGameStep(h_in, xsize, ysize, zsize);
      }
    }
  }
  
  // Free memory
  hipFree(d_in);
  hipFree(d_out);
  free(h_in);
}

// ------------------------------------------------------------------------------------------------
// Prints the usage message if a bad number of runtime arguments are passed.
// ------------------------------------------------------------------------------------------------
void printUsage() {
  printf("Arguments (separated by spaces):\n");
  printf("  MAX_ITERATIONS\n  SIZE_X\n  SIZE_Y\n  SIZE_Z\n  INITIAL_ALIVE_CHANCE (1-100)\n");
  printf("  ALIVE_THRESHOLD_LOW (inclusive)\n  ALIVE_THRESHOLD_HIGH (inclusive)\n");
  printf("  PRINT_ARRAY? (0=no, 1=yes)\n  WRITE_TO_FILE? (0=no, 1=yes)\n");
}

// ------------------------------------------------------------------------------------------------
// Main Method
// ------------------------------------------------------------------------------------------------
int main(int argc, char *argv[]) {
  
  // Ensure proper runtime argument count
  if(argc != 10) {
    printUsage();
    return EXIT_SUCCESS;
  }
  
  // Parse iteration count
  unsigned const int iterations = atoi(argv[1]);
  
  // Parse X-Size
  unsigned const int sizeX = atoi(argv[2]);
  
  // Parse Y-Size
  unsigned const int sizeY = atoi(argv[3]);
  
  // Parse Z-Size
  unsigned const int sizeZ = atoi(argv[4]);
  
  // Parse initial alive chance
  unsigned const int initChance = atoi(argv[5]);
  
  // Parse alive low threshold (inclusive)
  unsigned const int aliveLow = atoi(argv[6]);
  
  // Parse alive high threshold (inclusive)
  unsigned const int aliveHigh = atoi(argv[7]);
  
  // Parse whether or not to print the array
  unsigned const int printArray = atoi(argv[8]);
  
  // Parse whether or not to output to disk
  unsigned const int writeOut = atoi(argv[9]);
  
  // Print game information to the console
  printf("Starting %d iteration Game of Life (CUDA) with sizes x=%d, y=%d, z=%d\n", iterations,
         sizeX, sizeY, sizeZ);
  printf("  initial alive chance=%d, neighbors for alive=%d to %d\n", initChance, 
         aliveLow, aliveHigh);
  if(writeOut) {
    printf("  File output enabled.\n");
  }
  runLife(iterations, sizeX, sizeY, sizeZ, initChance, aliveLow, aliveHigh, printArray, writeOut);
  
  return EXIT_SUCCESS;
}
	.text
	.file	"Cuda3DLife.hip"
	.globl	_Z28__device_stub__lifeItrKernelPKcPcjjjjj # -- Begin function _Z28__device_stub__lifeItrKernelPKcPcjjjjj
	.type	_Z28__device_stub__lifeItrKernelPKcPcjjjjj,@function
_Z28__device_stub__lifeItrKernelPKcPcjjjjj: # @_Z28__device_stub__lifeItrKernelPKcPcjjjjj
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$144, %rsp
	.cfi_def_cfa_offset 192
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	40(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	32(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	12(%rsp), %rsi
	movl	%edx, (%rsi)
	leaq	8(%rsp), %rdx
	movl	%ecx, (%rdx)
	leaq	4(%rsp), %rcx
	movl	%r8d, (%rcx)
	movq	%rsp, %r8
	movl	%r9d, (%r8)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%rbx)
	movq	%r8, 40(%rbx)
	leaq	192(%rsp), %rax
	movq	%rax, 48(%rbx)
	leaq	64(%rsp), %r14
	leaq	48(%rsp), %r15
	leaq	24(%rsp), %r12
	leaq	16(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z13lifeItrKernelPKcPcjjjjj, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$160, %rsp
	.cfi_adjust_cfa_offset -160
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end0:
	.size	_Z28__device_stub__lifeItrKernelPKcPcjjjjj, .Lfunc_end0-_Z28__device_stub__lifeItrKernelPKcPcjjjjj
	.cfi_endproc
                                        # -- End function
	.globl	_Z33__device_stub__sumNeighborsKernelPKcPcjjj # -- Begin function _Z33__device_stub__sumNeighborsKernelPKcPcjjj
	.type	_Z33__device_stub__sumNeighborsKernelPKcPcjjj,@function
_Z33__device_stub__sumNeighborsKernelPKcPcjjj: # @_Z33__device_stub__sumNeighborsKernelPKcPcjjj
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$128, %rsp
	.cfi_def_cfa_offset 176
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	40(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	32(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	12(%rsp), %rsi
	movl	%edx, (%rsi)
	leaq	8(%rsp), %rdx
	movl	%ecx, (%rdx)
	leaq	4(%rsp), %rcx
	movl	%r8d, (%rcx)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%rbx)
	leaq	64(%rsp), %r14
	leaq	48(%rsp), %r15
	leaq	24(%rsp), %r12
	leaq	16(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z18sumNeighborsKernelPKcPcjjj, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$144, %rsp
	.cfi_adjust_cfa_offset -144
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end1:
	.size	_Z33__device_stub__sumNeighborsKernelPKcPcjjj, .Lfunc_end1-_Z33__device_stub__sumNeighborsKernelPKcPcjjj
	.cfi_endproc
                                        # -- End function
	.globl	_Z33__device_stub__setAliveDeadKernelPKcPcjjjjj # -- Begin function _Z33__device_stub__setAliveDeadKernelPKcPcjjjjj
	.type	_Z33__device_stub__setAliveDeadKernelPKcPcjjjjj,@function
_Z33__device_stub__setAliveDeadKernelPKcPcjjjjj: # @_Z33__device_stub__setAliveDeadKernelPKcPcjjjjj
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$144, %rsp
	.cfi_def_cfa_offset 192
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	40(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	32(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	12(%rsp), %rsi
	movl	%edx, (%rsi)
	leaq	8(%rsp), %rdx
	movl	%ecx, (%rdx)
	leaq	4(%rsp), %rcx
	movl	%r8d, (%rcx)
	movq	%rsp, %r8
	movl	%r9d, (%r8)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%rbx)
	movq	%r8, 40(%rbx)
	leaq	192(%rsp), %rax
	movq	%rax, 48(%rbx)
	leaq	64(%rsp), %r14
	leaq	48(%rsp), %r15
	leaq	24(%rsp), %r12
	leaq	16(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z18setAliveDeadKernelPKcPcjjjjj, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$160, %rsp
	.cfi_adjust_cfa_offset -160
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end2:
	.size	_Z33__device_stub__setAliveDeadKernelPKcPcjjjjj, .Lfunc_end2-_Z33__device_stub__setAliveDeadKernelPKcPcjjjjj
	.cfi_endproc
                                        # -- End function
	.globl	_Z11getArrIndexjjjjj            # -- Begin function _Z11getArrIndexjjjjj
	.type	_Z11getArrIndexjjjjj,@function
_Z11getArrIndexjjjjj:                   # @_Z11getArrIndexjjjjj
	.cfi_startproc
# %bb.0:
                                        # kill: def $esi killed $esi def $rsi
                                        # kill: def $edi killed $edi def $rdi
	imull	%ecx, %edi
	leal	(%rdi,%rsi), %eax
	imull	%r8d, %eax
	addl	%edx, %eax
	retq
.Lfunc_end3:
	.size	_Z11getArrIndexjjjjj, .Lfunc_end3-_Z11getArrIndexjjjjj
	.cfi_endproc
                                        # -- End function
	.globl	_Z12print3DArrayPcjjj           # -- Begin function _Z12print3DArrayPcjjj
	.type	_Z12print3DArrayPcjjj,@function
_Z12print3DArrayPcjjj:                  # @_Z12print3DArrayPcjjj
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	subq	$40, %rsp
	.cfi_def_cfa_offset 96
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movl	%edx, 16(%rsp)                  # 4-byte Spill
	movl	%esi, 24(%rsp)                  # 4-byte Spill
	testl	%esi, %esi
	je	.LBB4_9
# %bb.1:                                # %.lr.ph25
	movl	%ecx, %ebx
	movq	%rdi, %r15
	movl	%ecx, %eax
	movq	%rax, 32(%rsp)                  # 8-byte Spill
	movl	%ecx, %eax
	imull	16(%rsp), %eax                  # 4-byte Folded Reload
	movl	%eax, 20(%rsp)                  # 4-byte Spill
	movl	$0, 12(%rsp)                    # 4-byte Folded Spill
	xorl	%esi, %esi
.LBB4_2:                                # =>This Loop Header: Depth=1
                                        #     Child Loop BB4_4 Depth 2
                                        #       Child Loop BB4_6 Depth 3
	movl	$.L.str, %edi
	movl	%esi, 28(%rsp)                  # 4-byte Spill
	xorl	%eax, %eax
	callq	printf
	cmpl	$0, 16(%rsp)                    # 4-byte Folded Reload
	je	.LBB4_8
# %bb.3:                                # %.preheader.lr.ph
                                        #   in Loop: Header=BB4_2 Depth=1
	xorl	%r14d, %r14d
	movl	12(%rsp), %r12d                 # 4-byte Reload
.LBB4_4:                                # %.preheader
                                        #   Parent Loop BB4_2 Depth=1
                                        # =>  This Loop Header: Depth=2
                                        #       Child Loop BB4_6 Depth 3
	testl	%ebx, %ebx
	je	.LBB4_7
# %bb.5:                                # %.lr.ph
                                        #   in Loop: Header=BB4_4 Depth=2
	movq	32(%rsp), %r13                  # 8-byte Reload
	movl	%r12d, %ebp
.LBB4_6:                                #   Parent Loop BB4_2 Depth=1
                                        #     Parent Loop BB4_4 Depth=2
                                        # =>    This Inner Loop Header: Depth=3
	movslq	%ebp, %rbp
	movsbl	(%r15,%rbp), %esi
	movl	$.L.str.1, %edi
	xorl	%eax, %eax
	callq	printf
	incl	%ebp
	decq	%r13
	jne	.LBB4_6
.LBB4_7:                                # %._crit_edge
                                        #   in Loop: Header=BB4_4 Depth=2
	movl	$10, %edi
	callq	putchar@PLT
	incl	%r14d
	addl	%ebx, %r12d
	cmpl	16(%rsp), %r14d                 # 4-byte Folded Reload
	jne	.LBB4_4
.LBB4_8:                                # %._crit_edge22
                                        #   in Loop: Header=BB4_2 Depth=1
	movl	$10, %edi
	callq	putchar@PLT
	movl	28(%rsp), %esi                  # 4-byte Reload
	incl	%esi
	movl	12(%rsp), %eax                  # 4-byte Reload
	addl	20(%rsp), %eax                  # 4-byte Folded Reload
	movl	%eax, 12(%rsp)                  # 4-byte Spill
	cmpl	24(%rsp), %esi                  # 4-byte Folded Reload
	jne	.LBB4_2
.LBB4_9:                                # %._crit_edge26
	addq	$40, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end4:
	.size	_Z12print3DArrayPcjjj, .Lfunc_end4-_Z12print3DArrayPcjjj
	.cfi_endproc
                                        # -- End function
	.globl	_Z13randomizeGridPcjj           # -- Begin function _Z13randomizeGridPcjj
	.type	_Z13randomizeGridPcjj,@function
_Z13randomizeGridPcjj:                  # @_Z13randomizeGridPcjj
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movl	%edx, %ebx
	movl	%esi, %ebp
	movq	%rdi, %r14
	xorl	%edi, %edi
	callq	time
	movl	%eax, %edi
	callq	srand
	testl	%ebp, %ebp
	je	.LBB5_3
# %bb.1:                                # %.lr.ph.preheader
	movl	%ebp, %r15d
	xorl	%r12d, %r12d
.LBB5_2:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	callq	rand
	cltq
	imulq	$1374389535, %rax, %rcx         # imm = 0x51EB851F
	movq	%rcx, %rdx
	shrq	$63, %rdx
	sarq	$37, %rcx
	addl	%edx, %ecx
	imull	$100, %ecx, %ecx
	subl	%ecx, %eax
	cmpl	%ebx, %eax
	setbe	(%r14,%r12)
	incq	%r12
	cmpq	%r12, %r15
	jne	.LBB5_2
.LBB5_3:                                # %._crit_edge
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end5:
	.size	_Z13randomizeGridPcjj, .Lfunc_end5-_Z13randomizeGridPcjj
	.cfi_endproc
                                        # -- End function
	.globl	_Z12initGameFilejjjj            # -- Begin function _Z12initGameFilejjjj
	.type	_Z12initGameFilejjjj,@function
_Z12initGameFilejjjj:                   # @_Z12initGameFilejjjj
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movl	%ecx, %ebx
	movl	%edx, %ebp
	movl	%esi, %r14d
	movl	%edi, %r15d
	movl	$.L.str.3, %edi
	movl	$.L.str.4, %esi
	callq	fopen
	movq	%rax, %r12
	movl	$.L.str.5, %esi
	movq	%rax, %rdi
	movl	%r15d, %edx
	movl	%r14d, %ecx
	movl	%ebp, %r8d
	movl	%ebx, %r9d
	xorl	%eax, %eax
	callq	fprintf
	movq	%r12, %rdi
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	jmp	fclose                          # TAILCALL
.Lfunc_end6:
	.size	_Z12initGameFilejjjj, .Lfunc_end6-_Z12initGameFilejjjj
	.cfi_endproc
                                        # -- End function
	.globl	_Z13writeGameStepPcjjj          # -- Begin function _Z13writeGameStepPcjjj
	.type	_Z13writeGameStepPcjjj,@function
_Z13writeGameStepPcjjj:                 # @_Z13writeGameStepPcjjj
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	subq	$40, %rsp
	.cfi_def_cfa_offset 96
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movl	%ecx, %ebx
	movl	%edx, 12(%rsp)                  # 4-byte Spill
	movl	%esi, %ebp
	movq	%rdi, %r15
	movl	$.L.str.3, %edi
	movl	$.L.str.6, %esi
	callq	fopen
	movq	%rax, %r12
	movl	%ebp, 20(%rsp)                  # 4-byte Spill
	testl	%ebp, %ebp
	je	.LBB7_11
# %bb.1:                                # %.lr.ph33
	movl	%ebx, %eax
	movq	%rax, 32(%rsp)                  # 8-byte Spill
	movl	%ebx, %eax
	imull	12(%rsp), %eax                  # 4-byte Folded Reload
	movl	%eax, 16(%rsp)                  # 4-byte Spill
	movl	$0, 8(%rsp)                     # 4-byte Folded Spill
	xorl	%edx, %edx
	movl	%ebx, 28(%rsp)                  # 4-byte Spill
.LBB7_2:                                # =>This Loop Header: Depth=1
                                        #     Child Loop BB7_4 Depth 2
                                        #       Child Loop BB7_8 Depth 3
	movl	$.L.str.7, %esi
	movq	%r12, %rdi
	movl	%edx, 24(%rsp)                  # 4-byte Spill
	xorl	%eax, %eax
	callq	fprintf
	cmpl	$0, 12(%rsp)                    # 4-byte Folded Reload
	je	.LBB7_10
# %bb.3:                                # %.lr.ph29
                                        #   in Loop: Header=BB7_2 Depth=1
	xorl	%r13d, %r13d
	movl	8(%rsp), %r14d                  # 4-byte Reload
.LBB7_4:                                #   Parent Loop BB7_2 Depth=1
                                        # =>  This Loop Header: Depth=2
                                        #       Child Loop BB7_8 Depth 3
	testl	%r13d, %r13d
	je	.LBB7_6
# %bb.5:                                #   in Loop: Header=BB7_4 Depth=2
	movl	$44, %edi
	movq	%r12, %rsi
	callq	fputc@PLT
.LBB7_6:                                #   in Loop: Header=BB7_4 Depth=2
	testl	%ebx, %ebx
	je	.LBB7_9
# %bb.7:                                # %.lr.ph
                                        #   in Loop: Header=BB7_4 Depth=2
	movq	32(%rsp), %rbp                  # 8-byte Reload
	movl	%r14d, %ebx
.LBB7_8:                                #   Parent Loop BB7_2 Depth=1
                                        #     Parent Loop BB7_4 Depth=2
                                        # =>    This Inner Loop Header: Depth=3
	movslq	%ebx, %rbx
	movsbl	(%r15,%rbx), %edx
	movl	$.L.str.9, %esi
	movq	%r12, %rdi
	xorl	%eax, %eax
	callq	fprintf
	incl	%ebx
	decq	%rbp
	jne	.LBB7_8
.LBB7_9:                                # %._crit_edge
                                        #   in Loop: Header=BB7_4 Depth=2
	incl	%r13d
	movl	28(%rsp), %ebx                  # 4-byte Reload
	addl	%ebx, %r14d
	cmpl	12(%rsp), %r13d                 # 4-byte Folded Reload
	jne	.LBB7_4
.LBB7_10:                               # %._crit_edge30
                                        #   in Loop: Header=BB7_2 Depth=1
	movl	$10, %edi
	movq	%r12, %rsi
	callq	fputc@PLT
	movl	24(%rsp), %edx                  # 4-byte Reload
	incl	%edx
	movl	8(%rsp), %eax                   # 4-byte Reload
	addl	16(%rsp), %eax                  # 4-byte Folded Reload
	movl	%eax, 8(%rsp)                   # 4-byte Spill
	cmpl	20(%rsp), %edx                  # 4-byte Folded Reload
	jne	.LBB7_2
.LBB7_11:                               # %._crit_edge34
	movq	%r12, %rdi
	addq	$40, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	jmp	fclose                          # TAILCALL
.Lfunc_end7:
	.size	_Z13writeGameStepPcjjj, .Lfunc_end7-_Z13writeGameStepPcjjj
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function _Z7runLifejjjjjjjjj
.LCPI8_0:
	.quad	0x3fc0000000000000              # double 0.125
	.text
	.globl	_Z7runLifejjjjjjjjj
	.type	_Z7runLifejjjjjjjjj,@function
_Z7runLifejjjjjjjjj:                    # @_Z7runLifejjjjjjjjj
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	subq	$72, %rsp
	.cfi_def_cfa_offset 128
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movl	%r9d, 60(%rsp)                  # 4-byte Spill
	movl	%r8d, 8(%rsp)                   # 4-byte Spill
	movl	%ecx, %r14d
	movl	%edx, %r12d
	movl	%esi, %r15d
	movl	%edi, %ebp
	movl	%edx, %r13d
	imull	%esi, %r13d
	imull	%ecx, %r13d
	movl	%esi, %eax
	cvtsi2sd	%rax, %xmm0
	mulsd	.LCPI8_0(%rip), %xmm0
	callq	ceil@PLT
	cvttsd2si	%xmm0, %eax
	movq	%rax, 16(%rsp)                  # 8-byte Spill
	movl	%r12d, %eax
	xorps	%xmm0, %xmm0
	cvtsi2sd	%rax, %xmm0
	mulsd	.LCPI8_0(%rip), %xmm0
	callq	ceil@PLT
	cvttsd2si	%xmm0, %ebx
	movl	%r14d, %eax
	xorps	%xmm0, %xmm0
	cvtsi2sd	%rax, %xmm0
	mulsd	.LCPI8_0(%rip), %xmm0
	callq	ceil@PLT
	cvttsd2si	%xmm0, %ecx
	movl	$.L.str.10, %edi
	movq	16(%rsp), %rsi                  # 8-byte Reload
                                        # kill: def $esi killed $esi killed $rsi
	movq	%rbx, 32(%rsp)                  # 8-byte Spill
	movl	%ebx, %edx
	movl	%ecx, 68(%rsp)                  # 4-byte Spill
	xorl	%eax, %eax
	callq	printf
	movl	$.L.str.11, %edi
	movl	$8, %esi
	movl	$8, %edx
	movl	$8, %ecx
	xorl	%eax, %eax
	callq	printf
	movq	%r13, %rdi
	callq	malloc
	movq	%rax, %rbx
	movl	$.Lstr, %edi
	callq	puts@PLT
	movq	%rbx, %rdi
	movq	%r13, 40(%rsp)                  # 8-byte Spill
	movl	%r13d, %esi
	movl	8(%rsp), %edx                   # 4-byte Reload
	callq	_Z13randomizeGridPcjj
	cmpl	$0, 136(%rsp)
	je	.LBB8_2
# %bb.1:
	movl	$.Lstr.1, %edi
	callq	puts@PLT
	movq	%rbx, %rdi
	movl	%r15d, %esi
	movl	%r12d, %edx
	movl	%r14d, %ecx
	callq	_Z12print3DArrayPcjjj
.LBB8_2:
	cmpl	$0, 144(%rsp)
	je	.LBB8_4
# %bb.3:
	movl	%ebp, %edi
	movl	%r15d, %esi
	movl	%r12d, %edx
	movl	%r14d, %ecx
	callq	_Z12initGameFilejjjj
	movq	%rbx, %rdi
	movl	%r15d, %esi
	movl	%r12d, %edx
	movl	%r14d, %ecx
	callq	_Z13writeGameStepPcjjj
.LBB8_4:
	movl	$.L.str.14, %edi
	movq	%rbx, 8(%rsp)                   # 8-byte Spill
	movq	40(%rsp), %r13                  # 8-byte Reload
	movl	%r13d, %esi
	xorl	%eax, %eax
	callq	printf
	leaq	48(%rsp), %rbx
	movq	%rbx, %rdi
	movq	%r13, %rsi
	callq	hipMalloc
	leaq	24(%rsp), %rdi
	movq	%r13, %rsi
	callq	hipMalloc
	movq	(%rbx), %rdi
	movq	8(%rsp), %rsi                   # 8-byte Reload
	movq	%r13, %rdx
	movq	8(%rsp), %r13                   # 8-byte Reload
	movl	$1, %ecx
	callq	hipMemcpy
	testl	%ebp, %ebp
	je	.LBB8_16
# %bb.5:                                # %.lr.ph
	movq	32(%rsp), %rax                  # 8-byte Reload
	shlq	$32, %rax
	addq	16(%rsp), %rax                  # 8-byte Folded Reload
	movq	%rax, 32(%rsp)                  # 8-byte Spill
	movl	144(%rsp), %eax
	orl	136(%rsp), %eax
	movl	%eax, 16(%rsp)                  # 4-byte Spill
	xorl	%ebx, %ebx
	movl	%ebp, 64(%rsp)                  # 4-byte Spill
.LBB8_6:                                # =>This Inner Loop Header: Depth=1
	movl	$.L.str.15, %edi
	movl	%ebx, %esi
	xorl	%eax, %eax
	callq	printf
	callq	clock
	movq	%rax, %rbp
	movl	$512, %r8d                      # imm = 0x200
	movq	32(%rsp), %rdi                  # 8-byte Reload
	movl	68(%rsp), %esi                  # 4-byte Reload
	movabsq	$34359738376, %rdx              # imm = 0x800000008
	movl	$8, %ecx
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB8_8
# %bb.7:                                #   in Loop: Header=BB8_6 Depth=1
	movq	48(%rsp), %rdi
	movq	24(%rsp), %rsi
	movl	128(%rsp), %eax
	movl	%eax, (%rsp)
	movl	%r15d, %edx
	movl	%r12d, %ecx
	movl	%r14d, %r8d
	movl	60(%rsp), %r9d                  # 4-byte Reload
	callq	_Z28__device_stub__lifeItrKernelPKcPcjjjjj
.LBB8_8:                                #   in Loop: Header=BB8_6 Depth=1
	movl	%r15d, %r13d
	movl	%r12d, %r15d
	movl	%r14d, %r12d
	callq	hipDeviceSynchronize
	testl	%eax, %eax
	je	.LBB8_10
# %bb.9:                                #   in Loop: Header=BB8_6 Depth=1
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.16, %edi
	movq	%rax, %rsi
	xorl	%eax, %eax
	callq	printf
.LBB8_10:                               #   in Loop: Header=BB8_6 Depth=1
	callq	clock
	movq	%rax, %r14
	movq	48(%rsp), %rdi
	movq	24(%rsp), %rsi
	movq	40(%rsp), %rdx                  # 8-byte Reload
	movl	$3, %ecx
	callq	hipMemcpy
	subq	%rbp, %r14
	movl	$.L.str.17, %edi
	movq	%r14, %rsi
	xorl	%eax, %eax
	callq	printf
	cmpl	$0, 16(%rsp)                    # 4-byte Folded Reload
	movl	%r12d, %r14d
	movl	%r15d, %r12d
	movl	%r13d, %r15d
	movl	64(%rsp), %ebp                  # 4-byte Reload
	movq	8(%rsp), %r13                   # 8-byte Reload
	je	.LBB8_15
# %bb.11:                               #   in Loop: Header=BB8_6 Depth=1
	movq	24(%rsp), %rsi
	movq	%r13, %rdi
	movq	40(%rsp), %rdx                  # 8-byte Reload
	movl	$2, %ecx
	callq	hipMemcpy
	cmpl	$0, 136(%rsp)
	je	.LBB8_13
# %bb.12:                               #   in Loop: Header=BB8_6 Depth=1
	movq	%r13, %rdi
	movl	%r15d, %esi
	movl	%r12d, %edx
	movl	%r14d, %ecx
	callq	_Z12print3DArrayPcjjj
.LBB8_13:                               #   in Loop: Header=BB8_6 Depth=1
	cmpl	$0, 144(%rsp)
	je	.LBB8_15
# %bb.14:                               #   in Loop: Header=BB8_6 Depth=1
	movq	%r13, %rdi
	movl	%r15d, %esi
	movl	%r12d, %edx
	movl	%r14d, %ecx
	callq	_Z13writeGameStepPcjjj
.LBB8_15:                               #   in Loop: Header=BB8_6 Depth=1
	incl	%ebx
	cmpl	%ebx, %ebp
	jne	.LBB8_6
.LBB8_16:                               # %._crit_edge
	movq	48(%rsp), %rdi
	callq	hipFree
	movq	24(%rsp), %rdi
	callq	hipFree
	movq	%r13, %rdi
	callq	free
	addq	$72, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end8:
	.size	_Z7runLifejjjjjjjjj, .Lfunc_end8-_Z7runLifejjjjjjjjj
	.cfi_endproc
                                        # -- End function
	.globl	_Z10printUsagev                 # -- Begin function _Z10printUsagev
	.type	_Z10printUsagev,@function
_Z10printUsagev:                        # @_Z10printUsagev
	.cfi_startproc
# %bb.0:
	pushq	%rax
	.cfi_def_cfa_offset 16
	movl	$.Lstr.2, %edi
	callq	puts@PLT
	movl	$.Lstr.3, %edi
	callq	puts@PLT
	movl	$.Lstr.4, %edi
	callq	puts@PLT
	movl	$.Lstr.5, %edi
	popq	%rax
	.cfi_def_cfa_offset 8
	jmp	puts@PLT                        # TAILCALL
.Lfunc_end9:
	.size	_Z10printUsagev, .Lfunc_end9-_Z10printUsagev
	.cfi_endproc
                                        # -- End function
	.globl	main                            # -- Begin function main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	subq	$40, %rsp
	.cfi_def_cfa_offset 96
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	cmpl	$10, %edi
	jne	.LBB10_1
# %bb.2:
	movq	%rsi, %r13
	movq	8(%rsi), %rdi
	callq	atoi
	movl	%eax, %r15d
	movq	16(%r13), %rdi
	callq	atoi
	movl	%eax, %ebp
	movq	24(%r13), %rdi
	callq	atoi
	movl	%eax, %r12d
	movq	32(%r13), %rdi
	callq	atoi
	movl	%eax, 8(%rsp)                   # 4-byte Spill
	movq	40(%r13), %rdi
	callq	atoi
	movl	%eax, %ebx
	movq	48(%r13), %rdi
	callq	atoi
	movl	%eax, %r14d
	movq	56(%r13), %rdi
	callq	atoi
                                        # kill: def $eax killed $eax def $rax
	movq	%rax, 24(%rsp)                  # 8-byte Spill
	movq	64(%r13), %rdi
	callq	atoi
                                        # kill: def $eax killed $eax def $rax
	movq	%rax, 32(%rsp)                  # 8-byte Spill
	movq	72(%r13), %rdi
	callq	atoi
	movl	%eax, %r13d
	movl	$.L.str.22, %edi
	movl	%r15d, 16(%rsp)                 # 4-byte Spill
	movl	%r15d, %esi
	movl	%ebp, %r15d
	movl	%ebp, %edx
	movl	%r12d, 20(%rsp)                 # 4-byte Spill
	movl	%r12d, %ecx
	movl	8(%rsp), %r12d                  # 4-byte Reload
	movl	%r12d, %r8d
	xorl	%eax, %eax
	callq	printf
	movl	$.L.str.23, %edi
	movl	%ebx, 12(%rsp)                  # 4-byte Spill
	movl	%ebx, %esi
	movq	24(%rsp), %rbx                  # 8-byte Reload
	movl	%r14d, %ebp
	movl	%r14d, %edx
	movl	%ebx, %ecx
	xorl	%eax, %eax
	callq	printf
	testl	%r13d, %r13d
	je	.LBB10_4
# %bb.3:
	movl	$.Lstr.6, %edi
	callq	puts@PLT
.LBB10_4:
	subq	$8, %rsp
	.cfi_adjust_cfa_offset 8
	movl	24(%rsp), %edi                  # 4-byte Reload
	movl	%r15d, %esi
	movl	28(%rsp), %edx                  # 4-byte Reload
	movl	%r12d, %ecx
	movl	20(%rsp), %r8d                  # 4-byte Reload
	movl	%ebp, %r9d
	pushq	%r13
	.cfi_adjust_cfa_offset 8
	pushq	48(%rsp)                        # 8-byte Folded Reload
	.cfi_adjust_cfa_offset 8
	pushq	%rbx
	.cfi_adjust_cfa_offset 8
	callq	_Z7runLifejjjjjjjjj
	addq	$32, %rsp
	.cfi_adjust_cfa_offset -32
	jmp	.LBB10_5
.LBB10_1:
	callq	_Z10printUsagev
.LBB10_5:
	xorl	%eax, %eax
	addq	$40, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end10:
	.size	main, .Lfunc_end10-main
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_ctor
	.type	__hip_module_ctor,@function
__hip_module_ctor:                      # @__hip_module_ctor
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	subq	$32, %rsp
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -16
	movq	__hip_gpubin_handle(%rip), %rbx
	testq	%rbx, %rbx
	jne	.LBB11_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB11_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z13lifeItrKernelPKcPcjjjjj, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z18sumNeighborsKernelPKcPcjjj, %esi
	movl	$.L__unnamed_2, %edx
	movl	$.L__unnamed_2, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z18setAliveDeadKernelPKcPcjjjjj, %esi
	movl	$.L__unnamed_3, %edx
	movl	$.L__unnamed_3, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end11:
	.size	__hip_module_ctor, .Lfunc_end11-__hip_module_ctor
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_dtor
	.type	__hip_module_dtor,@function
__hip_module_dtor:                      # @__hip_module_dtor
	.cfi_startproc
# %bb.0:
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	je	.LBB12_2
# %bb.1:
	pushq	%rax
	.cfi_def_cfa_offset 16
	callq	__hipUnregisterFatBinary
	movq	$0, __hip_gpubin_handle(%rip)
	addq	$8, %rsp
	.cfi_def_cfa_offset 8
.LBB12_2:
	retq
.Lfunc_end12:
	.size	__hip_module_dtor, .Lfunc_end12-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_Z13lifeItrKernelPKcPcjjjjj,@object # @_Z13lifeItrKernelPKcPcjjjjj
	.section	.rodata,"a",@progbits
	.globl	_Z13lifeItrKernelPKcPcjjjjj
	.p2align	3, 0x0
_Z13lifeItrKernelPKcPcjjjjj:
	.quad	_Z28__device_stub__lifeItrKernelPKcPcjjjjj
	.size	_Z13lifeItrKernelPKcPcjjjjj, 8

	.type	_Z18sumNeighborsKernelPKcPcjjj,@object # @_Z18sumNeighborsKernelPKcPcjjj
	.globl	_Z18sumNeighborsKernelPKcPcjjj
	.p2align	3, 0x0
_Z18sumNeighborsKernelPKcPcjjj:
	.quad	_Z33__device_stub__sumNeighborsKernelPKcPcjjj
	.size	_Z18sumNeighborsKernelPKcPcjjj, 8

	.type	_Z18setAliveDeadKernelPKcPcjjjjj,@object # @_Z18setAliveDeadKernelPKcPcjjjjj
	.globl	_Z18setAliveDeadKernelPKcPcjjjjj
	.p2align	3, 0x0
_Z18setAliveDeadKernelPKcPcjjjjj:
	.quad	_Z33__device_stub__setAliveDeadKernelPKcPcjjjjj
	.size	_Z18setAliveDeadKernelPKcPcjjjjj, 8

	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"Dimension %d:\n"
	.size	.L.str, 15

	.type	.L.str.1,@object                # @.str.1
.L.str.1:
	.asciz	"%d "
	.size	.L.str.1, 4

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	"gol3DOutput.dat"
	.size	.L.str.3, 16

	.type	.L.str.4,@object                # @.str.4
.L.str.4:
	.asciz	"w+"
	.size	.L.str.4, 3

	.type	.L.str.5,@object                # @.str.5
.L.str.5:
	.asciz	"%d %d %d %d\n\n"
	.size	.L.str.5, 14

	.type	.L.str.6,@object                # @.str.6
.L.str.6:
	.asciz	"a"
	.size	.L.str.6, 2

	.type	.L.str.7,@object                # @.str.7
.L.str.7:
	.asciz	"%d:"
	.size	.L.str.7, 4

	.type	.L.str.9,@object                # @.str.9
.L.str.9:
	.asciz	"%d"
	.size	.L.str.9, 3

	.type	.L.str.10,@object               # @.str.10
.L.str.10:
	.asciz	"Grid dimension: %d,%d,%d\n"
	.size	.L.str.10, 26

	.type	.L.str.11,@object               # @.str.11
.L.str.11:
	.asciz	"Block dimension: %d,%d,%d\n"
	.size	.L.str.11, 27

	.type	.L.str.14,@object               # @.str.14
.L.str.14:
	.asciz	"Allocating %d bytes of memory on the GPU...\n"
	.size	.L.str.14, 45

	.type	.L.str.15,@object               # @.str.15
.L.str.15:
	.asciz	"Iteration %d "
	.size	.L.str.15, 14

	.type	.L.str.16,@object               # @.str.16
.L.str.16:
	.asciz	"Kernel lifeItr failed with error \"%s\".\n"
	.size	.L.str.16, 40

	.type	.L.str.17,@object               # @.str.17
.L.str.17:
	.asciz	"took %d ticks.\n"
	.size	.L.str.17, 16

	.type	.L.str.22,@object               # @.str.22
.L.str.22:
	.asciz	"Starting %d iteration Game of Life (CUDA) with sizes x=%d, y=%d, z=%d\n"
	.size	.L.str.22, 71

	.type	.L.str.23,@object               # @.str.23
.L.str.23:
	.asciz	"  initial alive chance=%d, neighbors for alive=%d to %d\n"
	.size	.L.str.23, 57

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z13lifeItrKernelPKcPcjjjjj"
	.size	.L__unnamed_1, 28

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"_Z18sumNeighborsKernelPKcPcjjj"
	.size	.L__unnamed_2, 31

	.type	.L__unnamed_3,@object           # @2
.L__unnamed_3:
	.asciz	"_Z18setAliveDeadKernelPKcPcjjjjj"
	.size	.L__unnamed_3, 33

	.type	__hip_fatbin_wrapper,@object    # @__hip_fatbin_wrapper
	.section	.hipFatBinSegment,"a",@progbits
	.p2align	3, 0x0
__hip_fatbin_wrapper:
	.long	1212764230                      # 0x48495046
	.long	1                               # 0x1
	.quad	__hip_fatbin
	.quad	0
	.size	__hip_fatbin_wrapper, 24

	.type	__hip_gpubin_handle,@object     # @__hip_gpubin_handle
	.local	__hip_gpubin_handle
	.comm	__hip_gpubin_handle,8,8
	.section	.init_array,"aw",@init_array
	.p2align	3, 0x0
	.quad	__hip_module_ctor
	.type	__hip_cuid_,@object             # @__hip_cuid_
	.bss
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               # 0x0
	.size	__hip_cuid_, 1

	.type	.Lstr,@object                   # @str
	.section	.rodata.str1.1,"aMS",@progbits,1
.Lstr:
	.asciz	"Randomizing initial game (could take a while)..."
	.size	.Lstr, 49

	.type	.Lstr.1,@object                 # @str.1
.Lstr.1:
	.asciz	"Initial grid:"
	.size	.Lstr.1, 14

	.type	.Lstr.2,@object                 # @str.2
.Lstr.2:
	.asciz	"Arguments (separated by spaces):"
	.size	.Lstr.2, 33

	.type	.Lstr.3,@object                 # @str.3
.Lstr.3:
	.asciz	"  MAX_ITERATIONS\n  SIZE_X\n  SIZE_Y\n  SIZE_Z\n  INITIAL_ALIVE_CHANCE (1-100)"
	.size	.Lstr.3, 75

	.type	.Lstr.4,@object                 # @str.4
.Lstr.4:
	.asciz	"  ALIVE_THRESHOLD_LOW (inclusive)\n  ALIVE_THRESHOLD_HIGH (inclusive)"
	.size	.Lstr.4, 69

	.type	.Lstr.5,@object                 # @str.5
.Lstr.5:
	.asciz	"  PRINT_ARRAY? (0=no, 1=yes)\n  WRITE_TO_FILE? (0=no, 1=yes)"
	.size	.Lstr.5, 60

	.type	.Lstr.6,@object                 # @str.6
.Lstr.6:
	.asciz	"  File output enabled."
	.size	.Lstr.6, 23

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym _Z28__device_stub__lifeItrKernelPKcPcjjjjj
	.addrsig_sym _Z33__device_stub__sumNeighborsKernelPKcPcjjj
	.addrsig_sym _Z33__device_stub__setAliveDeadKernelPKcPcjjjjj
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z13lifeItrKernelPKcPcjjjjj
	.addrsig_sym _Z18sumNeighborsKernelPKcPcjjj
	.addrsig_sym _Z18setAliveDeadKernelPKcPcjjjjj
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z13lifeItrKernelPKcPcjjjjj ; -- Begin function _Z13lifeItrKernelPKcPcjjjjj
	.globl	_Z13lifeItrKernelPKcPcjjjjj
	.p2align	8
	.type	_Z13lifeItrKernelPKcPcjjjjj,@function
_Z13lifeItrKernelPKcPcjjjjj:            ; @_Z13lifeItrKernelPKcPcjjjjj
; %bb.0:
	s_clause 0x1
	s_load_b64 s[2:3], s[0:1], 0x34
	s_load_b128 s[4:7], s[0:1], 0x10
	v_and_b32_e32 v1, 0x3ff, v0
	v_bfe_u32 v2, v0, 10, 10
	v_bfe_u32 v0, v0, 20, 10
	s_waitcnt lgkmcnt(0)
	s_and_b32 s16, s2, 0xffff
	s_lshr_b32 s17, s2, 16
	s_and_b32 s12, s3, 0xffff
	s_mul_i32 s13, s13, s16
	s_mul_i32 s14, s14, s17
	s_mul_i32 s15, s15, s12
	v_add_nc_u32_e32 v3, s13, v1
	v_add_nc_u32_e32 v4, s14, v2
	v_add_nc_u32_e32 v5, s15, v0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_cmp_gt_u32_e32 vcc_lo, s4, v3
	v_cmp_gt_u32_e64 s2, s5, v4
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cmp_gt_u32_e64 s3, s6, v5
	s_and_b32 s2, vcc_lo, s2
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	s_and_b32 s2, s2, s3
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_saveexec_b32 s3, s2
	s_cbranch_execz .LBB0_14
; %bb.1:
	s_clause 0x1
	s_load_b128 s[8:11], s[0:1], 0x0
	s_load_b32 s18, s[0:1], 0x20
	s_mul_i32 s3, s6, s5
	v_mul_lo_u32 v6, v4, s6
	v_mul_lo_u32 v7, v3, s3
	v_cvt_f32_u32_e32 v8, s17
	v_cvt_f32_u32_e32 v9, s12
	v_mad_u32_u24 v1, v1, s17, v2
	s_sub_i32 s0, 0, s16
	s_sub_i32 s1, 0, s17
	v_rcp_iflag_f32_e32 v8, v8
	v_rcp_iflag_f32_e32 v9, v9
	v_add3_u32 v6, v6, v7, v5
	v_cvt_f32_u32_e32 v7, s16
	s_sub_i32 s2, 0, s12
	v_mul_lo_u32 v1, v1, s12
	s_mov_b32 s19, 0
	s_add_i32 s20, s13, 1
	v_rcp_iflag_f32_e32 v7, v7
	s_add_i32 s21, s14, 1
	s_waitcnt lgkmcnt(0)
	global_load_u8 v14, v6, s[8:9]
	v_mul_f32_e32 v9, 0x4f7ffffe, v9
	s_add_i32 s22, s15, 1
	v_add3_u32 v0, v1, v0, 0
	v_dual_mul_f32 v8, 0x4f7ffffe, v8 :: v_dual_mul_f32 v7, 0x4f7ffffe, v7
	s_delay_alu instid0(VALU_DEP_1)
	v_cvt_u32_f32_e32 v12, v8
	v_add_nc_u32_e32 v8, -1, v4
	v_cvt_u32_f32_e32 v13, v9
	v_add_nc_u32_e32 v9, -1, v3
	v_cvt_u32_f32_e32 v2, v7
	v_mov_b32_e32 v7, 0
	v_mul_lo_u32 v11, s1, v12
	v_mul_lo_u32 v15, s2, v13
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_mul_lo_u32 v10, s0, v2
	v_mul_hi_u32 v17, v12, v11
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_mul_hi_u32 v15, v13, v15
	v_mul_hi_u32 v16, v2, v10
	v_add_nc_u32_e32 v10, -2, v5
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_nc_u32_e32 v12, v12, v17
	v_add_nc_u32_e32 v13, v13, v15
	s_delay_alu instid0(VALU_DEP_4)
	v_add_nc_u32_e32 v11, v2, v16
	s_waitcnt vmcnt(0)
	ds_store_b8 v0, v14
	s_waitcnt lgkmcnt(0)
	s_barrier
	buffer_gl0_inv
.LBB0_2:                                ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB0_3 Depth 2
                                        ;       Child Loop BB0_4 Depth 3
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_2)
	v_cmp_gt_u32_e32 vcc_lo, s4, v9
	s_mov_b32 s23, 0
	v_dual_mov_b32 v15, v8 :: v_dual_cndmask_b32 v0, 0, v9
	v_cmp_gt_i32_e32 vcc_lo, 0, v9
	v_cndmask_b32_e64 v1, v0, s4, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_mul_hi_u32 v0, v1, v11
	v_cmp_gt_u32_e64 s0, s13, v1
	v_cmp_le_u32_e64 s1, s20, v1
	s_or_b32 s24, s0, s1
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v0, v0, s16
	v_sub_nc_u32_e32 v0, v1, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_subrev_nc_u32_e32 v2, s16, v0
	v_cmp_le_u32_e32 vcc_lo, s16, v0
	v_cndmask_b32_e32 v0, v0, v2, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_subrev_nc_u32_e32 v2, s16, v0
	v_cmp_le_u32_e32 vcc_lo, s16, v0
	v_cndmask_b32_e32 v2, v0, v2, vcc_lo
	v_mul_lo_u32 v0, v1, s3
	v_cmp_eq_u32_e32 vcc_lo, v3, v1
	s_delay_alu instid0(VALU_DEP_3)
	v_mul_lo_u32 v14, v2, s17
.LBB0_3:                                ;   Parent Loop BB0_2 Depth=1
                                        ; =>  This Loop Header: Depth=2
                                        ;       Child Loop BB0_4 Depth 3
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_cmp_gt_u32_e64 s0, s5, v15
	s_mov_b32 s25, 0
	v_cndmask_b32_e64 v1, 0, v15, s0
	v_cmp_gt_i32_e64 s0, 0, v15
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e64 v16, v1, s5, s0
	v_mul_hi_u32 v1, v16, v12
	v_cmp_gt_u32_e64 s1, s14, v16
	v_cmp_le_u32_e64 s2, s21, v16
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	s_or_b32 s1, s24, s1
	v_mul_lo_u32 v1, v1, s17
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_or_b32 s1, s2, s1
	s_xor_b32 s26, s1, -1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_nc_u32_e32 v1, v16, v1
	v_subrev_nc_u32_e32 v2, s17, v1
	v_cmp_le_u32_e64 s0, s17, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e64 v1, v1, v2, s0
	v_subrev_nc_u32_e32 v2, s17, v1
	v_cmp_le_u32_e64 s0, s17, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_cndmask_b32_e64 v1, v1, v2, s0
	v_cmp_eq_u32_e64 s0, v4, v16
	v_add_nc_u32_e32 v17, v14, v1
	v_mad_u64_u32 v[1:2], null, v16, s6, v[0:1]
	v_mov_b32_e32 v16, v10
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_3)
	s_and_b32 s0, vcc_lo, s0
	v_mul_lo_u32 v2, v17, s12
	s_xor_b32 s2, s0, -1
.LBB0_4:                                ;   Parent Loop BB0_2 Depth=1
                                        ;     Parent Loop BB0_3 Depth=2
                                        ; =>    This Inner Loop Header: Depth=3
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v16, 1, v16
	v_cmp_gt_u32_e64 s0, s6, v16
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_cndmask_b32_e64 v17, 0, v16, s0
	v_cmp_gt_i32_e64 s0, 0, v16
	v_cndmask_b32_e64 v17, v17, s6, s0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cmp_ne_u32_e64 s0, v5, v17
	s_or_b32 s0, s2, s0
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_saveexec_b32 s27, s0
	s_cbranch_execz .LBB0_10
; %bb.5:                                ;   in Loop: Header=BB0_4 Depth=3
	v_cmp_le_u32_e64 s0, s15, v17
	v_cmp_gt_u32_e64 s1, s22, v17
	s_delay_alu instid0(VALU_DEP_2)
	s_and_b32 s0, s26, s0
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	s_and_b32 s0, s1, s0
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_saveexec_b32 s1, s0
	s_xor_b32 s1, exec_lo, s1
	s_cbranch_execz .LBB0_7
; %bb.6:                                ;   in Loop: Header=BB0_4 Depth=3
	v_mul_hi_u32 v18, v17, v13
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v18, v18, s12
	v_sub_nc_u32_e32 v17, v17, v18
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_subrev_nc_u32_e32 v18, s12, v17
	v_cmp_le_u32_e64 s0, s12, v17
	v_cndmask_b32_e64 v17, v17, v18, s0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_subrev_nc_u32_e32 v18, s12, v17
	v_cmp_le_u32_e64 s0, s12, v17
	v_cndmask_b32_e64 v17, v17, v18, s0
	s_delay_alu instid0(VALU_DEP_1)
	v_add3_u32 v17, v2, v17, 0
	ds_load_u8 v17, v17
	s_waitcnt lgkmcnt(0)
	v_add_nc_u16 v7, v17, v7
                                        ; implicit-def: $vgpr17
.LBB0_7:                                ; %Flow
                                        ;   in Loop: Header=BB0_4 Depth=3
	s_and_not1_saveexec_b32 s0, s1
	s_cbranch_execz .LBB0_9
; %bb.8:                                ;   in Loop: Header=BB0_4 Depth=3
	v_add_nc_u32_e32 v17, v1, v17
	global_load_u8 v17, v17, s[8:9]
	s_waitcnt vmcnt(0)
	v_add_nc_u16 v7, v17, v7
.LBB0_9:                                ; %Flow177
                                        ;   in Loop: Header=BB0_4 Depth=3
	s_or_b32 exec_lo, exec_lo, s0
.LBB0_10:                               ;   in Loop: Header=BB0_4 Depth=3
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	s_or_b32 exec_lo, exec_lo, s27
	v_cmp_gt_i32_e64 s0, v16, v5
	s_or_b32 s25, s0, s25
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s25
	s_cbranch_execnz .LBB0_4
; %bb.11:                               ;   in Loop: Header=BB0_3 Depth=2
	s_or_b32 exec_lo, exec_lo, s25
	v_add_nc_u32_e32 v1, 1, v15
	v_cmp_gt_i32_e64 s0, v15, v4
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mov_b32_e32 v15, v1
	s_or_b32 s23, s0, s23
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s23
	s_cbranch_execnz .LBB0_3
; %bb.12:                               ;   in Loop: Header=BB0_2 Depth=1
	s_or_b32 exec_lo, exec_lo, s23
	v_add_nc_u32_e32 v0, 1, v9
	v_cmp_gt_i32_e32 vcc_lo, v9, v3
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	v_mov_b32_e32 v9, v0
	s_or_b32 s19, vcc_lo, s19
	s_and_not1_b32 exec_lo, exec_lo, s19
	s_cbranch_execnz .LBB0_2
; %bb.13:                               ; %.sink.split
	s_or_b32 exec_lo, exec_lo, s19
	v_bfe_i32 v0, v7, 0, 8
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_cmp_le_u32_e32 vcc_lo, s7, v0
	v_cmp_ge_u32_e64 s0, s18, v0
	v_add_co_u32 v0, s1, s10, v6
	v_add_co_ci_u32_e64 v1, null, s11, 0, s1
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b32 s0, vcc_lo, s0
	v_cndmask_b32_e64 v2, 0, 1, s0
	global_store_b8 v[0:1], v2, off
.LBB0_14:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z13lifeItrKernelPKcPcjjjjj
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 296
		.amdhsa_user_sgpr_count 13
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 1
		.amdhsa_system_sgpr_workgroup_id_z 1
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 2
		.amdhsa_next_free_vgpr 19
		.amdhsa_next_free_sgpr 28
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end0:
	.size	_Z13lifeItrKernelPKcPcjjjjj, .Lfunc_end0-_Z13lifeItrKernelPKcPcjjjjj
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 1212
; NumSgprs: 30
; NumVgprs: 19
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 3
; VGPRBlocks: 2
; NumSGPRsForWavesPerEU: 30
; NumVGPRsForWavesPerEU: 19
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 13
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 1
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 2
	.text
	.protected	_Z18sumNeighborsKernelPKcPcjjj ; -- Begin function _Z18sumNeighborsKernelPKcPcjjj
	.globl	_Z18sumNeighborsKernelPKcPcjjj
	.p2align	8
	.type	_Z18sumNeighborsKernelPKcPcjjj,@function
_Z18sumNeighborsKernelPKcPcjjj:         ; @_Z18sumNeighborsKernelPKcPcjjj
; %bb.0:
	s_clause 0x1
	s_load_b64 s[2:3], s[0:1], 0x2c
	s_load_b128 s[4:7], s[0:1], 0x10
	v_and_b32_e32 v2, 0x3ff, v0
	v_bfe_u32 v3, v0, 10, 10
	v_bfe_u32 v4, v0, 20, 10
	s_waitcnt lgkmcnt(0)
	s_lshr_b32 s7, s2, 16
	s_and_b32 s2, s2, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1)
	v_mad_u64_u32 v[0:1], null, s13, s2, v[2:3]
	v_mad_u64_u32 v[1:2], null, s14, s7, v[3:4]
	s_and_b32 s2, s3, 0xffff
	s_add_i32 s3, s5, -1
	v_mad_u64_u32 v[2:3], null, s15, s2, v[4:5]
	s_add_i32 s2, s4, -1
	s_delay_alu instid0(VALU_DEP_3) | instid1(SALU_CYCLE_1)
	v_cmp_ge_u32_e32 vcc_lo, s2, v0
	s_delay_alu instid0(VALU_DEP_3)
	v_cmp_ge_u32_e64 s2, s3, v1
	s_add_i32 s3, s6, -1
	s_delay_alu instid0(VALU_DEP_3) | instid1(SALU_CYCLE_1)
	v_cmp_ge_u32_e64 s3, s3, v2
	s_delay_alu instid0(VALU_DEP_2)
	s_and_b32 s2, vcc_lo, s2
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	s_and_b32 s2, s2, s3
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_saveexec_b32 s3, s2
	s_cbranch_execz .LBB1_10
; %bb.1:
	s_load_b128 s[8:11], s[0:1], 0x0
	v_dual_mov_b32 v6, 0 :: v_dual_add_nc_u32 v9, -1, v0
	v_add_nc_u32_e32 v7, -1, v1
	v_add_nc_u32_e32 v8, -2, v2
	s_mul_i32 s1, s6, s5
	s_mov_b32 s2, 0
.LBB1_2:                                ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB1_3 Depth 2
                                        ;       Child Loop BB1_4 Depth 3
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_2)
	v_cmp_gt_u32_e32 vcc_lo, s4, v9
	s_mov_b32 s3, 0
	v_dual_mov_b32 v10, v7 :: v_dual_cndmask_b32 v3, 0, v9
	v_cmp_gt_i32_e32 vcc_lo, 0, v9
	v_cndmask_b32_e64 v4, v3, s4, vcc_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_mul_lo_u32 v3, v4, s1
	v_cmp_eq_u32_e32 vcc_lo, v0, v4
.LBB1_3:                                ;   Parent Loop BB1_2 Depth=1
                                        ; =>  This Loop Header: Depth=2
                                        ;       Child Loop BB1_4 Depth 3
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_cmp_gt_u32_e64 s0, s5, v10
	s_mov_b32 s7, 0
	v_cndmask_b32_e64 v4, 0, v10, s0
	v_cmp_gt_i32_e64 s0, 0, v10
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e64 v11, v4, s5, s0
	v_cmp_eq_u32_e64 s0, v1, v11
	v_mad_u64_u32 v[4:5], null, v11, s6, v[3:4]
	v_mov_b32_e32 v5, v8
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b32 s0, vcc_lo, s0
	s_xor_b32 s12, s0, -1
.LBB1_4:                                ;   Parent Loop BB1_2 Depth=1
                                        ;     Parent Loop BB1_3 Depth=2
                                        ; =>    This Inner Loop Header: Depth=3
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v5, 1, v5
	v_cmp_gt_u32_e64 s0, s6, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_cndmask_b32_e64 v11, 0, v5, s0
	v_cmp_gt_i32_e64 s0, 0, v5
	v_cndmask_b32_e64 v11, v11, s6, s0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cmp_ne_u32_e64 s0, v2, v11
	s_or_b32 s13, s12, s0
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_saveexec_b32 s0, s13
	s_cbranch_execz .LBB1_6
; %bb.5:                                ;   in Loop: Header=BB1_4 Depth=3
	v_add_nc_u32_e32 v11, v4, v11
	s_waitcnt lgkmcnt(0)
	global_load_u8 v11, v11, s[8:9]
	s_waitcnt vmcnt(0)
	v_add_nc_u16 v6, v11, v6
.LBB1_6:                                ;   in Loop: Header=BB1_4 Depth=3
	s_or_b32 exec_lo, exec_lo, s0
	v_cmp_gt_i32_e64 s0, v5, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_or_b32 s7, s0, s7
	s_and_not1_b32 exec_lo, exec_lo, s7
	s_cbranch_execnz .LBB1_4
; %bb.7:                                ;   in Loop: Header=BB1_3 Depth=2
	s_or_b32 exec_lo, exec_lo, s7
	v_add_nc_u32_e32 v4, 1, v10
	v_cmp_gt_i32_e64 s0, v10, v1
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mov_b32_e32 v10, v4
	s_or_b32 s3, s0, s3
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s3
	s_cbranch_execnz .LBB1_3
; %bb.8:                                ;   in Loop: Header=BB1_2 Depth=1
	s_or_b32 exec_lo, exec_lo, s3
	v_add_nc_u32_e32 v3, 1, v9
	v_cmp_gt_i32_e32 vcc_lo, v9, v0
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	v_mov_b32_e32 v9, v3
	s_or_b32 s2, vcc_lo, s2
	s_and_not1_b32 exec_lo, exec_lo, s2
	s_cbranch_execnz .LBB1_2
; %bb.9:
	s_or_b32 exec_lo, exec_lo, s2
	v_mul_lo_u32 v0, v0, s1
	v_mul_lo_u32 v1, v1, s6
	s_delay_alu instid0(VALU_DEP_1)
	v_add3_u32 v0, v1, v0, v2
	s_waitcnt lgkmcnt(0)
	global_store_b8 v0, v6, s[10:11]
.LBB1_10:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z18sumNeighborsKernelPKcPcjjj
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 288
		.amdhsa_user_sgpr_count 13
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 1
		.amdhsa_system_sgpr_workgroup_id_z 1
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 2
		.amdhsa_next_free_vgpr 12
		.amdhsa_next_free_sgpr 16
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end1:
	.size	_Z18sumNeighborsKernelPKcPcjjj, .Lfunc_end1-_Z18sumNeighborsKernelPKcPcjjj
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 580
; NumSgprs: 18
; NumVgprs: 12
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 12
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 13
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 1
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 2
	.text
	.protected	_Z18setAliveDeadKernelPKcPcjjjjj ; -- Begin function _Z18setAliveDeadKernelPKcPcjjjjj
	.globl	_Z18setAliveDeadKernelPKcPcjjjjj
	.p2align	8
	.type	_Z18setAliveDeadKernelPKcPcjjjjj,@function
_Z18setAliveDeadKernelPKcPcjjjjj:       ; @_Z18setAliveDeadKernelPKcPcjjjjj
; %bb.0:
	s_clause 0x1
	s_load_b64 s[2:3], s[0:1], 0x34
	s_load_b128 s[4:7], s[0:1], 0x10
	v_and_b32_e32 v2, 0x3ff, v0
	v_bfe_u32 v3, v0, 10, 10
	v_bfe_u32 v4, v0, 20, 10
	s_waitcnt lgkmcnt(0)
	s_lshr_b32 s8, s2, 16
	s_and_b32 s2, s2, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1)
	v_mad_u64_u32 v[0:1], null, s13, s2, v[2:3]
	v_mad_u64_u32 v[1:2], null, s14, s8, v[3:4]
	s_and_b32 s2, s3, 0xffff
	s_add_i32 s3, s5, -1
	v_mad_u64_u32 v[2:3], null, s15, s2, v[4:5]
	s_add_i32 s2, s4, -1
	s_delay_alu instid0(VALU_DEP_3) | instid1(SALU_CYCLE_1)
	v_cmp_ge_u32_e32 vcc_lo, s2, v0
	s_delay_alu instid0(VALU_DEP_3)
	v_cmp_ge_u32_e64 s2, s3, v1
	s_add_i32 s3, s6, -1
	s_delay_alu instid0(VALU_DEP_3) | instid1(SALU_CYCLE_1)
	v_cmp_ge_u32_e64 s3, s3, v2
	s_delay_alu instid0(VALU_DEP_2)
	s_and_b32 s2, vcc_lo, s2
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	s_and_b32 s2, s2, s3
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_saveexec_b32 s3, s2
	s_cbranch_execz .LBB2_2
; %bb.1:                                ; %.sink.split
	s_load_b128 s[8:11], s[0:1], 0x0
	v_mad_u64_u32 v[3:4], null, v0, s5, v[1:2]
	s_load_b32 s0, s[0:1], 0x20
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[0:1], null, v3, s6, v[2:3]
	v_ashrrev_i32_e32 v3, 31, v0
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v1, vcc_lo, s8, v0
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v2, vcc_lo, s9, v3, vcc_lo
	v_add_co_u32 v0, s1, s10, v0
	global_load_i8 v1, v[1:2], off
	s_waitcnt vmcnt(0)
	v_cmp_le_u32_e32 vcc_lo, s7, v1
	v_cmp_ge_u32_e64 s0, s0, v1
	v_add_co_ci_u32_e64 v1, s1, s11, v3, s1
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b32 s0, vcc_lo, s0
	v_cndmask_b32_e64 v2, 0, 1, s0
	global_store_b8 v[0:1], v2, off
.LBB2_2:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z18setAliveDeadKernelPKcPcjjjjj
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 296
		.amdhsa_user_sgpr_count 13
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 1
		.amdhsa_system_sgpr_workgroup_id_z 1
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 2
		.amdhsa_next_free_vgpr 6
		.amdhsa_next_free_sgpr 16
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end2:
	.size	_Z18setAliveDeadKernelPKcPcjjjjj, .Lfunc_end2-_Z18setAliveDeadKernelPKcPcjjjjj
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 304
; NumSgprs: 18
; NumVgprs: 6
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 6
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 13
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 1
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 2
	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_
	.amdgpu_metadata
---
amdhsa.kernels:
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .offset:         16
        .size:           4
        .value_kind:     by_value
      - .offset:         20
        .size:           4
        .value_kind:     by_value
      - .offset:         24
        .size:           4
        .value_kind:     by_value
      - .offset:         28
        .size:           4
        .value_kind:     by_value
      - .offset:         32
        .size:           4
        .value_kind:     by_value
      - .offset:         40
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         44
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         48
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         52
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         54
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         56
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         58
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         60
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         62
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         88
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         96
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         104
        .size:           2
        .value_kind:     hidden_grid_dims
      - .offset:         160
        .size:           4
        .value_kind:     hidden_dynamic_lds_size
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 296
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z13lifeItrKernelPKcPcjjjjj
    .private_segment_fixed_size: 0
    .sgpr_count:     30
    .sgpr_spill_count: 0
    .symbol:         _Z13lifeItrKernelPKcPcjjjjj.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     19
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .offset:         16
        .size:           4
        .value_kind:     by_value
      - .offset:         20
        .size:           4
        .value_kind:     by_value
      - .offset:         24
        .size:           4
        .value_kind:     by_value
      - .offset:         32
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         36
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         40
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         44
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         46
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         48
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         50
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         52
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         54
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         88
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         96
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 288
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z18sumNeighborsKernelPKcPcjjj
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z18sumNeighborsKernelPKcPcjjj.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     12
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .offset:         16
        .size:           4
        .value_kind:     by_value
      - .offset:         20
        .size:           4
        .value_kind:     by_value
      - .offset:         24
        .size:           4
        .value_kind:     by_value
      - .offset:         28
        .size:           4
        .value_kind:     by_value
      - .offset:         32
        .size:           4
        .value_kind:     by_value
      - .offset:         40
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         44
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         48
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         52
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         54
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         56
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         58
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         60
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         62
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         88
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         96
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         104
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 296
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z18setAliveDeadKernelPKcPcjjjjj
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z18setAliveDeadKernelPKcPcjjjjj.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     6
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
