/*
Please use "inp.txt" as input file and output/write your results of each question to a separate file named as "q1a.txt", "q1b.txt" etc. The output file should have the same format as the input file. 
You only need to submit three source code files, e.g. q1.cu, q2.cu and q3.cu and the input file "inp.txt". Don't submit any other files.
*/
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#define NUMBLOCKS 10

__global__ void convertTo1and0(int *array, int *b, int * idx){
	b[blockIdx.x] = array[blockIdx.x] % 2;
	idx[blockIdx.x] = array[blockIdx.x] % 2;
}

__global__ void PPRead(int * b, int * c, int d){
	if (blockIdx.x >= d) c[blockIdx.x] = b[blockIdx.x - d];
}


__global__ void PPWrite(int * b, int * c, int d){
	if (blockIdx.x >= d) b[blockIdx.x] += c[blockIdx.x];
}

__global__ void findOdds(int * array, int * idx, int * b, int * d){
	if (idx[blockIdx.x]){
		d[b[blockIdx.x] - 1] = array[blockIdx.x];
	}	
}

int main(void) {
	int numcomma = 0;
	char c;
	FILE* stream = fopen("inp.txt", "r");
	while(1){
		c = fgetc(stream);
		if (c == EOF)
			break;
		if (c == ',')
			numcomma ++;
	}
	fclose(stream);

	int array[numcomma+1];

	stream = fopen("inp.txt", "r");
	int i;
	for (i = 0; i <= numcomma; i ++){
		fscanf(stream, "%d,", &array[i]);
	}
	fclose(stream);


	int *d_array; // Holds device copy of array
	int size = sizeof(array);
	int *d_b; // Holds parallel prefix sum of d_idx
	int *d_c;
	int *d_idx; // 1 for odd number in array, 0 for even number in array
	int *d = (int *) malloc(size);
	int *d_d; // Holds final answer

	int *b = (int *) malloc(size);

	cudaMalloc((void **)&d_b, size);
	cudaMalloc((void **)&d_c, size);
	cudaMalloc((void **)&d_array, size);
	cudaMalloc((void **)&d_d, size);
	cudaMalloc((void **)&d_idx, size);
	cudaMemcpy(d_array, &array, size, cudaMemcpyHostToDevice);

	convertTo1and0<<<(numcomma+1), 1>>>(d_array, d_b, d_idx);
	
	for (int d = 1; d <= numcomma; d *= 2){
		PPRead<<<(numcomma+1), 1>>>(d_b, d_c, d);
		PPWrite<<<(numcomma+1), 1>>>(d_b, d_c, d);
	}	

	cudaMemcpy(b, d_b, size, cudaMemcpyDeviceToHost);
	int final_size = b[numcomma];

	findOdds<<<(numcomma+1), 1>>>(d_array, d_idx, d_b, d_d);

	cudaMemcpy(d, d_d, size, cudaMemcpyDeviceToHost);
	cudaFree(d_array); cudaFree(d_d); cudaFree(d_b), cudaFree(d_c); cudaFree(d_idx);

	FILE *q3 = fopen("q3.txt", "w+");
	for (i = 0; i < final_size; i ++){
		fprintf(q3, "%d", d[i]);
		if (i < final_size-1) fprintf(q3, ", ");
	}	

	free(d); free(b);
}





	.file	"tmpxft_0028677b_00000000-6_q3.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	_Z38__device_stub__Z14convertTo1and0PiS_S_PiS_S_
	.type	_Z38__device_stub__Z14convertTo1and0PiS_S_PiS_S_, @function
