#include <sys/types.h>
#include <sys/ioctl.h>
#include <sys/mman.h>
#include <linux/fb.h>
#include <fcntl.h>

#include <unistd.h>
#include <stdio.h>
#include <stdlib.h>
#include <time.h>

/*
   Conway's Game of Life running on nVidia GPUs.

   Each GPU thread updates the state of a single cell in the playing field, then the current generation is
   displayed in Linux's framebuffer (this means the program must be run outside X Window System).
   Since this program is just a demonstration, the data model is NOT decoupled from its representation.
*/



/**
   Cell drawing function

   Updates the state of a single cell and draws it on screen.
   Each GPU thread updates a single cell (the mapping is linear, so even if the field's size exceeds the
   number of computation units of the GPU, the pixels assigned to the same thread are far away from each other),
   based on the number of neighbors.
   Since the number of neighbors is only read (and this function is called AFTER the neighbors array is written)
   and each thread writes only its cell, there are no race conditions.

   @param field Array of cell states. 1: live cell, 0: dead cell.
   @param neighbors Array of neighbors' count (e.g. neighbors[42] = 4 means cell 42 has 4 neighbors).
   @param screen Framebuffer (the function assumes the buffer uses 32 bit colors).
   @param w Width of the framebuffer.
   @param h Height of the framebuffer.
*/
__global__ void updateCell(char *field, char *neighbors, char *screen, int w, int h)
{
  int i = threadIdx.x + blockIdx.x * blockDim.x; /* Pick the cell based on thread and block IDs. */
  
  if (i < w * h)
  {
    if ((field[i] != 0 && (neighbors[i] == 3 || neighbors[i] == 2)) || /* Survive rule: S23 */
	(field[i] == 0 && neighbors[i] == 3))                            /* Birth rule: B3    */
    {
      /* Live cell */
      field[i] = 1;
      screen[4 * i] =  0xff;     /* blue  */
      screen[4 * i + 1] =  0xff; /* green */
      screen[4 * i + 2] =  0xff; /* red   */
      screen[4 * i + 3] =  0xff; /* alpha */
    }
    else
    {
      /* Dead cell */
      field[i] = 0;
      screen[4 * i] =  0x00;
      screen[4 * i + 1] =  0x00;
      screen[4 * i + 2] =  0x00;
      screen[4 * i + 3] =  0x00;
    }
  }
}

/**
   Neighbors counting function

   Checks the eight neighbors' state and counts the live ones.
   Each GPU thread writes the neighbors count for a single cell while the field is not modified, so there
   are no race conditions.

   @param field Array of cell states. 1: live cell, 0: dead cell.
   @param neighbors Array of neighbors' count (e.g. neighbors[42] = 4 means cell 42 has 4 neighbors).
   @param w Width of the framebuffer.
   @param h Height of the framebuffer.
 */
__global__ void countNeighbors(char *field, char *neighbors, int w, int h)
{
  int i = threadIdx.x + blockIdx.x * blockDim.x;
  int x = i % w;
  int y = i / w * w; /* Rounding to the closest multiple of w. */
  int xp1 = (i + 1) % w;
  int xm1 = (i + w - 1) % w;
  int yp1 = ((i + w) % (h * w)) / w * w;
  int ym1 = ((i - w + h * w) % (h * w)) / w * w;
   
  neighbors[i] = 0;
  if (field[xm1 + ym1] != 0)
    neighbors[i]++;
  if (field[xm1 + y] != 0)
    neighbors[i]++;
  if (field[xm1 + yp1] != 0)
    neighbors[i]++;
  if (field[x + ym1] != 0)
    neighbors[i]++;
  if (field[x + yp1] != 0)
    neighbors[i]++;
  if (field[xp1 + ym1] != 0)
    neighbors[i]++;
  if (field[xp1 + y] != 0)
    neighbors[i]++;
  if (field[xp1 + yp1] != 0)
    neighbors[i]++;
}

/**
   Game of Life wrapper function.

   Initializes the playing field with a random initial state and iterates over Life's generations.
   This function takes care of synchronizing GPU threads so that no race conditions are triggered.

   @param screen Framebuffer.
   @param w Width of the framebuffer.
   @param h Height of the framebuffer.
 */
void life(char *screen, int w, int h)
{
  char field[w * h];
  
  /* Generate a random initial playing field. */
  srand(time(NULL));
  for (int i = 0; i < w * h; i++)
    field[i] = (rand() % 100 < 75)? 1: 0;

  /* Find the maximum number of threads for the current GPU. */
  struct cudaDeviceProp properties;
  cudaGetDeviceProperties(&properties, 0);
  char *d_field, *d_screen, *d_neighbors;
  int maxthds = properties.maxThreadsPerBlock;

  /*
     Copy memory to the GPU's internal RAM. Since this process is slow, it's done only once and the state
     of the playing field is effectively kept in GPU's memory, updating the main RAM only to display the
     result on screen.
  */
  cudaMalloc((void **) &d_field, sizeof(char) * w * h);
  cudaMalloc((void **) &d_neighbors, sizeof(char) * w * h);
  cudaMalloc((void **) &d_screen, sizeof(char) * w * h * 4);
  cudaMemcpy(d_field, field, sizeof(char) * w * h, cudaMemcpyHostToDevice);
  cudaMemcpy(d_screen, screen, sizeof(char) * w * h * 4, cudaMemcpyHostToDevice);
  
  
  while (1) /* Maybe a termination condition would have been more elegant... */
  {
    /*
       Count the neighbors, synchronize the threads (in order to avoid race conditions),
       then update the state of each cell.
    */
    countNeighbors<<<(w * h + maxthds - 1) / maxthds, maxthds>>>(d_field, d_neighbors, w, h);
    cudaDeviceSynchronize();
    updateCell<<<(w * h + maxthds - 1) / maxthds, maxthds>>>(d_field, d_neighbors, d_screen, w, h);

    /*
       Synchronize the threads again and finally copy the framebuffer from GPU's internal memory to main RAM
       (and therefore display the result).
    */
    cudaDeviceSynchronize();
    cudaMemcpy(screen, d_screen, sizeof(char) * w * h * 4, cudaMemcpyDeviceToHost);
  }
  
  
  /* Free the GPU's internal memory. Since the loop condition is never false, this cleanup is never performed. */
  cudaFree(d_field);
  cudaFree(d_neighbors);
  cudaFree(d_screen);
  
  /* printf("%s\n", cudaGetErrorString(cudaGetLastError())); Debug message. */
}


/* Semplice funzione main che controlla il framebuffer. Copiata da: https://stackoverflow.com/a/1830865

/**
   Program entrypoint.

   Initializes the main framebuffer (/dev/fb0), as suggested on: https://stackoverflow.com/a/1830865
   then passes it to the life() function.

   @return 0 if no error occurred, 1 otherwise.
 */
