/*
KAM PUI SO (ANTHONY)
CS 510 GPU
Homework 2

The Game of Life
Rules:

Any live cell with fewer than two live neighbours dies, as if caused by under-population.
Any live cell with two or three live neighbours lives on to the next generation.
Any live cell with more than three live neighbours dies, as if by overcrowding.
Any dead cell with exactly three live neighbours becomes a live cell, as if by reproduction.

*/

#include <sys/time.h>
#include <time.h>
#include <stdio.h>
#include <stdlib.h>
#include <string.h>

#define WIDTH 70
#define HEIGHT 47
#define MAX 2000
#define LIMIT 512
#define NEIGHBORS 8
#define ROW 0
#define COL 1

// global const
const int offsets[NEIGHBORS][2] = {{-1, 1},{0, 1},{1, 1},
                                   {-1, 0},       {1, 0},
                                   {-1,-1},{0,-1},{1,-1}};

__constant__ int offsets_dev[NEIGHBORS][2] = {{-1, 1},{0, 1},{1, 1},
                                              {-1, 0},       {1, 0},
                                              {-1,-1},{0,-1},{1,-1}};

/* The kernel that will execute on the GPU */
__global__ void step_kernel(int *board, int *result, int width, int height) {
    int n = width * height;
    int idx = blockDim.x * blockIdx.x + threadIdx.x;
    int num_neighbors = 0;
    int nx = 0;
    int ny = 0;
    int x = idx % width;
    int y = idx / width;
    int i = 0;

    for (i = 0; i < NEIGHBORS; i++) {
        // To make the board torroidal, we use modular arithmetic to
        // wrap neighbor coordinates around to the other side of the
        // board if they fall off.
        nx = (x + offsets_dev[i][ROW] + width) % width;
        ny = (y + offsets_dev[i][COL] + height) % height;
        if (board[ny * width + nx]) {
            num_neighbors++;
        }
    }

    // apply the Game of Life rules to this cell
    if (idx < n && ((board[idx] && num_neighbors==2) || num_neighbors==3))
        result[idx] = 1;
    else
        result[idx] = 0;

}

/* This function encapsulates the process of creating and tearing down the
 * environment used to execute our game of life iteration kernel. The steps of the
 * process are:
 *   1. Allocate memory on the device to hold our board vectors
 *   2. Copy the board vectors to device memory
 *   3. Execute the kernel
 *   4. Retrieve the result board vector from the device by copying it to the host
 *   5. Free memory on the device
 */
void step_dev(int *board, int *result, int width, int height) {
    // Step 1: Allocate memory
    int *board_dev, *result_dev;
    int n = width * height;

    // Since cudaMalloc does not return a pointer like C's traditional malloc
    // (it returns a success status instead), we provide as it's first argument
    // the address of our device pointer variable so that it can change the
    // value of our pointer to the correct device address.
    cudaMalloc((void **) &board_dev, sizeof(int) * n);
    cudaMalloc((void **) &result_dev, sizeof(int) * n);

    // Step 2: Copy the input vectors to the device
    cudaMemcpy(board_dev, board, sizeof(int) * n, cudaMemcpyHostToDevice);

    // Step 3: Invoke the kernel
    // We allocate enough blocks (each 512 threads long) in the grid to
    // accomodate all `n` elements in the vectors. The 512 long block size
    // is somewhat arbitrary, but with the constraint that we know the
    // hardware will support blocks of that size.
    dim3 dimGrid((n + LIMIT - 1) / LIMIT, 1, 1);
    dim3 dimBlock(LIMIT, 1, 1);
    step_kernel<<<dimGrid, dimBlock>>>(board_dev, result_dev, width, height);

    // Step 4: Retrieve the results
    cudaMemcpy(result, result_dev, sizeof(int) * n, cudaMemcpyDeviceToHost);

    // Step 5: Free device memory
    cudaFree(board_dev);
    cudaFree(result_dev);
}



/* The old-fashioned CPU-only way to step thru game of life*/
void step(int *current, int *next, int width, int height) {
    // coordinates of the cell we're currently evaluating
    int x, y;
    // offset index, neighbor coordinates, alive neighbor count
    int i, nx, ny, num_neighbors;

    // write the next board state
    for (y=0; y<height; y++) {
        for (x=0; x<width; x++) {
            // count this cell's alive neighbors
            num_neighbors = 0;
            for (i = 0; i < NEIGHBORS; i++) {
                // To make the board torroidal, we use modular arithmetic to
                // wrap neighbor coordinates around to the other side of the
                // board if they fall off.
                nx = (x + offsets[i][ROW] + width) % width;
                ny = (y + offsets[i][COL] + height) % height;
                if (current[ny * width + nx]) {
                    num_neighbors++;
                }
            }

            // apply the Game of Life rules to this cell
            next[y * width + x] = 0;
            if ((current[y * width + x] && num_neighbors==2) || num_neighbors==3) {
                next[y * width + x] = 1;
            }
        }
    }
}




// fill the board with random cells
void fill_board(int *board, int width, int height) {
    int i;
    for (i = 0; i < (width * height); i++)
        board[i] = rand() % 2;
}


// print board image
void print_board(int *board, int width, int height) {
    int x, y;
    for (y = 0; y<height; y++) {
        for (x = 0; x<width; x++) {
            char c = board[y * width + x] ? '#':' ';
            printf("%c", c);
        }
        printf("\n");
    }
    printf("-----\n");
}



// animate each cell
void animate(int *current, int *next, int width, int height) {
    struct timespec delay = {0, 0}; // 0.005 seconds
//    struct timespec delay = {0, 125000000}; // 0.125 seconds
//    struct timespec delay = {0, 250000000}; // 0.25 seconds
    struct timespec remaining;


//    while (1) {
    for (int i = 0; i < MAX; ++i) {
	printf("%d\n", i);
        print_board(current, width, height);
        step_dev(current, next, width, height);
        // Copy the next state, that step() just wrote into, to current state
        memcpy(current, next, sizeof(int) * width * height);
        // We sleep only because textual output is slow and the console needs
        // time to catch up. We don't sleep in the graphical X11 version.
        nanosleep(&delay, &remaining);
    }
}