_Z38__device_stub__Z14convertTo1and0PiS_S_PiS_S_:
.LFB2052:
	.cfi_startproc
	endbr64
	subq	$136, %rsp
	.cfi_def_cfa_offset 144
	movq	%rdi, 24(%rsp)
	leaq	40(%rsp), %rcx
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	32(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 120(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 96(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 112(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L2
	pushq	40(%rsp)
	.cfi_def_cfa_offset 152
	leaq	_Z14convertTo1and0PiS_S_(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 160
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	112(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 152
	popq	%rdx
	.cfi_def_cfa_offset 144
.L2:
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L4
	call	__stack_chk_fail@PLT
.L4:
	addq	$136, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2052:
	.size	_Z38__device_stub__Z14convertTo1and0PiS_S_PiS_S_, .-_Z38__device_stub__Z14convertTo1and0PiS_S_PiS_S_
	.globl	_Z14convertTo1and0PiS_S_
	.type	_Z14convertTo1and0PiS_S_, @function
_Z14convertTo1and0PiS_S_:
.LFB2053:
	.cfi_startproc
	endbr64
	jmp	_Z38__device_stub__Z14convertTo1and0PiS_S_PiS_S_
	.cfi_endproc
.LFE2053:
	.size	_Z14convertTo1and0PiS_S_, .-_Z14convertTo1and0PiS_S_
	.globl	_Z28__device_stub__Z6PPReadPiS_iPiS_i
	.type	_Z28__device_stub__Z6PPReadPiS_iPiS_i, @function
_Z28__device_stub__Z6PPReadPiS_iPiS_i:
.LFB2054:
	.cfi_startproc
	endbr64
	subq	$136, %rsp
	.cfi_def_cfa_offset 144
	movq	%rdi, 24(%rsp)
	leaq	40(%rsp), %rcx
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movl	%edx, 12(%rsp)
	leaq	32(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 120(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 96(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 112(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L8
	pushq	40(%rsp)
	.cfi_def_cfa_offset 152
	leaq	_Z6PPReadPiS_i(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 160
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	112(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 152
	popq	%rdx
	.cfi_def_cfa_offset 144
.L8:
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L10
	call	__stack_chk_fail@PLT
.L10:
	addq	$136, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2054:
	.size	_Z28__device_stub__Z6PPReadPiS_iPiS_i, .-_Z28__device_stub__Z6PPReadPiS_iPiS_i
	.globl	_Z6PPReadPiS_i
	.type	_Z6PPReadPiS_i, @function
_Z6PPReadPiS_i:
.LFB2055:
	.cfi_startproc
	endbr64
	jmp	_Z28__device_stub__Z6PPReadPiS_iPiS_i
	.cfi_endproc
.LFE2055:
	.size	_Z6PPReadPiS_i, .-_Z6PPReadPiS_i
	.globl	_Z29__device_stub__Z7PPWritePiS_iPiS_i
	.type	_Z29__device_stub__Z7PPWritePiS_iPiS_i, @function
_Z29__device_stub__Z7PPWritePiS_iPiS_i:
.LFB2056:
	.cfi_startproc
	endbr64
	subq	$136, %rsp
	.cfi_def_cfa_offset 144
	movq	%rdi, 24(%rsp)
	leaq	40(%rsp), %rcx
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movl	%edx, 12(%rsp)
	leaq	32(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 120(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 96(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 112(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L13
	pushq	40(%rsp)
	.cfi_def_cfa_offset 152
	leaq	_Z7PPWritePiS_i(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 160
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	112(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 152
	popq	%rdx
	.cfi_def_cfa_offset 144
.L13:
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L15
	call	__stack_chk_fail@PLT
.L15:
	addq	$136, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2056:
	.size	_Z29__device_stub__Z7PPWritePiS_iPiS_i, .-_Z29__device_stub__Z7PPWritePiS_iPiS_i
	.globl	_Z7PPWritePiS_i
	.type	_Z7PPWritePiS_i, @function
_Z7PPWritePiS_i:
.LFB2057:
	.cfi_startproc
	endbr64
	jmp	_Z29__device_stub__Z7PPWritePiS_iPiS_i
	.cfi_endproc
.LFE2057:
	.size	_Z7PPWritePiS_i, .-_Z7PPWritePiS_i
	.globl	_Z33__device_stub__Z8findOddsPiS_S_S_PiS_S_S_
	.type	_Z33__device_stub__Z8findOddsPiS_S_S_PiS_S_S_, @function
_Z33__device_stub__Z8findOddsPiS_S_S_PiS_S_S_:
.LFB2058:
	.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
	movq	%rdx, 8(%rsp)
	leaq	40(%rsp), %rdx
	movq	%rcx, (%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	8(%rsp), %rax
	movq	%rax, 120(%rsp)
	movq	%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	_Z8findOddsPiS_S_S_(%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	_Z33__device_stub__Z8findOddsPiS_S_S_PiS_S_S_, .-_Z33__device_stub__Z8findOddsPiS_S_S_PiS_S_S_
	.globl	_Z8findOddsPiS_S_S_
	.type	_Z8findOddsPiS_S_S_, @function
_Z8findOddsPiS_S_S_:
.LFB2059:
	.cfi_startproc
	endbr64
	jmp	_Z33__device_stub__Z8findOddsPiS_S_S_PiS_S_S_
	.cfi_endproc
.LFE2059:
	.size	_Z8findOddsPiS_S_S_, .-_Z8findOddsPiS_S_S_
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"r"
.LC1:
	.string	"inp.txt"
.LC2:
	.string	"%d,"
.LC3:
	.string	"w+"
.LC4:
	.string	"q3.txt"
.LC5:
	.string	"%d"
.LC6:
	.string	", "
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2027:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	movq	%rsp, %rbp
	.cfi_def_cfa_register 6
	pushq	%r15
	.cfi_offset 15, -24
	xorl	%r15d, %r15d
	pushq	%r14
	pushq	%r13
	.cfi_offset 14, -32
	.cfi_offset 13, -40
	leaq	.LC0(%rip), %r13
	pushq	%r12
	.cfi_offset 12, -48
	leaq	.LC1(%rip), %r12
	movq	%r13, %rsi
	pushq	%rbx
	movq	%r12, %rdi
	subq	$120, %rsp
	.cfi_offset 3, -56
	movq	%fs:40, %rax
	movq	%rax, -56(%rbp)
	xorl	%eax, %eax
	call	fopen@PLT
	movq	%rax, %r14
.L24:
	movq	%r14, %rdi
	leal	1(%r15), %ebx
	call	fgetc@PLT
	cmpb	$-1, %al
	je	.L25
	cmpb	$44, %al
	cmove	%ebx, %r15d
	jmp	.L24
.L25:
	movq	%r14, %rdi
	call	fclose@PLT
	movslq	%ebx, %rax
	movq	%rsp, %rcx
	salq	$2, %rax
	leaq	15(%rax), %rdx
	movq	%rax, -152(%rbp)
	movq	%rdx, %rax
	andq	$-4096, %rdx
	andq	$-16, %rax
	subq	%rdx, %rcx
.L28:
	cmpq	%rcx, %rsp
	je	.L29
	subq	$4096, %rsp
	orq	$0, 4088(%rsp)
	jmp	.L28
.L29:
	andl	$4095, %eax
	subq	%rax, %rsp
	testq	%rax, %rax
	je	.L30
	orq	$0, -8(%rsp,%rax)
.L30:
	movq	%r13, %rsi
	movq	%r12, %rdi
	movq	%rsp, -160(%rbp)
	xorl	%r14d, %r14d
	call	fopen@PLT
	movq	-160(%rbp), %r12
	movq	%rax, %r13
.L31:
	cmpl	%r14d, %r15d
	jl	.L47
	movq	%r12, %rdx
	leaq	.LC2(%rip), %rsi
	movq	%r13, %rdi
	xorl	%eax, %eax
	call	__isoc23_fscanf@PLT
	incl	%r14d
	addq	$4, %r12
	jmp	.L31
.L47:
	movq	%r13, %rdi
	call	fclose@PLT
	leal	0(,%rbx,4), %eax
	cltq
	movq	%rax, %rdi
	movq	%rax, -136(%rbp)
	call	malloc@PLT
	movq	-136(%rbp), %rdi
	movq	%rax, %r13
	call	malloc@PLT
	movq	-136(%rbp), %rsi
	leaq	-112(%rbp), %rdi
	movq	%rax, -144(%rbp)
	call	cudaMalloc@PLT
	movq	-136(%rbp), %rsi
	leaq	-104(%rbp), %rdi
	call	cudaMalloc@PLT
	movq	-136(%rbp), %rsi
	leaq	-120(%rbp), %rdi
	call	cudaMalloc@PLT
	movq	-136(%rbp), %rsi
	leaq	-88(%rbp), %rdi
	call	cudaMalloc@PLT
	movq	-136(%rbp), %rsi
	leaq	-96(%rbp), %rdi
	call	cudaMalloc@PLT
	movq	-136(%rbp), %rdx
	movq	-120(%rbp), %rdi
	movl	$1, %ecx
	movq	-160(%rbp), %rsi
	call	cudaMemcpy@PLT
	movl	%ebx, -80(%rbp)
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movabsq	$4294967297, %rdx
	movl	$1, %ecx
	movl	$1, -60(%rbp)
	movq	%rdx, -76(%rbp)
	movq	-80(%rbp), %rdi
	movl	-72(%rbp), %esi
	movq	%rdx, -68(%rbp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L33
	movq	-96(%rbp), %rdx
	movq	-112(%rbp), %rsi
	movq	-120(%rbp), %rdi
	call	_Z38__device_stub__Z14convertTo1and0PiS_S_PiS_S_
.L33:
	movabsq	$4294967297, %r12
	movl	$1, %r14d
.L34:
	cmpl	%r14d, %r15d
	jl	.L48
	movq	%r12, -76(%rbp)
	movl	-72(%rbp), %esi
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	%ebx, -80(%rbp)
	movq	-80(%rbp), %rdi
	movq	%r12, %rdx
	movl	$1, %ecx
	movq	%r12, -68(%rbp)
	movl	$1, -60(%rbp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L35
	movq	-104(%rbp), %rsi
	movq	-112(%rbp), %rdi
	movl	%r14d, %edx
	call	_Z28__device_stub__Z6PPReadPiS_iPiS_i
.L35:
	movq	%r12, -76(%rbp)
	movl	-72(%rbp), %esi
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	%ebx, -80(%rbp)
	movq	-80(%rbp), %rdi
	movq	%r12, %rdx
	movl	$1, %ecx
	movq	%r12, -68(%rbp)
	movl	$1, -60(%rbp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L36
	movq	-104(%rbp), %rsi
	movq	-112(%rbp), %rdi
	movl	%r14d, %edx
	call	_Z29__device_stub__Z7PPWritePiS_iPiS_i
.L36:
	addl	%r14d, %r14d
	jmp	.L34
.L48:
	movq	-136(%rbp), %rdx
	movq	-112(%rbp), %rsi
	movl	$2, %ecx
	movq	-144(%rbp), %rdi
	call	cudaMemcpy@PLT
	movl	%ebx, -80(%rbp)
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movq	-152(%rbp), %rcx
	movq	-144(%rbp), %rax
	movabsq	$4294967297, %rdx
	movl	$1, -60(%rbp)
	movq	%rdx, -76(%rbp)
	movq	-80(%rbp), %rdi
	movl	-72(%rbp), %esi
	movl	-4(%rax,%rcx), %r12d
	movl	$1, %ecx
	movq	%rdx, -68(%rbp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L38
	movq	-88(%rbp), %rcx
	movq	-112(%rbp), %rdx
	movq	-96(%rbp), %rsi
	movq	-120(%rbp), %rdi
	call	_Z33__device_stub__Z8findOddsPiS_S_S_PiS_S_S_
.L38:
	movq	-88(%rbp), %rsi
	movl	$2, %ecx
	movq	%r13, %rdi
	xorl	%r15d, %r15d
	movq	-136(%rbp), %rdx
	leaq	.LC5(%rip), %r14
	call	cudaMemcpy@PLT
	movq	-120(%rbp), %rdi
	call	cudaFree@PLT
	movq	-88(%rbp), %rdi
	call	cudaFree@PLT
	movq	-112(%rbp), %rdi
	call	cudaFree@PLT
	movq	-104(%rbp), %rdi
	call	cudaFree@PLT
	movq	-96(%rbp), %rdi
	call	cudaFree@PLT
	leaq	.LC3(%rip), %rsi
	leaq	.LC4(%rip), %rdi
	call	fopen@PLT
	movq	%rax, %rbx
.L39:
	cmpl	%r15d, %r12d
	jle	.L49
	movl	0(%r13,%r15,4), %ecx
	movq	%r14, %rdx
	movl	$2, %esi
	movq	%rbx, %rdi
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
	leal	-1(%r12), %eax
	cmpl	%r15d, %eax
	jle	.L40
	leaq	.LC6(%rip), %rdx
	movl	$2, %esi
	movq	%rbx, %rdi
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
.L40:
	incq	%r15
	jmp	.L39
.L49:
	movq	%r13, %rdi
	call	free@PLT
	movq	-144(%rbp), %rdi
	call	free@PLT
	movq	-56(%rbp), %rax
	subq	%fs:40, %rax
	je	.L42
	call	__stack_chk_fail@PLT
.L42:
	leaq	-40(%rbp), %rsp
	xorl	%eax, %eax
	popq	%rbx
	popq	%r12
	popq	%r13
	popq	%r14
	popq	%r15
	popq	%rbp
	.cfi_def_cfa 7, 8
	ret
	.cfi_endproc
.LFE2027:
	.size	main, .-main
	.section	.rodata.str1.1
.LC7:
	.string	"_Z8findOddsPiS_S_S_"
.LC8:
	.string	"_Z7PPWritePiS_i"
.LC9:
	.string	"_Z6PPReadPiS_i"
.LC10:
	.string	"_Z14convertTo1and0PiS_S_"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2061:
	.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	.LC7(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z8findOddsPiS_S_S_(%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	.LC8(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z7PPWritePiS_i(%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	.LC9(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z6PPReadPiS_i(%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	.LC10(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z14convertTo1and0PiS_S_(%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
.LFE2061:
	.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:

/*
Please use "inp.txt" as input file and output/write your results of each question to a separate file named as "q1a.txt", "q1b.txt" etc. The output file should have the same format as the input file. 
You only need to submit three source code files, e.g. q1.cu, q2.cu and q3.cu and the input file "inp.txt". Don't submit any other files.
*/

#include <hip/hip_runtime.h>
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#define NUMBLOCKS 10

__global__ void convertTo1and0(int *array, int *b, int * idx){
	b[blockIdx.x] = array[blockIdx.x] % 2;
	idx[blockIdx.x] = array[blockIdx.x] % 2;
}

__global__ void PPRead(int * b, int * c, int d){
	if (blockIdx.x >= d) c[blockIdx.x] = b[blockIdx.x - d];
}


__global__ void PPWrite(int * b, int * c, int d){
	if (blockIdx.x >= d) b[blockIdx.x] += c[blockIdx.x];
}

__global__ void findOdds(int * array, int * idx, int * b, int * d){
	if (idx[blockIdx.x]){
		d[b[blockIdx.x] - 1] = array[blockIdx.x];
	}	
}

int main(void) {
	int numcomma = 0;
	char c;
	FILE* stream = fopen("inp.txt", "r");
	while(1){
		c = fgetc(stream);
		if (c == EOF)
			break;
		if (c == ',')
			numcomma ++;
	}
	fclose(stream);

	int array[numcomma+1];

	stream = fopen("inp.txt", "r");
	int i;
	for (i = 0; i <= numcomma; i ++){
		fscanf(stream, "%d,", &array[i]);
	}
	fclose(stream);


	int *d_array; // Holds device copy of array
	int size = sizeof(array);
	int *d_b; // Holds parallel prefix sum of d_idx
	int *d_c;
	int *d_idx; // 1 for odd number in array, 0 for even number in array
	int *d = (int *) malloc(size);
	int *d_d; // Holds final answer

	int *b = (int *) malloc(size);

	hipMalloc((void **)&d_b, size);
	hipMalloc((void **)&d_c, size);
	hipMalloc((void **)&d_array, size);
	hipMalloc((void **)&d_d, size);
	hipMalloc((void **)&d_idx, size);
	hipMemcpy(d_array, &array, size, hipMemcpyHostToDevice);

	convertTo1and0<<<(numcomma+1), 1>>>(d_array, d_b, d_idx);
	
	for (int d = 1; d <= numcomma; d *= 2){
		PPRead<<<(numcomma+1), 1>>>(d_b, d_c, d);
		PPWrite<<<(numcomma+1), 1>>>(d_b, d_c, d);
	}	

	hipMemcpy(b, d_b, size, hipMemcpyDeviceToHost);
	int final_size = b[numcomma];

	findOdds<<<(numcomma+1), 1>>>(d_array, d_idx, d_b, d_d);

	hipMemcpy(d, d_d, size, hipMemcpyDeviceToHost);
	hipFree(d_array); hipFree(d_d); hipFree(d_b), hipFree(d_c); hipFree(d_idx);

	FILE *q3 = fopen("q3.txt", "w+");
	for (i = 0; i < final_size; i ++){
		fprintf(q3, "%d", d[i]);
		if (i < final_size-1) fprintf(q3, ", ");
	}	

	free(d); free(b);
}





	.text
	.file	"q3.hip"
	.globl	_Z29__device_stub__convertTo1and0PiS_S_ # -- Begin function _Z29__device_stub__convertTo1and0PiS_S_
	.type	_Z29__device_stub__convertTo1and0PiS_S_,@function
_Z29__device_stub__convertTo1and0PiS_S_: # @_Z29__device_stub__convertTo1and0PiS_S_
	.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), %rcx
	movq	%rsi, (%rcx)
	leaq	24(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%rbx)
	movq	%rsi, 16(%rbx)
	leaq	64(%rsp), %r14
	leaq	48(%rsp), %r15
	leaq	16(%rsp), %r12
	leaq	8(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z14convertTo1and0PiS_S_, %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	_Z29__device_stub__convertTo1and0PiS_S_, .Lfunc_end0-_Z29__device_stub__convertTo1and0PiS_S_
	.cfi_endproc
                                        # -- End function
	.globl	_Z21__device_stub__PPReadPiS_i  # -- Begin function _Z21__device_stub__PPReadPiS_i
	.type	_Z21__device_stub__PPReadPiS_i,@function
_Z21__device_stub__PPReadPiS_i:         # @_Z21__device_stub__PPReadPiS_i
	.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), %rcx
	movq	%rsi, (%rcx)
	leaq	12(%rsp), %rsi
	movl	%edx, (%rsi)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%rbx)
	movq	%rsi, 16(%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	$_Z6PPReadPiS_i, %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	_Z21__device_stub__PPReadPiS_i, .Lfunc_end1-_Z21__device_stub__PPReadPiS_i
	.cfi_endproc
                                        # -- End function
	.globl	_Z22__device_stub__PPWritePiS_i # -- Begin function _Z22__device_stub__PPWritePiS_i
	.type	_Z22__device_stub__PPWritePiS_i,@function
_Z22__device_stub__PPWritePiS_i:        # @_Z22__device_stub__PPWritePiS_i
	.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), %rcx
	movq	%rsi, (%rcx)
	leaq	12(%rsp), %rsi
	movl	%edx, (%rsi)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%rbx)
	movq	%rsi, 16(%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	$_Z7PPWritePiS_i, %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	_Z22__device_stub__PPWritePiS_i, .Lfunc_end2-_Z22__device_stub__PPWritePiS_i
	.cfi_endproc
                                        # -- End function
	.globl	_Z23__device_stub__findOddsPiS_S_S_ # -- Begin function _Z23__device_stub__findOddsPiS_S_S_
	.type	_Z23__device_stub__findOddsPiS_S_S_,@function
_Z23__device_stub__findOddsPiS_S_S_:    # @_Z23__device_stub__findOddsPiS_S_S_
	.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	24(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	16(%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	8(%rsp), %r12
	movq	%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	$_Z8findOddsPiS_S_S_, %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	_Z23__device_stub__findOddsPiS_S_S_, .Lfunc_end3-_Z23__device_stub__findOddsPiS_S_S_
	.cfi_endproc
                                        # -- End function
	.globl	main                            # -- Begin function main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset %rbp, -16
	movq	%rsp, %rbp
	.cfi_def_cfa_register %rbp
	pushq	%r15
	pushq	%r14
	pushq	%r13
	pushq	%r12
	pushq	%rbx
	subq	$72, %rsp
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	movl	$.L.str, %edi
	movl	$.L.str.1, %esi
	callq	fopen
	movq	%rax, %rbx
	xorl	%r15d, %r15d
	movl	$1, %r12d
.LBB4_1:                                # =>This Inner Loop Header: Depth=1
	movq	%rbx, %rdi
	callq	fgetc
	cmpb	$-1, %al
	je	.LBB4_4
# %bb.2:                                #   in Loop: Header=BB4_1 Depth=1
	movzbl	%al, %eax
	cmpl	$44, %eax
	jne	.LBB4_1
# %bb.3:                                #   in Loop: Header=BB4_1 Depth=1
	incl	%r15d
	incq	%r12
	jmp	.LBB4_1
.LBB4_4:                                # %.lr.ph.preheader
	movq	%rbx, %rdi
	callq	fclose
	movl	%r15d, %ecx
	leaq	19(,%rcx,4), %rax
	movq	%rcx, -104(%rbp)                # 8-byte Spill
	incq	%rcx
	movq	%rcx, -56(%rbp)                 # 8-byte Spill
	movabsq	$34359738352, %rcx              # imm = 0x7FFFFFFF0
	andq	%rax, %rcx
	movq	%rsp, %r13
	subq	%rcx, %r13
	movq	%r13, %rsp
	movl	$.L.str, %edi
	movl	$.L.str.1, %esi
	callq	fopen
	movq	%rax, %rbx
	movq	%r13, %r14
.LBB4_5:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movl	$.L.str.2, %esi
	movq	%rbx, %rdi
	movq	%r14, %rdx
	xorl	%eax, %eax
	callq	__isoc23_fscanf
	addq	$4, %r14
	decq	%r12
	jne	.LBB4_5
# %bb.6:                                # %._crit_edge
	movabsq	$4294967297, %r12               # imm = 0x100000001
	movq	%rbx, %rdi
	callq	fclose
	movq	-56(%rbp), %rax                 # 8-byte Reload
	leal	(,%rax,4), %eax
	movslq	%eax, %r14
	movq	%r14, %rdi
	callq	malloc
	movq	%rax, %rbx
	movq	%r14, %rdi
	callq	malloc
	movq	%rax, -96(%rbp)                 # 8-byte Spill
	leaq	-48(%rbp), %rdi
	movq	%r14, %rsi
	callq	hipMalloc
	leaq	-88(%rbp), %rdi
	movq	%r14, %rsi
	callq	hipMalloc
	leaq	-64(%rbp), %rdi
	movq	%r14, %rsi
	callq	hipMalloc
	leaq	-72(%rbp), %rdi
	movq	%r14, %rsi
	callq	hipMalloc
	leaq	-80(%rbp), %rdi
	movq	%r14, %rsi
	callq	hipMalloc
	leaq	-64(%rbp), %rax
	movq	(%rax), %rdi
	movq	%r13, %rsi
	movq	%r14, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	-56(%rbp), %rax                 # 8-byte Reload
	leaq	(%rax,%r12), %r13
	decq	%r13
	movq	%r13, %rdi
	movl	$1, %esi
	movq	%r12, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB4_8
# %bb.7:
	movq	-64(%rbp), %rdi
	movq	-48(%rbp), %rsi
	movq	-80(%rbp), %rdx
	callq	_Z29__device_stub__convertTo1and0PiS_S_
.LBB4_8:
	movq	%r14, -56(%rbp)                 # 8-byte Spill
	testl	%r15d, %r15d
	je	.LBB4_15
# %bb.9:                                # %.lr.ph97.preheader
	movl	$1, %r14d
.LBB4_10:                               # %.lr.ph97
                                        # =>This Inner Loop Header: Depth=1
	movq	%r13, %rdi
	movl	$1, %esi
	movq	%r12, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB4_12
# %bb.11:                               #   in Loop: Header=BB4_10 Depth=1
	movq	-48(%rbp), %rdi
	movq	-88(%rbp), %rsi
	movl	%r14d, %edx
	callq	_Z21__device_stub__PPReadPiS_i
.LBB4_12:                               #   in Loop: Header=BB4_10 Depth=1
	movq	%r13, %rdi
	movl	$1, %esi
	movq	%r12, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB4_14
# %bb.13:                               #   in Loop: Header=BB4_10 Depth=1
	movq	-48(%rbp), %rdi
	movq	-88(%rbp), %rsi
	movl	%r14d, %edx
	callq	_Z22__device_stub__PPWritePiS_i
.LBB4_14:                               #   in Loop: Header=BB4_10 Depth=1
	addl	%r14d, %r14d
	cmpl	%r15d, %r14d
	jle	.LBB4_10
.LBB4_15:                               # %._crit_edge98
	movq	-48(%rbp), %rsi
	movq	-96(%rbp), %r14                 # 8-byte Reload
	movq	%r14, %rdi
	movq	-56(%rbp), %rdx                 # 8-byte Reload
	movl	$2, %ecx
	callq	hipMemcpy
	movq	-104(%rbp), %rax                # 8-byte Reload
	movl	(%r14,%rax,4), %r15d
	movq	%r13, %rdi
	movl	$1, %esi
	movq	%r12, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB4_17
# %bb.16:
	movq	-64(%rbp), %rdi
	movq	-80(%rbp), %rsi
	movq	-48(%rbp), %rdx
	movq	-72(%rbp), %rcx
	callq	_Z23__device_stub__findOddsPiS_S_S_
.LBB4_17:
	movq	-72(%rbp), %rsi
	movq	%rbx, %rdi
	movq	-56(%rbp), %rdx                 # 8-byte Reload
	movl	$2, %ecx
	callq	hipMemcpy
	movq	-64(%rbp), %rdi
	callq	hipFree
	movq	-72(%rbp), %rdi
	callq	hipFree
	movq	-48(%rbp), %rdi
	callq	hipFree
	movq	-88(%rbp), %rdi
	callq	hipFree
	movq	-80(%rbp), %rdi
	callq	hipFree
	movl	$.L.str.3, %edi
	movl	$.L.str.4, %esi
	callq	fopen
	testl	%r15d, %r15d
	jle	.LBB4_22
# %bb.18:                               # %.lr.ph101
	movq	%rax, %r14
	leal	-1(%r15), %r12d
	xorl	%r13d, %r13d
.LBB4_19:                               # =>This Inner Loop Header: Depth=1
	movl	(%rbx,%r13,4), %edx
	movl	$.L.str.5, %esi
	movq	%r14, %rdi
	xorl	%eax, %eax
	callq	fprintf
	cmpq	%r12, %r13
	jae	.LBB4_21
# %bb.20:                               #   in Loop: Header=BB4_19 Depth=1
	movl	$.L.str.6, %edi
	movl	$2, %esi
	movl	$1, %edx
	movq	%r14, %rcx
	callq	fwrite@PLT
.LBB4_21:                               #   in Loop: Header=BB4_19 Depth=1
	incq	%r13
	cmpq	%r13, %r15
	jne	.LBB4_19
.LBB4_22:                               # %._crit_edge102
	movq	%rbx, %rdi
	callq	free
	movq	-96(%rbp), %rdi                 # 8-byte Reload
	callq	free
	xorl	%eax, %eax
	leaq	-40(%rbp), %rsp
	popq	%rbx
	popq	%r12
	popq	%r13
	popq	%r14
	popq	%r15
	popq	%rbp
	.cfi_def_cfa %rsp, 8
	retq
.Lfunc_end4:
	.size	main, .Lfunc_end4-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	.LBB5_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB5_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z14convertTo1and0PiS_S_, %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	$_Z6PPReadPiS_i, %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	$_Z7PPWritePiS_i, %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	$_Z8findOddsPiS_S_S_, %esi
	movl	$.L__unnamed_4, %edx
	movl	$.L__unnamed_4, %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_end5:
	.size	__hip_module_ctor, .Lfunc_end5-__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	.LBB6_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
.LBB6_2:
	retq
.Lfunc_end6:
	.size	__hip_module_dtor, .Lfunc_end6-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_Z14convertTo1and0PiS_S_,@object # @_Z14convertTo1and0PiS_S_
	.section	.rodata,"a",@progbits
	.globl	_Z14convertTo1and0PiS_S_
	.p2align	3, 0x0
_Z14convertTo1and0PiS_S_:
	.quad	_Z29__device_stub__convertTo1and0PiS_S_
	.size	_Z14convertTo1and0PiS_S_, 8

	.type	_Z6PPReadPiS_i,@object          # @_Z6PPReadPiS_i
	.globl	_Z6PPReadPiS_i
	.p2align	3, 0x0
_Z6PPReadPiS_i:
	.quad	_Z21__device_stub__PPReadPiS_i
	.size	_Z6PPReadPiS_i, 8

	.type	_Z7PPWritePiS_i,@object         # @_Z7PPWritePiS_i
	.globl	_Z7PPWritePiS_i
	.p2align	3, 0x0
_Z7PPWritePiS_i:
	.quad	_Z22__device_stub__PPWritePiS_i
	.size	_Z7PPWritePiS_i, 8

	.type	_Z8findOddsPiS_S_S_,@object     # @_Z8findOddsPiS_S_S_
	.globl	_Z8findOddsPiS_S_S_
	.p2align	3, 0x0
_Z8findOddsPiS_S_S_:
	.quad	_Z23__device_stub__findOddsPiS_S_S_
	.size	_Z8findOddsPiS_S_S_, 8

	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"inp.txt"
	.size	.L.str, 8

	.type	.L.str.1,@object                # @.str.1
.L.str.1:
	.asciz	"r"
	.size	.L.str.1, 2

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

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	"q3.txt"
	.size	.L.str.3, 7

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

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

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

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z14convertTo1and0PiS_S_"
	.size	.L__unnamed_1, 25

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"_Z6PPReadPiS_i"
	.size	.L__unnamed_2, 15

	.type	.L__unnamed_3,@object           # @2
.L__unnamed_3:
	.asciz	"_Z7PPWritePiS_i"
	.size	.L__unnamed_3, 16

	.type	.L__unnamed_4,@object           # @3
.L__unnamed_4:
	.asciz	"_Z8findOddsPiS_S_S_"
	.size	.L__unnamed_4, 20

	.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 _Z29__device_stub__convertTo1and0PiS_S_
	.addrsig_sym _Z21__device_stub__PPReadPiS_i
	.addrsig_sym _Z22__device_stub__PPWritePiS_i
	.addrsig_sym _Z23__device_stub__findOddsPiS_S_S_
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z14convertTo1and0PiS_S_
	.addrsig_sym _Z6PPReadPiS_i
	.addrsig_sym _Z7PPWritePiS_i
	.addrsig_sym _Z8findOddsPiS_S_S_
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z14convertTo1and0PiS_S_ ; -- Begin function _Z14convertTo1and0PiS_S_
	.globl	_Z14convertTo1and0PiS_S_
	.p2align	8
	.type	_Z14convertTo1and0PiS_S_,@function
_Z14convertTo1and0PiS_S_:               ; @_Z14convertTo1and0PiS_S_
; %bb.0:
	s_load_b128 s[4:7], s[0:1], 0x0
	s_mov_b32 s2, s15
	s_mov_b32 s3, 0
	s_load_b64 s[0:1], s[0:1], 0x10
	s_lshl_b64 s[2:3], s[2:3], 2
	s_waitcnt lgkmcnt(0)
	s_add_u32 s4, s4, s2
	s_addc_u32 s5, s5, s3
	s_load_b32 s8, s[4:5], 0x0
	s_waitcnt lgkmcnt(0)
	s_lshr_b32 s9, s8, 31
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_add_i32 s9, s8, s9
	s_and_b32 s9, s9, -2
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_sub_i32 s8, s8, s9
	v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v1, s8
	s_add_u32 s6, s6, s2
	s_addc_u32 s7, s7, s3
	s_add_u32 s0, s0, s2
	s_addc_u32 s1, s1, s3
	global_store_b32 v0, v1, s[6:7]
	global_load_b32 v1, v0, s[4:5]
	s_waitcnt vmcnt(0)
	v_lshrrev_b32_e32 v2, 31, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v2, v1, v2
	v_and_b32_e32 v2, -2, v2
	s_delay_alu instid0(VALU_DEP_1)
	v_sub_nc_u32_e32 v1, v1, v2
	global_store_b32 v0, v1, s[0:1]
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z14convertTo1and0PiS_S_
		.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 3
		.amdhsa_next_free_sgpr 16
		.amdhsa_reserve_vcc 0
		.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	_Z14convertTo1and0PiS_S_, .Lfunc_end0-_Z14convertTo1and0PiS_S_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 164
; NumSgprs: 16
; NumVgprs: 3
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 1
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 16
; NumVGPRsForWavesPerEU: 3
; 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	_Z6PPReadPiS_i          ; -- Begin function _Z6PPReadPiS_i
	.globl	_Z6PPReadPiS_i
	.p2align	8
	.type	_Z6PPReadPiS_i,@function
_Z6PPReadPiS_i:                         ; @_Z6PPReadPiS_i
; %bb.0:
	s_load_b32 s4, s[0:1], 0x10
	s_waitcnt lgkmcnt(0)
	s_cmp_lt_u32 s15, s4
	s_cbranch_scc1 .LBB1_2
; %bb.1:
	s_load_b128 s[8:11], s[0:1], 0x0
	s_mov_b32 s2, s15
	s_mov_b32 s3, 0
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_4) | instid1(SALU_CYCLE_1)
	s_lshl_b64 s[0:1], s[2:3], 2
	s_waitcnt lgkmcnt(0)
	s_add_u32 s0, s10, s0
	s_addc_u32 s1, s11, s1
	s_sub_i32 s2, s15, s4
	s_lshl_b64 s[2:3], s[2:3], 2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_add_u32 s2, s8, s2
	s_addc_u32 s3, s9, s3
	s_load_b32 s2, s[2:3], 0x0
	s_waitcnt lgkmcnt(0)
	v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v1, s2
	global_store_b32 v0, v1, s[0:1]
.LBB1_2:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z6PPReadPiS_i
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 20
		.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 2
		.amdhsa_next_free_sgpr 16
		.amdhsa_reserve_vcc 0
		.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	_Z6PPReadPiS_i, .Lfunc_end1-_Z6PPReadPiS_i
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 116
; NumSgprs: 16
; NumVgprs: 2
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 1
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 16
; NumVGPRsForWavesPerEU: 2
; 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	_Z7PPWritePiS_i         ; -- Begin function _Z7PPWritePiS_i
	.globl	_Z7PPWritePiS_i
	.p2align	8
	.type	_Z7PPWritePiS_i,@function
_Z7PPWritePiS_i:                        ; @_Z7PPWritePiS_i
; %bb.0:
	s_load_b32 s3, s[0:1], 0x10
	s_waitcnt lgkmcnt(0)
	s_cmp_lt_u32 s15, s3
	s_cbranch_scc1 .LBB2_2
; %bb.1:
	s_load_b128 s[4:7], s[0:1], 0x0
	s_mov_b32 s2, s15
	s_mov_b32 s3, 0
	s_delay_alu instid0(SALU_CYCLE_1)
	s_lshl_b64 s[0:1], s[2:3], 2
	s_waitcnt lgkmcnt(0)
	s_add_u32 s2, s4, s0
	s_addc_u32 s3, s5, s1
	s_add_u32 s0, s6, s0
	s_addc_u32 s1, s7, s1
	s_load_b32 s0, s[0:1], 0x0
	s_load_b32 s1, s[2:3], 0x0
	s_waitcnt lgkmcnt(0)
	s_add_i32 s0, s1, s0
	s_delay_alu instid0(SALU_CYCLE_1)
	v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v1, s0
	global_store_b32 v0, v1, s[2:3]
.LBB2_2:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z7PPWritePiS_i
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 20
		.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 2
		.amdhsa_next_free_sgpr 16
		.amdhsa_reserve_vcc 0
		.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	_Z7PPWritePiS_i, .Lfunc_end2-_Z7PPWritePiS_i
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 120
; NumSgprs: 16
; NumVgprs: 2
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 1
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 16
; NumVGPRsForWavesPerEU: 2
; 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	_Z8findOddsPiS_S_S_     ; -- Begin function _Z8findOddsPiS_S_S_
	.globl	_Z8findOddsPiS_S_S_
	.p2align	8
	.type	_Z8findOddsPiS_S_S_,@function
_Z8findOddsPiS_S_S_:                    ; @_Z8findOddsPiS_S_S_
; %bb.0:
	s_load_b256 s[0:7], s[0:1], 0x0
	s_mov_b32 s8, s15
	s_mov_b32 s9, 0
	s_delay_alu instid0(SALU_CYCLE_1)
	s_lshl_b64 s[8:9], s[8:9], 2
	s_waitcnt lgkmcnt(0)
	s_add_u32 s2, s2, s8
	s_addc_u32 s3, s3, s9
	s_load_b32 s2, s[2:3], 0x0
	s_waitcnt lgkmcnt(0)
	s_cmp_eq_u32 s2, 0
	s_cbranch_scc1 .LBB3_2
; %bb.1:
	s_add_u32 s0, s0, s8
	s_addc_u32 s1, s1, s9
	s_add_u32 s2, s4, s8
	s_addc_u32 s3, s5, s9
	s_load_b32 s2, s[2:3], 0x0
	s_load_b32 s4, s[0:1], 0x0
	v_mov_b32_e32 v0, 0
	s_waitcnt lgkmcnt(0)
	s_ashr_i32 s3, s2, 31
	v_mov_b32_e32 v1, s4
	s_lshl_b64 s[0:1], s[2:3], 2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_add_u32 s0, s6, s0
	s_addc_u32 s1, s7, s1
	global_store_b32 v0, v1, s[0:1] offset:-4
.LBB3_2:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z8findOddsPiS_S_S_
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 32
		.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 2
		.amdhsa_next_free_sgpr 16
		.amdhsa_reserve_vcc 0
		.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	_Z8findOddsPiS_S_S_, .Lfunc_end3-_Z8findOddsPiS_S_S_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 140
; NumSgprs: 16
; NumVgprs: 2
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 1
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 16
; NumVGPRsForWavesPerEU: 2
; Occupancy: 16
; WaveLimiterHint : 1
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_
	.amdgpu_metadata
---
amdhsa.kernels:
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
    .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:           _Z14convertTo1and0PiS_S_
    .private_segment_fixed_size: 0
    .sgpr_count:     16
    .sgpr_spill_count: 0
    .symbol:         _Z14convertTo1and0PiS_S_.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     3
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .offset:         16
        .size:           4
        .value_kind:     by_value
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 20
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z6PPReadPiS_i
    .private_segment_fixed_size: 0
    .sgpr_count:     16
    .sgpr_spill_count: 0
    .symbol:         _Z6PPReadPiS_i.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     2
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .offset:         16
        .size:           4
        .value_kind:     by_value
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 20
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z7PPWritePiS_i
    .private_segment_fixed_size: 0
    .sgpr_count:     16
    .sgpr_spill_count: 0
    .symbol:         _Z7PPWritePiS_i.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     2
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         24
        .size:           8
        .value_kind:     global_buffer
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 32
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z8findOddsPiS_S_S_
    .private_segment_fixed_size: 0
    .sgpr_count:     16
    .sgpr_spill_count: 0
    .symbol:         _Z8findOddsPiS_S_S_.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     2
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
