/*  LA-CC-16080
    Copyright © 2016 Priscilla Kelly and Los Alamos National Laboratory. All Rights Reserved.

    Redistribution and use in source and binary forms, with or without modification, are permitted provided that the following conditions are met:

    1. Redistributions of source code must retain the above copyright notice, this list of conditions and the following disclaimer.

    2. Redistributions in binary form must reproduce the above copyright notice, this list of conditions and the following disclaimer in the documentation and/or other materials provided with the distribution.

    3. The name of the author may not be used to endorse or promote products derived from this software without specific prior written permission.

    THIS SOFTWARE IS PROVIDED BY Priscilla Kelly and Los Alamos National Laboratory "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE AUTHOR BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.

   	Priscilla Kelly <priscilla.noreen@gmail.com>
*/
#include "stdio.h"
#include "stdlib.h"
#define maxThread 32
/*******************************************/
/* Cuda kernel to apply the rules of Life  */
/*******************************************/
__global__ void applyRules(int row,int col,int *update, int *hold) {
	int threadMax = blockDim.x;
	int blockID = blockIdx.x;
	int threadID = threadIdx.x;
	int linID;
	if (blockID == 1) {
		linID = threadID;
	} else {
		linID = blockID*threadMax+threadID;
	}
	int elements = (row-2)*(col-2);
	
	int i = linID%(col-2);
	int j = linID/(row-2);
	int loc = col + i*col + j + 1;
	if (linID < elements) {
		int liveCells = 0;   

		int n, s, e, w, nw, ne, sw, se; // location in halo

		n = loc-col;
		nw = n-1;
		ne = n+1;
		w = loc-1;
		e = loc+1;
		s = loc+col;
    	sw = s-1;
    	se = s+1; 

		liveCells = hold[nw] + hold[n] + hold[ne]
				  + hold[w]            + hold[e]
		    	  + hold[sw] + hold[s] + hold[se]; 	
	
		// Apply Rules
		if (hold[loc] == 0) {
			if (liveCells == 3) {
				update[loc] = 1; // reproduction
			} else {
				update[loc] = 0; // remain dead
			}
		} else {  
			if (liveCells < 2){ 
				update[loc] = 0; // under population
			} else {
				if (liveCells < 4) {
					update[loc] = 1; // survivor
				} else {
					 update[loc] = 0; // over population
				}
			}
		}
	}
}
/*******************************************/
/* Cuda kernel to upload N/S halo elements */
/*******************************************/
__global__ void add_NS_Halo(int row,int col,int *haloMat,int *subMat) {
	int b = blockIdx.x;
	int t = threadIdx.x;
	// add North portion	
	if (b == 0) {
		subMat[t] = haloMat[t];
	}
	// add South portion
	if (b == 1) {
		int subLoc = col*(row-1)+t;
		int haloLoc = (row+col)+t;
		subMat[subLoc] = haloMat[haloLoc];
	}
}
/*******************************************/
/* Cuda kernel to upload E/W halo elements */
/*******************************************/
__global__ void add_EW_Halo(int row,int col,int *haloMat,int *subMat) {
	int b = blockIdx.x;
	int t = threadIdx.x;
	// add East portion	
	if (b == 0) {
		int subLoc = col-1+(col)*t;
		int haloLoc = col+t;
		subMat[subLoc] = haloMat[haloLoc];
	}
	// add the West portion
	if (b == 1) {
		int subLoc = (col)*t;
		int haloLoc = 2*row+col+t;
		subMat[subLoc] = haloMat[haloLoc];
	}
}

/*******************************************/
/* Cuda kernel to get N/S halo elements */
/*******************************************/
__global__ void get_NS_Halo(int row,int col,int *haloMat,int *subMat) {
	int b = blockIdx.x;
	int t = threadIdx.x;
	// add North portion	
	if (b == 0) {
		haloMat[t]=subMat[t+col]; 
	}
	// add South portion
	if (b == 1) {
		int subLoc = col*(row-2)+t;
		int haloLoc = (row+col)+t;
		haloMat[haloLoc]=subMat[subLoc];
	}
}
/*******************************************/
/* Cuda kernel to get E/W halo elements */
/*******************************************/
__global__ void get_EW_Halo(int row,int col,int *haloMat,int *subMat) {
	int b = blockIdx.x;
	int t = threadIdx.x;
	// add East portion	
	if (b == 0) {
		int subLoc = (col-2)+col*t;
		int haloLoc = col+t; 
		haloMat[haloLoc]=subMat[subLoc];
	}
	// add the West portion
	if (b == 1) {
		int subLoc = 1+col*t;
		int haloLoc = col+2*row+t;
		haloMat[haloLoc]=subMat[subLoc] ;
	}
}
/***************************************/
/* External c subroutine for CUDA      */
/***************************************/
extern "C" void call_cuda_applyRules(int flag,int rows, int cols,int *halo, int *halo_dev, int *update, int *hold) {

	/**************************************************/
	/* Get the values to exchange over MPI */
	/**************************************************/
	if (flag == 0) {
	int haloSize = sizeof(int)*2*(rows+cols);
	cudaError_t err = cudaSuccess;
		
	// Add the North and South rows to hold
	get_NS_Halo<<<2, cols>>>(rows,cols,halo_dev,hold);
	err = cudaGetLastError();	
	if (err != cudaSuccess) {
		fprintf(stderr, "Failed to launch Kernel (error code %s)!\n", cudaGetErrorString(err));
		exit(EXIT_FAILURE);
	}
	
	cudaDeviceSynchronize();
	
	// Add the East and West columns to hold
	get_EW_Halo<<<2, rows>>>(rows,cols,halo_dev,hold);
	err = cudaGetLastError();	
	if (err != cudaSuccess) {
		fprintf(stderr, "Failed to launch Kernel (error code %s)!\n", cudaGetErrorString(err));
		exit(EXIT_FAILURE);
	}
	
	cudaDeviceSynchronize();
	
	err = cudaMemcpy(halo,halo_dev,haloSize,cudaMemcpyDeviceToHost);	
	return;
	}

	/*****************************************************/
	/* Update hold with halo, then apply rules to update */
	/*****************************************************/
	if (flag == 1) {
	int haloSize = sizeof(int)*2*(rows+cols);
	cudaError_t err = cudaSuccess;
	
	// Copy updated halo to GPU
	err = cudaMemcpy(halo_dev,halo,haloSize,cudaMemcpyHostToDevice);	
	
	// Add the North and South rows to hold
	add_NS_Halo<<<2, cols>>>(rows,cols,halo_dev,hold);
	err = cudaGetLastError();	
	if (err != cudaSuccess) {
		fprintf(stderr, "Failed to launch Kernel (error code %s)!\n", cudaGetErrorString(err));
		exit(EXIT_FAILURE);
	}
	
	cudaDeviceSynchronize();

	// Add the East and West columns to hold
	add_EW_Halo<<<2, rows>>>(rows,cols,halo_dev,hold);
	err = cudaGetLastError();	
	if (err != cudaSuccess) {
		fprintf(stderr, "Failed to launch Kernel (error code %s)!\n", cudaGetErrorString(err));
		exit(EXIT_FAILURE);
	}
	cudaDeviceSynchronize();
	// Apply the Rules

	int N = (rows-2)*(cols-2);
	int threadCnt;
	int blockCnt;
	
	//threads per block
	threadCnt = maxThread;

	//blocks
	int check = N/maxThread;
	if (check == 0) {
		blockCnt = 1;
	} else {
		blockCnt = check + 1;
	}
	applyRules<<<blockCnt, threadCnt>>>(rows,cols,update,hold);
	if (err != cudaSuccess) {
		fprintf(stderr, "Failed to copy halo to device (error code %s)!\n", cudaGetErrorString(err));
		exit(EXIT_FAILURE);
	}
	cudaDeviceSynchronize();

	return;
	}
}

	.file	"tmpxft_0027922e_00000000-6_applyRules_CUDA.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2030:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2030:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z34__device_stub__Z10applyRulesiiPiS_iiPiS_
	.type	_Z34__device_stub__Z10applyRulesiiPiS_iiPiS_, @function
