/*
This program is written to find the nearest neighbour of each point in 3 deminsional
space by implementing the brute force algorithm.
The brute force approach can easily be converted into a embarassingly parallel algorithm for
the GPU where there is no interaction between the threads.
Benchmarking is done to compare the CPU and GPU computational approaches to the problem.
Also this program tries to use shared memory to measure if there is any performance gain due to
the use of shared memory between threads.
*/

/*
Note that there is a considerable dependency of the ratio of execution times of the CPU and GPU on the 
hardware which is being used to execute the run the program.
*/

// Importing the required headers
#include<stdio.h>
#include<stdlib.h>
#include<math.h>
#include<cuda.h>
#include<time.h>

struct position
{
    int x,y,z;      //odd number of parameters in the structure helps reducing bank conflicts in shared memory(if used).
};

__device__ const int blocksize = 128;

// Returns the duration from start to end times in sec
double time_elapsed(struct timespec *start, struct timespec *end) 
{
	double t;
	t = (end->tv_sec - start->tv_sec); // diff in seconds
	t += (end->tv_nsec - start->tv_nsec) * 0.000000001; //diff in nanoseconds
	return t;
}

// GPU Kernel
__global__ void GPU_Find(struct position *points, int *nearest, int n)
{
    int id = threadIdx.x + blockIdx.x * blocksize;
    if(id >= n) return;

    __shared__ struct position sharedpoints[blocksize];

    int dtc = 1<<22;
    int temp;
    int iofc = -1;
    struct position thispoint;
    thispoint.x = points[id].x;
    thispoint.y = points[id].y;
    thispoint.z = points[id].z;

    for(int j = 0; j < gridDim.x; j++)
    {
        int a = threadIdx.x + j * blocksize;

        if(a < n)
        {
            sharedpoints[threadIdx.x].x = points[a].x;
            sharedpoints[threadIdx.x].y = points[a].y;
            sharedpoints[threadIdx.x].z = points[a].z;
        }
        __syncthreads();

        int *ptr = &sharedpoints[0].x;

        for(int i = 0; i < blocksize; i++)
        {
            temp = (thispoint.x - ptr[0]) * (thispoint.x - ptr[0]);
            temp += (thispoint.y - ptr[1]) * (thispoint.y - ptr[1]);
            temp += (thispoint.z - ptr[2]) * (thispoint.z - ptr[2]);
            ptr += 3;

            {
                int index = i + j*blocksize;
                if(temp < dtc && (index < n) && (index != id))
                {
                    dtc = temp;
                    iofc = index;
                }
            }
        }

        __syncthreads();
    }

    nearest[id] = iofc;

    return;
}

// CPU Function
void CPU_Find(struct position *points, int *nearest, int n)
{
    int min;       //All the distances are going to be smaller than this.
    int temp;

    for(int i = 0; i < n; i++)
    {
        min = 1<<22;
        for(int j = 0; j < n; j++)
        {
            if(i == j) continue;

            temp = (points[i].x - points[j].x)*(points[i].x - points[j].x); 
            temp += (points[i].y - points[j].y)*(points[i].y - points[j].y);
            temp += (points[i].z - points[j].z)*(points[i].z - points[j].z);

            //temp = (int)sqrt(temp);

            if(temp < min)
            {
                min = temp;
                nearest[i] = j;
            }
        }
    }

    return;
}

// Code execution begins here
int main()
{
    struct timespec start1, end1;
    struct timespec start2, end2;

    int n;

    printf("Enter the value of n: ");
    scanf("%d", &n);

    struct position *points;
    int *nearest1;
    int *nearest2;

    cudaMallocManaged(&points, n*sizeof(struct position));
    cudaMallocManaged(&nearest1, n*sizeof(int));
    cudaMallocManaged(&nearest2, n*sizeof(int));

    for(int i = 0; i < n; i++)
    {
        points[i].x = rand()%10000;
        points[i].y = rand()%10000;
        points[i].z = rand()%10000;
        nearest1[i] = -1;
        nearest2[i] = -1;
    }

    clock_gettime(CLOCK_REALTIME, &start1); //start timestamp
	GPU_Find<<<(n/128+1), 128>>>(points, nearest1, n);
	cudaDeviceSynchronize();
    clock_gettime(CLOCK_REALTIME, &end1);	//end timestamp
    
    clock_gettime(CLOCK_REALTIME, &start2); //start timestamp
	CPU_Find(points, nearest2, n);
    clock_gettime(CLOCK_REALTIME, &end2);	//end timestamp
    
    printf("\nTime taken by GPU is: %lf\n", time_elapsed(&start1, &end1));	 //print result for GPU
    printf("Time taken by CPU is: %lf\n", time_elapsed(&start2, &end2));	 //print result for CPU

    cudaFree(points);
    cudaFree(nearest1);
    cudaFree(nearest2);

    return 0;
}