// main function
int main(void) {
    // variable
    int width = WIDTH;
    int height = HEIGHT;
    int n = width * height;
    int *current = (int *) malloc(n* sizeof(int));
    int *next = (int *) malloc(n * sizeof(int));

    // initialize the global "current"
    fill_board(current, width, height);
    animate(current, next, width, height);

    // free memory
    free(current);
    free(next);

    return 0;
}






	.file	"tmpxft_0033e71a_00000000-6_hw2_notile.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2035:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2035:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z4stepPiS_ii
	.type	_Z4stepPiS_ii, @function
_Z4stepPiS_ii:
.LFB2028:
	.cfi_startproc
	endbr64
	pushq	%r14
	.cfi_def_cfa_offset 16
	.cfi_offset 14, -16
	movl	%ecx, %r8d
	xorl	%r11d, %r11d
	leaq	64+_ZL7offsets(%rip), %r14
	pushq	%r13
	.cfi_def_cfa_offset 24
	.cfi_offset 13, -24
	xorl	%r13d, %r13d
	pushq	%r12
	.cfi_def_cfa_offset 32
	.cfi_offset 12, -32
	pushq	%rbp
	.cfi_def_cfa_offset 40
	.cfi_offset 6, -40
	movq	%rsi, %rbp
	movl	%edx, %esi
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset 3, -48
	movq	%rdi, %rbx
.L3:
	cmpl	%r8d, %r11d
	jge	.L2
	movslq	%r13d, %rcx
	xorl	%r10d, %r10d
	salq	$2, %rcx
.L11:
	cmpl	%esi, %r10d
	jge	.L19
	leaq	_ZL7offsets(%rip), %r9
	xorl	%r12d, %r12d
.L5:
	movl	4(%r9), %eax
	addl	%r11d, %eax
	addl	%r8d, %eax
	cltd
	idivl	%r8d
	movl	(%r9), %eax
	addl	%r10d, %eax
	addl	%esi, %eax
	imull	%esi, %edx
	movl	%edx, %edi
	cltd
	idivl	%esi
	addl	%edx, %edi
	movslq	%edi, %rdi
	cmpl	$0, (%rbx,%rdi,4)
	je	.L4
	incl	%r12d
.L4:
	addq	$8, %r9
	cmpq	%r9, %r14
	jne	.L5
	xorl	%eax, %eax
	movl	%eax, 0(%rbp,%rcx)
	cmpl	$0, (%rbx,%rcx)
	je	.L12
	cmpl	$2, %r12d
	je	.L6
.L12:
	cmpl	$3, %r12d
	jne	.L8
.L6:
	movl	$1, 0(%rbp,%rcx)
.L8:
	incl	%r10d
	addq	$4, %rcx
	jmp	.L11
.L19:
	incl	%r11d
	addl	%esi, %r13d
	jmp	.L3
.L2:
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%rbp
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r13
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2028:
	.size	_Z4stepPiS_ii, .-_Z4stepPiS_ii
	.globl	_Z10fill_boardPiii
	.type	_Z10fill_boardPiii, @function
_Z10fill_boardPiii:
.LFB2029:
	.cfi_startproc
	endbr64
	imull	%edx, %esi
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	movl	$2, %r13d
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	movq	%rdi, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	movl	%esi, %ebp
	xorl	%ebx, %ebx
	pushq	%rcx
	.cfi_def_cfa_offset 48
.L21:
	cmpl	%ebx, %ebp
	jle	.L24
	call	rand@PLT
	cltd
	idivl	%r13d
	movl	%edx, (%r12,%rbx,4)
	incq	%rbx
	jmp	.L21
.L24:
	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	_Z10fill_boardPiii, .-_Z10fill_boardPiii
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"%c"
.LC1:
	.string	"\n"
.LC2:
	.string	"-----\n"
	.text
	.globl	_Z11print_boardPiii
	.type	_Z11print_boardPiii, @function
_Z11print_boardPiii:
.LFB2030:
	.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
	movl	%edx, %r14d
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	xorl	%r13d, %r13d
	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
	movq	%rdi, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movl	%esi, %ebx
	subq	$24, %rsp
	.cfi_def_cfa_offset 80
.L26:
	cmpl	%r14d, %r13d
	jge	.L29
	movslq	%r12d, %rax
	xorl	%r15d, %r15d
	leaq	0(%rbp,%rax,4), %r8
.L30:
	cmpl	%r15d, %ebx
	jle	.L33
	cmpl	$1, (%r8,%r15,4)
	movl	$2, %edi
	movq	%r8, 8(%rsp)
	leaq	.LC0(%rip), %rsi
	sbbl	%edx, %edx
	xorl	%eax, %eax
	incq	%r15
	andl	$-3, %edx
	addl	$35, %edx
	movsbl	%dl, %edx
	call	__printf_chk@PLT
	movq	8(%rsp), %r8
	jmp	.L30
.L33:
	leaq	.LC1(%rip), %rsi
	xorl	%eax, %eax
	incl	%r13d
	addl	%ebx, %r12d
	movl	$2, %edi
	call	__printf_chk@PLT
	jmp	.L26
.L29:
	addq	$24, %rsp
	.cfi_def_cfa_offset 56
	leaq	.LC2(%rip), %rsi
	movl	$2, %edi
	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
	jmp	__printf_chk@PLT
	.cfi_endproc
.LFE2030:
	.size	_Z11print_boardPiii, .-_Z11print_boardPiii
	.globl	_Z35__device_stub__Z11step_kernelPiS_iiPiS_ii
	.type	_Z35__device_stub__Z11step_kernelPiS_iiPiS_ii, @function