_Z34__device_stub__Z10applyRulesiiPiS_iiPiS_:
.LFB2052:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movl	%edi, 28(%rsp)
	leaq	56(%rsp), %rdi
	movl	%esi, 24(%rsp)
	leaq	68(%rsp), %rsi
	movq	%rdx, 16(%rsp)
	leaq	40(%rsp), %rdx
	movq	%rcx, 8(%rsp)
	leaq	48(%rsp), %rcx
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	28(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 104(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	16(%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	.L2
	pushq	48(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z10applyRulesiiPiS_(%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
.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
.LFE2052:
	.size	_Z34__device_stub__Z10applyRulesiiPiS_iiPiS_, .-_Z34__device_stub__Z10applyRulesiiPiS_iiPiS_
	.globl	_Z10applyRulesiiPiS_
	.type	_Z10applyRulesiiPiS_, @function
_Z10applyRulesiiPiS_:
.LFB2053:
	.cfi_startproc
	endbr64
	jmp	_Z34__device_stub__Z10applyRulesiiPiS_iiPiS_
	.cfi_endproc
.LFE2053:
	.size	_Z10applyRulesiiPiS_, .-_Z10applyRulesiiPiS_
	.globl	_Z35__device_stub__Z11add_NS_HaloiiPiS_iiPiS_
	.type	_Z35__device_stub__Z11add_NS_HaloiiPiS_iiPiS_, @function
_Z35__device_stub__Z11add_NS_HaloiiPiS_iiPiS_:
.LFB2054:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movl	%edi, 28(%rsp)
	leaq	56(%rsp), %rdi
	movl	%esi, 24(%rsp)
	leaq	68(%rsp), %rsi
	movq	%rdx, 16(%rsp)
	leaq	40(%rsp), %rdx
	movq	%rcx, 8(%rsp)
	leaq	48(%rsp), %rcx
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	28(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 104(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	16(%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	_Z11add_NS_HaloiiPiS_(%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
.LFE2054:
	.size	_Z35__device_stub__Z11add_NS_HaloiiPiS_iiPiS_, .-_Z35__device_stub__Z11add_NS_HaloiiPiS_iiPiS_
	.globl	_Z11add_NS_HaloiiPiS_
	.type	_Z11add_NS_HaloiiPiS_, @function
_Z11add_NS_HaloiiPiS_:
.LFB2055:
	.cfi_startproc
	endbr64
	jmp	_Z35__device_stub__Z11add_NS_HaloiiPiS_iiPiS_
	.cfi_endproc
.LFE2055:
	.size	_Z11add_NS_HaloiiPiS_, .-_Z11add_NS_HaloiiPiS_
	.globl	_Z35__device_stub__Z11add_EW_HaloiiPiS_iiPiS_
	.type	_Z35__device_stub__Z11add_EW_HaloiiPiS_iiPiS_, @function
_Z35__device_stub__Z11add_EW_HaloiiPiS_iiPiS_:
.LFB2056:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movl	%edi, 28(%rsp)
	leaq	56(%rsp), %rdi
	movl	%esi, 24(%rsp)
	leaq	68(%rsp), %rsi
	movq	%rdx, 16(%rsp)
	leaq	40(%rsp), %rdx
	movq	%rcx, 8(%rsp)
	leaq	48(%rsp), %rcx
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	28(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 104(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	16(%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	.L13
	pushq	48(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z11add_EW_HaloiiPiS_(%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
.L13:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L15
	call	__stack_chk_fail@PLT
.L15:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2056:
	.size	_Z35__device_stub__Z11add_EW_HaloiiPiS_iiPiS_, .-_Z35__device_stub__Z11add_EW_HaloiiPiS_iiPiS_
	.globl	_Z11add_EW_HaloiiPiS_
	.type	_Z11add_EW_HaloiiPiS_, @function
_Z11add_EW_HaloiiPiS_:
.LFB2057:
	.cfi_startproc
	endbr64
	jmp	_Z35__device_stub__Z11add_EW_HaloiiPiS_iiPiS_
	.cfi_endproc
.LFE2057:
	.size	_Z11add_EW_HaloiiPiS_, .-_Z11add_EW_HaloiiPiS_
	.globl	_Z35__device_stub__Z11get_NS_HaloiiPiS_iiPiS_
	.type	_Z35__device_stub__Z11get_NS_HaloiiPiS_iiPiS_, @function
_Z35__device_stub__Z11get_NS_HaloiiPiS_iiPiS_:
.LFB2058:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movl	%edi, 28(%rsp)
	leaq	56(%rsp), %rdi
	movl	%esi, 24(%rsp)
	leaq	68(%rsp), %rsi
	movq	%rdx, 16(%rsp)
	leaq	40(%rsp), %rdx
	movq	%rcx, 8(%rsp)
	leaq	48(%rsp), %rcx
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	28(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 104(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	16(%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	.L18
	pushq	48(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z11get_NS_HaloiiPiS_(%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
.L18:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L20
	call	__stack_chk_fail@PLT
.L20:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2058:
	.size	_Z35__device_stub__Z11get_NS_HaloiiPiS_iiPiS_, .-_Z35__device_stub__Z11get_NS_HaloiiPiS_iiPiS_
	.globl	_Z11get_NS_HaloiiPiS_
	.type	_Z11get_NS_HaloiiPiS_, @function
_Z11get_NS_HaloiiPiS_:
.LFB2059:
	.cfi_startproc
	endbr64
	jmp	_Z35__device_stub__Z11get_NS_HaloiiPiS_iiPiS_
	.cfi_endproc
.LFE2059:
	.size	_Z11get_NS_HaloiiPiS_, .-_Z11get_NS_HaloiiPiS_
	.globl	_Z35__device_stub__Z11get_EW_HaloiiPiS_iiPiS_
	.type	_Z35__device_stub__Z11get_EW_HaloiiPiS_iiPiS_, @function
_Z35__device_stub__Z11get_EW_HaloiiPiS_iiPiS_:
.LFB2060:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movl	%edi, 28(%rsp)
	leaq	56(%rsp), %rdi
	movl	%esi, 24(%rsp)
	leaq	68(%rsp), %rsi
	movq	%rdx, 16(%rsp)
	leaq	40(%rsp), %rdx
	movq	%rcx, 8(%rsp)
	leaq	48(%rsp), %rcx
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	28(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 104(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	16(%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	.L23
	pushq	48(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z11get_EW_HaloiiPiS_(%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
.L23:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L25
	call	__stack_chk_fail@PLT
.L25:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2060:
	.size	_Z35__device_stub__Z11get_EW_HaloiiPiS_iiPiS_, .-_Z35__device_stub__Z11get_EW_HaloiiPiS_iiPiS_
	.globl	_Z11get_EW_HaloiiPiS_
	.type	_Z11get_EW_HaloiiPiS_, @function
_Z11get_EW_HaloiiPiS_:
.LFB2061:
	.cfi_startproc
	endbr64
	jmp	_Z35__device_stub__Z11get_EW_HaloiiPiS_iiPiS_
	.cfi_endproc
.LFE2061:
	.size	_Z11get_EW_HaloiiPiS_, .-_Z11get_EW_HaloiiPiS_
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"Failed to launch Kernel (error code %s)!\n"
	.text
	.globl	call_cuda_applyRules
	.type	call_cuda_applyRules, @function
call_cuda_applyRules:
.LFB2027:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movq	%rcx, %r15
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	movq	%r8, %r14
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movl	%edx, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movl	%esi, %ebx
	subq	$40, %rsp
	.cfi_def_cfa_offset 96
	movq	96(%rsp), %r12
	testl	%edi, %edi
	jne	.L29
	leal	(%rsi,%rdx), %r13d
	movl	%edx, 20(%rsp)
	movl	$2147483649, %edi
	xorl	%r9d, %r9d
	addq	%rdi, %rdi
	xorl	%r8d, %r8d
	movl	$1, %esi
	sall	$3, %r13d
	movabsq	$4294967297, %rax
	movq	%rdi, 8(%rsp)
	movq	%rax, 24(%rsp)
	movl	28(%rsp), %ecx
	movq	20(%rsp), %rdx
	movl	$1, 16(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L30
	movq	%r12, %rcx
	movq	%r14, %rdx
	movl	%ebp, %esi
	movl	%ebx, %edi
	call	_Z35__device_stub__Z11get_NS_HaloiiPiS_iiPiS_
.L30:
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	jne	.L42
	call	cudaDeviceSynchronize@PLT
	movl	%ebx, 20(%rsp)
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movabsq	$4294967297, %rax
	movl	$2147483649, %edi
	movl	$1, %esi
	movl	$1, 16(%rsp)
	movq	%rax, 24(%rsp)
	movl	28(%rsp), %ecx
	addq	%rdi, %rdi
	movq	20(%rsp), %rdx
	movq	%rdi, 8(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L32
	movq	%r12, %rcx
	movq	%r14, %rdx
	movl	%ebp, %esi
	movl	%ebx, %edi
	call	_Z35__device_stub__Z11get_EW_HaloiiPiS_iiPiS_
.L32:
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	jne	.L42
	call	cudaDeviceSynchronize@PLT
	addq	$40, %rsp
	.cfi_remember_state
	.cfi_def_cfa_offset 56
	movslq	%r13d, %rdx
	movq	%r14, %rsi
	popq	%rbx
	.cfi_def_cfa_offset 48
	movq	%r15, %rdi
	popq	%rbp
	.cfi_def_cfa_offset 40
	movl	$2, %ecx
	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	cudaMemcpy@PLT
.L29:
	.cfi_restore_state
	movl	%edi, %eax
	decl	%eax
	jne	.L28
	leal	(%rsi,%rdx), %edx
	movq	%r8, %rdi
	movl	$1, %ecx
	movq	%r15, %rsi
	sall	$3, %edx
	movq	%r9, %r13
	movslq	%edx, %rdx
	call	cudaMemcpy@PLT
	movl	%ebp, 20(%rsp)
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movabsq	$4294967297, %rax
	movl	$2147483649, %edi
	movl	$1, %esi
	movl	$1, 16(%rsp)
	movq	%rax, 24(%rsp)
	movl	28(%rsp), %ecx
	addq	%rdi, %rdi
	movq	20(%rsp), %rdx
	movq	%rdi, 8(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L35
	movq	%r12, %rcx
	movq	%r14, %rdx
	movl	%ebp, %esi
	movl	%ebx, %edi
	call	_Z35__device_stub__Z11add_NS_HaloiiPiS_iiPiS_
.L35:
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	jne	.L42
	call	cudaDeviceSynchronize@PLT
	movl	%ebx, 20(%rsp)
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movabsq	$4294967297, %rax
	movl	$2147483649, %edi
	movl	$1, %esi
	movl	$1, 16(%rsp)
	movq	%rax, 24(%rsp)
	movl	28(%rsp), %ecx
	addq	%rdi, %rdi
	movq	20(%rsp), %rdx
	movq	%rdi, 8(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L37
	movq	%r12, %rcx
	movq	%r14, %rdx
	movl	%ebp, %esi
	movl	%ebx, %edi
	call	_Z35__device_stub__Z11add_EW_HaloiiPiS_iiPiS_
.L37:
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L38
.L42:
	call	cudaGetErrorString@PLT
	movq	stderr(%rip), %rdi
	movl	$2, %esi
	leaq	.LC0(%rip), %rdx
	movq	%rax, %rcx
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
	movl	$1, %edi
	call	exit@PLT
.L38:
	call	cudaDeviceSynchronize@PLT
	leal	-2(%rbp), %edx
	leal	-2(%rbx), %eax
	movl	$32, %esi
	imull	%edx, %eax
	movl	$134217729, %ecx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	salq	$5, %rcx
	movl	$1, 28(%rsp)
	movq	%rcx, 20(%rsp)
	cltd
	idivl	%esi
	movq	%rcx, %rdx
	movl	$1, %ecx
	incl	%eax
	movl	%eax, 8(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 12(%rsp)
	movq	8(%rsp), %rdi
	movl	16(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L39
	movq	%r12, %rcx
	movq	%r13, %rdx
	movl	%ebp, %esi
	movl	%ebx, %edi
	call	_Z34__device_stub__Z10applyRulesiiPiS_iiPiS_
.L39:
	addq	$40, %rsp
	.cfi_remember_state
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	jmp	cudaDeviceSynchronize@PLT
.L28:
	.cfi_restore_state
	addq	$40, %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	call_cuda_applyRules, .-call_cuda_applyRules
	.section	.rodata.str1.1
.LC1:
	.string	"_Z11get_EW_HaloiiPiS_"
.LC2:
	.string	"_Z11get_NS_HaloiiPiS_"
.LC3:
	.string	"_Z11add_EW_HaloiiPiS_"
.LC4:
	.string	"_Z11add_NS_HaloiiPiS_"
.LC5:
	.string	"_Z10applyRulesiiPiS_"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2063:
	.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	.LC1(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z11get_EW_HaloiiPiS_(%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	.LC2(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z11get_NS_HaloiiPiS_(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC3(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z11add_EW_HaloiiPiS_(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC4(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z11add_NS_HaloiiPiS_(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC5(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z10applyRulesiiPiS_(%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
.LFE2063:
	.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:

/*  LA-CC-16080
    Copyright © 2016 Priscilla Kelly and Los Alamos National Laboratory. All Rights Reserved.

    Redistribution and use in source and binary forms, with or without modification, are permitted provided that the following conditions are met:

    1. Redistributions of source code must retain the above copyright notice, this list of conditions and the following disclaimer.

    2. Redistributions in binary form must reproduce the above copyright notice, this list of conditions and the following disclaimer in the documentation and/or other materials provided with the distribution.

    3. The name of the author may not be used to endorse or promote products derived from this software without specific prior written permission.

    THIS SOFTWARE IS PROVIDED BY Priscilla Kelly and Los Alamos National Laboratory "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE AUTHOR BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.

   	Priscilla Kelly <priscilla.noreen@gmail.com>
*/

#include <hip/hip_runtime.h>
#include "stdio.h"
#include "stdlib.h"
#define maxThread 32
/*******************************************/
/* Cuda kernel to apply the rules of Life  */
/*******************************************/
__global__ void applyRules(int row,int col,int *update, int *hold) {
	int threadMax = blockDim.x;
	int blockID = blockIdx.x;
	int threadID = threadIdx.x;
	int linID;
	if (blockID == 1) {
		linID = threadID;
	} else {
		linID = blockID*threadMax+threadID;
	}
	int elements = (row-2)*(col-2);
	
	int i = linID%(col-2);
	int j = linID/(row-2);
	int loc = col + i*col + j + 1;
	if (linID < elements) {
		int liveCells = 0;   

		int n, s, e, w, nw, ne, sw, se; // location in halo

		n = loc-col;
		nw = n-1;
		ne = n+1;
		w = loc-1;
		e = loc+1;
		s = loc+col;
    	sw = s-1;
    	se = s+1; 

		liveCells = hold[nw] + hold[n] + hold[ne]
				  + hold[w]            + hold[e]
		    	  + hold[sw] + hold[s] + hold[se]; 	
	
		// Apply Rules
		if (hold[loc] == 0) {
			if (liveCells == 3) {
				update[loc] = 1; // reproduction
			} else {
				update[loc] = 0; // remain dead
			}
		} else {  
			if (liveCells < 2){ 
				update[loc] = 0; // under population
			} else {
				if (liveCells < 4) {
					update[loc] = 1; // survivor
				} else {
					 update[loc] = 0; // over population
				}
			}
		}
	}
}
/*******************************************/
/* Cuda kernel to upload N/S halo elements */
/*******************************************/
__global__ void add_NS_Halo(int row,int col,int *haloMat,int *subMat) {
	int b = blockIdx.x;
	int t = threadIdx.x;
	// add North portion	
	if (b == 0) {
		subMat[t] = haloMat[t];
	}
	// add South portion
	if (b == 1) {
		int subLoc = col*(row-1)+t;
		int haloLoc = (row+col)+t;
		subMat[subLoc] = haloMat[haloLoc];
	}
}
/*******************************************/
/* Cuda kernel to upload E/W halo elements */
/*******************************************/
__global__ void add_EW_Halo(int row,int col,int *haloMat,int *subMat) {
	int b = blockIdx.x;
	int t = threadIdx.x;
	// add East portion	
	if (b == 0) {
		int subLoc = col-1+(col)*t;
		int haloLoc = col+t;
		subMat[subLoc] = haloMat[haloLoc];
	}
	// add the West portion
	if (b == 1) {
		int subLoc = (col)*t;
		int haloLoc = 2*row+col+t;
		subMat[subLoc] = haloMat[haloLoc];
	}
}

/*******************************************/
/* Cuda kernel to get N/S halo elements */
/*******************************************/
__global__ void get_NS_Halo(int row,int col,int *haloMat,int *subMat) {
	int b = blockIdx.x;
	int t = threadIdx.x;
	// add North portion	
	if (b == 0) {
		haloMat[t]=subMat[t+col]; 
	}
	// add South portion
	if (b == 1) {
		int subLoc = col*(row-2)+t;
		int haloLoc = (row+col)+t;
		haloMat[haloLoc]=subMat[subLoc];
	}
}
/*******************************************/
/* Cuda kernel to get E/W halo elements */
/*******************************************/
__global__ void get_EW_Halo(int row,int col,int *haloMat,int *subMat) {
	int b = blockIdx.x;
	int t = threadIdx.x;
	// add East portion	
	if (b == 0) {
		int subLoc = (col-2)+col*t;
		int haloLoc = col+t; 
		haloMat[haloLoc]=subMat[subLoc];
	}
	// add the West portion
	if (b == 1) {
		int subLoc = 1+col*t;
		int haloLoc = col+2*row+t;
		haloMat[haloLoc]=subMat[subLoc] ;
	}
}
/***************************************/
/* External c subroutine for CUDA      */
/***************************************/
extern "C" void call_cuda_applyRules(int flag,int rows, int cols,int *halo, int *halo_dev, int *update, int *hold) {

	/**************************************************/
	/* Get the values to exchange over MPI */
	/**************************************************/
	if (flag == 0) {
	int haloSize = sizeof(int)*2*(rows+cols);
	hipError_t err = hipSuccess;
		
	// Add the North and South rows to hold
	get_NS_Halo<<<2, cols>>>(rows,cols,halo_dev,hold);
	err = hipGetLastError();	
	if (err != hipSuccess) {
		fprintf(stderr, "Failed to launch Kernel (error code %s)!\n", hipGetErrorString(err));
		exit(EXIT_FAILURE);
	}
	
	hipDeviceSynchronize();
	
	// Add the East and West columns to hold
	get_EW_Halo<<<2, rows>>>(rows,cols,halo_dev,hold);
	err = hipGetLastError();	
	if (err != hipSuccess) {
		fprintf(stderr, "Failed to launch Kernel (error code %s)!\n", hipGetErrorString(err));
		exit(EXIT_FAILURE);
	}
	
	hipDeviceSynchronize();
	
	err = hipMemcpy(halo,halo_dev,haloSize,hipMemcpyDeviceToHost);	
	return;
	}

	/*****************************************************/
	/* Update hold with halo, then apply rules to update */
	/*****************************************************/
	if (flag == 1) {
	int haloSize = sizeof(int)*2*(rows+cols);
	hipError_t err = hipSuccess;
	
	// Copy updated halo to GPU
	err = hipMemcpy(halo_dev,halo,haloSize,hipMemcpyHostToDevice);	
	
	// Add the North and South rows to hold
	add_NS_Halo<<<2, cols>>>(rows,cols,halo_dev,hold);
	err = hipGetLastError();	
	if (err != hipSuccess) {
		fprintf(stderr, "Failed to launch Kernel (error code %s)!\n", hipGetErrorString(err));
		exit(EXIT_FAILURE);
	}
	
	hipDeviceSynchronize();

	// Add the East and West columns to hold
	add_EW_Halo<<<2, rows>>>(rows,cols,halo_dev,hold);
	err = hipGetLastError();	
	if (err != hipSuccess) {
		fprintf(stderr, "Failed to launch Kernel (error code %s)!\n", hipGetErrorString(err));
		exit(EXIT_FAILURE);
	}
	hipDeviceSynchronize();
	// Apply the Rules

	int N = (rows-2)*(cols-2);
	int threadCnt;
	int blockCnt;
	
	//threads per block
	threadCnt = maxThread;

	//blocks
	int check = N/maxThread;
	if (check == 0) {
		blockCnt = 1;
	} else {
		blockCnt = check + 1;
	}
	applyRules<<<blockCnt, threadCnt>>>(rows,cols,update,hold);
	if (err != hipSuccess) {
		fprintf(stderr, "Failed to copy halo to device (error code %s)!\n", hipGetErrorString(err));
		exit(EXIT_FAILURE);
	}
	hipDeviceSynchronize();

	return;
	}
}

	.text
	.file	"applyRules_CUDA.hip"
	.globl	_Z25__device_stub__applyRulesiiPiS_ # -- Begin function _Z25__device_stub__applyRulesiiPiS_
	.type	_Z25__device_stub__applyRulesiiPiS_,@function
_Z25__device_stub__applyRulesiiPiS_:    # @_Z25__device_stub__applyRulesiiPiS_
	.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	12(%rsp), %rax
	movl	%edi, (%rax)
	leaq	8(%rsp), %rdi
	movl	%esi, (%rdi)
	leaq	40(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	32(%rsp), %rdx
	movq	%rcx, (%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	$_Z10applyRulesiiPiS_, %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	_Z25__device_stub__applyRulesiiPiS_, .Lfunc_end0-_Z25__device_stub__applyRulesiiPiS_
	.cfi_endproc
                                        # -- End function
	.globl	_Z26__device_stub__add_NS_HaloiiPiS_ # -- Begin function _Z26__device_stub__add_NS_HaloiiPiS_
	.type	_Z26__device_stub__add_NS_HaloiiPiS_,@function
_Z26__device_stub__add_NS_HaloiiPiS_:   # @_Z26__device_stub__add_NS_HaloiiPiS_
	.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	12(%rsp), %rax
	movl	%edi, (%rax)
	leaq	8(%rsp), %rdi
	movl	%esi, (%rdi)
	leaq	40(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	32(%rsp), %rdx
	movq	%rcx, (%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	$_Z11add_NS_HaloiiPiS_, %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	_Z26__device_stub__add_NS_HaloiiPiS_, .Lfunc_end1-_Z26__device_stub__add_NS_HaloiiPiS_
	.cfi_endproc
                                        # -- End function
	.globl	_Z26__device_stub__add_EW_HaloiiPiS_ # -- Begin function _Z26__device_stub__add_EW_HaloiiPiS_
	.type	_Z26__device_stub__add_EW_HaloiiPiS_,@function
_Z26__device_stub__add_EW_HaloiiPiS_:   # @_Z26__device_stub__add_EW_HaloiiPiS_
	.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	12(%rsp), %rax
	movl	%edi, (%rax)
	leaq	8(%rsp), %rdi
	movl	%esi, (%rdi)
	leaq	40(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	32(%rsp), %rdx
	movq	%rcx, (%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	$_Z11add_EW_HaloiiPiS_, %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_end2:
	.size	_Z26__device_stub__add_EW_HaloiiPiS_, .Lfunc_end2-_Z26__device_stub__add_EW_HaloiiPiS_
	.cfi_endproc
                                        # -- End function
	.globl	_Z26__device_stub__get_NS_HaloiiPiS_ # -- Begin function _Z26__device_stub__get_NS_HaloiiPiS_
	.type	_Z26__device_stub__get_NS_HaloiiPiS_,@function
_Z26__device_stub__get_NS_HaloiiPiS_:   # @_Z26__device_stub__get_NS_HaloiiPiS_
	.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	12(%rsp), %rax
	movl	%edi, (%rax)
	leaq	8(%rsp), %rdi
	movl	%esi, (%rdi)
	leaq	40(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	32(%rsp), %rdx
	movq	%rcx, (%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	$_Z11get_NS_HaloiiPiS_, %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_end3:
	.size	_Z26__device_stub__get_NS_HaloiiPiS_, .Lfunc_end3-_Z26__device_stub__get_NS_HaloiiPiS_
	.cfi_endproc
                                        # -- End function
	.globl	_Z26__device_stub__get_EW_HaloiiPiS_ # -- Begin function _Z26__device_stub__get_EW_HaloiiPiS_
	.type	_Z26__device_stub__get_EW_HaloiiPiS_,@function
_Z26__device_stub__get_EW_HaloiiPiS_:   # @_Z26__device_stub__get_EW_HaloiiPiS_
	.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	12(%rsp), %rax
	movl	%edi, (%rax)
	leaq	8(%rsp), %rdi
	movl	%esi, (%rdi)
	leaq	40(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	32(%rsp), %rdx
	movq	%rcx, (%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	$_Z11get_EW_HaloiiPiS_, %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_end4:
	.size	_Z26__device_stub__get_EW_HaloiiPiS_, .Lfunc_end4-_Z26__device_stub__get_EW_HaloiiPiS_
	.cfi_endproc
                                        # -- End function
	.globl	call_cuda_applyRules            # -- Begin function call_cuda_applyRules
	.type	call_cuda_applyRules,@function
call_cuda_applyRules:                   # @call_cuda_applyRules
	.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
	pushq	%rax
	.cfi_def_cfa_offset 64
	.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	%r8, %r12
	movq	%rcx, %r13
	movl	%edx, %ebx
	movl	%esi, %r14d
	movq	64(%rsp), %r15
	cmpl	$1, %edi
	je	.LBB5_9
# %bb.1:
	testl	%edi, %edi
	jne	.LBB5_21
# %bb.2:
	movabsq	$4294967296, %rbp               # imm = 0x100000000
	movl	%ebx, %edx
	btsq	$32, %rdx
	leaq	2(%rbp), %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB5_4
# %bb.3:
	movl	%r14d, %edi
	movl	%ebx, %esi
	movq	%r12, %rdx
	movq	%r15, %rcx
	callq	_Z26__device_stub__get_NS_HaloiiPiS_
.LBB5_4:
	callq	hipGetLastError
	testl	%eax, %eax
	jne	.LBB5_22
# %bb.5:
	callq	hipDeviceSynchronize
	movl	%r14d, %edx
	orq	%rbp, %rdx
	addq	$2, %rbp
	movq	%rbp, %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB5_7
# %bb.6:
	movl	%r14d, %edi
	movl	%ebx, %esi
	movq	%r12, %rdx
	movq	%r15, %rcx
	callq	_Z26__device_stub__get_EW_HaloiiPiS_
.LBB5_7:
	callq	hipGetLastError
	testl	%eax, %eax
	jne	.LBB5_22
# %bb.8:
	addl	%r14d, %ebx
	shll	$3, %ebx
	callq	hipDeviceSynchronize
	movslq	%ebx, %rdx
	movq	%r13, %rdi
	movq	%r12, %rsi
	movl	$2, %ecx
	addq	$8, %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	hipMemcpy                       # TAILCALL
.LBB5_9:
	.cfi_def_cfa_offset 64
	movq	%r9, (%rsp)                     # 8-byte Spill
	movabsq	$4294967296, %rbp               # imm = 0x100000000
	leal	(%rbx,%r14), %eax
	shll	$3, %eax
	movslq	%eax, %rdx
	movq	%r12, %rdi
	movq	%r13, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movl	%ebx, %edx
	btsq	$32, %rdx
	leaq	2(%rbp), %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB5_11
# %bb.10:
	movl	%r14d, %edi
	movl	%ebx, %esi
	movq	%r12, %rdx
	movq	%r15, %rcx
	callq	_Z26__device_stub__add_NS_HaloiiPiS_
.LBB5_11:
	callq	hipGetLastError
	testl	%eax, %eax
	jne	.LBB5_22
# %bb.12:
	callq	hipDeviceSynchronize
	movl	%r14d, %edx
	orq	%rbp, %rdx
	leaq	2(%rbp), %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB5_14
# %bb.13:
	movl	%r14d, %edi
	movl	%ebx, %esi
	movq	%r12, %rdx
	movq	%r15, %rcx
	callq	_Z26__device_stub__add_EW_HaloiiPiS_
.LBB5_14:
	callq	hipGetLastError
	testl	%eax, %eax
	jne	.LBB5_22
# %bb.15:
	callq	hipDeviceSynchronize
	leal	-2(%r14), %ecx
	leal	-2(%rbx), %eax
	imull	%ecx, %eax
	leal	31(%rax), %edi
	cmpl	$63, %edi
	jae	.LBB5_17
# %bb.16:
	leaq	1(%rbp), %rdi
	jmp	.LBB5_18
.LBB5_21:
	addq	$8, %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
.LBB5_17:
	.cfi_def_cfa_offset 64
	testl	%eax, %eax
	cmovnsl	%eax, %edi
	sarl	$5, %edi
	incl	%edi
	orq	%rbp, %rdi
.LBB5_18:
	addq	$32, %rbp
	movl	$1, %esi
	movq	%rbp, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB5_20
# %bb.19:
	movl	%r14d, %edi
	movl	%ebx, %esi
	movq	(%rsp), %rdx                    # 8-byte Reload
	movq	%r15, %rcx
	callq	_Z25__device_stub__applyRulesiiPiS_
.LBB5_20:
	addq	$8, %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	hipDeviceSynchronize            # TAILCALL
.LBB5_22:
	.cfi_def_cfa_offset 64
	movq	stderr(%rip), %rbx
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str, %esi
	movq	%rbx, %rdi
	movq	%rax, %rdx
	xorl	%eax, %eax
	callq	fprintf
	movl	$1, %edi
	callq	exit
.Lfunc_end5:
	.size	call_cuda_applyRules, .Lfunc_end5-call_cuda_applyRules
	.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	.LBB6_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB6_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z10applyRulesiiPiS_, %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	$_Z11add_NS_HaloiiPiS_, %esi
	movl	$.L__unnamed_2, %edx
	movl	$.L__unnamed_2, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z11add_EW_HaloiiPiS_, %esi
	movl	$.L__unnamed_3, %edx
	movl	$.L__unnamed_3, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z11get_NS_HaloiiPiS_, %esi
	movl	$.L__unnamed_4, %edx
	movl	$.L__unnamed_4, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z11get_EW_HaloiiPiS_, %esi
	movl	$.L__unnamed_5, %edx
	movl	$.L__unnamed_5, %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_end6:
	.size	__hip_module_ctor, .Lfunc_end6-__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	.LBB7_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
.LBB7_2:
	retq
.Lfunc_end7:
	.size	__hip_module_dtor, .Lfunc_end7-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_Z10applyRulesiiPiS_,@object    # @_Z10applyRulesiiPiS_
	.section	.rodata,"a",@progbits
	.globl	_Z10applyRulesiiPiS_
	.p2align	3, 0x0
_Z10applyRulesiiPiS_:
	.quad	_Z25__device_stub__applyRulesiiPiS_
	.size	_Z10applyRulesiiPiS_, 8

	.type	_Z11add_NS_HaloiiPiS_,@object   # @_Z11add_NS_HaloiiPiS_
	.globl	_Z11add_NS_HaloiiPiS_
	.p2align	3, 0x0
_Z11add_NS_HaloiiPiS_:
	.quad	_Z26__device_stub__add_NS_HaloiiPiS_
	.size	_Z11add_NS_HaloiiPiS_, 8

	.type	_Z11add_EW_HaloiiPiS_,@object   # @_Z11add_EW_HaloiiPiS_
	.globl	_Z11add_EW_HaloiiPiS_
	.p2align	3, 0x0
_Z11add_EW_HaloiiPiS_:
	.quad	_Z26__device_stub__add_EW_HaloiiPiS_
	.size	_Z11add_EW_HaloiiPiS_, 8

	.type	_Z11get_NS_HaloiiPiS_,@object   # @_Z11get_NS_HaloiiPiS_
	.globl	_Z11get_NS_HaloiiPiS_
	.p2align	3, 0x0
_Z11get_NS_HaloiiPiS_:
	.quad	_Z26__device_stub__get_NS_HaloiiPiS_
	.size	_Z11get_NS_HaloiiPiS_, 8

	.type	_Z11get_EW_HaloiiPiS_,@object   # @_Z11get_EW_HaloiiPiS_
	.globl	_Z11get_EW_HaloiiPiS_
	.p2align	3, 0x0
_Z11get_EW_HaloiiPiS_:
	.quad	_Z26__device_stub__get_EW_HaloiiPiS_
	.size	_Z11get_EW_HaloiiPiS_, 8

	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"Failed to launch Kernel (error code %s)!\n"
	.size	.L.str, 42

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z10applyRulesiiPiS_"
	.size	.L__unnamed_1, 21

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"_Z11add_NS_HaloiiPiS_"
	.size	.L__unnamed_2, 22

	.type	.L__unnamed_3,@object           # @2
.L__unnamed_3:
	.asciz	"_Z11add_EW_HaloiiPiS_"
	.size	.L__unnamed_3, 22

	.type	.L__unnamed_4,@object           # @3
.L__unnamed_4:
	.asciz	"_Z11get_NS_HaloiiPiS_"
	.size	.L__unnamed_4, 22

	.type	.L__unnamed_5,@object           # @4
.L__unnamed_5:
	.asciz	"_Z11get_EW_HaloiiPiS_"
	.size	.L__unnamed_5, 22

	.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__applyRulesiiPiS_
	.addrsig_sym _Z26__device_stub__add_NS_HaloiiPiS_
	.addrsig_sym _Z26__device_stub__add_EW_HaloiiPiS_
	.addrsig_sym _Z26__device_stub__get_NS_HaloiiPiS_
	.addrsig_sym _Z26__device_stub__get_EW_HaloiiPiS_
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z10applyRulesiiPiS_
	.addrsig_sym _Z11add_NS_HaloiiPiS_
	.addrsig_sym _Z11add_EW_HaloiiPiS_
	.addrsig_sym _Z11get_NS_HaloiiPiS_
	.addrsig_sym _Z11get_EW_HaloiiPiS_
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z10applyRulesiiPiS_    ; -- Begin function _Z10applyRulesiiPiS_
	.globl	_Z10applyRulesiiPiS_
	.p2align	8
	.type	_Z10applyRulesiiPiS_,@function
_Z10applyRulesiiPiS_:                   ; @_Z10applyRulesiiPiS_
; %bb.0:
	s_clause 0x1
	s_load_b32 s4, s[0:1], 0x24
	s_load_b64 s[2:3], s[0:1], 0x0
	s_waitcnt lgkmcnt(0)
	s_and_b32 s4, s4, 0xffff
	s_cmp_lg_u32 s15, 1
	s_mul_i32 s15, s15, s4
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_3) | instid1(SALU_CYCLE_1)
	s_cselect_b32 s4, s15, 0
	s_add_i32 s2, s2, -2
	v_add_nc_u32_e32 v0, s4, v0
	s_add_i32 s4, s3, -2
	s_mul_i32 s5, s4, s2
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_cmp_gt_i32_e32 vcc_lo, s5, v0
	s_and_saveexec_b32 s5, vcc_lo
	s_cbranch_execz .LBB0_16
; %bb.1:
	s_ashr_i32 s5, s4, 31
	v_ashrrev_i32_e32 v4, 31, v0
	s_add_i32 s4, s4, s5
	s_delay_alu instid0(SALU_CYCLE_1)
	s_xor_b32 s4, s4, s5
	s_ashr_i32 s5, s2, 31
	v_cvt_f32_u32_e32 v1, s4
	s_add_i32 s2, s2, s5
	v_add_nc_u32_e32 v0, v0, v4
	s_xor_b32 s6, s2, s5
	s_sub_i32 s2, 0, s4
	v_rcp_iflag_f32_e32 v1, v1
	v_cvt_f32_u32_e32 v2, s6
	v_xor_b32_e32 v0, v0, v4
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_iflag_f32_e32 v2, v2
	s_waitcnt_depctr 0xfff
	v_dual_mul_f32 v1, 0x4f7ffffe, v1 :: v_dual_mul_f32 v2, 0x4f7ffffe, v2
	v_cvt_u32_f32_e32 v1, v1
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cvt_u32_f32_e32 v2, v2
	v_mul_lo_u32 v3, s2, v1
	s_sub_i32 s2, 0, s6
	s_delay_alu instid0(VALU_DEP_2) | instid1(SALU_CYCLE_1)
	v_mul_lo_u32 v5, s2, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_hi_u32 v3, v1, v3
	v_add_nc_u32_e32 v1, v1, v3
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_hi_u32 v3, v2, v5
	v_mul_hi_u32 v1, v0, v1
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v2, v2, v3
	v_mul_lo_u32 v1, v1, s4
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_hi_u32 v2, v0, v2
	v_sub_nc_u32_e32 v1, v0, v1
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_lo_u32 v3, v2, s6
	v_subrev_nc_u32_e32 v5, s4, v1
	v_cmp_le_u32_e32 vcc_lo, s4, v1
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_sub_nc_u32_e32 v0, v0, v3
	v_add_nc_u32_e32 v3, 1, v2
	v_cndmask_b32_e32 v1, v1, v5, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_subrev_nc_u32_e32 v6, s6, v0
	v_cmp_le_u32_e32 vcc_lo, s6, v0
	v_subrev_nc_u32_e32 v5, s4, v1
	v_cmp_le_u32_e64 s2, s4, v1
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_cndmask_b32_e32 v0, v0, v6, vcc_lo
	v_cndmask_b32_e32 v2, v2, v3, vcc_lo
	v_cndmask_b32_e64 v1, v1, v5, s2
	v_xor_b32_e32 v5, s5, v4
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_cmp_le_u32_e32 vcc_lo, s6, v0
	v_add_nc_u32_e32 v3, 1, v2
	s_delay_alu instid0(VALU_DEP_4)
	v_xor_b32_e32 v1, v1, v4
	s_load_b128 s[4:7], s[0:1], 0x8
	s_mov_b32 s0, s3
	s_ashr_i32 s1, s3, 31
	v_cndmask_b32_e32 v0, v2, v3, vcc_lo
	v_sub_nc_u32_e32 v1, v1, v4
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_xor_b32_e32 v0, v0, v5
	v_mul_lo_u32 v1, s3, v1
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_nc_u32_e32 v0, v0, v5
	v_add3_u32 v0, v1, s3, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_ashrrev_i32_e32 v1, 31, v0
	v_sub_co_u32 v4, vcc_lo, v0, s0
	v_add_nc_u32_e32 v2, s3, v0
	v_subrev_co_ci_u32_e32 v5, vcc_lo, s1, v1, vcc_lo
	v_lshlrev_b64 v[0:1], 2, v[0:1]
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_ashrrev_i32_e32 v3, 31, v2
	s_lshl_b64 s[0:1], s[0:1], 2
	v_lshlrev_b64 v[4:5], 2, v[4:5]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_lshlrev_b64 v[6:7], 2, v[2:3]
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v2, vcc_lo, s6, v4
	s_delay_alu instid0(VALU_DEP_3)
	v_add_co_ci_u32_e32 v3, vcc_lo, s7, v5, vcc_lo
	v_add_co_u32 v8, vcc_lo, s6, v0
	v_add_co_ci_u32_e32 v9, vcc_lo, s7, v1, vcc_lo
	global_load_b96 v[2:4], v[2:3], off
	v_add_co_u32 v10, vcc_lo, s6, v6
	v_add_co_ci_u32_e32 v11, vcc_lo, s7, v7, vcc_lo
	global_load_b96 v[5:7], v[8:9], off
	v_add_co_u32 v8, vcc_lo, v8, s0
	v_add_co_ci_u32_e32 v9, vcc_lo, s1, v9, vcc_lo
	s_clause 0x1
	global_load_b32 v10, v[10:11], off
	global_load_b64 v[8:9], v[8:9], off offset:4
	s_mov_b32 s0, exec_lo
	s_waitcnt vmcnt(3)
	v_add_nc_u32_e32 v2, v3, v2
	s_waitcnt vmcnt(2)
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_add3_u32 v2, v2, v4, v5
	s_waitcnt vmcnt(1)
	v_add3_u32 v2, v2, v7, v10
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_1)
	v_add3_u32 v2, v2, v8, v9
	v_cmpx_ne_u32_e32 0, v6
	s_xor_b32 s1, exec_lo, s0
	s_cbranch_execz .LBB0_11
; %bb.2:
	s_mov_b32 s0, exec_lo
	v_cmpx_lt_i32_e32 1, v2
	s_xor_b32 s2, exec_lo, s0
	s_cbranch_execz .LBB0_8
; %bb.3:
	v_add_co_u32 v0, s0, s4, v0
	s_delay_alu instid0(VALU_DEP_1)
	v_add_co_ci_u32_e64 v1, s0, s5, v1, s0
	s_mov_b32 s0, exec_lo
	v_cmpx_lt_u32_e32 3, v2
	s_xor_b32 s0, exec_lo, s0
	s_cbranch_execz .LBB0_5
; %bb.4:
	v_mov_b32_e32 v2, 0
	global_store_b32 v[0:1], v2, off offset:4
                                        ; implicit-def: $vgpr0_vgpr1
.LBB0_5:                                ; %Flow
	s_and_not1_saveexec_b32 s0, s0
	s_cbranch_execz .LBB0_7
; %bb.6:
	v_mov_b32_e32 v2, 1
	global_store_b32 v[0:1], v2, off offset:4
.LBB0_7:                                ; %Flow75
	s_or_b32 exec_lo, exec_lo, s0
                                        ; implicit-def: $vgpr0_vgpr1
.LBB0_8:                                ; %Flow76
	s_and_not1_saveexec_b32 s0, s2
	s_cbranch_execz .LBB0_10
; %bb.9:
	v_add_co_u32 v0, vcc_lo, s4, v0
	v_add_co_ci_u32_e32 v1, vcc_lo, s5, v1, vcc_lo
	v_mov_b32_e32 v2, 0
	global_store_b32 v[0:1], v2, off offset:4
.LBB0_10:                               ; %Flow77
	s_or_b32 exec_lo, exec_lo, s0
                                        ; implicit-def: $vgpr2
                                        ; implicit-def: $vgpr0_vgpr1
.LBB0_11:                               ; %Flow80
	s_and_not1_saveexec_b32 s0, s1
	s_cbranch_execz .LBB0_16
; %bb.12:
	v_add_co_u32 v0, s0, s4, v0
	s_delay_alu instid0(VALU_DEP_1)
	v_add_co_ci_u32_e64 v1, s0, s5, v1, s0
	s_mov_b32 s0, exec_lo
	v_cmpx_ne_u32_e32 3, v2
	s_xor_b32 s0, exec_lo, s0
	s_cbranch_execz .LBB0_14
; %bb.13:
	v_mov_b32_e32 v2, 0
	global_store_b32 v[0:1], v2, off offset:4
                                        ; implicit-def: $vgpr0_vgpr1
.LBB0_14:                               ; %Flow78
	s_and_not1_saveexec_b32 s0, s0
	s_cbranch_execz .LBB0_16
; %bb.15:
	v_mov_b32_e32 v2, 1
	global_store_b32 v[0:1], v2, off offset:4
.LBB0_16:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z10applyRulesiiPiS_
		.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 12
		.amdhsa_next_free_sgpr 16
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end0:
	.size	_Z10applyRulesiiPiS_, .Lfunc_end0-_Z10applyRulesiiPiS_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 860
; NumSgprs: 18
; NumVgprs: 12
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 12
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 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	_Z11add_NS_HaloiiPiS_   ; -- Begin function _Z11add_NS_HaloiiPiS_
	.globl	_Z11add_NS_HaloiiPiS_
	.p2align	8
	.type	_Z11add_NS_HaloiiPiS_,@function
_Z11add_NS_HaloiiPiS_:                  ; @_Z11add_NS_HaloiiPiS_
; %bb.0:
	s_cmp_lt_i32 s15, 1
	s_cbranch_scc1 .LBB1_3
; %bb.1:                                ; %LeafBlock30
	s_cmp_eq_u32 s15, 1
	s_mov_b32 s3, 0
	s_cbranch_scc0 .LBB1_4
; %bb.2:
	s_load_b64 s[4:5], s[0:1], 0x0
	s_waitcnt lgkmcnt(0)
	s_add_i32 s2, s4, -1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[3:4], null, s2, s5, v[0:1]
	v_add3_u32 v1, s5, s4, v0
	s_mov_b32 s2, -1
	v_ashrrev_i32_e32 v2, 31, v1
	s_delay_alu instid0(VALU_DEP_3)
	v_ashrrev_i32_e32 v4, 31, v3
	s_branch .LBB1_5
.LBB1_3:
	s_mov_b32 s3, -1
.LBB1_4:                                ; %Flow
	s_mov_b32 s2, 0
                                        ; implicit-def: $vgpr1_vgpr2
                                        ; implicit-def: $vgpr3_vgpr4
.LBB1_5:                                ; %Flow
	s_and_b32 vcc_lo, exec_lo, s3
	s_cbranch_vccz .LBB1_10
; %bb.6:                                ; %LeafBlock
	s_cmp_eq_u32 s15, 0
	s_cbranch_scc1 .LBB1_8
; %bb.7:
                                        ; implicit-def: $vgpr1_vgpr2
	s_branch .LBB1_9
.LBB1_8:
	v_mov_b32_e32 v1, 0
	s_mov_b32 s2, -1
	s_delay_alu instid0(VALU_DEP_1)
	v_mov_b32_e32 v2, v1
	v_mov_b32_e32 v1, v0
.LBB1_9:                                ; %Flow34
	s_delay_alu instid0(VALU_DEP_1)
	v_dual_mov_b32 v4, v2 :: v_dual_mov_b32 v3, v1
.LBB1_10:                               ; %Flow33
	s_and_b32 vcc_lo, exec_lo, s2
	s_cbranch_vccz .LBB1_12
; %bb.11:                               ; %.sink.split
	s_load_b128 s[0:3], s[0:1], 0x8
	v_lshlrev_b64 v[0:1], 2, v[1:2]
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v0, vcc_lo, s0, v0
	v_add_co_ci_u32_e32 v1, vcc_lo, s1, v1, vcc_lo
	global_load_b32 v2, v[0:1], off
	v_lshlrev_b64 v[0:1], 2, v[3:4]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v0, vcc_lo, s2, v0
	v_add_co_ci_u32_e32 v1, vcc_lo, s3, v1, vcc_lo
	s_waitcnt vmcnt(0)
	global_store_b32 v[0:1], v2, off
.LBB1_12:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z11add_NS_HaloiiPiS_
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 24
		.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 5
		.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	_Z11add_NS_HaloiiPiS_, .Lfunc_end1-_Z11add_NS_HaloiiPiS_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 236
; NumSgprs: 18
; NumVgprs: 5
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 5
; 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	_Z11add_EW_HaloiiPiS_   ; -- Begin function _Z11add_EW_HaloiiPiS_
	.globl	_Z11add_EW_HaloiiPiS_
	.p2align	8
	.type	_Z11add_EW_HaloiiPiS_,@function
_Z11add_EW_HaloiiPiS_:                  ; @_Z11add_EW_HaloiiPiS_
; %bb.0:
	s_load_b64 s[2:3], s[0:1], 0x0
	s_cmp_lt_i32 s15, 1
	s_cbranch_scc1 .LBB2_3
; %bb.1:                                ; %LeafBlock37
	s_cmp_eq_u32 s15, 1
	s_mov_b32 s4, 0
	s_cbranch_scc0 .LBB2_4
; %bb.2:
	s_waitcnt lgkmcnt(0)
	s_lshl_b32 s2, s2, 1
	v_mul_lo_u32 v1, v0, s3
	v_add3_u32 v2, s2, s3, v0
	s_mov_b32 s2, -1
	s_branch .LBB2_5
.LBB2_3:
	s_mov_b32 s4, -1
.LBB2_4:                                ; %Flow
	s_waitcnt lgkmcnt(0)
	s_mov_b32 s2, 0
                                        ; implicit-def: $vgpr1
                                        ; implicit-def: $vgpr2
.LBB2_5:                                ; %Flow
	s_and_b32 vcc_lo, exec_lo, s4
	s_cbranch_vccz .LBB2_9
; %bb.6:                                ; %LeafBlock
	s_cmp_eq_u32 s15, 0
	s_cbranch_scc1 .LBB2_8
; %bb.7:
                                        ; implicit-def: $vgpr1
                                        ; implicit-def: $vgpr2
	s_branch .LBB2_9
.LBB2_8:
	v_mul_lo_u32 v1, v0, s3
	v_add_nc_u32_e32 v2, s3, v0
	s_mov_b32 s2, -1
	s_delay_alu instid0(VALU_DEP_2)
	v_add3_u32 v1, s3, -1, v1
.LBB2_9:                                ; %Flow40
	s_and_b32 vcc_lo, exec_lo, s2
	s_cbranch_vccz .LBB2_11
; %bb.10:                               ; %.sink.split
	s_load_b128 s[0:3], s[0:1], 0x8
	v_ashrrev_i32_e32 v3, 31, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[2:3], 2, v[2:3]
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v2, vcc_lo, s0, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_add_co_ci_u32_e32 v3, vcc_lo, s1, v3, vcc_lo
	global_load_b32 v3, v[2:3], off
	v_ashrrev_i32_e32 v2, 31, v1
	v_lshlrev_b64 v[0:1], 2, v[1:2]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v0, vcc_lo, s2, v0
	v_add_co_ci_u32_e32 v1, vcc_lo, s3, v1, vcc_lo
	s_waitcnt vmcnt(0)
	global_store_b32 v[0:1], v3, off
.LBB2_11:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z11add_EW_HaloiiPiS_
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 24
		.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 4
		.amdhsa_next_free_sgpr 16
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end2:
	.size	_Z11add_EW_HaloiiPiS_, .Lfunc_end2-_Z11add_EW_HaloiiPiS_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 232
; NumSgprs: 18
; NumVgprs: 4
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 4
; 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	_Z11get_NS_HaloiiPiS_   ; -- Begin function _Z11get_NS_HaloiiPiS_
	.globl	_Z11get_NS_HaloiiPiS_
	.p2align	8
	.type	_Z11get_NS_HaloiiPiS_,@function
_Z11get_NS_HaloiiPiS_:                  ; @_Z11get_NS_HaloiiPiS_
; %bb.0:
	s_load_b64 s[2:3], s[0:1], 0x0
	s_cmp_lt_i32 s15, 1
	s_cbranch_scc1 .LBB3_3
; %bb.1:                                ; %LeafBlock31
	s_cmp_eq_u32 s15, 1
	s_mov_b32 s4, 0
	s_cbranch_scc0 .LBB3_4
; %bb.2:
	s_waitcnt lgkmcnt(0)
	v_add3_u32 v4, s2, s3, v0
	s_add_i32 s5, s2, -2
	s_mov_b32 s2, -1
	v_mad_u64_u32 v[2:3], null, s5, s3, v[0:1]
	s_delay_alu instid0(VALU_DEP_2)
	v_ashrrev_i32_e32 v5, 31, v4
	s_branch .LBB3_5
.LBB3_3:
	s_mov_b32 s4, -1
.LBB3_4:                                ; %Flow
	s_waitcnt lgkmcnt(0)
	s_mov_b32 s2, 0
                                        ; implicit-def: $vgpr2
                                        ; implicit-def: $vgpr4_vgpr5
.LBB3_5:                                ; %Flow
	s_and_b32 vcc_lo, exec_lo, s4
	s_cbranch_vccz .LBB3_9
; %bb.6:                                ; %LeafBlock
	s_cmp_eq_u32 s15, 0
	s_cbranch_scc1 .LBB3_8
; %bb.7:
                                        ; implicit-def: $vgpr2
                                        ; implicit-def: $vgpr4_vgpr5
	s_branch .LBB3_9
.LBB3_8:
	v_dual_mov_b32 v1, 0 :: v_dual_add_nc_u32 v2, s3, v0
	s_mov_b32 s2, -1
	s_delay_alu instid0(VALU_DEP_1)
	v_dual_mov_b32 v5, v1 :: v_dual_mov_b32 v4, v0
.LBB3_9:                                ; %Flow34
	s_and_b32 vcc_lo, exec_lo, s2
	s_cbranch_vccz .LBB3_11
; %bb.10:                               ; %.sink.split
	s_load_b128 s[0:3], s[0:1], 0x8
	v_ashrrev_i32_e32 v3, 31, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[0:1], 2, v[2:3]
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v0, vcc_lo, s2, v0
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_add_co_ci_u32_e32 v1, vcc_lo, s3, v1, vcc_lo
	global_load_b32 v2, v[0:1], off
	v_lshlrev_b64 v[0:1], 2, v[4:5]
	v_add_co_u32 v0, vcc_lo, s0, v0
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v1, vcc_lo, s1, v1, vcc_lo
	s_waitcnt vmcnt(0)
	global_store_b32 v[0:1], v2, off
.LBB3_11:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z11get_NS_HaloiiPiS_
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 24
		.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 6
		.amdhsa_next_free_sgpr 16
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end3:
	.size	_Z11get_NS_HaloiiPiS_, .Lfunc_end3-_Z11get_NS_HaloiiPiS_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 232
; NumSgprs: 18
; NumVgprs: 6
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 6
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 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	_Z11get_EW_HaloiiPiS_   ; -- Begin function _Z11get_EW_HaloiiPiS_
	.globl	_Z11get_EW_HaloiiPiS_
	.p2align	8
	.type	_Z11get_EW_HaloiiPiS_,@function
_Z11get_EW_HaloiiPiS_:                  ; @_Z11get_EW_HaloiiPiS_
; %bb.0:
	s_clause 0x1
	s_load_b64 s[4:5], s[0:1], 0x0
	s_load_b128 s[0:3], s[0:1], 0x8
	s_cmp_lt_i32 s15, 1
	s_cbranch_scc1 .LBB4_3
; %bb.1:                                ; %LeafBlock36
	s_cmp_eq_u32 s15, 1
	s_mov_b32 s6, 0
	s_cbranch_scc0 .LBB4_4
; %bb.2:
	s_waitcnt lgkmcnt(0)
	v_mul_lo_u32 v1, v0, s5
	s_lshl_b32 s4, s4, 1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v2, 31, v1
	v_lshlrev_b64 v[1:2], 2, v[1:2]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v3, vcc_lo, s2, v1
	v_add_co_ci_u32_e32 v4, vcc_lo, s3, v2, vcc_lo
	v_add3_u32 v1, s4, s5, v0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_u32 v2, vcc_lo, v3, 4
	v_add_co_ci_u32_e32 v3, vcc_lo, 0, v4, vcc_lo
	s_mov_b32 s4, -1
	s_branch .LBB4_5
.LBB4_3:
	s_mov_b32 s6, -1
.LBB4_4:                                ; %Flow
	s_waitcnt lgkmcnt(0)
	s_mov_b32 s4, 0
                                        ; implicit-def: $vgpr1
                                        ; implicit-def: $vgpr2_vgpr3
.LBB4_5:                                ; %Flow
	s_and_b32 vcc_lo, exec_lo, s6
	s_cbranch_vccz .LBB4_9
; %bb.6:                                ; %LeafBlock
	s_cmp_eq_u32 s15, 0
	s_cbranch_scc1 .LBB4_8
; %bb.7:
                                        ; implicit-def: $vgpr1
                                        ; implicit-def: $vgpr2_vgpr3
	s_branch .LBB4_9
.LBB4_8:
	v_mul_lo_u32 v1, v0, s5
	s_mov_b32 s6, s5
	s_ashr_i32 s4, s5, 31
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_ashrrev_i32_e32 v2, 31, v1
	v_add_co_u32 v1, vcc_lo, s6, v1
	v_add_co_ci_u32_e32 v2, vcc_lo, s4, v2, vcc_lo
	s_mov_b32 s4, -1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[1:2], 2, v[1:2]
	v_add_co_u32 v3, vcc_lo, s2, v1
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_add_co_ci_u32_e32 v4, vcc_lo, s3, v2, vcc_lo
	v_add_nc_u32_e32 v1, s5, v0
	v_add_co_u32 v2, vcc_lo, v3, -8
	s_delay_alu instid0(VALU_DEP_3)
	v_add_co_ci_u32_e32 v3, vcc_lo, -1, v4, vcc_lo
.LBB4_9:                                ; %Flow39
	s_and_b32 vcc_lo, exec_lo, s4
	s_cbranch_vccz .LBB4_11
; %bb.10:                               ; %.sink.split
	global_load_b32 v3, v[2:3], off
	v_ashrrev_i32_e32 v2, 31, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[0:1], 2, v[1:2]
	v_add_co_u32 v0, vcc_lo, s0, v0
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v1, vcc_lo, s1, v1, vcc_lo
	s_waitcnt vmcnt(0)
	global_store_b32 v[0:1], v3, off
.LBB4_11:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z11get_EW_HaloiiPiS_
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 24
		.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 5
		.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_end4:
	.size	_Z11get_EW_HaloiiPiS_, .Lfunc_end4-_Z11get_EW_HaloiiPiS_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 312
; NumSgprs: 18
; NumVgprs: 5
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 5
; 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:
      - .offset:         0
        .size:           4
        .value_kind:     by_value
      - .offset:         4
        .size:           4
        .value_kind:     by_value
      - .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:     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:           _Z10applyRulesiiPiS_
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z10applyRulesiiPiS_.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     12
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .offset:         0
        .size:           4
        .value_kind:     by_value
      - .offset:         4
        .size:           4
        .value_kind:     by_value
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 24
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z11add_NS_HaloiiPiS_
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z11add_NS_HaloiiPiS_.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     5
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .offset:         0
        .size:           4
        .value_kind:     by_value
      - .offset:         4
        .size:           4
        .value_kind:     by_value
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 24
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z11add_EW_HaloiiPiS_
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z11add_EW_HaloiiPiS_.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     4
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .offset:         0
        .size:           4
        .value_kind:     by_value
      - .offset:         4
        .size:           4
        .value_kind:     by_value
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 24
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z11get_NS_HaloiiPiS_
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z11get_NS_HaloiiPiS_.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     6
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .offset:         0
        .size:           4
        .value_kind:     by_value
      - .offset:         4
        .size:           4
        .value_kind:     by_value
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 24
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z11get_EW_HaloiiPiS_
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z11get_EW_HaloiiPiS_.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     5
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