/*
The results obtained by the CPU and GPU may differ. Why so?
*/

	.file	"tmpxft_00312e92_00000000-6_NearestNeighbour2.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2032:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2032:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z12time_elapsedP8timespecS0_
	.type	_Z12time_elapsedP8timespecS0_, @function
_Z12time_elapsedP8timespecS0_:
.LFB2027:
	.cfi_startproc
	endbr64
	movq	(%rsi), %rax
	subq	(%rdi), %rax
	cvtsi2sdq	%rax, %xmm1
	movq	8(%rsi), %rax
	subq	8(%rdi), %rax
	cvtsi2sdq	%rax, %xmm0
	mulsd	.LC0(%rip), %xmm0
	addsd	%xmm1, %xmm0
	ret
	.cfi_endproc
.LFE2027:
	.size	_Z12time_elapsedP8timespecS0_, .-_Z12time_elapsedP8timespecS0_
	.globl	_Z8CPU_FindP8positionPii
	.type	_Z8CPU_FindP8positionPii, @function
_Z8CPU_FindP8positionPii:
.LFB2028:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	movq	%rsi, %r11
	movq	%rdi, %r10
	movl	%edx, %r9d
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	movq	%rdi, %rsi
	xorl	%ecx, %ecx
.L4:
	movl	%ecx, %ebp
	cmpl	%ecx, %r9d
	jle	.L11
	movq	%r10, %r8
	xorl	%edi, %edi
	movl	$4194304, %ebx
.L6:
	cmpl	%ebp, %edi
	je	.L5
	movl	(%rsi), %eax
	subl	(%r8), %eax
	imull	%eax, %eax
	movl	%eax, %edx
	movl	4(%rsi), %eax
	subl	4(%r8), %eax
	imull	%eax, %eax
	addl	%eax, %edx
	movl	8(%rsi), %eax
	subl	8(%r8), %eax
	imull	%eax, %eax
	addl	%edx, %eax
	cmpl	%ebx, %eax
	jge	.L5
	movl	%edi, (%r11,%rcx,4)
	movl	%eax, %ebx
.L5:
	incl	%edi
	addq	$12, %r8
	cmpl	%edi, %r9d
	jne	.L6
	incq	%rcx
	addq	$12, %rsi
	jmp	.L4
.L11:
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2028:
	.size	_Z8CPU_FindP8positionPii, .-_Z8CPU_FindP8positionPii
	.globl	_Z38__device_stub__Z8GPU_FindP8positionPiiP8positionPii
	.type	_Z38__device_stub__Z8GPU_FindP8positionPiiP8positionPii, @function