_Z35__device_stub__Z11step_kernelPiS_iiPiS_ii:
.LFB2057:
	.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	.L34
	pushq	48(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z11step_kernelPiS_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
.L34:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L36
	call	__stack_chk_fail@PLT
.L36:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2057:
	.size	_Z35__device_stub__Z11step_kernelPiS_iiPiS_ii, .-_Z35__device_stub__Z11step_kernelPiS_iiPiS_ii
	.globl	_Z11step_kernelPiS_ii
	.type	_Z11step_kernelPiS_ii, @function
_Z11step_kernelPiS_ii:
.LFB2058:
	.cfi_startproc
	endbr64
	jmp	_Z35__device_stub__Z11step_kernelPiS_iiPiS_ii
	.cfi_endproc
.LFE2058:
	.size	_Z11step_kernelPiS_ii, .-_Z11step_kernelPiS_ii
	.globl	_Z8step_devPiS_ii
	.type	_Z8step_devPiS_ii, @function
_Z8step_devPiS_ii:
.LFB2027:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movq	%rdi, %r15
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	movl	%ecx, %r14d
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	movq	%rsi, %r13
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movl	%edx, %r12d
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movl	%edx, %ebx
	imull	%ecx, %ebx
	subq	$56, %rsp
	.cfi_def_cfa_offset 112
	movq	%fs:40, %rax
	movq	%rax, 40(%rsp)
	xorl	%eax, %eax
	movslq	%ebx, %rbp
	movq	%rsp, %rdi
	salq	$2, %rbp
	movq	%rbp, %rsi
	call	cudaMalloc@PLT
	leaq	8(%rsp), %rdi
	movq	%rbp, %rsi
	call	cudaMalloc@PLT
	movq	(%rsp), %rdi
	movq	%rbp, %rdx
	movq	%r15, %rsi
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	movl	$512, %ecx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	leal	511(%rbx), %eax
	cltd
	idivl	%ecx
	movl	$8388609, %edx
	movl	$1, %ecx
	salq	$9, %rdx
	movl	%eax, 16(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 20(%rsp)
	movq	16(%rsp), %rdi
	movl	24(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L40
	movq	8(%rsp), %rsi
	movq	(%rsp), %rdi
	movl	%r14d, %ecx
	movl	%r12d, %edx
	call	_Z35__device_stub__Z11step_kernelPiS_iiPiS_ii
.L40:
	movq	8(%rsp), %rsi
	movl	$2, %ecx
	movq	%rbp, %rdx
	movq	%r13, %rdi
	call	cudaMemcpy@PLT
	movq	(%rsp), %rdi
	call	cudaFree@PLT
	movq	8(%rsp), %rdi
	call	cudaFree@PLT
	movq	40(%rsp), %rax
	subq	%fs:40, %rax
	je	.L41
	call	__stack_chk_fail@PLT
.L41:
	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
.LFE2027:
	.size	_Z8step_devPiS_ii, .-_Z8step_devPiS_ii
	.section	.rodata.str1.1
.LC3:
	.string	"%d\n"
	.text
	.globl	_Z7animatePiS_ii
	.type	_Z7animatePiS_ii, @function
_Z7animatePiS_ii:
.LFB2031:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	leaq	.LC3(%rip), %r15
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	movl	%ecx, %r14d
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movq	%rsi, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movq	%rdi, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movslq	%edx, %rbx
	movq	%rbx, %r13
	subq	$72, %rsp
	.cfi_def_cfa_offset 128
	movq	%fs:40, %rax
	movq	%rax, 56(%rsp)
	xorl	%eax, %eax
	xorl	%eax, %eax
	movq	%rax, 24(%rsp)
	movq	%rax, 32(%rsp)
	movslq	%ecx, %rax
	imulq	%rax, %rbx
	leaq	0(,%rbx,4), %rax
	xorl	%ebx, %ebx
	movq	%rax, 8(%rsp)
.L44:
	movl	%ebx, %edx
	movq	%r15, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	%r14d, %edx
	movl	%r13d, %esi
	movq	%rbp, %rdi
	call	_Z11print_boardPiii
	movl	%r14d, %ecx
	movl	%r13d, %edx
	movq	%r12, %rsi
	movq	%rbp, %rdi
	incl	%ebx
	call	_Z8step_devPiS_ii
	movq	8(%rsp), %rcx
	movq	%rbp, %rdi
	movq	%r12, %rsi
	rep movsb
	leaq	40(%rsp), %rsi
	leaq	24(%rsp), %rdi
	call	nanosleep@PLT
	cmpl	$2000, %ebx
	jne	.L44
	movq	56(%rsp), %rax
	subq	%fs:40, %rax
	je	.L45
	call	__stack_chk_fail@PLT
.L45:
	addq	$72, %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
.LFE2031:
	.size	_Z7animatePiS_ii, .-_Z7animatePiS_ii
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2032:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	movl	$13160, %edi
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	pushq	%rax
	.cfi_def_cfa_offset 32
	call	malloc@PLT
	movl	$13160, %edi
	movq	%rax, %rbx
	call	malloc@PLT
	movq	%rbx, %rdi
	movl	$47, %edx
	movl	$70, %esi
	movq	%rax, %rbp
	call	_Z10fill_boardPiii
	movl	$70, %edx
	movq	%rbp, %rsi
	movq	%rbx, %rdi
	movl	$47, %ecx
	call	_Z7animatePiS_ii
	movq	%rbx, %rdi
	call	free@PLT
	movq	%rbp, %rdi
	call	free@PLT
	popq	%rdx
	.cfi_def_cfa_offset 24
	xorl	%eax, %eax
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2032:
	.size	main, .-main
	.section	.rodata.str1.1
.LC4:
	.string	"_Z11step_kernelPiS_ii"
.LC5:
	.string	"offsets_dev"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2060:
	.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	_Z11step_kernelPiS_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	%r8d, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC5(%rip), %rdx
	movl	$64, %r9d
	leaq	_ZL11offsets_dev(%rip), %rsi
	pushq	$1
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	call	__cudaRegisterVar@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	popq	%rax
	.cfi_def_cfa_offset 24
	popq	%rdx
	.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
.LFE2060:
	.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
	.local	_ZL11offsets_dev
	.comm	_ZL11offsets_dev,64,32
	.section	.rodata
	.align 32
	.type	_ZL7offsets, @object
	.size	_ZL7offsets, 64
_ZL7offsets:
	.long	-1
	.long	1
	.long	0
	.long	1
	.long	1
	.long	1
	.long	-1
	.long	0
	.long	1
	.long	0
	.long	-1
	.long	-1
	.long	0
	.long	-1
	.long	1
	.long	-1
	.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:

/*
KAM PUI SO (ANTHONY)
CS 510 GPU
Homework 2

The Game of Life
Rules:

Any live cell with fewer than two live neighbours dies, as if caused by under-population.
Any live cell with two or three live neighbours lives on to the next generation.
Any live cell with more than three live neighbours dies, as if by overcrowding.
Any dead cell with exactly three live neighbours becomes a live cell, as if by reproduction.

*/


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

#define WIDTH 70
#define HEIGHT 47
#define MAX 2000
#define LIMIT 512
#define NEIGHBORS 8
#define ROW 0
#define COL 1

// global const
const int offsets[NEIGHBORS][2] = {{-1, 1},{0, 1},{1, 1},
                                   {-1, 0},       {1, 0},
                                   {-1,-1},{0,-1},{1,-1}};

__constant__ int offsets_dev[NEIGHBORS][2] = {{-1, 1},{0, 1},{1, 1},
                                              {-1, 0},       {1, 0},
                                              {-1,-1},{0,-1},{1,-1}};

/* The kernel that will execute on the GPU */
__global__ void step_kernel(int *board, int *result, int width, int height) {
    int n = width * height;
    int idx = blockDim.x * blockIdx.x + threadIdx.x;
    int num_neighbors = 0;
    int nx = 0;
    int ny = 0;
    int x = idx % width;
    int y = idx / width;
    int i = 0;

    for (i = 0; i < NEIGHBORS; i++) {
        // To make the board torroidal, we use modular arithmetic to
        // wrap neighbor coordinates around to the other side of the
        // board if they fall off.
        nx = (x + offsets_dev[i][ROW] + width) % width;
        ny = (y + offsets_dev[i][COL] + height) % height;
        if (board[ny * width + nx]) {
            num_neighbors++;
        }
    }

    // apply the Game of Life rules to this cell
    if (idx < n && ((board[idx] && num_neighbors==2) || num_neighbors==3))
        result[idx] = 1;
    else
        result[idx] = 0;

}

/* This function encapsulates the process of creating and tearing down the
 * environment used to execute our game of life iteration kernel. The steps of the
 * process are:
 *   1. Allocate memory on the device to hold our board vectors
 *   2. Copy the board vectors to device memory
 *   3. Execute the kernel
 *   4. Retrieve the result board vector from the device by copying it to the host
 *   5. Free memory on the device
 */
void step_dev(int *board, int *result, int width, int height) {
    // Step 1: Allocate memory
    int *board_dev, *result_dev;
    int n = width * height;

    // Since cudaMalloc does not return a pointer like C's traditional malloc
    // (it returns a success status instead), we provide as it's first argument
    // the address of our device pointer variable so that it can change the
    // value of our pointer to the correct device address.
    hipMalloc((void **) &board_dev, sizeof(int) * n);
    hipMalloc((void **) &result_dev, sizeof(int) * n);

    // Step 2: Copy the input vectors to the device
    hipMemcpy(board_dev, board, sizeof(int) * n, hipMemcpyHostToDevice);

    // Step 3: Invoke the kernel
    // We allocate enough blocks (each 512 threads long) in the grid to
    // accomodate all `n` elements in the vectors. The 512 long block size
    // is somewhat arbitrary, but with the constraint that we know the
    // hardware will support blocks of that size.
    dim3 dimGrid((n + LIMIT - 1) / LIMIT, 1, 1);
    dim3 dimBlock(LIMIT, 1, 1);
    step_kernel<<<dimGrid, dimBlock>>>(board_dev, result_dev, width, height);

    // Step 4: Retrieve the results
    hipMemcpy(result, result_dev, sizeof(int) * n, hipMemcpyDeviceToHost);

    // Step 5: Free device memory
    hipFree(board_dev);
    hipFree(result_dev);
}



/* The old-fashioned CPU-only way to step thru game of life*/
void step(int *current, int *next, int width, int height) {
    // coordinates of the cell we're currently evaluating
    int x, y;
    // offset index, neighbor coordinates, alive neighbor count
    int i, nx, ny, num_neighbors;

    // write the next board state
    for (y=0; y<height; y++) {
        for (x=0; x<width; x++) {
            // count this cell's alive neighbors
            num_neighbors = 0;
            for (i = 0; i < NEIGHBORS; i++) {
                // To make the board torroidal, we use modular arithmetic to
                // wrap neighbor coordinates around to the other side of the
                // board if they fall off.
                nx = (x + offsets[i][ROW] + width) % width;
                ny = (y + offsets[i][COL] + height) % height;
                if (current[ny * width + nx]) {
                    num_neighbors++;
                }
            }

            // apply the Game of Life rules to this cell
            next[y * width + x] = 0;
            if ((current[y * width + x] && num_neighbors==2) || num_neighbors==3) {
                next[y * width + x] = 1;
            }
        }
    }
}




// fill the board with random cells
void fill_board(int *board, int width, int height) {
    int i;
    for (i = 0; i < (width * height); i++)
        board[i] = rand() % 2;
}


// print board image
void print_board(int *board, int width, int height) {
    int x, y;
    for (y = 0; y<height; y++) {
        for (x = 0; x<width; x++) {
            char c = board[y * width + x] ? '#':' ';
            printf("%c", c);
        }
        printf("\n");
    }
    printf("-----\n");
}



// animate each cell
void animate(int *current, int *next, int width, int height) {
    struct timespec delay = {0, 0}; // 0.005 seconds
//    struct timespec delay = {0, 125000000}; // 0.125 seconds
//    struct timespec delay = {0, 250000000}; // 0.25 seconds
    struct timespec remaining;


//    while (1) {
    for (int i = 0; i < MAX; ++i) {
	printf("%d\n", i);
        print_board(current, width, height);
        step_dev(current, next, width, height);
        // Copy the next state, that step() just wrote into, to current state
        memcpy(current, next, sizeof(int) * width * height);
        // We sleep only because textual output is slow and the console needs
        // time to catch up. We don't sleep in the graphical X11 version.
        nanosleep(&delay, &remaining);
    }
}



// main function
int main(void) {
    // variable
    int width = WIDTH;
    int height = HEIGHT;
    int n = width * height;
    int *current = (int *) malloc(n* sizeof(int));
    int *next = (int *) malloc(n * sizeof(int));

    // initialize the global "current"
    fill_board(current, width, height);
    animate(current, next, width, height);

    // free memory
    free(current);
    free(next);

    return 0;
}






	.text
	.file	"hw2_notile.hip"
	.globl	_Z26__device_stub__step_kernelPiS_ii # -- Begin function _Z26__device_stub__step_kernelPiS_ii
	.type	_Z26__device_stub__step_kernelPiS_ii,@function
_Z26__device_stub__step_kernelPiS_ii:   # @_Z26__device_stub__step_kernelPiS_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	$_Z11step_kernelPiS_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_end0:
	.size	_Z26__device_stub__step_kernelPiS_ii, .Lfunc_end0-_Z26__device_stub__step_kernelPiS_ii
	.cfi_endproc
                                        # -- End function
	.globl	_Z8step_devPiS_ii               # -- Begin function _Z8step_devPiS_ii
	.type	_Z8step_devPiS_ii,@function
_Z8step_devPiS_ii:                      # @_Z8step_devPiS_ii
	.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	$24, %rsp
	.cfi_def_cfa_offset 80
	.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, %ebp
	movl	%edx, %r15d
	movq	%rsi, 16(%rsp)                  # 8-byte Spill
	movq	%rdi, %r13
	movl	%ecx, %eax
	imull	%edx, %eax
	movslq	%eax, %r12
	leaq	(,%r12,4), %r14
	leaq	8(%rsp), %rbx
	movq	%rbx, %rdi
	movq	%r14, %rsi
	callq	hipMalloc
	movq	%rsp, %rdi
	movq	%r14, %rsi
	callq	hipMalloc
	movq	(%rbx), %rdi
	movq	%r13, %rsi
	movq	%r14, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	leal	511(%r12), %eax
	addl	$1022, %r12d                    # imm = 0x3FE
	testl	%eax, %eax
	cmovnsl	%eax, %r12d
	sarl	$9, %r12d
	btsq	$32, %r12
	movabsq	$4294967296, %rdx               # imm = 0x100000000
	orq	$512, %rdx                      # imm = 0x200
	movq	%r12, %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB1_2
# %bb.1:
	movq	8(%rsp), %rdi
	movq	(%rsp), %rsi
	movl	%r15d, %edx
	movl	%ebp, %ecx
	callq	_Z26__device_stub__step_kernelPiS_ii
.LBB1_2:
	movq	(%rsp), %rsi
	movq	16(%rsp), %rdi                  # 8-byte Reload
	movq	%r14, %rdx
	movl	$2, %ecx
	callq	hipMemcpy
	movq	8(%rsp), %rdi
	callq	hipFree
	movq	(%rsp), %rdi
	callq	hipFree
	addq	$24, %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_end1:
	.size	_Z8step_devPiS_ii, .Lfunc_end1-_Z8step_devPiS_ii
	.cfi_endproc
                                        # -- End function
	.globl	_Z4stepPiS_ii                   # -- Begin function _Z4stepPiS_ii
	.type	_Z4stepPiS_ii,@function
_Z4stepPiS_ii:                          # @_Z4stepPiS_ii
	.cfi_startproc
# %bb.0:
                                        # kill: def $ecx killed $ecx def $rcx
	testl	%ecx, %ecx
	jle	.LBB2_9
# %bb.1:                                # %.preheader42.lr.ph
	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
	.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, %r8d
	movl	%edx, %r10d
	movl	%ecx, %eax
	movq	%rax, -8(%rsp)                  # 8-byte Spill
	xorl	%ebx, %ebx
.LBB2_2:                                # %.preheader42
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB2_4 Depth 2
                                        #       Child Loop BB2_5 Depth 3
	testl	%r8d, %r8d
	jle	.LBB2_7
# %bb.3:                                # %.preheader.lr.ph
                                        #   in Loop: Header=BB2_2 Depth=1
	movl	%ebx, %r14d
	imull	%r8d, %r14d
	leal	(%rbx,%rcx), %ebp
	xorl	%r15d, %r15d
.LBB2_4:                                # %.preheader
                                        #   Parent Loop BB2_2 Depth=1
                                        # =>  This Loop Header: Depth=2
                                        #       Child Loop BB2_5 Depth 3
	leal	(%r15,%r8), %r13d
	xorl	%r12d, %r12d
	movq	$-64, %r11
.LBB2_5:                                #   Parent Loop BB2_2 Depth=1
                                        #     Parent Loop BB2_4 Depth=2
                                        # =>    This Inner Loop Header: Depth=3
	movl	_ZL7offsets+64(%r11), %eax
	addl	%r13d, %eax
	cltd
	idivl	%r8d
	movl	%edx, %r9d
	movl	_ZL7offsets+68(%r11), %eax
	addl	%ebp, %eax
	cltd
	idivl	%ecx
	imull	%r8d, %edx
	addl	%r9d, %edx
	movslq	%edx, %rax
	cmpl	$1, (%rdi,%rax,4)
	sbbl	$-1, %r12d
	addq	$8, %r11
	jne	.LBB2_5
# %bb.6:                                #   in Loop: Header=BB2_4 Depth=2
	leaq	(%r15,%r14), %rax
	movl	$0, (%rsi,%rax,4)
	cmpl	$0, (%rdi,%rax,4)
	setne	%dl
	cmpl	$2, %r12d
	sete	%r9b
	andb	%dl, %r9b
	cmpl	$3, %r12d
	sete	%dl
	orb	%r9b, %dl
	movzbl	%dl, %edx
	movl	%edx, (%rsi,%rax,4)
	incq	%r15
	cmpq	%r10, %r15
	jne	.LBB2_4
.LBB2_7:                                # %._crit_edge
                                        #   in Loop: Header=BB2_2 Depth=1
	incq	%rbx
	cmpq	-8(%rsp), %rbx                  # 8-byte Folded Reload
	jne	.LBB2_2
# %bb.8:
	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
	.cfi_restore %rbx
	.cfi_restore %r12
	.cfi_restore %r13
	.cfi_restore %r14
	.cfi_restore %r15
	.cfi_restore %rbp
.LBB2_9:                                # %._crit_edge47
	retq
.Lfunc_end2:
	.size	_Z4stepPiS_ii, .Lfunc_end2-_Z4stepPiS_ii
	.cfi_endproc
                                        # -- End function
	.globl	_Z10fill_boardPiii              # -- Begin function _Z10fill_boardPiii
	.type	_Z10fill_boardPiii,@function
_Z10fill_boardPiii:                     # @_Z10fill_boardPiii
	.cfi_startproc
# %bb.0:
	imull	%edx, %esi
	testl	%esi, %esi
	jle	.LBB3_4
# %bb.1:                                # %.lr.ph.preheader
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset %rbx, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movq	%rdi, %rbx
	movl	%esi, %r14d
	xorl	%r15d, %r15d
.LBB3_2:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	callq	rand
	movl	%eax, %ecx
	shrl	$31, %ecx
	addl	%eax, %ecx
	andl	$-2, %ecx
	subl	%ecx, %eax
	movl	%eax, (%rbx,%r15,4)
	incq	%r15
	cmpq	%r15, %r14
	jne	.LBB3_2
# %bb.3:
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	.cfi_restore %rbx
	.cfi_restore %r14
	.cfi_restore %r15
.LBB3_4:                                # %._crit_edge
	retq
.Lfunc_end3:
	.size	_Z10fill_boardPiii, .Lfunc_end3-_Z10fill_boardPiii
	.cfi_endproc
                                        # -- End function
	.globl	_Z11print_boardPiii             # -- Begin function _Z11print_boardPiii
	.type	_Z11print_boardPiii,@function
_Z11print_boardPiii:                    # @_Z11print_boardPiii
	.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	$24, %rsp
	.cfi_def_cfa_offset 80
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movq	%rdi, 8(%rsp)                   # 8-byte Spill
	testl	%edx, %edx
	jle	.LBB4_6
# %bb.1:                                # %.preheader.lr.ph
	movl	%esi, %ebx
	movl	%edx, %eax
	movq	%rax, 16(%rsp)                  # 8-byte Spill
	movl	%esi, %r12d
	xorl	%r13d, %r13d
	xorl	%ebp, %ebp
.LBB4_2:                                # %.preheader
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB4_4 Depth 2
	testl	%ebx, %ebx
	jle	.LBB4_5
# %bb.3:                                # %.lr.ph
                                        #   in Loop: Header=BB4_2 Depth=1
	movl	%r13d, %eax
	movq	8(%rsp), %rcx                   # 8-byte Reload
	leaq	(%rcx,%rax,4), %r14
	xorl	%r15d, %r15d
.LBB4_4:                                #   Parent Loop BB4_2 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	xorl	%eax, %eax
	cmpl	$0, (%r14,%r15,4)
	setne	%al
	leal	(%rax,%rax,2), %edi
	addl	$32, %edi
	callq	putchar@PLT
	incq	%r15
	cmpq	%r15, %r12
	jne	.LBB4_4
.LBB4_5:                                # %._crit_edge
                                        #   in Loop: Header=BB4_2 Depth=1
	movl	$10, %edi
	callq	putchar@PLT
	incq	%rbp
	addl	%ebx, %r13d
	cmpq	16(%rsp), %rbp                  # 8-byte Folded Reload
	jne	.LBB4_2
.LBB4_6:                                # %._crit_edge15
	movl	$.Lstr, %edi
	addq	$24, %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	puts@PLT                        # TAILCALL
.Lfunc_end4:
	.size	_Z11print_boardPiii, .Lfunc_end4-_Z11print_boardPiii
	.cfi_endproc
                                        # -- End function
	.globl	_Z7animatePiS_ii                # -- Begin function _Z7animatePiS_ii
	.type	_Z7animatePiS_ii,@function
_Z7animatePiS_ii:                       # @_Z7animatePiS_ii
	.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, %ebp
	movq	%rsi, %r14
	movq	%rdi, %r15
	xorps	%xmm0, %xmm0
	movq	%rsp, %rax
	movaps	%xmm0, (%rax)
	movslq	%edx, %rax
	movslq	%ecx, %r13
	imulq	%rax, %r13
	shlq	$2, %r13
	xorl	%r12d, %r12d
.LBB5_1:                                # =>This Inner Loop Header: Depth=1
	movl	$.L.str.3, %edi
	movl	%r12d, %esi
	xorl	%eax, %eax
	callq	printf
	movq	%r15, %rdi
	movl	%ebp, %esi
	movl	%ebx, %edx
	callq	_Z11print_boardPiii
	movq	%r15, %rdi
	movq	%r14, %rsi
	movl	%ebp, %edx
	movl	%ebx, %ecx
	callq	_Z8step_devPiS_ii
	movq	%r15, %rdi
	movq	%r14, %rsi
	movq	%r13, %rdx
	callq	memcpy@PLT
	movq	%rsp, %rdi
	leaq	24(%rsp), %rsi
	callq	nanosleep
	incl	%r12d
	cmpl	$2000, %r12d                    # imm = 0x7D0
	jne	.LBB5_1
# %bb.2:
	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_end5:
	.size	_Z7animatePiS_ii, .Lfunc_end5-_Z7animatePiS_ii
	.cfi_endproc
                                        # -- End function
	.globl	main                            # -- Begin function main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset %rbx, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movl	$13160, %edi                    # imm = 0x3368
	callq	malloc
	movq	%rax, %rbx
	movl	$13160, %edi                    # imm = 0x3368
	callq	malloc
	movq	%rax, %r14
	xorl	%r15d, %r15d
.LBB6_1:                                # %.lr.ph.i
                                        # =>This Inner Loop Header: Depth=1
	callq	rand
	movl	%eax, %ecx
	shrl	$31, %ecx
	addl	%eax, %ecx
	andl	$-2, %ecx
	subl	%ecx, %eax
	movl	%eax, (%rbx,%r15,4)
	incq	%r15
	cmpq	$3290, %r15                     # imm = 0xCDA
	jne	.LBB6_1
# %bb.2:                                # %_Z10fill_boardPiii.exit
	movq	%rbx, %rdi
	movq	%r14, %rsi
	movl	$70, %edx
	movl	$47, %ecx
	callq	_Z7animatePiS_ii
	movq	%rbx, %rdi
	callq	free
	movq	%r14, %rdi
	callq	free
	xorl	%eax, %eax
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end6:
	.size	main, .Lfunc_end6-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	.LBB7_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB7_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z11step_kernelPiS_ii, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$0, 8(%rsp)
	movl	$1, (%rsp)
	movl	$offsets_dev, %esi
	movl	$.L__unnamed_2, %edx
	movl	$.L__unnamed_2, %ecx
	movl	$64, %r9d
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	callq	__hipRegisterVar
	movl	$__hip_module_dtor, %edi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end7:
	.size	__hip_module_ctor, .Lfunc_end7-__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	.LBB8_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
.LBB8_2:
	retq
.Lfunc_end8:
	.size	__hip_module_dtor, .Lfunc_end8-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	offsets_dev,@object             # @offsets_dev
	.local	offsets_dev
	.comm	offsets_dev,64,16
	.type	_Z11step_kernelPiS_ii,@object   # @_Z11step_kernelPiS_ii
	.section	.rodata,"a",@progbits
	.globl	_Z11step_kernelPiS_ii
	.p2align	3, 0x0
_Z11step_kernelPiS_ii:
	.quad	_Z26__device_stub__step_kernelPiS_ii
	.size	_Z11step_kernelPiS_ii, 8

	.type	_ZL7offsets,@object             # @_ZL7offsets
	.p2align	4, 0x0
_ZL7offsets:
	.long	4294967295                      # 0xffffffff
	.long	1                               # 0x1
	.long	0                               # 0x0
	.long	1                               # 0x1
	.long	1                               # 0x1
	.long	1                               # 0x1
	.long	4294967295                      # 0xffffffff
	.long	0                               # 0x0
	.long	1                               # 0x1
	.long	0                               # 0x0
	.zero	8,255
	.long	0                               # 0x0
	.long	4294967295                      # 0xffffffff
	.long	1                               # 0x1
	.long	4294967295                      # 0xffffffff
	.size	_ZL7offsets, 64

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

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z11step_kernelPiS_ii"
	.size	.L__unnamed_1, 22

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"offsets_dev"
	.size	.L__unnamed_2, 12

	.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	"-----"
	.size	.Lstr, 6

	.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 _Z26__device_stub__step_kernelPiS_ii
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym offsets_dev
	.addrsig_sym _Z11step_kernelPiS_ii
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z11step_kernelPiS_ii   ; -- Begin function _Z11step_kernelPiS_ii
	.globl	_Z11step_kernelPiS_ii
	.p2align	8
	.type	_Z11step_kernelPiS_ii,@function
_Z11step_kernelPiS_ii:                  ; @_Z11step_kernelPiS_ii
; %bb.0:
	s_load_b64 s[2:3], s[0:1], 0x10
	s_mov_b64 s[8:9], 0
	s_waitcnt lgkmcnt(0)
	s_ashr_i32 s4, s2, 31
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_add_i32 s5, s2, s4
	s_xor_b32 s10, s5, s4
	s_load_b32 s5, s[0:1], 0x24
	v_cvt_f32_u32_e32 v1, s10
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_4) | instid1(VALU_DEP_1)
	v_rcp_iflag_f32_e32 v1, v1
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v1, 0x4f7ffffe, v1
	s_waitcnt lgkmcnt(0)
	s_and_b32 s5, s5, 0xffff
	v_cvt_u32_f32_e32 v3, v1
	v_mad_u64_u32 v[1:2], null, s15, s5, v[0:1]
	s_sub_i32 s5, 0, s10
	s_delay_alu instid0(VALU_DEP_2) | instid1(SALU_CYCLE_1)
	v_mul_lo_u32 v0, s5, v3
	s_ashr_i32 s5, s3, 31
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	s_add_i32 s6, s3, s5
	v_ashrrev_i32_e32 v2, 31, v1
	s_xor_b32 s11, s6, s5
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_cvt_f32_u32_e32 v6, s11
	v_mul_hi_u32 v0, v3, v0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_nc_u32_e32 v4, v1, v2
	v_rcp_iflag_f32_e32 v6, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_xor_b32_e32 v4, v4, v2
	v_add_nc_u32_e32 v0, v3, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_1)
	v_mul_hi_u32 v3, v4, v0
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v6, 0x4f7ffffe, v6
	v_mul_lo_u32 v5, v3, s10
	v_sub_nc_u32_e32 v4, v4, v5
	v_add_nc_u32_e32 v5, 1, v3
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_cmp_le_u32_e32 vcc_lo, s10, v4
	v_subrev_nc_u32_e32 v7, s10, v4
	v_cndmask_b32_e32 v3, v3, v5, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_nc_u32_e32 v5, 1, v3
	v_cndmask_b32_e32 v4, v4, v7, vcc_lo
	v_xor_b32_e32 v7, s4, v2
	s_sub_i32 s4, 0, s11
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_cmp_le_u32_e32 vcc_lo, s10, v4
	v_cvt_u32_f32_e32 v4, v6
	v_cndmask_b32_e32 v3, v3, v5, vcc_lo
	v_mul_lo_u32 v5, s4, v4
	s_load_b128 s[4:7], s[0:1], 0x0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_xor_b32_e32 v3, v3, v7
	v_mul_hi_u32 v5, v4, v5
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_nc_u32_e32 v6, v3, v7
	v_mul_lo_u32 v3, v6, s2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_sub_nc_u32_e32 v7, v1, v3
	v_add_nc_u32_e32 v3, v4, v5
	v_dual_mov_b32 v4, 0 :: v_dual_add_nc_u32 v5, s3, v6
	s_delay_alu instid0(VALU_DEP_3)
	v_add_nc_u32_e32 v6, s2, v7
.LBB0_1:                                ; =>This Inner Loop Header: Depth=1
	s_getpc_b64 s[0:1]
	s_add_u32 s0, s0, offsets_dev@rel32@lo+4
	s_addc_u32 s1, s1, offsets_dev@rel32@hi+12
	s_add_u32 s0, s8, s0
	s_addc_u32 s1, s9, s1
	s_add_u32 s8, s8, 8
	s_load_b64 s[0:1], s[0:1], 0x0
	s_addc_u32 s9, s9, 0
	s_cmp_lg_u32 s8, 64
	s_waitcnt lgkmcnt(0)
	v_add_nc_u32_e32 v8, s1, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v10, 31, v8
	v_add_nc_u32_e32 v8, v8, v10
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v8, v8, v10
	v_mul_hi_u32 v12, v8, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v12, v12, s11
	v_sub_nc_u32_e32 v8, v8, v12
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_subrev_nc_u32_e32 v12, s11, v8
	v_cmp_le_u32_e32 vcc_lo, s11, v8
	v_dual_cndmask_b32 v8, v8, v12 :: v_dual_add_nc_u32 v7, s0, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_ashrrev_i32_e32 v9, 31, v7
	v_subrev_nc_u32_e32 v12, s11, v8
	v_cmp_le_u32_e32 vcc_lo, s11, v8
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_cndmask_b32 v8, v8, v12 :: v_dual_add_nc_u32 v7, v7, v9
	v_xor_b32_e32 v7, v7, v9
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_xor_b32_e32 v8, v8, v10
	v_mul_hi_u32 v11, v7, v0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_sub_nc_u32_e32 v10, v8, v10
	v_mul_lo_u32 v11, v11, s10
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_nc_u32_e32 v7, v7, v11
	v_subrev_nc_u32_e32 v11, s10, v7
	v_cmp_le_u32_e64 s0, s10, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e64 v7, v7, v11, s0
	v_subrev_nc_u32_e32 v11, s10, v7
	v_cmp_le_u32_e64 s0, s10, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e64 v7, v7, v11, s0
	v_xor_b32_e32 v7, v7, v9
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_nc_u32_e32 v7, v7, v9
	v_mad_u64_u32 v[8:9], null, v10, s2, v[7:8]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v9, 31, v8
	v_lshlrev_b64 v[7:8], 2, v[8:9]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v7, vcc_lo, s4, v7
	v_add_co_ci_u32_e32 v8, vcc_lo, s5, v8, vcc_lo
	global_load_b32 v7, v[7:8], off
	s_waitcnt vmcnt(0)
	v_cmp_ne_u32_e32 vcc_lo, 0, v7
	v_add_co_ci_u32_e32 v4, vcc_lo, 0, v4, vcc_lo
	s_cbranch_scc1 .LBB0_1
; %bb.2:
	v_lshlrev_b64 v[2:3], 2, v[1:2]
	v_mov_b32_e32 v0, 0
	s_mul_i32 s0, s3, s2
	s_mov_b32 s2, exec_lo
	v_cmpx_gt_i32_e64 s0, v1
	s_cbranch_execz .LBB0_4
; %bb.3:
	v_add_co_u32 v0, vcc_lo, s4, v2
	v_add_co_ci_u32_e32 v1, vcc_lo, s5, v3, vcc_lo
	v_cmp_eq_u32_e32 vcc_lo, 2, v4
	v_cmp_eq_u32_e64 s1, 3, v4
	global_load_b32 v0, v[0:1], off
	s_waitcnt vmcnt(0)
	v_cmp_ne_u32_e64 s0, 0, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b32 s0, s0, vcc_lo
	s_or_b32 s0, s0, s1
	s_delay_alu instid0(SALU_CYCLE_1)
	v_cndmask_b32_e64 v0, 0, 1, s0
.LBB0_4:                                ; %._crit_edge
	s_or_b32 exec_lo, exec_lo, s2
	v_add_co_u32 v1, vcc_lo, s6, v2
	v_add_co_ci_u32_e32 v2, vcc_lo, s7, v3, vcc_lo
	global_store_b32 v[1:2], v0, off
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z11step_kernelPiS_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_end0:
	.size	_Z11step_kernelPiS_ii, .Lfunc_end0-_Z11step_kernelPiS_ii
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 784
; 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
	.protected	offsets_dev             ; @offsets_dev
	.type	offsets_dev,@object
	.data
	.globl	offsets_dev
	.p2align	4, 0x0
offsets_dev:
	.long	4294967295                      ; 0xffffffff
	.long	1                               ; 0x1
	.long	0                               ; 0x0
	.long	1                               ; 0x1
	.long	1                               ; 0x1
	.long	1                               ; 0x1
	.long	4294967295                      ; 0xffffffff
	.long	0                               ; 0x0
	.long	1                               ; 0x1
	.long	0                               ; 0x0
	.zero	8,255
	.long	0                               ; 0x0
	.long	4294967295                      ; 0xffffffff
	.long	1                               ; 0x1
	.long	4294967295                      ; 0xffffffff
	.size	offsets_dev, 64

	.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 offsets_dev
	.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:     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:           _Z11step_kernelPiS_ii
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z11step_kernelPiS_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