int main()
{
  struct fb_var_screeninfo screen_info;
  struct fb_fix_screeninfo fixed_info;
  char *buffer = NULL;
  size_t buflen;
  int fd = -1;
  int r = 1;

  fd = open("/dev/fb0", O_RDWR);
  if (fd >= 0)
  {
    if (!ioctl(fd, FBIOGET_VSCREENINFO, &screen_info) &&
        !ioctl(fd, FBIOGET_FSCREENINFO, &fixed_info))
    {
      buflen = screen_info.yres_virtual * fixed_info.line_length;
      buffer = (char *) mmap(NULL, buflen,
                       PROT_READ|PROT_WRITE, MAP_SHARED, fd, 0);
      if (buffer != MAP_FAILED)
      {
	life(buffer, screen_info.xres_virtual, screen_info.yres_virtual); /* Start the game. */
	r = 0;
      }
      else
      {
	perror("mmap");
      }
    }
    else
      {
	perror("ioctl");
      }
   }
   else
   {
       perror("open");
   }

   if (buffer && buffer != MAP_FAILED)
     munmap(buffer, buflen);
   if (fd >= 0)
     close(fd);

   return r;
}

	.file	"tmpxft_0032622d_00000000-6_life.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2048:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2048:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z36__device_stub__Z10updateCellPcS_S_iiPcS_S_ii
	.type	_Z36__device_stub__Z10updateCellPcS_S_iiPcS_S_ii, @function
