/*
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 THREADNUM 16
#define THREADS_PER_BLOCK 1024

__global__ void min(int *array, int *answer, int n){
	int index = (threadIdx.x + blockIdx.x * blockDim.x)*2;

	int d, val;
	for (d = n; d >= 1;  d = d/2){
		if (index < d){
			val = array[index];
			if (array[index+1] < val)
				val = array[index+1];
		}
		__syncthreads();
		if (index < d){
			array[index/2] = val;
		}
		__syncthreads();
	}
	*answer = array[0];

/*




	int chunk_size = n/blockDim.x;
	int i, localmin = 10000;
	__shared__ int min[THREADNUM];
	if (threadIdx.x < blockDim.x-1){
		for (i = threadIdx.x * chunk_size; i < threadIdx.x * chunk_size + chunk_size; i++){
			if(array[i] < localmin)
				localmin = array[i];
		}
	}
	else{
		for (i = threadIdx.x * chunk_size; i < n; i++){
			if(array[i] < localmin)
				localmin = array[i];
		}
	}
	min[threadIdx.x] = localmin;
	__syncthreads();
	if(threadIdx.x == 0){
		int globalmin = 10000;
		for(i = 0; i < blockDim.x; i ++)
			if(min[i] < globalmin)
				globalmin = min[i];
		*answer = globalmin;
	}
*/
}

__global__ void last_digit(int *array, int *b){
	b[blockIdx.x] = array[blockIdx.x] % 10;
}

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;
	int answer;
	int *d_answer;    			
	int size = sizeof(array);
	int *b = (int *) malloc(size);
	int *d_b;

	// Allocate space for device copies of array
	cudaMalloc((void **)&d_array, size);
	cudaMalloc((void **)&d_answer, sizeof(int));
	cudaMalloc((void **)&d_b, size);

	cudaMemcpy(d_array, &array, size, cudaMemcpyHostToDevice);

	min<<<(numcomma + THREADS_PER_BLOCK)/THREADS_PER_BLOCK,THREADS_PER_BLOCK>>>(d_array, d_answer, numcomma+1);
	cudaMemcpy(&answer, d_answer, sizeof(int), cudaMemcpyDeviceToHost);

	last_digit<<<(numcomma+1), 1>>>(d_array, d_b);
	cudaMemcpy(b, d_b, size, cudaMemcpyDeviceToHost);

	cudaFree(d_answer); cudaFree(d_array); cudaFree(d_b);

	FILE *q1a = fopen("q1a.txt", "w+");
	fprintf(q1a, "Min: %d\n", answer);
	
	FILE *q1b = fopen("q1b.txt", "w+");
	for (i = 0; i <= numcomma; i ++){
		fprintf(q1b, "%d", b[i]);
		if (i < numcomma) fprintf(q1b, ", ");
	}
	free(b);
}





	.file	"tmpxft_0031c674_00000000-6_q1.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	_Z25__device_stub__Z3minPiS_iPiS_i
	.type	_Z25__device_stub__Z3minPiS_iPiS_i, @function
_Z25__device_stub__Z3minPiS_iPiS_i:
.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
	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	.L2
	pushq	40(%rsp)
	.cfi_def_cfa_offset 152
	leaq	_Z3minPiS_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
.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	_Z25__device_stub__Z3minPiS_iPiS_i, .-_Z25__device_stub__Z3minPiS_iPiS_i
	.globl	_Z3minPiS_i
	.type	_Z3minPiS_i, @function
_Z3minPiS_i:
.LFB2053:
	.cfi_startproc
	endbr64
	jmp	_Z25__device_stub__Z3minPiS_iPiS_i
	.cfi_endproc
.LFE2053:
	.size	_Z3minPiS_i, .-_Z3minPiS_i
	.globl	_Z32__device_stub__Z10last_digitPiS_PiS_
	.type	_Z32__device_stub__Z10last_digitPiS_PiS_, @function