_Z38__device_stub__Z8GPU_FindP8positionPiiP8positionPii:
.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	.L12
	pushq	40(%rsp)
	.cfi_def_cfa_offset 152
	leaq	_Z8GPU_FindP8positionPii(%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
.L12:
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L14
	call	__stack_chk_fail@PLT
.L14:
	addq	$136, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2054:
	.size	_Z38__device_stub__Z8GPU_FindP8positionPiiP8positionPii, .-_Z38__device_stub__Z8GPU_FindP8positionPiiP8positionPii
	.globl	_Z8GPU_FindP8positionPii
	.type	_Z8GPU_FindP8positionPii, @function
_Z8GPU_FindP8positionPii:
.LFB2055:
	.cfi_startproc
	endbr64
	jmp	_Z38__device_stub__Z8GPU_FindP8positionPiiP8positionPii
	.cfi_endproc
.LFE2055:
	.size	_Z8GPU_FindP8positionPii, .-_Z8GPU_FindP8positionPii
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC1:
	.string	"Enter the value of n: "
.LC2:
	.string	"%d"
.LC3:
	.string	"\nTime taken by GPU is: %lf\n"
.LC4:
	.string	"Time taken by CPU is: %lf\n"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2029:
	.cfi_startproc
	endbr64
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	leaq	.LC1(%rip), %rsi
	movl	$2, %edi
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	movl	$10000, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	xorl	%ebx, %ebx
	subq	$120, %rsp
	.cfi_def_cfa_offset 160
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	call	__printf_chk@PLT
	leaq	12(%rsp), %rsi
	leaq	.LC2(%rip), %rdi
	xorl	%eax, %eax
	call	__isoc23_scanf@PLT
	movslq	12(%rsp), %rsi
	leaq	16(%rsp), %rdi
	movl	$1, %edx
	imulq	$12, %rsi, %rsi
	call	cudaMallocManaged@PLT
	movslq	12(%rsp), %rsi
	leaq	24(%rsp), %rdi
	movl	$1, %edx
	salq	$2, %rsi
	call	cudaMallocManaged@PLT
	movslq	12(%rsp), %rsi
	leaq	32(%rsp), %rdi
	movl	$1, %edx
	salq	$2, %rsi
	call	cudaMallocManaged@PLT
.L18:
	cmpl	%ebx, 12(%rsp)
	jle	.L23
	call	rand@PLT
	imulq	$12, %rbx, %r12
	cltd
	idivl	%ebp
	movq	16(%rsp), %rax
	movl	%edx, (%rax,%r12)
	call	rand@PLT
	cltd
	idivl	%ebp
	movq	16(%rsp), %rax
	movl	%edx, 4(%rax,%r12)
	call	rand@PLT
	cltd
	idivl	%ebp
	movq	16(%rsp), %rax
	movl	%edx, 8(%rax,%r12)
	movq	24(%rsp), %rax
	movl	$-1, (%rax,%rbx,4)
	movq	32(%rsp), %rax
	movl	$-1, (%rax,%rbx,4)
	incq	%rbx
	jmp	.L18
.L23:
	leaq	40(%rsp), %r12
	xorl	%edi, %edi
	movq	%r12, %rsi
	call	clock_gettime@PLT
	movl	12(%rsp), %eax
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$128, %esi
	movl	$33554433, %ecx
	movl	$1, 96(%rsp)
	cltd
	salq	$7, %rcx
	idivl	%esi
	movq	%rcx, 88(%rsp)
	movq	%rcx, %rdx
	movl	$1, %ecx
	incl	%eax
	movl	%eax, 72(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 76(%rsp)
	movq	72(%rsp), %rdi
	movl	80(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L20
	movl	12(%rsp), %edx
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	call	_Z38__device_stub__Z8GPU_FindP8positionPiiP8positionPii
.L20:
	call	cudaDeviceSynchronize@PLT
	leaq	56(%rsp), %r13
	xorl	%edi, %edi
	leaq	72(%rsp), %rbx
	movq	%r13, %rsi
	leaq	88(%rsp), %rbp
	call	clock_gettime@PLT
	xorl	%edi, %edi
	movq	%rbx, %rsi
	call	clock_gettime@PLT
	movl	12(%rsp), %edx
	movq	32(%rsp), %rsi
	movq	16(%rsp), %rdi
	call	_Z8CPU_FindP8positionPii
	xorl	%edi, %edi
	movq	%rbp, %rsi
	call	clock_gettime@PLT
	movq	%r13, %rsi
	movq	%r12, %rdi
	call	_Z12time_elapsedP8timespecS0_
	leaq	.LC3(%rip), %rsi
	movl	$2, %edi
	movb	$1, %al
	call	__printf_chk@PLT
	movq	%rbp, %rsi
	movq	%rbx, %rdi
	call	_Z12time_elapsedP8timespecS0_
	leaq	.LC4(%rip), %rsi
	movl	$2, %edi
	movb	$1, %al
	call	__printf_chk@PLT
	movq	16(%rsp), %rdi
	call	cudaFree@PLT
	movq	24(%rsp), %rdi
	call	cudaFree@PLT
	movq	32(%rsp), %rdi
	call	cudaFree@PLT
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L21
	call	__stack_chk_fail@PLT
.L21:
	addq	$120, %rsp
	.cfi_def_cfa_offset 40
	xorl	%eax, %eax
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%rbp
	.cfi_def_cfa_offset 24
	popq	%r12
	.cfi_def_cfa_offset 16
	popq	%r13
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2029:
	.size	main, .-main
	.section	.rodata.str1.1
.LC5:
	.string	"_Z8GPU_FindP8positionPii"
.LC6:
	.string	"blocksize"
	.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	.LC5(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z8GPU_FindP8positionPii(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC6(%rip), %rdx
	movl	$4, %r9d
	leaq	_ZL9blocksize(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	call	__cudaRegisterVar@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	popq	%rax
	.cfi_def_cfa_offset 24
	popq	%rdx
	.cfi_def_cfa_offset 16
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.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
	.section	.rodata
	.align 4
	.type	_ZL9blocksize, @object
	.size	_ZL9blocksize, 4
_ZL9blocksize:
	.long	128
	.section	.rodata.cst8,"aM",@progbits,8
	.align 8
.LC0:
	.long	-400107883
	.long	1041313291
	.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:

/*
This program is written to find the nearest neighbour of each point in 3 deminsional
space by implementing the brute force algorithm.
The brute force approach can easily be converted into a embarassingly parallel algorithm for
the GPU where there is no interaction between the threads.
Benchmarking is done to compare the CPU and GPU computational approaches to the problem.
Also this program tries to use shared memory to measure if there is any performance gain due to
the use of shared memory between threads.
*/

/*
Note that there is a considerable dependency of the ratio of execution times of the CPU and GPU on the 
hardware which is being used to execute the run the program.
*/

// Importing the required headers
#include<stdio.h>
#include<stdlib.h>
#include<math.h>
#include<hip/hip_runtime.h>
#include<time.h>

struct position
{
    int x,y,z;      //odd number of parameters in the structure helps reducing bank conflicts in shared memory(if used).
};

__device__ const int blocksize = 128;

// Returns the duration from start to end times in sec
double time_elapsed(struct timespec *start, struct timespec *end) 
{
	double t;
	t = (end->tv_sec - start->tv_sec); // diff in seconds
	t += (end->tv_nsec - start->tv_nsec) * 0.000000001; //diff in nanoseconds
	return t;
}

// GPU Kernel
__global__ void GPU_Find(struct position *points, int *nearest, int n)
{
    int id = threadIdx.x + blockIdx.x * blocksize;
    if(id >= n) return;

    __shared__ struct position sharedpoints[blocksize];

    int dtc = 1<<22;
    int temp;
    int iofc = -1;
    struct position thispoint;
    thispoint.x = points[id].x;
    thispoint.y = points[id].y;
    thispoint.z = points[id].z;

    for(int j = 0; j < gridDim.x; j++)
    {
        int a = threadIdx.x + j * blocksize;

        if(a < n)
        {
            sharedpoints[threadIdx.x].x = points[a].x;
            sharedpoints[threadIdx.x].y = points[a].y;
            sharedpoints[threadIdx.x].z = points[a].z;
        }
        __syncthreads();

        int *ptr = &sharedpoints[0].x;

        for(int i = 0; i < blocksize; i++)
        {
            temp = (thispoint.x - ptr[0]) * (thispoint.x - ptr[0]);
            temp += (thispoint.y - ptr[1]) * (thispoint.y - ptr[1]);
            temp += (thispoint.z - ptr[2]) * (thispoint.z - ptr[2]);
            ptr += 3;

            {
                int index = i + j*blocksize;
                if(temp < dtc && (index < n) && (index != id))
                {
                    dtc = temp;
                    iofc = index;
                }
            }
        }

        __syncthreads();
    }

    nearest[id] = iofc;

    return;
}

// CPU Function
void CPU_Find(struct position *points, int *nearest, int n)
{
    int min;       //All the distances are going to be smaller than this.
    int temp;

    for(int i = 0; i < n; i++)
    {
        min = 1<<22;
        for(int j = 0; j < n; j++)
        {
            if(i == j) continue;

            temp = (points[i].x - points[j].x)*(points[i].x - points[j].x); 
            temp += (points[i].y - points[j].y)*(points[i].y - points[j].y);
            temp += (points[i].z - points[j].z)*(points[i].z - points[j].z);

            //temp = (int)sqrt(temp);

            if(temp < min)
            {
                min = temp;
                nearest[i] = j;
            }
        }
    }

    return;
}

// Code execution begins here
int main()
{
    struct timespec start1, end1;
    struct timespec start2, end2;

    int n;

    printf("Enter the value of n: ");
    scanf("%d", &n);

    struct position *points;
    int *nearest1;
    int *nearest2;

    hipMallocManaged(&points, n*sizeof(struct position));
    hipMallocManaged(&nearest1, n*sizeof(int));
    hipMallocManaged(&nearest2, n*sizeof(int));

    for(int i = 0; i < n; i++)
    {
        points[i].x = rand()%10000;
        points[i].y = rand()%10000;
        points[i].z = rand()%10000;
        nearest1[i] = -1;
        nearest2[i] = -1;
    }

    clock_gettime(CLOCK_REALTIME, &start1); //start timestamp
	GPU_Find<<<(n/128+1), 128>>>(points, nearest1, n);
	hipDeviceSynchronize();
    clock_gettime(CLOCK_REALTIME, &end1);	//end timestamp
    
    clock_gettime(CLOCK_REALTIME, &start2); //start timestamp
	CPU_Find(points, nearest2, n);
    clock_gettime(CLOCK_REALTIME, &end2);	//end timestamp
    
    printf("\nTime taken by GPU is: %lf\n", time_elapsed(&start1, &end1));	 //print result for GPU
    printf("Time taken by CPU is: %lf\n", time_elapsed(&start2, &end2));	 //print result for CPU

    hipFree(points);
    hipFree(nearest1);
    hipFree(nearest2);

    return 0;
}

/*
The results obtained by the CPU and GPU may differ. Why so?
*/

	.text
	.file	"NearestNeighbour2.hip"
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function _Z12time_elapsedP8timespecS0_
.LCPI0_0:
	.quad	0x3e112e0be826d695              # double 1.0000000000000001E-9
	.text
	.globl	_Z12time_elapsedP8timespecS0_
	.type	_Z12time_elapsedP8timespecS0_,@function
_Z12time_elapsedP8timespecS0_:          # @_Z12time_elapsedP8timespecS0_
	.cfi_startproc
# %bb.0:
	movq	(%rsi), %rax
	movq	8(%rsi), %rcx
	subq	(%rdi), %rax
	cvtsi2sd	%rax, %xmm1
	subq	8(%rdi), %rcx
	cvtsi2sd	%rcx, %xmm0
	mulsd	.LCPI0_0(%rip), %xmm0
	addsd	%xmm1, %xmm0
	retq
.Lfunc_end0:
	.size	_Z12time_elapsedP8timespecS0_, .Lfunc_end0-_Z12time_elapsedP8timespecS0_
	.cfi_endproc
                                        # -- End function
	.globl	_Z23__device_stub__GPU_FindP8positionPii # -- Begin function _Z23__device_stub__GPU_FindP8positionPii
	.type	_Z23__device_stub__GPU_FindP8positionPii,@function
_Z23__device_stub__GPU_FindP8positionPii: # @_Z23__device_stub__GPU_FindP8positionPii
	.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	$_Z8GPU_FindP8positionPii, %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	_Z23__device_stub__GPU_FindP8positionPii, .Lfunc_end1-_Z23__device_stub__GPU_FindP8positionPii
	.cfi_endproc
                                        # -- End function
	.globl	_Z8CPU_FindP8positionPii        # -- Begin function _Z8CPU_FindP8positionPii
	.type	_Z8CPU_FindP8positionPii,@function
_Z8CPU_FindP8positionPii:               # @_Z8CPU_FindP8positionPii
	.cfi_startproc
# %bb.0:
	testl	%edx, %edx
	jle	.LBB2_9
# %bb.1:                                # %.preheader.lr.ph
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset %rbx, -32
	.cfi_offset %r14, -24
	.cfi_offset %rbp, -16
	movl	%edx, %eax
	leaq	8(%rdi), %rcx
	xorl	%edx, %edx
.LBB2_2:                                # %.lr.ph
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB2_3 Depth 2
	leaq	(%rdx,%rdx,2), %r8
	leaq	(%rdi,%r8,4), %r8
	movl	$4194304, %r11d                 # imm = 0x400000
	movq	%rcx, %r9
	xorl	%r10d, %r10d
.LBB2_3:                                #   Parent Loop BB2_2 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	cmpq	%r10, %rdx
	je	.LBB2_6
# %bb.4:                                #   in Loop: Header=BB2_3 Depth=2
	movl	(%r8), %ebp
	subl	-8(%r9), %ebp
	imull	%ebp, %ebp
	movl	4(%r8), %r14d
	subl	-4(%r9), %r14d
	imull	%r14d, %r14d
	movl	8(%r8), %ebx
	subl	(%r9), %ebx
	addl	%ebp, %r14d
	imull	%ebx, %ebx
	addl	%r14d, %ebx
	cmpl	%r11d, %ebx
	jge	.LBB2_6
# %bb.5:                                #   in Loop: Header=BB2_3 Depth=2
	movl	%r10d, (%rsi,%rdx,4)
	movl	%ebx, %r11d
.LBB2_6:                                #   in Loop: Header=BB2_3 Depth=2
	incq	%r10
	addq	$12, %r9
	cmpq	%r10, %rax
	jne	.LBB2_3
# %bb.7:                                # %._crit_edge
                                        #   in Loop: Header=BB2_2 Depth=1
	incq	%rdx
	cmpq	%rax, %rdx
	jne	.LBB2_2
# %bb.8:
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	.cfi_restore %rbx
	.cfi_restore %r14
	.cfi_restore %rbp
.LBB2_9:                                # %._crit_edge46
	retq
.Lfunc_end2:
	.size	_Z8CPU_FindP8positionPii, .Lfunc_end2-_Z8CPU_FindP8positionPii
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function main
.LCPI3_0:
	.quad	0x3e112e0be826d695              # double 1.0000000000000001E-9
	.text
	.globl	main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	subq	$104, %rsp
	.cfi_def_cfa_offset 144
	.cfi_offset %rbx, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movl	$.L.str, %edi
	xorl	%eax, %eax
	callq	printf
	leaq	12(%rsp), %rbx
	movl	$.L.str.1, %edi
	movq	%rbx, %rsi
	xorl	%eax, %eax
	callq	__isoc23_scanf
	movslq	(%rbx), %rax
	shlq	$2, %rax
	leaq	(%rax,%rax,2), %rsi
	leaq	16(%rsp), %rdi
	movl	$1, %edx
	callq	hipMallocManaged
	movslq	(%rbx), %rsi
	shlq	$2, %rsi
	leaq	32(%rsp), %rdi
	movl	$1, %edx
	callq	hipMallocManaged
	movslq	(%rbx), %rsi
	shlq	$2, %rsi
	leaq	24(%rsp), %rdi
	movl	$1, %edx
	callq	hipMallocManaged
	cmpl	$0, (%rbx)
	jle	.LBB3_3
# %bb.1:                                # %.lr.ph.preheader
	xorl	%ebx, %ebx
	movl	$-1, %ebp
	xorl	%r14d, %r14d
.LBB3_2:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	callq	rand
	cltq
	imulq	$1759218605, %rax, %rcx         # imm = 0x68DB8BAD
	movq	%rcx, %rdx
	shrq	$63, %rdx
	sarq	$44, %rcx
	addl	%edx, %ecx
	imull	$10000, %ecx, %ecx              # imm = 0x2710
	subl	%ecx, %eax
	movq	16(%rsp), %rcx
	movl	%eax, (%rcx,%rbx)
	callq	rand
	cltq
	imulq	$1759218605, %rax, %rcx         # imm = 0x68DB8BAD
	movq	%rcx, %rdx
	shrq	$63, %rdx
	sarq	$44, %rcx
	addl	%edx, %ecx
	imull	$10000, %ecx, %ecx              # imm = 0x2710
	subl	%ecx, %eax
	movq	16(%rsp), %rcx
	movl	%eax, 4(%rcx,%rbx)
	callq	rand
	cltq
	imulq	$1759218605, %rax, %rcx         # imm = 0x68DB8BAD
	movq	%rcx, %rdx
	shrq	$63, %rdx
	sarq	$44, %rcx
	addl	%edx, %ecx
	imull	$10000, %ecx, %ecx              # imm = 0x2710
	subl	%ecx, %eax
	movq	16(%rsp), %rcx
	movl	%eax, 8(%rcx,%rbx)
	movq	32(%rsp), %rax
	movl	%ebp, (%rax,%r14,4)
	movq	24(%rsp), %rax
	movl	%ebp, (%rax,%r14,4)
	incq	%r14
	movslq	12(%rsp), %rax
	addq	$12, %rbx
	cmpq	%rax, %r14
	jl	.LBB3_2
.LBB3_3:                                # %._crit_edge
	leaq	40(%rsp), %rsi
	xorl	%edi, %edi
	callq	clock_gettime
	movl	12(%rsp), %eax
	leal	127(%rax), %edi
	testl	%eax, %eax
	cmovnsl	%eax, %edi
	sarl	$7, %edi
	incl	%edi
	btsq	$32, %rdi
	movabsq	$4294967296, %rdx               # imm = 0x100000000
	orq	$128, %rdx
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB3_5
# %bb.4:
	movq	16(%rsp), %rdi
	movq	32(%rsp), %rsi
	movl	12(%rsp), %edx
	callq	_Z23__device_stub__GPU_FindP8positionPii
.LBB3_5:
	callq	hipDeviceSynchronize
	leaq	88(%rsp), %r14
	xorl	%edi, %edi
	movq	%r14, %rsi
	callq	clock_gettime
	leaq	72(%rsp), %rbx
	xorl	%edi, %edi
	movq	%rbx, %rsi
	callq	clock_gettime
	movq	16(%rsp), %rdi
	movq	24(%rsp), %rsi
	movl	12(%rsp), %edx
	callq	_Z8CPU_FindP8positionPii
	leaq	56(%rsp), %r15
	xorl	%edi, %edi
	movq	%r15, %rsi
	callq	clock_gettime
	movq	(%r14), %rax
	movq	8(%r14), %rcx
	subq	40(%rsp), %rax
	cvtsi2sd	%rax, %xmm1
	subq	48(%rsp), %rcx
	cvtsi2sd	%rcx, %xmm0
	mulsd	.LCPI3_0(%rip), %xmm0
	addsd	%xmm1, %xmm0
	movl	$.L.str.2, %edi
	movb	$1, %al
	callq	printf
	movq	(%r15), %rax
	movq	8(%r15), %rcx
	subq	(%rbx), %rax
	xorps	%xmm1, %xmm1
	cvtsi2sd	%rax, %xmm1
	subq	8(%rbx), %rcx
	xorps	%xmm0, %xmm0
	cvtsi2sd	%rcx, %xmm0
	mulsd	.LCPI3_0(%rip), %xmm0
	addsd	%xmm1, %xmm0
	movl	$.L.str.3, %edi
	movb	$1, %al
	callq	printf
	movq	16(%rsp), %rdi
	callq	hipFree
	movq	32(%rsp), %rdi
	callq	hipFree
	movq	24(%rsp), %rdi
	callq	hipFree
	xorl	%eax, %eax
	addq	$104, %rsp
	.cfi_def_cfa_offset 40
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end3:
	.size	main, .Lfunc_end3-main
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_ctor
	.type	__hip_module_ctor,@function
__hip_module_ctor:                      # @__hip_module_ctor
	.cfi_startproc
# %bb.0:
	subq	$40, %rsp
	.cfi_def_cfa_offset 48
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	jne	.LBB4_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rdi
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB4_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z8GPU_FindP8positionPii, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$40, %rsp
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end4:
	.size	__hip_module_ctor, .Lfunc_end4-__hip_module_ctor
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_dtor
	.type	__hip_module_dtor,@function
__hip_module_dtor:                      # @__hip_module_dtor
	.cfi_startproc
# %bb.0:
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	je	.LBB5_2
# %bb.1:
	pushq	%rax
	.cfi_def_cfa_offset 16
	callq	__hipUnregisterFatBinary
	movq	$0, __hip_gpubin_handle(%rip)
	addq	$8, %rsp
	.cfi_def_cfa_offset 8
.LBB5_2:
	retq
.Lfunc_end5:
	.size	__hip_module_dtor, .Lfunc_end5-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_Z8GPU_FindP8positionPii,@object # @_Z8GPU_FindP8positionPii
	.section	.rodata,"a",@progbits
	.globl	_Z8GPU_FindP8positionPii
	.p2align	3, 0x0
_Z8GPU_FindP8positionPii:
	.quad	_Z23__device_stub__GPU_FindP8positionPii
	.size	_Z8GPU_FindP8positionPii, 8

	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"Enter the value of n: "
	.size	.L.str, 23

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

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	"\nTime taken by GPU is: %lf\n"
	.size	.L.str.2, 28

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	"Time taken by CPU is: %lf\n"
	.size	.L.str.3, 27

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

	.type	__hip_fatbin_wrapper,@object    # @__hip_fatbin_wrapper
	.section	.hipFatBinSegment,"a",@progbits
	.p2align	3, 0x0
__hip_fatbin_wrapper:
	.long	1212764230                      # 0x48495046
	.long	1                               # 0x1
	.quad	__hip_fatbin
	.quad	0
	.size	__hip_fatbin_wrapper, 24

	.type	__hip_gpubin_handle,@object     # @__hip_gpubin_handle
	.local	__hip_gpubin_handle
	.comm	__hip_gpubin_handle,8,8
	.section	.init_array,"aw",@init_array
	.p2align	3, 0x0
	.quad	__hip_module_ctor
	.type	__hip_cuid_,@object             # @__hip_cuid_
	.bss
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               # 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym _Z23__device_stub__GPU_FindP8positionPii
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z8GPU_FindP8positionPii
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z8GPU_FindP8positionPii ; -- Begin function _Z8GPU_FindP8positionPii
	.globl	_Z8GPU_FindP8positionPii
	.p2align	8
	.type	_Z8GPU_FindP8positionPii,@function
_Z8GPU_FindP8positionPii:               ; @_Z8GPU_FindP8positionPii
; %bb.0:
	s_load_b32 s2, s[0:1], 0x10
	v_lshl_add_u32 v4, s15, 7, v0
	s_mov_b32 s3, exec_lo
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_gt_i32_e64 s2, v4
	s_cbranch_execz .LBB0_10
; %bb.1:
	s_clause 0x1
	s_load_b32 s3, s[0:1], 0x18
	s_load_b128 s[4:7], s[0:1], 0x0
	v_ashrrev_i32_e32 v5, 31, v4
	s_mov_b32 s1, 0
	s_waitcnt lgkmcnt(0)
	s_cmp_eq_u32 s3, 0
	s_cbranch_scc1 .LBB0_8
; %bb.2:                                ; %.lr.ph
	v_mad_i64_i32 v[1:2], null, v4, 12, s[4:5]
	v_mul_u32_u24_e32 v7, 12, v0
	v_mad_u32_u24 v8, v0, 12, 8
	v_dual_mov_b32 v9, 0x400000 :: v_dual_mov_b32 v6, -1
	s_mov_b32 s8, 0
	global_load_b96 v[1:3], v[1:2], off
.LBB0_3:                                ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB0_6 Depth 2
	v_lshl_add_u32 v10, s8, 7, v0
	s_mov_b32 s0, exec_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_gt_i32_e64 s2, v10
	s_cbranch_execz .LBB0_5
; %bb.4:                                ;   in Loop: Header=BB0_3 Depth=1
	v_mad_i64_i32 v[11:12], null, v10, 12, s[4:5]
	global_load_b96 v[10:12], v[11:12], off
	s_waitcnt vmcnt(0)
	ds_store_2addr_b32 v7, v10, v11 offset1:1
	ds_store_b32 v8, v12
.LBB0_5:                                ;   in Loop: Header=BB0_3 Depth=1
	s_or_b32 exec_lo, exec_lo, s0
	s_movk_i32 s9, 0xfa00
	s_mov_b32 s10, s1
	s_waitcnt vmcnt(0) lgkmcnt(0)
	s_barrier
	buffer_gl0_inv
.LBB0_6:                                ;   Parent Loop BB0_3 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_add_i32 s0, s9, 0x600
	v_mov_b32_e32 v12, s9
	v_mov_b32_e32 v10, s0
	s_cmp_ge_i32 s10, s2
	v_cmp_eq_u32_e64 s0, s10, v4
	s_cselect_b32 s11, -1, 0
	s_add_i32 s9, s9, 12
	ds_load_2addr_b32 v[10:11], v10 offset1:1
	ds_load_b32 v12, v12 offset:1544
	s_waitcnt lgkmcnt(1)
	v_sub_nc_u32_e32 v10, v1, v10
	v_sub_nc_u32_e32 v11, v2, v11
	s_waitcnt lgkmcnt(0)
	v_sub_nc_u32_e32 v12, v3, v12
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_mul_lo_u32 v10, v10, v10
	v_mul_lo_u32 v11, v11, v11
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v12, v12, v12
	v_add3_u32 v10, v11, v10, v12
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	v_cmp_ge_i32_e32 vcc_lo, v10, v9
	s_or_b32 s11, vcc_lo, s11
	s_or_b32 vcc_lo, s11, s0
	v_cndmask_b32_e32 v9, v10, v9, vcc_lo
	v_cndmask_b32_e32 v6, s10, v6, vcc_lo
	s_add_i32 s10, s10, 1
	s_cmp_eq_u32 s9, 0
	s_cbranch_scc0 .LBB0_6
; %bb.7:                                ;   in Loop: Header=BB0_3 Depth=1
	s_add_i32 s8, s8, 1
	s_addk_i32 s1, 0x80
	s_cmp_eq_u32 s8, s3
	s_barrier
	buffer_gl0_inv
	s_cbranch_scc0 .LBB0_3
	s_branch .LBB0_9
.LBB0_8:
	v_mov_b32_e32 v6, -1
.LBB0_9:                                ; %Flow87
	v_lshlrev_b64 v[0:1], 2, v[4:5]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v0, vcc_lo, s6, v0
	v_add_co_ci_u32_e32 v1, vcc_lo, s7, v1, vcc_lo
	global_store_b32 v[0:1], v6, off
.LBB0_10:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z8GPU_FindP8positionPii
		.amdhsa_group_segment_fixed_size 1536
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 280
		.amdhsa_user_sgpr_count 15
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 13
		.amdhsa_next_free_sgpr 16
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end0:
	.size	_Z8GPU_FindP8positionPii, .Lfunc_end0-_Z8GPU_FindP8positionPii
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 444
; NumSgprs: 18
; NumVgprs: 13
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 1536 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 13
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

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

	.end_amdgpu_metadata