_Z36__device_stub__Z10updateCellPcS_S_iiPcS_S_ii:
.LFB2070:
	.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
	movq	%rdx, 8(%rsp)
	leaq	32(%rsp), %rdx
	movl	%ecx, 4(%rsp)
	leaq	40(%rsp), %rcx
	movl	%r8d, (%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	8(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 120(%rsp)
	movq	%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	.L2
	pushq	40(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z10updateCellPcS_S_ii(%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
.L2:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L4
	call	__stack_chk_fail@PLT
.L4:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2070:
	.size	_Z36__device_stub__Z10updateCellPcS_S_iiPcS_S_ii, .-_Z36__device_stub__Z10updateCellPcS_S_iiPcS_S_ii
	.globl	_Z10updateCellPcS_S_ii
	.type	_Z10updateCellPcS_S_ii, @function
_Z10updateCellPcS_S_ii:
.LFB2071:
	.cfi_startproc
	endbr64
	jmp	_Z36__device_stub__Z10updateCellPcS_S_iiPcS_S_ii
	.cfi_endproc
.LFE2071:
	.size	_Z10updateCellPcS_S_ii, .-_Z10updateCellPcS_S_ii
	.globl	_Z38__device_stub__Z14countNeighborsPcS_iiPcS_ii
	.type	_Z38__device_stub__Z14countNeighborsPcS_iiPcS_ii, @function
_Z38__device_stub__Z14countNeighborsPcS_iiPcS_ii:
.LFB2072:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movq	%rdi, 24(%rsp)
	leaq	56(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	68(%rsp), %rsi
	movl	%edx, 12(%rsp)
	leaq	40(%rsp), %rdx
	movl	%ecx, 8(%rsp)
	leaq	48(%rsp), %rcx
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 104(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 128(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 56(%rsp)
	movq	%rax, 68(%rsp)
	movl	$1, 76(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L8
	pushq	48(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z14countNeighborsPcS_ii(%rip), %rdi
	pushq	48(%rsp)
	.cfi_def_cfa_offset 176
	movq	84(%rsp), %rcx
	movl	92(%rsp), %r8d
	movq	72(%rsp), %rsi
	movl	80(%rsp), %edx
	leaq	120(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 168
	popq	%rdx
	.cfi_def_cfa_offset 160
.L8:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L10
	call	__stack_chk_fail@PLT
.L10:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2072:
	.size	_Z38__device_stub__Z14countNeighborsPcS_iiPcS_ii, .-_Z38__device_stub__Z14countNeighborsPcS_iiPcS_ii
	.globl	_Z14countNeighborsPcS_ii
	.type	_Z14countNeighborsPcS_ii, @function
_Z14countNeighborsPcS_ii:
.LFB2073:
	.cfi_startproc
	endbr64
	jmp	_Z38__device_stub__Z14countNeighborsPcS_iiPcS_ii
	.cfi_endproc
.LFE2073:
	.size	_Z14countNeighborsPcS_ii, .-_Z14countNeighborsPcS_ii
	.globl	_Z4lifePcii
	.type	_Z4lifePcii, @function
_Z4lifePcii:
.LFB2044:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	movq	%rsp, %rbp
	.cfi_def_cfa_register 6
	pushq	%r15
	.cfi_offset 15, -24
	movl	%esi, %r15d
	pushq	%r14
	imull	%edx, %r15d
	pushq	%r13
	.cfi_offset 14, -32
	.cfi_offset 13, -40
	movl	%edx, %r13d
	pushq	%r12
	.cfi_offset 12, -48
	movl	%esi, %r12d
	pushq	%rbx
	subq	$1128, %rsp
	.cfi_offset 3, -56
	movq	%rdi, -1160(%rbp)
	movq	%rsp, %rcx
	movq	%fs:40, %rax
	movq	%rax, -56(%rbp)
	xorl	%eax, %eax
	movslq	%r15d, %rax
	addq	$15, %rax
	movq	%rax, %rdx
	andq	$-4096, %rax
	andq	$-16, %rdx
	subq	%rax, %rcx
.L14:
	cmpq	%rcx, %rsp
	je	.L15
	subq	$4096, %rsp
	orq	$0, 4088(%rsp)
	jmp	.L14
.L15:
	movq	%rdx, %rax
	andl	$4095, %eax
	subq	%rax, %rsp
	testq	%rax, %rax
	je	.L16
	orq	$0, -8(%rsp,%rax)
.L16:
	xorl	%edi, %edi
	movq	%rsp, -1152(%rbp)
	xorl	%r14d, %r14d
	movl	$100, %ebx
	call	time@PLT
	movl	%eax, %edi
	call	srand@PLT
.L17:
	cmpl	%r14d, %r15d
	jle	.L27
	call	rand@PLT
	cltd
	idivl	%ebx
	movq	-1152(%rbp), %rax
	cmpl	$74, %edx
	setle	(%rax,%r14)
	incq	%r14
	jmp	.L17
.L27:
	leaq	-1088(%rbp), %rdi
	xorl	%esi, %esi
	movslq	%r12d, %rbx
	call	cudaGetDeviceProperties_v2@PLT
	movslq	%r13d, %rax
	leaq	-1136(%rbp), %rdi
	movl	-768(%rbp), %r14d
	imulq	%rax, %rbx
	movq	%rbx, %rsi
	call	cudaMalloc@PLT
	movq	%rbx, %rsi
	leaq	-1120(%rbp), %rdi
	call	cudaMalloc@PLT
	leaq	0(,%rbx,4), %rax
	leaq	-1128(%rbp), %rdi
	movq	%rax, %rsi
	movq	%rax, -1144(%rbp)
	call	cudaMalloc@PLT
	movq	-1152(%rbp), %rsi
	movq	%rbx, %rdx
	movq	-1136(%rbp), %rdi
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	movq	-1144(%rbp), %rdx
	movl	$1, %ecx
	movq	-1160(%rbp), %rsi
	movq	-1128(%rbp), %rdi
	call	cudaMemcpy@PLT
	leal	-1(%r15,%r14), %eax
	movabsq	$4294967297, %r15
	cltd
	idivl	%r14d
	movl	%eax, %ebx
.L21:
	movq	%r15, -1096(%rbp)
	movl	-1092(%rbp), %ecx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movq	%r15, -1108(%rbp)
	movl	-1104(%rbp), %esi
	movl	%r14d, -1100(%rbp)
	movq	-1100(%rbp), %rdx
	movl	%ebx, -1112(%rbp)
	movq	-1112(%rbp), %rdi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L19
	movq	-1120(%rbp), %rsi
	movq	-1136(%rbp), %rdi
	movl	%r13d, %ecx
	movl	%r12d, %edx
	call	_Z38__device_stub__Z14countNeighborsPcS_iiPcS_ii
.L19:
	call	cudaDeviceSynchronize@PLT
	movq	%r15, -1096(%rbp)
	xorl	%r9d, %r9d
	movl	-1092(%rbp), %ecx
	movq	%r15, -1108(%rbp)
	movl	-1104(%rbp), %esi
	xorl	%r8d, %r8d
	movl	%r14d, -1100(%rbp)
	movq	-1100(%rbp), %rdx
	movl	%ebx, -1112(%rbp)
	movq	-1112(%rbp), %rdi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L20
	movq	-1128(%rbp), %rdx
	movq	-1120(%rbp), %rsi
	movl	%r13d, %r8d
	movl	%r12d, %ecx
	movq	-1136(%rbp), %rdi
	call	_Z36__device_stub__Z10updateCellPcS_S_iiPcS_S_ii
.L20:
	call	cudaDeviceSynchronize@PLT
	movq	-1144(%rbp), %rdx
	movl	$2, %ecx
	movq	-1128(%rbp), %rsi
	movq	-1160(%rbp), %rdi
	call	cudaMemcpy@PLT
	jmp	.L21
	.cfi_endproc
.LFE2044:
	.size	_Z4lifePcii, .-_Z4lifePcii
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"/dev/fb0"
.LC1:
	.string	"ioctl"
.LC2:
	.string	"mmap"
.LC3:
	.string	"open"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2045:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	movl	$2, %esi
	leaq	.LC0(%rip), %rdi
	subq	$256, %rsp
	.cfi_def_cfa_offset 272
	movq	%fs:40, %rax
	movq	%rax, 248(%rsp)
	xorl	%eax, %eax
	call	__open_2@PLT
	testl	%eax, %eax
	js	.L29
	movl	%eax, %ebx
	movl	%eax, %edi
	leaq	88(%rsp), %rdx
	xorl	%eax, %eax
	movl	$17920, %esi
	call	ioctl@PLT
	testl	%eax, %eax
	jne	.L30
	leaq	8(%rsp), %rdx
	movl	$17922, %esi
	movl	%ebx, %edi
	call	ioctl@PLT
	testl	%eax, %eax
	jne	.L30
	movl	100(%rsp), %esi
	imull	56(%rsp), %esi
	xorl	%edi, %edi
	xorl	%r9d, %r9d
	movl	%ebx, %r8d
	movl	$1, %ecx
	movl	$3, %edx
	call	mmap@PLT
	movq	%rax, %rdi
	cmpq	$-1, %rax
	jne	.L31
	jmp	.L41
.L30:
	leaq	.LC1(%rip), %rdi
	jmp	.L40
.L31:
	movq	248(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L38
	movl	100(%rsp), %edx
	movl	96(%rsp), %esi
	call	_Z4lifePcii
.L41:
	leaq	.LC2(%rip), %rdi
.L40:
	call	perror@PLT
	movl	%ebx, %edi
	call	close@PLT
	jmp	.L35
.L29:
	leaq	.LC3(%rip), %rdi
	call	perror@PLT
.L35:
	movq	248(%rsp), %rax
	subq	%fs:40, %rax
	je	.L36
.L38:
	call	__stack_chk_fail@PLT
.L36:
	addq	$256, %rsp
	.cfi_def_cfa_offset 16
	movl	$1, %eax
	popq	%rbx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2045:
	.size	main, .-main
	.section	.rodata.str1.1
.LC4:
	.string	"_Z14countNeighborsPcS_ii"
.LC5:
	.string	"_Z10updateCellPcS_S_ii"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2075:
	.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	.LC4(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z14countNeighborsPcS_ii(%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	.LC5(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z10updateCellPcS_S_ii(%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
.LFE2075:
	.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
	.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:


#include <hip/hip_runtime.h>
#include <sys/types.h>
#include <sys/ioctl.h>
#include <sys/mman.h>
#include <linux/fb.h>
#include <fcntl.h>

#include <unistd.h>
#include <stdio.h>
#include <stdlib.h>
#include <time.h>

/*
   Conway's Game of Life running on nVidia GPUs.

   Each GPU thread updates the state of a single cell in the playing field, then the current generation is
   displayed in Linux's framebuffer (this means the program must be run outside X Window System).
   Since this program is just a demonstration, the data model is NOT decoupled from its representation.
*/



/**
   Cell drawing function

   Updates the state of a single cell and draws it on screen.
   Each GPU thread updates a single cell (the mapping is linear, so even if the field's size exceeds the
   number of computation units of the GPU, the pixels assigned to the same thread are far away from each other),
   based on the number of neighbors.
   Since the number of neighbors is only read (and this function is called AFTER the neighbors array is written)
   and each thread writes only its cell, there are no race conditions.

   @param field Array of cell states. 1: live cell, 0: dead cell.
   @param neighbors Array of neighbors' count (e.g. neighbors[42] = 4 means cell 42 has 4 neighbors).
   @param screen Framebuffer (the function assumes the buffer uses 32 bit colors).
   @param w Width of the framebuffer.
   @param h Height of the framebuffer.
*/
__global__ void updateCell(char *field, char *neighbors, char *screen, int w, int h)
{
  int i = threadIdx.x + blockIdx.x * blockDim.x; /* Pick the cell based on thread and block IDs. */
  
  if (i < w * h)
  {
    if ((field[i] != 0 && (neighbors[i] == 3 || neighbors[i] == 2)) || /* Survive rule: S23 */
	(field[i] == 0 && neighbors[i] == 3))                            /* Birth rule: B3    */
    {
      /* Live cell */
      field[i] = 1;
      screen[4 * i] =  0xff;     /* blue  */
      screen[4 * i + 1] =  0xff; /* green */
      screen[4 * i + 2] =  0xff; /* red   */
      screen[4 * i + 3] =  0xff; /* alpha */
    }
    else
    {
      /* Dead cell */
      field[i] = 0;
      screen[4 * i] =  0x00;
      screen[4 * i + 1] =  0x00;
      screen[4 * i + 2] =  0x00;
      screen[4 * i + 3] =  0x00;
    }
  }
}

/**
   Neighbors counting function

   Checks the eight neighbors' state and counts the live ones.
   Each GPU thread writes the neighbors count for a single cell while the field is not modified, so there
   are no race conditions.

   @param field Array of cell states. 1: live cell, 0: dead cell.
   @param neighbors Array of neighbors' count (e.g. neighbors[42] = 4 means cell 42 has 4 neighbors).
   @param w Width of the framebuffer.
   @param h Height of the framebuffer.
 */
__global__ void countNeighbors(char *field, char *neighbors, int w, int h)
{
  int i = threadIdx.x + blockIdx.x * blockDim.x;
  int x = i % w;
  int y = i / w * w; /* Rounding to the closest multiple of w. */
  int xp1 = (i + 1) % w;
  int xm1 = (i + w - 1) % w;
  int yp1 = ((i + w) % (h * w)) / w * w;
  int ym1 = ((i - w + h * w) % (h * w)) / w * w;
   
  neighbors[i] = 0;
  if (field[xm1 + ym1] != 0)
    neighbors[i]++;
  if (field[xm1 + y] != 0)
    neighbors[i]++;
  if (field[xm1 + yp1] != 0)
    neighbors[i]++;
  if (field[x + ym1] != 0)
    neighbors[i]++;
  if (field[x + yp1] != 0)
    neighbors[i]++;
  if (field[xp1 + ym1] != 0)
    neighbors[i]++;
  if (field[xp1 + y] != 0)
    neighbors[i]++;
  if (field[xp1 + yp1] != 0)
    neighbors[i]++;
}

/**
   Game of Life wrapper function.

   Initializes the playing field with a random initial state and iterates over Life's generations.
   This function takes care of synchronizing GPU threads so that no race conditions are triggered.

   @param screen Framebuffer.
   @param w Width of the framebuffer.
   @param h Height of the framebuffer.
 */
void life(char *screen, int w, int h)
{
  char field[w * h];
  
  /* Generate a random initial playing field. */
  srand(time(NULL));
  for (int i = 0; i < w * h; i++)
    field[i] = (rand() % 100 < 75)? 1: 0;

  /* Find the maximum number of threads for the current GPU. */
  struct hipDeviceProp_t properties;
  hipGetDeviceProperties(&properties, 0);
  char *d_field, *d_screen, *d_neighbors;
  int maxthds = properties.maxThreadsPerBlock;

  /*
     Copy memory to the GPU's internal RAM. Since this process is slow, it's done only once and the state
     of the playing field is effectively kept in GPU's memory, updating the main RAM only to display the
     result on screen.
  */
  hipMalloc((void **) &d_field, sizeof(char) * w * h);
  hipMalloc((void **) &d_neighbors, sizeof(char) * w * h);
  hipMalloc((void **) &d_screen, sizeof(char) * w * h * 4);
  hipMemcpy(d_field, field, sizeof(char) * w * h, hipMemcpyHostToDevice);
  hipMemcpy(d_screen, screen, sizeof(char) * w * h * 4, hipMemcpyHostToDevice);
  
  
  while (1) /* Maybe a termination condition would have been more elegant... */
  {
    /*
       Count the neighbors, synchronize the threads (in order to avoid race conditions),
       then update the state of each cell.
    */
    countNeighbors<<<(w * h + maxthds - 1) / maxthds, maxthds>>>(d_field, d_neighbors, w, h);
    hipDeviceSynchronize();
    updateCell<<<(w * h + maxthds - 1) / maxthds, maxthds>>>(d_field, d_neighbors, d_screen, w, h);

    /*
       Synchronize the threads again and finally copy the framebuffer from GPU's internal memory to main RAM
       (and therefore display the result).
    */
    hipDeviceSynchronize();
    hipMemcpy(screen, d_screen, sizeof(char) * w * h * 4, hipMemcpyDeviceToHost);
  }
  
  
  /* Free the GPU's internal memory. Since the loop condition is never false, this cleanup is never performed. */
  hipFree(d_field);
  hipFree(d_neighbors);
  hipFree(d_screen);
  
  /* printf("%s\n", cudaGetErrorString(cudaGetLastError())); Debug message. */
}


/* Semplice funzione main che controlla il framebuffer. Copiata da: https://stackoverflow.com/a/1830865

/**
   Program entrypoint.

   Initializes the main framebuffer (/dev/fb0), as suggested on: https://stackoverflow.com/a/1830865
   then passes it to the life() function.

   @return 0 if no error occurred, 1 otherwise.
 */
int main()
{
  struct fb_var_screeninfo screen_info;
  struct fb_fix_screeninfo fixed_info;
  char *buffer = NULL;
  size_t buflen;
  int fd = -1;
  int r = 1;

  fd = open("/dev/fb0", O_RDWR);
  if (fd >= 0)
  {
    if (!ioctl(fd, FBIOGET_VSCREENINFO, &screen_info) &&
        !ioctl(fd, FBIOGET_FSCREENINFO, &fixed_info))
    {
      buflen = screen_info.yres_virtual * fixed_info.line_length;
      buffer = (char *) mmap(NULL, buflen,
                       PROT_READ|PROT_WRITE, MAP_SHARED, fd, 0);
      if (buffer != MAP_FAILED)
      {
	life(buffer, screen_info.xres_virtual, screen_info.yres_virtual); /* Start the game. */
	r = 0;
      }
      else
      {
	perror("mmap");
      }
    }
    else
      {
	perror("ioctl");
      }
   }
   else
   {
       perror("open");
   }

   if (buffer && buffer != MAP_FAILED)
     munmap(buffer, buflen);
   if (fd >= 0)
     close(fd);

   return r;
}

	.text
	.file	"life.hip"
	.globl	_Z25__device_stub__updateCellPcS_S_ii # -- Begin function _Z25__device_stub__updateCellPcS_S_ii
	.type	_Z25__device_stub__updateCellPcS_S_ii,@function
_Z25__device_stub__updateCellPcS_S_ii:  # @_Z25__device_stub__updateCellPcS_S_ii
	.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	24(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	4(%rsp), %rdx
	movl	%ecx, (%rdx)
	movq	%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	16(%rsp), %r12
	leaq	8(%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	$_Z10updateCellPcS_S_ii, %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_end0:
	.size	_Z25__device_stub__updateCellPcS_S_ii, .Lfunc_end0-_Z25__device_stub__updateCellPcS_S_ii
	.cfi_endproc
                                        # -- End function
	.globl	_Z29__device_stub__countNeighborsPcS_ii # -- Begin function _Z29__device_stub__countNeighborsPcS_ii
	.type	_Z29__device_stub__countNeighborsPcS_ii,@function
_Z29__device_stub__countNeighborsPcS_ii: # @_Z29__device_stub__countNeighborsPcS_ii
	.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	$112, %rsp
	.cfi_def_cfa_offset 160
	.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	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%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	$_Z14countNeighborsPcS_ii, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$128, %rsp
	.cfi_adjust_cfa_offset -128
	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	_Z29__device_stub__countNeighborsPcS_ii, .Lfunc_end1-_Z29__device_stub__countNeighborsPcS_ii
	.cfi_endproc
                                        # -- End function
	.globl	_Z4lifePcii                     # -- Begin function _Z4lifePcii
	.type	_Z4lifePcii,@function
_Z4lifePcii:                            # @_Z4lifePcii
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset %rbp, -16
	movq	%rsp, %rbp
	.cfi_def_cfa_register %rbp
	pushq	%r15
	pushq	%r14
	pushq	%r13
	pushq	%r12
	pushq	%rbx
	subq	$1528, %rsp                     # imm = 0x5F8
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	movq	%rdi, %r15
	movl	%edx, -44(%rbp)                 # 4-byte Spill
	movl	%edx, %ebx
	movl	%esi, %r14d
	imull	%esi, %ebx
	leaq	15(%rbx), %rax
	andq	$-16, %rax
	movq	%rsp, %r13
	subq	%rax, %r13
	movq	%r13, %rsp
	xorl	%edi, %edi
	callq	time
	movl	%eax, %edi
	callq	srand
	testl	%ebx, %ebx
	jle	.LBB2_3
# %bb.1:                                # %.lr.ph.preheader
	xorl	%r12d, %r12d
.LBB2_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	$75, %eax
	setl	(%r13,%r12)
	incq	%r12
	cmpq	%r12, %rbx
	jne	.LBB2_2
.LBB2_3:                                # %._crit_edge
	leaq	-1560(%rbp), %r12
	movq	%r12, %rdi
	xorl	%esi, %esi
	callq	hipGetDevicePropertiesR0600
	movl	320(%r12), %r12d
	movslq	%r14d, %rax
	movslq	-44(%rbp), %rsi                 # 4-byte Folded Reload
	imulq	%rax, %rsi
	movq	%rsi, -56(%rbp)                 # 8-byte Spill
	leaq	-72(%rbp), %rdi
	callq	hipMalloc
	leaq	-80(%rbp), %rdi
	movq	-56(%rbp), %rsi                 # 8-byte Reload
	callq	hipMalloc
	movq	%r13, -88(%rbp)                 # 8-byte Spill
	movq	-56(%rbp), %rax                 # 8-byte Reload
	leaq	(,%rax,4), %r13
	leaq	-64(%rbp), %rdi
	movq	%r13, %rsi
	callq	hipMalloc
	leaq	-72(%rbp), %rax
	movq	(%rax), %rdi
	movq	-88(%rbp), %rsi                 # 8-byte Reload
	movq	-56(%rbp), %rdx                 # 8-byte Reload
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	-64(%rbp), %rax
	movq	(%rax), %rdi
	movq	%r15, %rsi
	movq	%r13, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	leal	(%rbx,%r12), %eax
	decl	%eax
	cltd
	idivl	%r12d
	movl	%eax, %ebx
	btsq	$32, %rbx
	btsq	$32, %r12
.LBB2_4:                                # =>This Inner Loop Header: Depth=1
	movq	%rbx, %rdi
	movl	$1, %esi
	movq	%r12, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB2_6
# %bb.5:                                #   in Loop: Header=BB2_4 Depth=1
	movq	-72(%rbp), %rdi
	movq	-80(%rbp), %rsi
	movl	%r14d, %edx
	movl	-44(%rbp), %ecx                 # 4-byte Reload
	callq	_Z29__device_stub__countNeighborsPcS_ii
.LBB2_6:                                #   in Loop: Header=BB2_4 Depth=1
	callq	hipDeviceSynchronize
	movq	%rbx, %rdi
	movl	$1, %esi
	movq	%r12, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB2_8
# %bb.7:                                #   in Loop: Header=BB2_4 Depth=1
	movq	-72(%rbp), %rdi
	movq	-80(%rbp), %rsi
	movq	-64(%rbp), %rdx
	movl	%r14d, %ecx
	movl	-44(%rbp), %r8d                 # 4-byte Reload
	callq	_Z25__device_stub__updateCellPcS_S_ii
.LBB2_8:                                #   in Loop: Header=BB2_4 Depth=1
	callq	hipDeviceSynchronize
	movq	-64(%rbp), %rsi
	movq	%r15, %rdi
	movq	%r13, %rdx
	movl	$2, %ecx
	callq	hipMemcpy
	jmp	.LBB2_4
.Lfunc_end2:
	.size	_Z4lifePcii, .Lfunc_end2-_Z4lifePcii
	.cfi_endproc
                                        # -- End function
	.globl	main                            # -- Begin function main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%r14
	.cfi_def_cfa_offset 16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	subq	$248, %rsp
	.cfi_def_cfa_offset 272
	.cfi_offset %rbx, -24
	.cfi_offset %r14, -16
	movl	$.L.str, %edi
	movl	$2, %esi
	xorl	%eax, %eax
	callq	open
	movl	%eax, %ebx
	testl	%eax, %eax
	js	.LBB3_6
# %bb.1:
	leaq	88(%rsp), %rdx
	movl	$17920, %esi                    # imm = 0x4600
	movl	%ebx, %edi
	xorl	%eax, %eax
	callq	ioctl
	testl	%eax, %eax
	jne	.LBB3_5
# %bb.2:
	leaq	8(%rsp), %rdx
	movl	$17922, %esi                    # imm = 0x4602
	movl	%ebx, %edi
	xorl	%eax, %eax
	callq	ioctl
	testl	%eax, %eax
	jne	.LBB3_5
# %bb.3:
	movl	56(%rsp), %r14d
	imull	100(%rsp), %r14d
	xorl	%edi, %edi
	movq	%r14, %rsi
	movl	$3, %edx
	movl	$1, %ecx
	movl	%ebx, %r8d
	xorl	%r9d, %r9d
	callq	mmap
	cmpq	$-1, %rax
	jne	.LBB3_13
# %bb.4:
	movl	$.L.str.1, %edi
	callq	perror
	movq	$-1, %rdi
	jmp	.LBB3_8
.LBB3_6:
	movl	$.L.str.3, %edi
	jmp	.LBB3_7
.LBB3_5:
	movl	$.L.str.2, %edi
.LBB3_7:
	callq	perror
	xorl	%edi, %edi
                                        # implicit-def: $r14
.LBB3_8:
	leaq	1(%rdi), %rax
	cmpq	$2, %rax
	jb	.LBB3_10
# %bb.9:
	movq	%r14, %rsi
	callq	munmap
.LBB3_10:
	testl	%ebx, %ebx
	js	.LBB3_12
# %bb.11:
	movl	%ebx, %edi
	callq	close
.LBB3_12:
	movl	$1, %eax
	addq	$248, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	retq
.LBB3_13:
	.cfi_def_cfa_offset 272
	movl	96(%rsp), %esi
	movl	100(%rsp), %edx
	movq	%rax, %rdi
	callq	_Z4lifePcii
.Lfunc_end3:
	.size	main, .Lfunc_end3-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	.LBB4_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB4_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z10updateCellPcS_S_ii, %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	$_Z14countNeighborsPcS_ii, %esi
	movl	$.L__unnamed_2, %edx
	movl	$.L__unnamed_2, %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_end4:
	.size	__hip_module_ctor, .Lfunc_end4-__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	.LBB5_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
.LBB5_2:
	retq
.Lfunc_end5:
	.size	__hip_module_dtor, .Lfunc_end5-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_Z10updateCellPcS_S_ii,@object  # @_Z10updateCellPcS_S_ii
	.section	.rodata,"a",@progbits
	.globl	_Z10updateCellPcS_S_ii
	.p2align	3, 0x0
_Z10updateCellPcS_S_ii:
	.quad	_Z25__device_stub__updateCellPcS_S_ii
	.size	_Z10updateCellPcS_S_ii, 8

	.type	_Z14countNeighborsPcS_ii,@object # @_Z14countNeighborsPcS_ii
	.globl	_Z14countNeighborsPcS_ii
	.p2align	3, 0x0
_Z14countNeighborsPcS_ii:
	.quad	_Z29__device_stub__countNeighborsPcS_ii
	.size	_Z14countNeighborsPcS_ii, 8

	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"/dev/fb0"
	.size	.L.str, 9

	.type	.L.str.1,@object                # @.str.1
.L.str.1:
	.asciz	"mmap"
	.size	.L.str.1, 5

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

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	"open"
	.size	.L.str.3, 5

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z10updateCellPcS_S_ii"
	.size	.L__unnamed_1, 23

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"_Z14countNeighborsPcS_ii"
	.size	.L__unnamed_2, 25

	.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

	.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 _Z25__device_stub__updateCellPcS_S_ii
	.addrsig_sym _Z29__device_stub__countNeighborsPcS_ii
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z10updateCellPcS_S_ii
	.addrsig_sym _Z14countNeighborsPcS_ii
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z10updateCellPcS_S_ii  ; -- Begin function _Z10updateCellPcS_S_ii
	.globl	_Z10updateCellPcS_S_ii
	.p2align	8
	.type	_Z10updateCellPcS_S_ii,@function
_Z10updateCellPcS_S_ii:                 ; @_Z10updateCellPcS_S_ii
; %bb.0:
	s_clause 0x1
	s_load_b32 s4, s[0:1], 0x2c
	s_load_b64 s[2:3], s[0:1], 0x18
	s_waitcnt lgkmcnt(0)
	s_and_b32 s4, s4, 0xffff
	s_mul_i32 s2, s3, s2
	v_mad_u64_u32 v[1:2], null, s15, s4, v[0:1]
	s_delay_alu instid0(VALU_DEP_1)
	v_cmp_gt_i32_e32 vcc_lo, s2, v1
	s_and_saveexec_b32 s2, vcc_lo
	s_cbranch_execz .LBB0_8
; %bb.1:
	s_load_b128 s[4:7], s[0:1], 0x0
	v_ashrrev_i32_e32 v0, 31, v1
	s_load_b64 s[0:1], s[0:1], 0x10
	s_mov_b32 s2, 0
	s_mov_b32 s3, exec_lo
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v2, vcc_lo, s4, v1
	v_add_co_ci_u32_e32 v3, vcc_lo, s5, v0, vcc_lo
	v_add_co_u32 v4, vcc_lo, s6, v1
	v_add_co_ci_u32_e32 v5, vcc_lo, s7, v0, vcc_lo
	global_load_u8 v6, v[2:3], off
                                        ; implicit-def: $sgpr5
                                        ; implicit-def: $sgpr4
	global_load_u8 v5, v[4:5], off
	s_waitcnt vmcnt(1)
	v_cmpx_ne_u16_e32 0, v6
	s_xor_b32 s3, exec_lo, s3
	s_cbranch_execz .LBB0_3
; %bb.2:
	s_waitcnt vmcnt(0)
	v_and_b32_e32 v0, 0xfe, v5
	s_mov_b32 s4, 1
	s_movk_i32 s5, 0xff
                                        ; implicit-def: $vgpr5
	s_delay_alu instid0(VALU_DEP_1)
	v_cmp_ne_u16_e32 vcc_lo, 2, v0
	s_and_b32 s2, vcc_lo, exec_lo
.LBB0_3:                                ; %Flow
	s_or_saveexec_b32 s3, s3
	v_mov_b32_e32 v4, s5
	v_mov_b32_e32 v0, s4
	s_xor_b32 exec_lo, exec_lo, s3
	s_cbranch_execz .LBB0_5
; %bb.4:                                ; %.critedge
	s_waitcnt vmcnt(0)
	v_and_b32_e32 v0, 0xff, v5
	v_mov_b32_e32 v4, 0xff
	s_and_not1_b32 s2, s2, exec_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(SALU_CYCLE_1)
	v_cmp_ne_u16_e32 vcc_lo, 3, v0
	v_mov_b32_e32 v0, 1
	s_and_b32 s4, vcc_lo, exec_lo
	s_or_b32 s2, s2, s4
.LBB0_5:                                ; %Flow68
	s_or_b32 exec_lo, exec_lo, s3
	s_and_saveexec_b32 s3, s2
; %bb.6:
	v_mov_b32_e32 v0, 0
	v_mov_b32_e32 v4, 0
; %bb.7:                                ; %.sink.split
	s_or_b32 exec_lo, exec_lo, s3
	v_lshlrev_b32_e32 v1, 2, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_3)
	v_ashrrev_i32_e32 v6, 31, v1
	s_waitcnt vmcnt(0)
	v_add_co_u32 v5, vcc_lo, s0, v1
	v_perm_b32 v1, v4, v4, 0x4040404
	v_add_co_ci_u32_e32 v6, vcc_lo, s1, v6, vcc_lo
	global_store_b8 v[2:3], v0, off
	global_store_b32 v[5:6], v1, off
.LBB0_8:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z10updateCellPcS_S_ii
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 288
		.amdhsa_user_sgpr_count 15
		.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 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 7
		.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_end0:
	.size	_Z10updateCellPcS_S_ii, .Lfunc_end0-_Z10updateCellPcS_S_ii
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 332
; NumSgprs: 18
; NumVgprs: 7
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 7
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.text
	.protected	_Z14countNeighborsPcS_ii ; -- Begin function _Z14countNeighborsPcS_ii
	.globl	_Z14countNeighborsPcS_ii
	.p2align	8
	.type	_Z14countNeighborsPcS_ii,@function
_Z14countNeighborsPcS_ii:               ; @_Z14countNeighborsPcS_ii
; %bb.0:
	s_load_b64 s[4:5], s[0:1], 0x10
	s_waitcnt lgkmcnt(0)
	s_mul_i32 s2, s5, s4
	s_load_b32 s5, s[0:1], 0x24
	s_ashr_i32 s3, s2, 31
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_2) | instid1(SALU_CYCLE_1)
	s_add_i32 s6, s2, s3
	s_sub_i32 s2, s2, s4
	s_xor_b32 s3, s6, s3
	v_cvt_f32_u32_e32 v1, s3
	s_delay_alu instid0(VALU_DEP_1)
	v_rcp_iflag_f32_e32 v1, v1
	s_waitcnt lgkmcnt(0)
	s_and_b32 s5, s5, 0xffff
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v3, 0x4f7ffffe, v1
	v_mad_u64_u32 v[1:2], null, s15, s5, v[0:1]
	s_sub_i32 s5, 0, s3
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v2, s2, v1
	s_ashr_i32 s2, s4, 31
	v_add_nc_u32_e32 v9, s4, v1
	v_ashrrev_i32_e32 v4, 31, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v2, v2, v4
	v_cvt_u32_f32_e32 v0, v3
	v_xor_b32_e32 v2, v2, v4
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	v_mul_lo_u32 v3, s5, v0
	s_add_i32 s5, s4, s2
	s_xor_b32 s2, s5, s2
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	s_sub_i32 s5, 0, s2
	v_mul_hi_u32 v3, v0, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v7, v0, v3
	v_mul_hi_u32 v0, v2, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v0, v0, s3
	v_sub_nc_u32_e32 v0, v2, v0
	v_cvt_f32_u32_e32 v2, s2
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_subrev_nc_u32_e32 v3, s3, v0
	v_cmp_le_u32_e32 vcc_lo, s3, v0
	v_rcp_iflag_f32_e32 v2, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v0, v0, v3, vcc_lo
	v_subrev_nc_u32_e32 v3, s3, v0
	v_cmp_le_u32_e32 vcc_lo, s3, v0
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v2, 0x4f7ffffe, v2
	v_cndmask_b32_e32 v0, v0, v3, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cvt_u32_f32_e32 v2, v2
	v_xor_b32_e32 v0, v0, v4
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_mul_lo_u32 v3, s5, v2
	s_load_b128 s[4:7], s[0:1], 0x0
	s_mov_b32 s1, exec_lo
	v_sub_nc_u32_e32 v0, v0, v4
	v_add_nc_u32_e32 v4, -1, v9
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_mul_hi_u32 v3, v2, v3
	v_ashrrev_i32_e32 v5, 31, v0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_ashrrev_i32_e32 v6, 31, v4
	v_add_nc_u32_e32 v8, v0, v5
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_add_nc_u32_e32 v10, v4, v6
	v_add_nc_u32_e32 v4, v2, v3
	v_xor_b32_e32 v2, v8, v5
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_xor_b32_e32 v3, v10, v6
	v_mul_hi_u32 v8, v2, v4
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_hi_u32 v10, v3, v4
	v_mul_lo_u32 v8, v8, s2
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_lo_u32 v10, v10, s2
	v_sub_nc_u32_e32 v2, v2, v8
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_sub_nc_u32_e32 v3, v3, v10
	v_subrev_nc_u32_e32 v8, s2, v2
	v_cmp_le_u32_e32 vcc_lo, s2, v2
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_subrev_nc_u32_e32 v10, s2, v3
	v_cndmask_b32_e32 v2, v2, v8, vcc_lo
	v_cmp_le_u32_e32 vcc_lo, s2, v3
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_subrev_nc_u32_e32 v8, s2, v2
	v_cndmask_b32_e32 v3, v3, v10, vcc_lo
	v_cmp_le_u32_e32 vcc_lo, s2, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_subrev_nc_u32_e32 v10, s2, v3
	v_cndmask_b32_e32 v2, v2, v8, vcc_lo
	v_cmp_le_u32_e32 vcc_lo, s2, v3
	v_ashrrev_i32_e32 v8, 31, v1
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_xor_b32_e32 v2, v2, v5
	v_cndmask_b32_e32 v3, v3, v10, vcc_lo
	v_sub_nc_u32_e32 v2, v5, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_xor_b32_e32 v3, v3, v6
	v_add_nc_u32_e32 v5, v0, v2
	s_delay_alu instid0(VALU_DEP_2)
	v_sub_nc_u32_e32 v10, v3, v6
	v_mov_b32_e32 v6, 0
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v2, vcc_lo, s6, v1
	v_add_co_ci_u32_e32 v3, vcc_lo, s7, v8, vcc_lo
	v_add_nc_u32_e32 v0, v5, v10
	global_store_b8 v[2:3], v6, off
	v_add_nc_u32_e32 v6, v1, v8
	v_ashrrev_i32_e32 v12, 31, v0
	v_add_co_u32 v11, vcc_lo, s4, v0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_xor_b32_e32 v6, v6, v8
	v_add_co_ci_u32_e32 v12, vcc_lo, s5, v12, vcc_lo
	global_load_u8 v0, v[11:12], off
	v_mul_hi_u32 v11, v6, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v11, v11, s2
	v_sub_nc_u32_e32 v6, v6, v11
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_subrev_nc_u32_e32 v11, s2, v6
	v_cmp_le_u32_e32 vcc_lo, s2, v6
	v_cndmask_b32_e32 v6, v6, v11, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_subrev_nc_u32_e32 v11, s2, v6
	v_cmp_le_u32_e32 vcc_lo, s2, v6
	v_cndmask_b32_e32 v6, v6, v11, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v6, v6, v8
	v_sub_nc_u32_e32 v8, v6, v8
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_nc_u32_e32 v6, v1, v8
	v_add_nc_u32_e32 v11, v6, v10
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v12, 31, v11
	v_add_co_u32 v11, s0, s4, v11
	v_add_co_ci_u32_e64 v12, s0, s5, v12, s0
	s_waitcnt vmcnt(0)
	v_cmp_ne_u16_e32 vcc_lo, 0, v0
	v_cndmask_b32_e64 v0, 0, 1, vcc_lo
	global_store_b8 v[2:3], v0, off
	global_load_u8 v11, v[11:12], off
	s_waitcnt vmcnt(0)
	v_cmpx_ne_u16_e32 0, v11
	s_cbranch_execz .LBB1_2
; %bb.1:
	v_cndmask_b32_e64 v0, 1, 2, vcc_lo
	global_store_b8 v[2:3], v0, off
.LBB1_2:
	s_or_b32 exec_lo, exec_lo, s1
	v_ashrrev_i32_e32 v11, 31, v9
	s_mov_b32 s0, exec_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v9, v9, v11
	v_xor_b32_e32 v9, v9, v11
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_hi_u32 v7, v9, v7
	v_mul_lo_u32 v7, v7, s3
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_nc_u32_e32 v7, v9, v7
	v_subrev_nc_u32_e32 v9, s3, v7
	v_cmp_le_u32_e32 vcc_lo, s3, v7
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v7, v7, v9, vcc_lo
	v_subrev_nc_u32_e32 v9, s3, v7
	v_cmp_le_u32_e32 vcc_lo, s3, v7
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v7, v7, v9, vcc_lo
	v_xor_b32_e32 v7, v7, v11
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_nc_u32_e32 v7, v7, v11
	v_ashrrev_i32_e32 v9, 31, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v11, v7, v9
	v_xor_b32_e32 v11, v11, v9
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_hi_u32 v12, v11, v4
	v_mul_lo_u32 v12, v12, s2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_nc_u32_e32 v11, v11, v12
	v_subrev_nc_u32_e32 v12, s2, v11
	v_cmp_le_u32_e32 vcc_lo, s2, v11
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v11, v11, v12, vcc_lo
	v_subrev_nc_u32_e32 v12, s2, v11
	v_cmp_le_u32_e32 vcc_lo, s2, v11
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v11, v11, v12, vcc_lo
	v_xor_b32_e32 v11, v11, v9
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_nc_u32_e32 v9, v9, v11
	v_add_nc_u32_e32 v7, v7, v9
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v9, v7, v10
	v_ashrrev_i32_e32 v10, 31, v9
	v_add_co_u32 v9, vcc_lo, s4, v9
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v10, vcc_lo, s5, v10, vcc_lo
	global_load_u8 v9, v[9:10], off
	s_waitcnt vmcnt(0)
	v_cmpx_ne_u16_e32 0, v9
	s_cbranch_execz .LBB1_4
; %bb.3:
	v_add_nc_u16 v0, v0, 1
	global_store_b8 v[2:3], v0, off
.LBB1_4:
	s_or_b32 exec_lo, exec_lo, s0
	v_add_nc_u32_e32 v9, v5, v8
	s_mov_b32 s0, exec_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_ashrrev_i32_e32 v10, 31, v9
	v_add_co_u32 v9, vcc_lo, s4, v9
	v_add_co_ci_u32_e32 v10, vcc_lo, s5, v10, vcc_lo
	global_load_u8 v9, v[9:10], off
	s_waitcnt vmcnt(0)
	v_cmpx_ne_u16_e32 0, v9
	s_cbranch_execz .LBB1_6
; %bb.5:
	v_add_nc_u16 v0, v0, 1
	global_store_b8 v[2:3], v0, off
.LBB1_6:
	s_or_b32 exec_lo, exec_lo, s0
	v_add_nc_u32_e32 v8, v7, v8
	s_mov_b32 s0, exec_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_ashrrev_i32_e32 v9, 31, v8
	v_add_co_u32 v8, vcc_lo, s4, v8
	v_add_co_ci_u32_e32 v9, vcc_lo, s5, v9, vcc_lo
	global_load_u8 v8, v[8:9], off
	s_waitcnt vmcnt(0)
	v_cmpx_ne_u16_e32 0, v8
	s_cbranch_execz .LBB1_8
; %bb.7:
	v_add_nc_u16 v0, v0, 1
	global_store_b8 v[2:3], v0, off
.LBB1_8:
	s_or_b32 exec_lo, exec_lo, s0
	v_add_nc_u32_e32 v1, 1, v1
	s_mov_b32 s0, exec_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v8, 31, v1
	v_add_nc_u32_e32 v1, v1, v8
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v1, v1, v8
	v_mul_hi_u32 v4, v1, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v4, v4, s2
	v_sub_nc_u32_e32 v1, v1, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_subrev_nc_u32_e32 v4, s2, v1
	v_cmp_le_u32_e32 vcc_lo, s2, v1
	v_cndmask_b32_e32 v1, v1, v4, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_subrev_nc_u32_e32 v4, s2, v1
	v_cmp_le_u32_e32 vcc_lo, s2, v1
	v_cndmask_b32_e32 v1, v1, v4, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v1, v1, v8
	v_sub_nc_u32_e32 v1, v1, v8
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v4, v5, v1
	v_ashrrev_i32_e32 v5, 31, v4
	v_add_co_u32 v4, vcc_lo, s4, v4
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v5, vcc_lo, s5, v5, vcc_lo
	global_load_u8 v4, v[4:5], off
	s_waitcnt vmcnt(0)
	v_cmpx_ne_u16_e32 0, v4
	s_cbranch_execz .LBB1_10
; %bb.9:
	v_add_nc_u16 v0, v0, 1
	global_store_b8 v[2:3], v0, off
.LBB1_10:
	s_or_b32 exec_lo, exec_lo, s0
	v_add_nc_u32_e32 v4, v6, v1
	s_mov_b32 s0, exec_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_ashrrev_i32_e32 v5, 31, v4
	v_add_co_u32 v4, vcc_lo, s4, v4
	v_add_co_ci_u32_e32 v5, vcc_lo, s5, v5, vcc_lo
	global_load_u8 v4, v[4:5], off
	s_waitcnt vmcnt(0)
	v_cmpx_ne_u16_e32 0, v4
	s_cbranch_execz .LBB1_12
; %bb.11:
	v_add_nc_u16 v0, v0, 1
	global_store_b8 v[2:3], v0, off
.LBB1_12:
	s_or_b32 exec_lo, exec_lo, s0
	v_add_nc_u32_e32 v1, v7, v1
	s_mov_b32 s0, exec_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_ashrrev_i32_e32 v5, 31, v1
	v_add_co_u32 v4, vcc_lo, s4, v1
	v_add_co_ci_u32_e32 v5, vcc_lo, s5, v5, vcc_lo
	global_load_u8 v1, v[4:5], off
	s_waitcnt vmcnt(0)
	v_cmpx_ne_u16_e32 0, v1
	s_cbranch_execz .LBB1_14
; %bb.13:
	v_add_nc_u16 v0, v0, 1
	global_store_b8 v[2:3], v0, off
.LBB1_14:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z14countNeighborsPcS_ii
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 280
		.amdhsa_user_sgpr_count 15
		.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 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 13
		.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	_Z14countNeighborsPcS_ii, .Lfunc_end1-_Z14countNeighborsPcS_ii
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 1424
; NumSgprs: 18
; NumVgprs: 13
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 13
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.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
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .offset:         24
        .size:           4
        .value_kind:     by_value
      - .offset:         28
        .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:           _Z10updateCellPcS_S_ii
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z10updateCellPcS_S_ii.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     7
    .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:     hidden_block_count_x
      - .offset:         28
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         32
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         36
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         38
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         40
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         42
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         44
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         46
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         64
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         88
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 280
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z14countNeighborsPcS_ii
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z14countNeighborsPcS_ii.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     13
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