_Z32__device_stub__Z10last_digitPiS_PiS_:
.LFB2054:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movq	%rdi, 8(%rsp)
	leaq	32(%rsp), %rcx
	leaq	24(%rsp), %rdx
	movq	%rsi, (%rsp)
	leaq	40(%rsp), %rdi
	leaq	52(%rsp), %rsi
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rax
	movl	$1, 48(%rsp)
	movq	%rax, 88(%rsp)
	movq	%rsp, %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 40(%rsp)
	movq	%rax, 52(%rsp)
	movl	$1, 60(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L8
	pushq	32(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z10last_digitPiS_(%rip), %rdi
	pushq	32(%rsp)
	.cfi_def_cfa_offset 144
	movq	68(%rsp), %rcx
	movl	76(%rsp), %r8d
	movq	56(%rsp), %rsi
	movl	64(%rsp), %edx
	leaq	104(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L8:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L10
	call	__stack_chk_fail@PLT
.L10:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2054:
	.size	_Z32__device_stub__Z10last_digitPiS_PiS_, .-_Z32__device_stub__Z10last_digitPiS_PiS_
	.globl	_Z10last_digitPiS_
	.type	_Z10last_digitPiS_, @function
_Z10last_digitPiS_:
.LFB2055:
	.cfi_startproc
	endbr64
	jmp	_Z32__device_stub__Z10last_digitPiS_PiS_
	.cfi_endproc
.LFE2055:
	.size	_Z10last_digitPiS_, .-_Z10last_digitPiS_
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"r"
.LC1:
	.string	"inp.txt"
.LC2:
	.string	"%d,"
.LC3:
	.string	"w+"
.LC4:
	.string	"q1a.txt"
.LC5:
	.string	"Min: %d\n"
.LC6:
	.string	"q1b.txt"
.LC7:
	.string	"%d"
.LC8:
	.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
	pushq	%r14
	pushq	%r13
	.cfi_offset 15, -24
	.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
	.cfi_offset 3, -56
	xorl	%ebx, %ebx
	subq	$88, %rsp
	movq	%fs:40, %rax
	movq	%rax, -56(%rbp)
	xorl	%eax, %eax
	call	fopen@PLT
	movq	%rax, %r15
.L14:
	movq	%r15, %rdi
	leal	1(%rbx), %r14d
	call	fgetc@PLT
	cmpb	$-1, %al
	je	.L15
	cmpb	$44, %al
	cmove	%r14d, %ebx
	jmp	.L14
.L15:
	movq	%r15, %rdi
	call	fclose@PLT
	movslq	%r14d, %rax
	movq	%rsp, %rcx
	leaq	15(,%rax,4), %rdx
	movq	%rdx, %rax
	andq	$-4096, %rdx
	andq	$-16, %rax
	subq	%rdx, %rcx
.L18:
	cmpq	%rcx, %rsp
	je	.L19
	subq	$4096, %rsp
	orq	$0, 4088(%rsp)
	jmp	.L18
.L19:
	andl	$4095, %eax
	subq	%rax, %rsp
	testq	%rax, %rax
	je	.L20
	orq	$0, -8(%rsp,%rax)
.L20:
	movq	%r12, %rdi
	movq	%r13, %rsi
	movq	%rsp, -120(%rbp)
	xorl	%r15d, %r15d
	call	fopen@PLT
	movq	-120(%rbp), %r12
	leaq	.LC2(%rip), %r13
	movq	%rax, %rdi
.L21:
	cmpl	%r15d, %ebx
	jl	.L33
	movq	%r12, %rdx
	movq	%r13, %rsi
	xorl	%eax, %eax
	movq	%rdi, -128(%rbp)
	call	__isoc23_fscanf@PLT
	incl	%r15d
	movq	-128(%rbp), %rdi
	addq	$4, %r12
	jmp	.L21
.L33:
	call	fclose@PLT
	leal	0(,%r14,4), %r12d
	movslq	%r12d, %r12
	movq	%r12, %rdi
	call	malloc@PLT
	leaq	-104(%rbp), %rdi
	movq	%r12, %rsi
	movq	%rax, %r13
	call	cudaMalloc@PLT
	leaq	-96(%rbp), %rdi
	movl	$4, %esi
	call	cudaMalloc@PLT
	leaq	-88(%rbp), %rdi
	movq	%r12, %rsi
	call	cudaMalloc@PLT
	movq	-120(%rbp), %rsi
	movq	-104(%rbp), %rdi
	movq	%r12, %rdx
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	movl	$1024, %esi
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	leal	1024(%rbx), %eax
	movl	$4194305, %ecx
	movl	$1, -60(%rbp)
	cltd
	salq	$10, %rcx
	idivl	%esi
	movq	%rcx, -68(%rbp)
	movq	%rcx, %rdx
	movl	$1, %ecx
	movl	%eax, -80(%rbp)
	movabsq	$4294967297, %rax
	movq	%rax, -76(%rbp)
	movq	-80(%rbp), %rdi
	movl	-72(%rbp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L23
	movq	-96(%rbp), %rsi
	movq	-104(%rbp), %rdi
	movl	%r14d, %edx
	call	_Z25__device_stub__Z3minPiS_iPiS_i
.L23:
	movq	-96(%rbp), %rsi
	leaq	-108(%rbp), %rdi
	movl	$2, %ecx
	movl	$4, %edx
	call	cudaMemcpy@PLT
	movl	%r14d, -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	.L24
	movq	-88(%rbp), %rsi
	movq	-104(%rbp), %rdi
	call	_Z32__device_stub__Z10last_digitPiS_PiS_
.L24:
	movq	-88(%rbp), %rsi
	movq	%r12, %rdx
	movl	$2, %ecx
	movq	%r13, %rdi
	leaq	.LC3(%rip), %r12
	leaq	.LC7(%rip), %r15
	call	cudaMemcpy@PLT
	movq	-96(%rbp), %rdi
	call	cudaFree@PLT
	movq	-104(%rbp), %rdi
	call	cudaFree@PLT
	movq	-88(%rbp), %rdi
	call	cudaFree@PLT
	movq	%r12, %rsi
	leaq	.LC4(%rip), %rdi
	call	fopen@PLT
	movl	-108(%rbp), %ecx
	movl	$2, %esi
	leaq	.LC5(%rip), %rdx
	movq	%rax, %rdi
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
	movq	%r12, %rsi
	leaq	.LC6(%rip), %rdi
	xorl	%r12d, %r12d
	call	fopen@PLT
	movq	%rax, %r14
.L25:
	cmpl	%r12d, %ebx
	jl	.L34
	movl	0(%r13,%r12,4), %ecx
	xorl	%eax, %eax
	movq	%r15, %rdx
	movl	$2, %esi
	movq	%r14, %rdi
	call	__fprintf_chk@PLT
	cmpl	%r12d, %ebx
	jle	.L26
	leaq	.LC8(%rip), %rdx
	movl	$2, %esi
	movq	%r14, %rdi
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
.L26:
	incq	%r12
	jmp	.L25
.L34:
	movq	%r13, %rdi
	call	free@PLT
	movq	-56(%rbp), %rax
	subq	%fs:40, %rax
	je	.L28
	call	__stack_chk_fail@PLT
.L28:
	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
.LC9:
	.string	"_Z10last_digitPiS_"
.LC10:
	.string	"_Z3minPiS_i"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2057:
	.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	.LC9(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z10last_digitPiS_(%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	.LC10(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z3minPiS_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
	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
.LFE2057:
	.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 THREADNUM 16
#define THREADS_PER_BLOCK 1024

__global__ void min(int *array, int *answer, int n){
	int index = (threadIdx.x + blockIdx.x * blockDim.x)*2;

	int d, val;
	for (d = n; d >= 1;  d = d/2){
		if (index < d){
			val = array[index];
			if (array[index+1] < val)
				val = array[index+1];
		}
		__syncthreads();
		if (index < d){
			array[index/2] = val;
		}
		__syncthreads();
	}
	*answer = array[0];

/*




	int chunk_size = n/blockDim.x;
	int i, localmin = 10000;
	__shared__ int min[THREADNUM];
	if (threadIdx.x < blockDim.x-1){
		for (i = threadIdx.x * chunk_size; i < threadIdx.x * chunk_size + chunk_size; i++){
			if(array[i] < localmin)
				localmin = array[i];
		}
	}
	else{
		for (i = threadIdx.x * chunk_size; i < n; i++){
			if(array[i] < localmin)
				localmin = array[i];
		}
	}
	min[threadIdx.x] = localmin;
	__syncthreads();
	if(threadIdx.x == 0){
		int globalmin = 10000;
		for(i = 0; i < blockDim.x; i ++)
			if(min[i] < globalmin)
				globalmin = min[i];
		*answer = globalmin;
	}
*/
}

__global__ void last_digit(int *array, int *b){
	b[blockIdx.x] = array[blockIdx.x] % 10;
}

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;
	int answer;
	int *d_answer;    			
	int size = sizeof(array);
	int *b = (int *) malloc(size);
	int *d_b;

	// Allocate space for device copies of array
	hipMalloc((void **)&d_array, size);
	hipMalloc((void **)&d_answer, sizeof(int));
	hipMalloc((void **)&d_b, size);

	hipMemcpy(d_array, &array, size, hipMemcpyHostToDevice);

	min<<<(numcomma + THREADS_PER_BLOCK)/THREADS_PER_BLOCK,THREADS_PER_BLOCK>>>(d_array, d_answer, numcomma+1);
	hipMemcpy(&answer, d_answer, sizeof(int), hipMemcpyDeviceToHost);

	last_digit<<<(numcomma+1), 1>>>(d_array, d_b);
	hipMemcpy(b, d_b, size, hipMemcpyDeviceToHost);

	hipFree(d_answer); hipFree(d_array); hipFree(d_b);

	FILE *q1a = fopen("q1a.txt", "w+");
	fprintf(q1a, "Min: %d\n", answer);
	
	FILE *q1b = fopen("q1b.txt", "w+");
	for (i = 0; i <= numcomma; i ++){
		fprintf(q1b, "%d", b[i]);
		if (i < numcomma) fprintf(q1b, ", ");
	}
	free(b);
}





	.text
	.file	"q1.hip"
	.globl	_Z18__device_stub__minPiS_i     # -- Begin function _Z18__device_stub__minPiS_i
	.type	_Z18__device_stub__minPiS_i,@function
_Z18__device_stub__minPiS_i:            # @_Z18__device_stub__minPiS_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	$_Z3minPiS_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_end0:
	.size	_Z18__device_stub__minPiS_i, .Lfunc_end0-_Z18__device_stub__minPiS_i
	.cfi_endproc
                                        # -- End function
	.globl	_Z25__device_stub__last_digitPiS_ # -- Begin function _Z25__device_stub__last_digitPiS_
	.type	_Z25__device_stub__last_digitPiS_,@function
_Z25__device_stub__last_digitPiS_:      # @_Z25__device_stub__last_digitPiS_
	.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	$80, %rsp
	.cfi_def_cfa_offset 128
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	24(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	16(%rsp), %rcx
	movq	%rsi, (%rcx)
	leaq	64(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%rbx)
	leaq	48(%rsp), %r14
	leaq	32(%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	$_Z10last_digitPiS_, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$96, %rsp
	.cfi_adjust_cfa_offset -96
	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	_Z25__device_stub__last_digitPiS_, .Lfunc_end1-_Z25__device_stub__last_digitPiS_
	.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	$56, %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	%r14d, %r14d
	movl	$1, %r15d
.LBB2_1:                                # =>This Inner Loop Header: Depth=1
	movq	%rbx, %rdi
	callq	fgetc
	cmpb	$-1, %al
	je	.LBB2_4
# %bb.2:                                #   in Loop: Header=BB2_1 Depth=1
	movzbl	%al, %eax
	cmpl	$44, %eax
	jne	.LBB2_1
# %bb.3:                                #   in Loop: Header=BB2_1 Depth=1
	incl	%r14d
	incq	%r15
	jmp	.LBB2_1
.LBB2_4:                                # %.lr.ph.preheader
	movq	%rbx, %rdi
	callq	fclose
	movl	%r14d, %ecx
	leaq	19(,%rcx,4), %rax
	movq	%rcx, -96(%rbp)                 # 8-byte Spill
	incq	%rcx
	movq	%rcx, -88(%rbp)                 # 8-byte Spill
	movabsq	$34359738352, %rcx              # imm = 0x7FFFFFFF0
	andq	%rax, %rcx
	movq	%rsp, %r12
	subq	%rcx, %r12
	movq	%r12, %rsp
	movl	$.L.str, %edi
	movl	$.L.str.1, %esi
	callq	fopen
	movq	%rax, %rbx
	movq	%r15, %r13
	movq	%r12, -56(%rbp)                 # 8-byte Spill
.LBB2_5:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movl	$.L.str.2, %esi
	movq	%rbx, %rdi
	movq	%r12, %rdx
	xorl	%eax, %eax
	callq	__isoc23_fscanf
	addq	$4, %r12
	decq	%r13
	jne	.LBB2_5
# %bb.6:                                # %._crit_edge
	movabsq	$4294967296, %r12               # imm = 0x100000000
	movq	%rbx, %rdi
	callq	fclose
	movq	-88(%rbp), %rax                 # 8-byte Reload
	leal	(,%rax,4), %eax
	movslq	%eax, %r13
	movq	%r13, %rdi
	callq	malloc
	movq	%rax, %rbx
	leaq	-48(%rbp), %rdi
	movq	%r13, %rsi
	callq	hipMalloc
	leaq	-72(%rbp), %rdi
	movl	$4, %esi
	callq	hipMalloc
	leaq	-64(%rbp), %rdi
	movq	%r13, %rsi
	callq	hipMalloc
	leaq	-48(%rbp), %rax
	movq	(%rax), %rdi
	movq	-56(%rbp), %rsi                 # 8-byte Reload
	movq	%r13, -56(%rbp)                 # 8-byte Spill
	movq	%r13, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	addl	$1024, %r14d                    # imm = 0x400
	shrl	$10, %r14d
	btsq	$32, %r14
	leaq	1024(%r12), %rdx
	movq	%r14, %rdi
	movq	-88(%rbp), %r14                 # 8-byte Reload
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB2_8
# %bb.7:
	movq	-48(%rbp), %rdi
	movq	-72(%rbp), %rsi
	movl	%r14d, %edx
	callq	_Z18__device_stub__minPiS_i
.LBB2_8:
	movq	-72(%rbp), %rsi
	leaq	-76(%rbp), %rdi
	movl	$4, %edx
	movl	$2, %ecx
	callq	hipMemcpy
	orq	%r12, %r14
	incq	%r12
	movq	%r14, %rdi
	movl	$1, %esi
	movq	%r12, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	movq	-96(%rbp), %r13                 # 8-byte Reload
	jne	.LBB2_10
# %bb.9:
	movq	-48(%rbp), %rdi
	movq	-64(%rbp), %rsi
	callq	_Z25__device_stub__last_digitPiS_
.LBB2_10:                               # %.lr.ph59.preheader
	movq	-64(%rbp), %rsi
	movq	%rbx, %rdi
	movq	-56(%rbp), %rdx                 # 8-byte Reload
	movl	$2, %ecx
	callq	hipMemcpy
	movq	-72(%rbp), %rdi
	callq	hipFree
	movq	-48(%rbp), %rdi
	callq	hipFree
	movq	-64(%rbp), %rdi
	callq	hipFree
	movl	$.L.str.3, %edi
	movl	$.L.str.4, %esi
	callq	fopen
	movl	-76(%rbp), %edx
	movl	$.L.str.5, %esi
	movq	%rax, %rdi
	xorl	%eax, %eax
	callq	fprintf
	movl	$.L.str.6, %edi
	movl	$.L.str.4, %esi
	callq	fopen
	movq	%rax, %r14
	xorl	%r12d, %r12d
.LBB2_11:                               # %.lr.ph59
                                        # =>This Inner Loop Header: Depth=1
	movl	(%rbx,%r12,4), %edx
	movl	$.L.str.7, %esi
	movq	%r14, %rdi
	xorl	%eax, %eax
	callq	fprintf
	cmpq	%r13, %r12
	jae	.LBB2_13
# %bb.12:                               #   in Loop: Header=BB2_11 Depth=1
	movl	$.L.str.8, %edi
	movl	$2, %esi
	movl	$1, %edx
	movq	%r14, %rcx
	callq	fwrite@PLT
.LBB2_13:                               #   in Loop: Header=BB2_11 Depth=1
	incq	%r12
	decq	%r15
	jne	.LBB2_11
# %bb.14:                               # %._crit_edge60
	movq	%rbx, %rdi
	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_end2:
	.size	main, .Lfunc_end2-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	.LBB3_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB3_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z3minPiS_i, %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	$_Z10last_digitPiS_, %esi
	movl	$.L__unnamed_2, %edx
	movl	$.L__unnamed_2, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end3:
	.size	__hip_module_ctor, .Lfunc_end3-__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	.LBB4_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
.LBB4_2:
	retq
.Lfunc_end4:
	.size	__hip_module_dtor, .Lfunc_end4-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_Z3minPiS_i,@object             # @_Z3minPiS_i
	.section	.rodata,"a",@progbits
	.globl	_Z3minPiS_i
	.p2align	3, 0x0
_Z3minPiS_i:
	.quad	_Z18__device_stub__minPiS_i
	.size	_Z3minPiS_i, 8

	.type	_Z10last_digitPiS_,@object      # @_Z10last_digitPiS_
	.globl	_Z10last_digitPiS_
	.p2align	3, 0x0
_Z10last_digitPiS_:
	.quad	_Z25__device_stub__last_digitPiS_
	.size	_Z10last_digitPiS_, 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	"q1a.txt"
	.size	.L.str.3, 8

	.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	"Min: %d\n"
	.size	.L.str.5, 9

	.type	.L.str.6,@object                # @.str.6
.L.str.6:
	.asciz	"q1b.txt"
	.size	.L.str.6, 8

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

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

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z3minPiS_i"
	.size	.L__unnamed_1, 12

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"_Z10last_digitPiS_"
	.size	.L__unnamed_2, 19

	.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 _Z18__device_stub__minPiS_i
	.addrsig_sym _Z25__device_stub__last_digitPiS_
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z3minPiS_i
	.addrsig_sym _Z10last_digitPiS_
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z3minPiS_i             ; -- Begin function _Z3minPiS_i
	.globl	_Z3minPiS_i
	.p2align	8
	.type	_Z3minPiS_i,@function
_Z3minPiS_i:                            ; @_Z3minPiS_i
; %bb.0:
	s_clause 0x1
	s_load_b32 s2, s[0:1], 0x10
	s_load_b128 s[4:7], s[0:1], 0x0
	s_waitcnt lgkmcnt(0)
	s_cmp_lt_i32 s2, 1
	s_cbranch_scc1 .LBB0_7
; %bb.1:                                ; %.lr.ph
	s_load_b32 s0, s[0:1], 0x24
	s_waitcnt lgkmcnt(0)
	s_and_b32 s0, s0, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[1:2], null, s15, s0, v[0:1]
	v_lshlrev_b32_e32 v0, 1, v1
	v_bfe_i32 v2, v1, 0, 31
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_ashrrev_i32_e32 v1, 31, v0
	v_ashrrev_i32_e32 v3, 31, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_lshlrev_b64 v[4:5], 2, v[0:1]
	v_lshlrev_b64 v[6:7], 2, v[2:3]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_u32 v1, vcc_lo, s4, v4
	v_add_co_ci_u32_e32 v2, vcc_lo, s5, v5, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_co_u32 v3, vcc_lo, s4, v6
	v_add_co_ci_u32_e32 v4, vcc_lo, s5, v7, vcc_lo
                                        ; implicit-def: $vgpr5
.LBB0_2:                                ; =>This Inner Loop Header: Depth=1
	v_cmp_gt_i32_e32 vcc_lo, s2, v0
	s_and_saveexec_b32 s0, vcc_lo
	s_cbranch_execz .LBB0_4
; %bb.3:                                ;   in Loop: Header=BB0_2 Depth=1
	global_load_b64 v[5:6], v[1:2], off
	s_waitcnt vmcnt(0)
	v_min_i32_e32 v5, v6, v5
.LBB0_4:                                ;   in Loop: Header=BB0_2 Depth=1
	s_or_b32 exec_lo, exec_lo, s0
	s_barrier
	buffer_gl0_inv
	s_and_saveexec_b32 s0, vcc_lo
	s_cbranch_execz .LBB0_6
; %bb.5:                                ;   in Loop: Header=BB0_2 Depth=1
	global_store_b32 v[3:4], v5, off
.LBB0_6:                                ;   in Loop: Header=BB0_2 Depth=1
	s_or_b32 exec_lo, exec_lo, s0
	s_lshr_b32 s0, s2, 1
	s_cmp_gt_u32 s2, 1
	s_mov_b32 s2, s0
	s_waitcnt_vscnt null, 0x0
	s_barrier
	buffer_gl0_inv
	s_cbranch_scc1 .LBB0_2
.LBB0_7:                                ; %._crit_edge
	v_mov_b32_e32 v0, 0
	global_load_b32 v1, v0, s[4:5]
	s_waitcnt vmcnt(0)
	global_store_b32 v0, v1, s[6:7]
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z3minPiS_i
		.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 8
		.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	_Z3minPiS_i, .Lfunc_end0-_Z3minPiS_i
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 272
; NumSgprs: 18
; NumVgprs: 8
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 8
; 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	_Z10last_digitPiS_      ; -- Begin function _Z10last_digitPiS_
	.globl	_Z10last_digitPiS_
	.p2align	8
	.type	_Z10last_digitPiS_,@function
_Z10last_digitPiS_:                     ; @_Z10last_digitPiS_
; %bb.0:
	s_load_b128 s[0:3], s[0:1], 0x0
	s_mov_b32 s4, s15
	s_mov_b32 s5, 0
	s_delay_alu instid0(SALU_CYCLE_1)
	s_lshl_b64 s[4:5], s[4:5], 2
	s_waitcnt lgkmcnt(0)
	s_add_u32 s0, s0, s4
	s_addc_u32 s1, s1, s5
	s_load_b32 s0, s[0:1], 0x0
	s_waitcnt lgkmcnt(0)
	s_mul_hi_i32 s1, s0, 0x66666667
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	s_lshr_b32 s6, s1, 31
	s_ashr_i32 s1, s1, 2
	s_add_i32 s1, s1, s6
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_mul_i32 s1, s1, 10
	s_sub_i32 s0, s0, s1
	s_delay_alu instid0(SALU_CYCLE_1)
	v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v1, s0
	s_add_u32 s0, s2, s4
	s_addc_u32 s1, s3, s5
	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 _Z10last_digitPiS_
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 16
		.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	_Z10last_digitPiS_, .Lfunc_end1-_Z10last_digitPiS_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 124
; 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
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

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