#include <stdio.h>

/*
Generic swap function in C, I got this off of stack overflow.
*/
#define swap(x,y) do { \
		unsigned char swap_temp[sizeof(x) == sizeof(y) ? (signed)sizeof(x) : -1]; \
		memcpy(swap_temp, &y, sizeof(x)); \
		memcpy(&y, &x, sizeof(x)); \
		memcpy(&x, swap_temp, sizeof(x)); \
	} while(0)


#define N 41
#define MAX 100

#define BLOCKS 4 // Number of blocks/subproblems that we're solving for the overall array.
#define THREADS 10 // Number of subproblems inside a single block that we're solving.


/*
Creates and returns a pointer to an array sorted in decreasing order. If size = N,
then the returned array is A[0..N-1] with A[0] = N, A[1] = N-1, ..., A[N-1] = 1
*/
int* CreateUnsortedArray(int size);

/*
This creates and returns a pointer to a randomly generated array.
*/
int* CreateRandomArray(int size);

/*
Implementation of insertion sort. I copied and pasted this online. The idea is that
each sub-array inside a block will be sorted by insertion sort. Sorts the array's elements
in the range [first, last]
*/
__host__ __device__ void InsertionSort(int* array, int first, int last);

/*
Prints the elements of an array with a description of them attached in the range [first, last]
*/
__host__ __device__ void PrintArray(const char* descp, int* array, int first, int last);


/*
Merges the elements of the two sub-arrays corresponding to indices [leftFirst, leftLast] and
[rightFirst, rightLast], respectively into a single, sorted array [leftFirst, rightLast].
Note that rightFirst = leftLast + 1 for this function to work. The merge code was also
obtained online.
*/
__host__ __device__ void Merge(int *array, int *temp, int leftFirst, int leftLast, int rightFirst, int rightLast);


/*
This function does only the "merging" part of the "MergeSort" function below.
It is intended to be called after each block in the array has been sorted,
so that this would do a pair-wise merge of the blocks into a single array.
*/
__global__ void DoMergeOnly(int *array, int *temp, int size);

/*
Sorts the array using a method that behaves like merge sort.
Note that temp is passed on to avoid dynamically allocating
memory on the GPU.

The idea here is that the array is paritioned into blocks,
each block holds a single sub array. Each block is then
partitioned into threads, where the threads themselves are sub-arrays
of the block.

The thread arrays are first sorted using insertion sort (since they're
the smallest "unit" of the array). Then, each pair of adjacent
threads are merged together until a single, merged array remains;
this array is the current block sorted in ascending order.

After the blocks each contain sorted subarrays. the function "DoMergeOnly"
is called on a single block with the same number of threads as there were
blocks in the preceding step. It merges the sub arrays in the blocks together
and the final array is then sorted in ascending order.

Note: If the number of blocks does not evenly divide the array size,
then the remaining elements are added on to the last block. Same with
the threads for the size of the array in each block.
*/
__global__ void MergeSort(int *array, int *temp, int size);

int main(int argc, char* argv[]) {

	// Create the test array
	int* a = CreateUnsortedArray(N);
	int *dev_a, *dev_temp;
	cudaMalloc((void**)&dev_a, N*sizeof(int));
	cudaMemcpy(dev_a, a, N*sizeof(int), cudaMemcpyHostToDevice);

	// Allocate the temporary array (to avoid doing it on the GPU)
	cudaMalloc((void**)&dev_temp, N*sizeof(int));

	// Print the array's contents
	PrintArray("Before MergeSort:", a, 0, N-1);
	printf("\n");

	// Do the first step of merge sort, where the array contains BLOCKS
	// number of sorted subarrays after this statement
	MergeSort<<<BLOCKS, THREADS>>>(dev_a, dev_temp, N);

	// Output the intermediate step (to verify that the sub arrays are sorted
	cudaMemcpy(a, dev_a, N*sizeof(int), cudaMemcpyDeviceToHost);
	PrintArray("Intermediate:", a, 0, N-1);
	printf("\n");

 	// Now do a pair-wise merge on the subarrays so that the final array is sorted.
	DoMergeOnly<<<1, BLOCKS>>>(dev_a, dev_temp, N);

	// Output the final array
	cudaMemcpy(a, dev_a, N*sizeof(int), cudaMemcpyDeviceToHost);
	PrintArray("After MergeSort:", a, 0, N-1);
	printf("\n");


	// Free the memory
	free(a); cudaFree(dev_a); cudaFree(dev_temp);

	return 0;
}


int* CreateUnsortedArray(int size) {
	int* array =(int*)malloc(size*sizeof(int));
	for (int i = 0; i < size; ++i) {
		array[i] = size-i;
	}

	return array;
}

int* CreateRandomArray(int size) {
	int* array = (int*)malloc(size*sizeof(int));
	for (int i = 0; i < size; ++i) {
		array[i] = rand() % MAX;
	}

	return array;
}



__host__ __device__ void InsertionSort(int* array, int first, int last) {
	for (int i = first; i <= last; ++i) {
		int j = i;
		while (j > first && array[j] < array[j-1]) {
			swap(array[j], array[j-1]);
			--j;
		}
	}
}

__host__ __device__ void PrintArray(const char* descp, int* array, int first, int last) {
	printf("%s\n", descp);
	for (int i = first; i <= last; ++i) {
		printf("%d ", array[i]);
	}
	printf("\n");
}

__host__ __device__ void Merge(int *array, int *temp, int leftFirst, int leftLast, int rightFirst, int rightLast) {
	int i, j, k;
	i = leftFirst;
	k = leftFirst;
	j = rightFirst;

	while (i <= leftLast && j <= rightLast) {
		if (array[i] < array[j]) {
			temp[k++] = array[i++];
		} else {
			temp[k++] = array[j++];
		}
	}

	while (i <= leftLast) {
		temp[k++] = array[i++];
	}

	while (j <= rightLast) {
		temp[k++] = array[j++];
	}

	for (i = leftFirst; i <= rightLast; ++i) {
		array[i] = temp[i];
	}
}

__global__ void MergeSort(int *array, int *temp, int size) {
	int elemPerBlock = size/gridDim.x; // Number of array elements per sub array of the block
	int blockFirst = blockIdx.x*elemPerBlock; // The starting point of the array for this block

	// Ending point of the sub array. Note if we're at the last block, we simply set this
	// To be N - 1, where N = size
	int blockLast = (blockIdx.x == (gridDim.x - 1) ? size - 1 : blockFirst + elemPerBlock - 1); 

	// Number of array elements in the sub array for a single thread
	int elemPerThread = elemPerBlock/blockDim.x;
	int threadFirst = blockFirst + threadIdx.x*elemPerThread; // Same logic as blockFirst, save for threads

	// Same logic as blockLast, save now we're doing it for threads
	int threadLast = (threadIdx.x == (blockDim.x - 1) ? blockLast : threadFirst + elemPerThread - 1);

	InsertionSort(array, threadFirst, threadLast); // Sort the subarrays in each thread by insertion sort
	__syncthreads(); // Wait until all threads are finished

	//Now we merge pair by pair
	int numThreads = (blockDim.x+2-1)/2; // Initial pair is the number of threads per block over 2
	while (numThreads > 1) { // The greater than 1 is because we may not have our threads as a power of 2, so we take the ceiling but ceiling never reaches 0
		if (threadIdx.x < numThreads) { // Is a valid pair that we are considering
			int startId = threadIdx.x*2;
			threadFirst = blockFirst + startId*elemPerThread; // Start of the first pair
			int splitPoint = threadFirst+elemPerThread; // End location of the first pair

			if (threadIdx.x == (numThreads - 1)) { // If it's the last thread, we want the end to be the last
				threadLast = blockLast;
			} else { // Otherwise, we traverse 2*M elements, where M is the number of elements per pair to the end
				threadLast = threadFirst+2*elemPerThread-1;
			}
			Merge(array, temp, threadFirst, splitPoint-1, splitPoint, threadLast);
		}
		__syncthreads(); // Wait until all threads are done
		numThreads = (numThreads+2-1)/2;
		elemPerThread *= 2;
	}


	if (threadIdx.x == 0) { //Finish off the merge. We did not set test condition to 0 above because we always took the ceiling. So we address single thread case here.
		int splitPoint = blockFirst+elemPerThread; // Calculate the split point

		//Last merge
		Merge(array, temp, blockFirst, splitPoint-1, splitPoint, blockLast);
	}
}

// This code is the same as the latter steps in the above. I did not write this
// in a separate function at the time because I wanted to test my code first
// and avoid repeating the first 5-6 lines of computation in the threads.
__global__ void DoMergeOnly(int *array, int *temp, int size) {
	int elemPerBlock = size;
	int blockFirst = 0;
	int blockLast = size-1;

	int elemPerThread = elemPerBlock/blockDim.x;

	//Now we merge pair by pair
	int numThreads = (blockDim.x+2-1)/2;
	while (numThreads > 1) {
		if (threadIdx.x < numThreads) {
			int startId = threadIdx.x*2;
			int threadFirst = blockFirst + startId*elemPerThread;
			int splitPoint = threadFirst+elemPerThread;

			int threadLast;
			if (threadIdx.x == (numThreads - 1)) {
				threadLast = blockLast;
			} else {
				threadLast = threadFirst+2*elemPerThread-1;
			}

			Merge(array, temp, threadFirst, splitPoint-1, splitPoint, threadLast);
		}
		__syncthreads();
		numThreads = (numThreads+2-1)/2;
		elemPerThread *= 2;
	}

	if (threadIdx.x == 0) { //Finish off the merge
		int splitPoint = blockFirst+elemPerThread;
		Merge(array, temp, blockFirst, splitPoint-1, splitPoint, blockLast);
	}
}

	.file	"tmpxft_00383e56_00000000-6_MergeSort.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2035:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2035:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z19CreateUnsortedArrayi
	.type	_Z19CreateUnsortedArrayi, @function
_Z19CreateUnsortedArrayi:
.LFB2028:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	movslq	%edi, %rdi
	movq	%rdi, %rbx
	salq	$2, %rdi
	call	malloc@PLT
	xorl	%edx, %edx
.L3:
	cmpl	%edx, %ebx
	jle	.L7
	movl	%ebx, %ecx
	subl	%edx, %ecx
	movl	%ecx, (%rax,%rdx,4)
	incq	%rdx
	jmp	.L3
.L7:
	popq	%rbx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2028:
	.size	_Z19CreateUnsortedArrayi, .-_Z19CreateUnsortedArrayi
	.globl	_Z17CreateRandomArrayi
	.type	_Z17CreateRandomArrayi, @function
_Z17CreateRandomArrayi:
.LFB2029:
	.cfi_startproc
	endbr64
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	movslq	%edi, %rdi
	movl	$100, %r13d
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	movq	%rdi, %r12
	salq	$2, %rdi
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	xorl	%ebp, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	pushq	%rcx
	.cfi_def_cfa_offset 48
	call	malloc@PLT
	movq	%rax, %rbx
.L9:
	cmpl	%ebp, %r12d
	jle	.L12
	call	rand@PLT
	cltd
	idivl	%r13d
	movl	%edx, (%rbx,%rbp,4)
	incq	%rbp
	jmp	.L9
.L12:
	popq	%rdx
	.cfi_def_cfa_offset 40
	movq	%rbx, %rax
	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	_Z17CreateRandomArrayi, .-_Z17CreateRandomArrayi
	.globl	_Z13InsertionSortPiii
	.type	_Z13InsertionSortPiii, @function
_Z13InsertionSortPiii:
.LFB2030:
	.cfi_startproc
	endbr64
	movslq	%esi, %r8
	movl	%esi, %ecx
.L14:
	cmpl	%edx, %ecx
	jg	.L20
	movq	%r8, %rax
.L18:
	cmpl	%eax, %esi
	jl	.L15
.L17:
	incl	%ecx
	incq	%r8
	jmp	.L14
.L15:
	movl	(%rdi,%rax,4), %r10d
	decq	%rax
	movl	(%rdi,%rax,4), %r9d
	cmpl	%r9d, %r10d
	jge	.L17
	movl	%r10d, (%rdi,%rax,4)
	movl	%r9d, 4(%rdi,%rax,4)
	jmp	.L18
.L20:
	ret
	.cfi_endproc
.LFE2030:
	.size	_Z13InsertionSortPiii, .-_Z13InsertionSortPiii
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"%s\n"
.LC1:
	.string	"%d "
.LC2:
	.string	"\n"
	.text
	.globl	_Z10PrintArrayPKcPiii
	.type	_Z10PrintArrayPKcPiii, @function
_Z10PrintArrayPKcPiii:
.LFB2031:
	.cfi_startproc
	endbr64
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	xorl	%eax, %eax
	leaq	.LC1(%rip), %r13
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	movq	%rsi, %r12
	leaq	.LC0(%rip), %rsi
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	movl	%ecx, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	movslq	%edx, %rbx
	movq	%rdi, %rdx
	movl	$2, %edi
	pushq	%r8
	.cfi_def_cfa_offset 48
	call	__printf_chk@PLT
.L22:
	cmpl	%ebx, %ebp
	jl	.L25
	movl	(%r12,%rbx,4), %edx
	movq	%r13, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	incq	%rbx
	call	__printf_chk@PLT
	jmp	.L22
.L25:
	popq	%rdx
	.cfi_def_cfa_offset 40
	leaq	.LC2(%rip), %rsi
	popq	%rbx
	.cfi_def_cfa_offset 32
	movl	$2, %edi
	popq	%rbp
	.cfi_def_cfa_offset 24
	xorl	%eax, %eax
	popq	%r12
	.cfi_def_cfa_offset 16
	popq	%r13
	.cfi_def_cfa_offset 8
	jmp	__printf_chk@PLT
	.cfi_endproc
.LFE2031:
	.size	_Z10PrintArrayPKcPiii, .-_Z10PrintArrayPKcPiii
	.globl	_Z5MergePiS_iiii
	.type	_Z5MergePiS_iiii, @function
_Z5MergePiS_iiii:
.LFB2032:
	.cfi_startproc
	endbr64
	pushq	%r14
	.cfi_def_cfa_offset 16
	.cfi_offset 14, -16
	movslq	%edx, %r11
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	movq	%r11, %rax
.L27:
	cmpl	%ecx, %edx
	movslq	%edx, %r10
	jg	.L38
	cmpl	%r9d, %r8d
	jg	.L38
	movslq	%r8d, %rbx
	movl	(%rdi,%r10,4), %r10d
	movl	(%rdi,%rbx,4), %ebx
	cmpl	%ebx, %r10d
	jge	.L28
	incl	%edx
	jmp	.L29
.L28:
	incl	%r8d
	movl	%ebx, %r10d
.L29:
	movl	%r10d, (%rsi,%rax,4)
	incq	%rax
	jmp	.L27
.L38:
	movl	%eax, %ebx
	cltq
.L32:
	cmpl	%r10d, %ecx
	jl	.L43
	movl	(%rdi,%r10,4), %r14d
	incq	%r10
	movl	%r14d, (%rsi,%rax,4)
	incq	%rax
	jmp	.L32
.L43:
	incl	%ecx
	movslq	%r8d, %r8
	movl	%ecx, %eax
	subl	%edx, %eax
	cmpl	%edx, %ecx
	movl	$0, %edx
	cmovl	%edx, %eax
	addl	%ebx, %eax
	cltq
.L34:
	cmpl	%r8d, %r9d
	jl	.L36
	movl	(%rdi,%r8,4), %edx
	incq	%r8
	movl	%edx, (%rsi,%rax,4)
	incq	%rax
	jmp	.L34
.L36:
	cmpl	%r11d, %r9d
	jl	.L44
	movl	(%rsi,%r11,4), %eax
	movl	%eax, (%rdi,%r11,4)
	incq	%r11
	jmp	.L36
.L44:
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2032:
	.size	_Z5MergePiS_iiii, .-_Z5MergePiS_iiii
	.globl	_Z31__device_stub__Z9MergeSortPiS_iPiS_i
	.type	_Z31__device_stub__Z9MergeSortPiS_iPiS_i, @function
_Z31__device_stub__Z9MergeSortPiS_iPiS_i:
.LFB2057:
	.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	.L45
	pushq	40(%rsp)
	.cfi_def_cfa_offset 152
	leaq	_Z9MergeSortPiS_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
.L45:
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L47
	call	__stack_chk_fail@PLT
.L47:
	addq	$136, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2057:
	.size	_Z31__device_stub__Z9MergeSortPiS_iPiS_i, .-_Z31__device_stub__Z9MergeSortPiS_iPiS_i
	.globl	_Z9MergeSortPiS_i
	.type	_Z9MergeSortPiS_i, @function
_Z9MergeSortPiS_i:
.LFB2058:
	.cfi_startproc
	endbr64
	jmp	_Z31__device_stub__Z9MergeSortPiS_iPiS_i
	.cfi_endproc
.LFE2058:
	.size	_Z9MergeSortPiS_i, .-_Z9MergeSortPiS_i
	.globl	_Z34__device_stub__Z11DoMergeOnlyPiS_iPiS_i
	.type	_Z34__device_stub__Z11DoMergeOnlyPiS_iPiS_i, @function
_Z34__device_stub__Z11DoMergeOnlyPiS_iPiS_i:
.LFB2059:
	.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	.L50
	pushq	40(%rsp)
	.cfi_def_cfa_offset 152
	leaq	_Z11DoMergeOnlyPiS_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
.L50:
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L52
	call	__stack_chk_fail@PLT
.L52:
	addq	$136, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2059:
	.size	_Z34__device_stub__Z11DoMergeOnlyPiS_iPiS_i, .-_Z34__device_stub__Z11DoMergeOnlyPiS_iPiS_i
	.globl	_Z11DoMergeOnlyPiS_i
	.type	_Z11DoMergeOnlyPiS_i, @function
_Z11DoMergeOnlyPiS_i:
.LFB2060:
	.cfi_startproc
	endbr64
	jmp	_Z34__device_stub__Z11DoMergeOnlyPiS_iPiS_i
	.cfi_endproc
.LFE2060:
	.size	_Z11DoMergeOnlyPiS_i, .-_Z11DoMergeOnlyPiS_i
	.section	.rodata.str1.1
.LC3:
	.string	"Before MergeSort:"
.LC4:
	.string	"Intermediate:"
.LC5:
	.string	"After MergeSort:"
	.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
	movl	$41, %edi
	leaq	.LC2(%rip), %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	subq	$56, %rsp
	.cfi_def_cfa_offset 80
	movq	%fs:40, %rax
	movq	%rax, 40(%rsp)
	xorl	%eax, %eax
	call	_Z19CreateUnsortedArrayi
	movq	%rsp, %rdi
	movl	$164, %esi
	movq	%rax, %rbx
	call	cudaMalloc@PLT
	movq	(%rsp), %rdi
	movl	$1, %ecx
	movq	%rbx, %rsi
	movl	$164, %edx
	call	cudaMemcpy@PLT
	leaq	8(%rsp), %rdi
	movl	$164, %esi
	call	cudaMalloc@PLT
	xorl	%edx, %edx
	movl	$40, %ecx
	movq	%rbx, %rsi
	leaq	.LC3(%rip), %rdi
	call	_Z10PrintArrayPKcPiii
	movq	%rbp, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	$2147483653, %edx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$1073741825, %edi
	addq	%rdx, %rdx
	movl	$1, %ecx
	movl	$1, %esi
	salq	$2, %rdi
	movq	%rdx, 28(%rsp)
	movl	$1, 36(%rsp)
	movq	%rdi, 16(%rsp)
	movl	$1, 24(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L56
	movq	8(%rsp), %rsi
	movq	(%rsp), %rdi
	movl	$41, %edx
	call	_Z31__device_stub__Z9MergeSortPiS_iPiS_i
.L56:
	movq	(%rsp), %rsi
	movl	$2, %ecx
	movl	$164, %edx
	movq	%rbx, %rdi
	call	cudaMemcpy@PLT
	xorl	%edx, %edx
	movl	$40, %ecx
	movq	%rbx, %rsi
	leaq	.LC4(%rip), %rdi
	call	_Z10PrintArrayPKcPiii
	movq	%rbp, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	$1073741825, %edx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	salq	$2, %rdx
	movl	$1, %ecx
	movabsq	$4294967297, %rdi
	movl	$1, %esi
	movq	%rdx, 28(%rsp)
	movl	$1, 36(%rsp)
	movq	%rdi, 16(%rsp)
	movl	$1, 24(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L57
	movq	8(%rsp), %rsi
	movq	(%rsp), %rdi
	movl	$41, %edx
	call	_Z34__device_stub__Z11DoMergeOnlyPiS_iPiS_i
.L57:
	movq	(%rsp), %rsi
	movl	$2, %ecx
	movl	$164, %edx
	movq	%rbx, %rdi
	call	cudaMemcpy@PLT
	xorl	%edx, %edx
	movl	$40, %ecx
	movq	%rbx, %rsi
	leaq	.LC5(%rip), %rdi
	call	_Z10PrintArrayPKcPiii
	movq	%rbp, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	%rbx, %rdi
	call	free@PLT
	movq	(%rsp), %rdi
	call	cudaFree@PLT
	movq	8(%rsp), %rdi
	call	cudaFree@PLT
	movq	40(%rsp), %rax
	subq	%fs:40, %rax
	je	.L58
	call	__stack_chk_fail@PLT
.L58:
	addq	$56, %rsp
	.cfi_def_cfa_offset 24
	xorl	%eax, %eax
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2027:
	.size	main, .-main
	.section	.rodata.str1.1
.LC6:
	.string	"_Z11DoMergeOnlyPiS_i"
.LC7:
	.string	"_Z9MergeSortPiS_i"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2062:
	.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	.LC6(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z11DoMergeOnlyPiS_i(%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	.LC7(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z9MergeSortPiS_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
.LFE2062:
	.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:


#include <hip/hip_runtime.h>
#include <stdio.h>

/*
Generic swap function in C, I got this off of stack overflow.
*/
#define swap(x,y) do { \
		unsigned char swap_temp[sizeof(x) == sizeof(y) ? (signed)sizeof(x) : -1]; \
		memcpy(swap_temp, &y, sizeof(x)); \
		memcpy(&y, &x, sizeof(x)); \
		memcpy(&x, swap_temp, sizeof(x)); \
	} while(0)


#define N 41
#define MAX 100

#define BLOCKS 4 // Number of blocks/subproblems that we're solving for the overall array.
#define THREADS 10 // Number of subproblems inside a single block that we're solving.


/*
Creates and returns a pointer to an array sorted in decreasing order. If size = N,
then the returned array is A[0..N-1] with A[0] = N, A[1] = N-1, ..., A[N-1] = 1
*/
int* CreateUnsortedArray(int size);

/*
This creates and returns a pointer to a randomly generated array.
*/
int* CreateRandomArray(int size);

/*
Implementation of insertion sort. I copied and pasted this online. The idea is that
each sub-array inside a block will be sorted by insertion sort. Sorts the array's elements
in the range [first, last]
*/
__host__ __device__ void InsertionSort(int* array, int first, int last);

/*
Prints the elements of an array with a description of them attached in the range [first, last]
*/
__host__ __device__ void PrintArray(const char* descp, int* array, int first, int last);


/*
Merges the elements of the two sub-arrays corresponding to indices [leftFirst, leftLast] and
[rightFirst, rightLast], respectively into a single, sorted array [leftFirst, rightLast].
Note that rightFirst = leftLast + 1 for this function to work. The merge code was also
obtained online.
*/
__host__ __device__ void Merge(int *array, int *temp, int leftFirst, int leftLast, int rightFirst, int rightLast);


/*
This function does only the "merging" part of the "MergeSort" function below.
It is intended to be called after each block in the array has been sorted,
so that this would do a pair-wise merge of the blocks into a single array.
*/
__global__ void DoMergeOnly(int *array, int *temp, int size);

/*
Sorts the array using a method that behaves like merge sort.
Note that temp is passed on to avoid dynamically allocating
memory on the GPU.

The idea here is that the array is paritioned into blocks,
each block holds a single sub array. Each block is then
partitioned into threads, where the threads themselves are sub-arrays
of the block.

The thread arrays are first sorted using insertion sort (since they're
the smallest "unit" of the array). Then, each pair of adjacent
threads are merged together until a single, merged array remains;
this array is the current block sorted in ascending order.

After the blocks each contain sorted subarrays. the function "DoMergeOnly"
is called on a single block with the same number of threads as there were
blocks in the preceding step. It merges the sub arrays in the blocks together
and the final array is then sorted in ascending order.

Note: If the number of blocks does not evenly divide the array size,
then the remaining elements are added on to the last block. Same with
the threads for the size of the array in each block.
*/
__global__ void MergeSort(int *array, int *temp, int size);

int main(int argc, char* argv[]) {

	// Create the test array
	int* a = CreateUnsortedArray(N);
	int *dev_a, *dev_temp;
	hipMalloc((void**)&dev_a, N*sizeof(int));
	hipMemcpy(dev_a, a, N*sizeof(int), hipMemcpyHostToDevice);

	// Allocate the temporary array (to avoid doing it on the GPU)
	hipMalloc((void**)&dev_temp, N*sizeof(int));

	// Print the array's contents
	PrintArray("Before MergeSort:", a, 0, N-1);
	printf("\n");

	// Do the first step of merge sort, where the array contains BLOCKS
	// number of sorted subarrays after this statement
	MergeSort<<<BLOCKS, THREADS>>>(dev_a, dev_temp, N);

	// Output the intermediate step (to verify that the sub arrays are sorted
	hipMemcpy(a, dev_a, N*sizeof(int), hipMemcpyDeviceToHost);
	PrintArray("Intermediate:", a, 0, N-1);
	printf("\n");

 	// Now do a pair-wise merge on the subarrays so that the final array is sorted.
	DoMergeOnly<<<1, BLOCKS>>>(dev_a, dev_temp, N);

	// Output the final array
	hipMemcpy(a, dev_a, N*sizeof(int), hipMemcpyDeviceToHost);
	PrintArray("After MergeSort:", a, 0, N-1);
	printf("\n");


	// Free the memory
	free(a); hipFree(dev_a); hipFree(dev_temp);

	return 0;
}


int* CreateUnsortedArray(int size) {
	int* array =(int*)malloc(size*sizeof(int));
	for (int i = 0; i < size; ++i) {
		array[i] = size-i;
	}

	return array;
}

int* CreateRandomArray(int size) {
	int* array = (int*)malloc(size*sizeof(int));
	for (int i = 0; i < size; ++i) {
		array[i] = rand() % MAX;
	}

	return array;
}



__host__ __device__ void InsertionSort(int* array, int first, int last) {
	for (int i = first; i <= last; ++i) {
		int j = i;
		while (j > first && array[j] < array[j-1]) {
			swap(array[j], array[j-1]);
			--j;
		}
	}
}

__host__ __device__ void PrintArray(const char* descp, int* array, int first, int last) {
	printf("%s\n", descp);
	for (int i = first; i <= last; ++i) {
		printf("%d ", array[i]);
	}
	printf("\n");
}

__host__ __device__ void Merge(int *array, int *temp, int leftFirst, int leftLast, int rightFirst, int rightLast) {
	int i, j, k;
	i = leftFirst;
	k = leftFirst;
	j = rightFirst;

	while (i <= leftLast && j <= rightLast) {
		if (array[i] < array[j]) {
			temp[k++] = array[i++];
		} else {
			temp[k++] = array[j++];
		}
	}

	while (i <= leftLast) {
		temp[k++] = array[i++];
	}

	while (j <= rightLast) {
		temp[k++] = array[j++];
	}

	for (i = leftFirst; i <= rightLast; ++i) {
		array[i] = temp[i];
	}
}

__global__ void MergeSort(int *array, int *temp, int size) {
	int elemPerBlock = size/gridDim.x; // Number of array elements per sub array of the block
	int blockFirst = blockIdx.x*elemPerBlock; // The starting point of the array for this block

	// Ending point of the sub array. Note if we're at the last block, we simply set this
	// To be N - 1, where N = size
	int blockLast = (blockIdx.x == (gridDim.x - 1) ? size - 1 : blockFirst + elemPerBlock - 1); 

	// Number of array elements in the sub array for a single thread
	int elemPerThread = elemPerBlock/blockDim.x;
	int threadFirst = blockFirst + threadIdx.x*elemPerThread; // Same logic as blockFirst, save for threads

	// Same logic as blockLast, save now we're doing it for threads
	int threadLast = (threadIdx.x == (blockDim.x - 1) ? blockLast : threadFirst + elemPerThread - 1);

	InsertionSort(array, threadFirst, threadLast); // Sort the subarrays in each thread by insertion sort
	__syncthreads(); // Wait until all threads are finished

	//Now we merge pair by pair
	int numThreads = (blockDim.x+2-1)/2; // Initial pair is the number of threads per block over 2
	while (numThreads > 1) { // The greater than 1 is because we may not have our threads as a power of 2, so we take the ceiling but ceiling never reaches 0
		if (threadIdx.x < numThreads) { // Is a valid pair that we are considering
			int startId = threadIdx.x*2;
			threadFirst = blockFirst + startId*elemPerThread; // Start of the first pair
			int splitPoint = threadFirst+elemPerThread; // End location of the first pair

			if (threadIdx.x == (numThreads - 1)) { // If it's the last thread, we want the end to be the last
				threadLast = blockLast;
			} else { // Otherwise, we traverse 2*M elements, where M is the number of elements per pair to the end
				threadLast = threadFirst+2*elemPerThread-1;
			}
			Merge(array, temp, threadFirst, splitPoint-1, splitPoint, threadLast);
		}
		__syncthreads(); // Wait until all threads are done
		numThreads = (numThreads+2-1)/2;
		elemPerThread *= 2;
	}


	if (threadIdx.x == 0) { //Finish off the merge. We did not set test condition to 0 above because we always took the ceiling. So we address single thread case here.
		int splitPoint = blockFirst+elemPerThread; // Calculate the split point

		//Last merge
		Merge(array, temp, blockFirst, splitPoint-1, splitPoint, blockLast);
	}
}

// This code is the same as the latter steps in the above. I did not write this
// in a separate function at the time because I wanted to test my code first
// and avoid repeating the first 5-6 lines of computation in the threads.
__global__ void DoMergeOnly(int *array, int *temp, int size) {
	int elemPerBlock = size;
	int blockFirst = 0;
	int blockLast = size-1;

	int elemPerThread = elemPerBlock/blockDim.x;

	//Now we merge pair by pair
	int numThreads = (blockDim.x+2-1)/2;
	while (numThreads > 1) {
		if (threadIdx.x < numThreads) {
			int startId = threadIdx.x*2;
			int threadFirst = blockFirst + startId*elemPerThread;
			int splitPoint = threadFirst+elemPerThread;

			int threadLast;
			if (threadIdx.x == (numThreads - 1)) {
				threadLast = blockLast;
			} else {
				threadLast = threadFirst+2*elemPerThread-1;
			}

			Merge(array, temp, threadFirst, splitPoint-1, splitPoint, threadLast);
		}
		__syncthreads();
		numThreads = (numThreads+2-1)/2;
		elemPerThread *= 2;
	}

	if (threadIdx.x == 0) { //Finish off the merge
		int splitPoint = blockFirst+elemPerThread;
		Merge(array, temp, blockFirst, splitPoint-1, splitPoint, blockLast);
	}
}

	.text
	.file	"MergeSort.hip"
	.globl	main                            # -- Begin function main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	subq	$16, %rsp
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movl	$164, %edi
	callq	malloc
	movq	%rax, %rbx
	movl	$41, %eax
	movq	%rbx, %rcx
.LBB0_1:                                # %.lr.ph.i
                                        # =>This Inner Loop Header: Depth=1
	movl	%eax, (%rcx)
	addq	$4, %rcx
	decq	%rax
	jne	.LBB0_1
# %bb.2:                                # %_Z19CreateUnsortedArrayi.exit
	movabsq	$4294967300, %r14               # imm = 0x100000004
	movq	%rsp, %r15
	movl	$164, %esi
	movq	%r15, %rdi
	callq	hipMalloc
	movq	(%r15), %rdi
	movl	$164, %edx
	movq	%rbx, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	8(%rsp), %rdi
	movl	$164, %esi
	callq	hipMalloc
	movl	$.L.str, %edi
	movq	%rbx, %rsi
	xorl	%edx, %edx
	movl	$40, %ecx
	callq	_Z10PrintArrayPKcPiii
	movl	$10, %edi
	callq	putchar@PLT
	leaq	6(%r14), %rdx
	movq	%r14, %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_4
# %bb.3:
	movq	(%rsp), %rdi
	movq	8(%rsp), %rsi
	movl	$41, %edx
	callq	_Z24__device_stub__MergeSortPiS_i
.LBB0_4:
	movq	(%rsp), %rsi
	movl	$164, %edx
	movq	%rbx, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movl	$.L.str.2, %edi
	movq	%rbx, %rsi
	xorl	%edx, %edx
	movl	$40, %ecx
	callq	_Z10PrintArrayPKcPiii
	movl	$10, %edi
	callq	putchar@PLT
	leaq	-3(%r14), %rdi
	movl	$1, %esi
	movq	%r14, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_6
# %bb.5:
	movq	(%rsp), %rdi
	movq	8(%rsp), %rsi
	movl	$41, %edx
	callq	_Z26__device_stub__DoMergeOnlyPiS_i
.LBB0_6:
	movq	(%rsp), %rsi
	movl	$164, %edx
	movq	%rbx, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movl	$.L.str.3, %edi
	movq	%rbx, %rsi
	xorl	%edx, %edx
	movl	$40, %ecx
	callq	_Z10PrintArrayPKcPiii
	movl	$10, %edi
	callq	putchar@PLT
	movq	%rbx, %rdi
	callq	free
	movq	(%rsp), %rdi
	callq	hipFree
	movq	8(%rsp), %rdi
	callq	hipFree
	xorl	%eax, %eax
	addq	$16, %rsp
	.cfi_def_cfa_offset 32
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end0:
	.size	main, .Lfunc_end0-main
	.cfi_endproc
                                        # -- End function
	.globl	_Z19CreateUnsortedArrayi        # -- Begin function _Z19CreateUnsortedArrayi
	.type	_Z19CreateUnsortedArrayi,@function
_Z19CreateUnsortedArrayi:               # @_Z19CreateUnsortedArrayi
	.cfi_startproc
# %bb.0:
	pushq	%r14
	.cfi_def_cfa_offset 16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	pushq	%rax
	.cfi_def_cfa_offset 32
	.cfi_offset %rbx, -24
	.cfi_offset %r14, -16
	movl	%edi, %ebx
	movslq	%edi, %r14
	leaq	(,%r14,4), %rdi
	callq	malloc
	testl	%r14d, %r14d
	jle	.LBB1_3
# %bb.1:                                # %.lr.ph.preheader
	movl	%ebx, %ecx
	movq	%rax, %rdx
.LBB1_2:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movl	%ecx, (%rdx)
	addq	$4, %rdx
	decq	%rcx
	jne	.LBB1_2
.LBB1_3:                                # %._crit_edge
	addq	$8, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end1:
	.size	_Z19CreateUnsortedArrayi, .Lfunc_end1-_Z19CreateUnsortedArrayi
	.cfi_endproc
                                        # -- End function
	.globl	_Z10PrintArrayPKcPiii           # -- Begin function _Z10PrintArrayPKcPiii
	.type	_Z10PrintArrayPKcPiii,@function
_Z10PrintArrayPKcPiii:                  # @_Z10PrintArrayPKcPiii
	.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
	pushq	%rax
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movl	%ecx, %ebx
	movl	%edx, %ebp
	movq	%rsi, %r14
	callq	puts@PLT
	cmpl	%ebx, %ebp
	jg	.LBB2_3
# %bb.1:                                # %.lr.ph.preheader
	movslq	%ebp, %rax
	leaq	(%r14,%rax,4), %r14
	subl	%ebp, %ebx
	incl	%ebx
	xorl	%r15d, %r15d
.LBB2_2:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movl	(%r14,%r15,4), %esi
	movl	$.L.str.5, %edi
	xorl	%eax, %eax
	callq	printf
	incq	%r15
	cmpl	%r15d, %ebx
	jne	.LBB2_2
.LBB2_3:                                # %._crit_edge
	movl	$10, %edi
	addq	$8, %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
	jmp	putchar@PLT                     # TAILCALL
.Lfunc_end2:
	.size	_Z10PrintArrayPKcPiii, .Lfunc_end2-_Z10PrintArrayPKcPiii
	.cfi_endproc
                                        # -- End function
	.globl	_Z24__device_stub__MergeSortPiS_i # -- Begin function _Z24__device_stub__MergeSortPiS_i
	.type	_Z24__device_stub__MergeSortPiS_i,@function
_Z24__device_stub__MergeSortPiS_i:      # @_Z24__device_stub__MergeSortPiS_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	$_Z9MergeSortPiS_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_end3:
	.size	_Z24__device_stub__MergeSortPiS_i, .Lfunc_end3-_Z24__device_stub__MergeSortPiS_i
	.cfi_endproc
                                        # -- End function
	.globl	_Z26__device_stub__DoMergeOnlyPiS_i # -- Begin function _Z26__device_stub__DoMergeOnlyPiS_i
	.type	_Z26__device_stub__DoMergeOnlyPiS_i,@function
_Z26__device_stub__DoMergeOnlyPiS_i:    # @_Z26__device_stub__DoMergeOnlyPiS_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	$_Z11DoMergeOnlyPiS_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_end4:
	.size	_Z26__device_stub__DoMergeOnlyPiS_i, .Lfunc_end4-_Z26__device_stub__DoMergeOnlyPiS_i
	.cfi_endproc
                                        # -- End function
	.globl	_Z17CreateRandomArrayi          # -- Begin function _Z17CreateRandomArrayi
	.type	_Z17CreateRandomArrayi,@function
_Z17CreateRandomArrayi:                 # @_Z17CreateRandomArrayi
	.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
	pushq	%rax
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movl	%edi, %ebp
	movslq	%edi, %r14
	leaq	(,%r14,4), %rdi
	callq	malloc
	movq	%rax, %rbx
	testl	%r14d, %r14d
	jle	.LBB5_3
# %bb.1:                                # %.lr.ph.preheader
	movl	%ebp, %r14d
	xorl	%r15d, %r15d
.LBB5_2:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	callq	rand
	cltq
	imulq	$1374389535, %rax, %rcx         # imm = 0x51EB851F
	movq	%rcx, %rdx
	shrq	$63, %rdx
	sarq	$37, %rcx
	addl	%edx, %ecx
	imull	$100, %ecx, %ecx
	subl	%ecx, %eax
	movl	%eax, (%rbx,%r15,4)
	incq	%r15
	cmpq	%r15, %r14
	jne	.LBB5_2
.LBB5_3:                                # %._crit_edge
	movq	%rbx, %rax
	addq	$8, %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_end5:
	.size	_Z17CreateRandomArrayi, .Lfunc_end5-_Z17CreateRandomArrayi
	.cfi_endproc
                                        # -- End function
	.globl	_Z13InsertionSortPiii           # -- Begin function _Z13InsertionSortPiii
	.type	_Z13InsertionSortPiii,@function
_Z13InsertionSortPiii:                  # @_Z13InsertionSortPiii
	.cfi_startproc
# %bb.0:
	cmpl	%edx, %esi
	jg	.LBB6_6
# %bb.1:                                # %.preheader.preheader
	movslq	%esi, %rax
	movq	%rax, %rcx
.LBB6_2:                                # %.preheader
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB6_3 Depth 2
	movq	%rcx, %r8
	cmpl	%esi, %ecx
	jle	.LBB6_5
.LBB6_3:                                # %.lr.ph
                                        #   Parent Loop BB6_2 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movl	-4(%rdi,%r8,4), %r9d
	movl	(%rdi,%r8,4), %r10d
	cmpl	%r9d, %r10d
	jge	.LBB6_5
# %bb.4:                                #   in Loop: Header=BB6_3 Depth=2
	movl	%r10d, -4(%rdi,%r8,4)
	movl	%r9d, (%rdi,%r8,4)
	decq	%r8
	cmpq	%rax, %r8
	jg	.LBB6_3
.LBB6_5:                                # %.critedge
                                        #   in Loop: Header=BB6_2 Depth=1
	movq	%rcx, %r8
	incq	%r8
	cmpl	%edx, %ecx
	movq	%r8, %rcx
	jl	.LBB6_2
.LBB6_6:                                # %._crit_edge
	retq
.Lfunc_end6:
	.size	_Z13InsertionSortPiii, .Lfunc_end6-_Z13InsertionSortPiii
	.cfi_endproc
                                        # -- End function
	.globl	_Z5MergePiS_iiii                # -- Begin function _Z5MergePiS_iiii
	.type	_Z5MergePiS_iiii,@function
_Z5MergePiS_iiii:                       # @_Z5MergePiS_iiii
	.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
	.cfi_offset %rbx, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
                                        # kill: def $r9d killed $r9d def $r9
                                        # kill: def $r8d killed $r8d def $r8
                                        # kill: def $ecx killed $ecx def $rcx
	cmpl	%ecx, %edx
	jg	.LBB7_1
# %bb.13:
	cmpl	%r9d, %r8d
	jg	.LBB7_1
# %bb.14:                               # %.lr.ph.preheader
	movslq	%edx, %rax
	leaq	(%rsi,%rax,4), %r11
	movl	%edx, %eax
	movl	%edx, %r10d
.LBB7_15:                               # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movslq	%r10d, %r10
	movl	(%rdi,%r10,4), %ebx
	movslq	%r8d, %r8
	movl	(%rdi,%r8,4), %ebp
	xorl	%r14d, %r14d
	xorl	%r15d, %r15d
	cmpl	%ebp, %ebx
	setge	%r14b
	setl	%r15b
	cmovll	%ebx, %ebp
	addl	%r15d, %r10d
	addl	%r14d, %r8d
	movl	%ebp, (%r11)
	incl	%eax
	cmpl	%ecx, %r10d
	jg	.LBB7_2
# %bb.16:                               # %.lr.ph
                                        #   in Loop: Header=BB7_15 Depth=1
	addq	$4, %r11
	cmpl	%r9d, %r8d
	jle	.LBB7_15
	jmp	.LBB7_2
.LBB7_1:
	movl	%edx, %r10d
	movl	%edx, %eax
.LBB7_2:                                # %.preheader51
	cmpl	%ecx, %r10d
	jg	.LBB7_6
# %bb.3:                                # %.lr.ph60.preheader
	movslq	%eax, %rbx
	movslq	%r10d, %r11
	leaq	(%rdi,%r11,4), %r11
	subl	%r10d, %ecx
	incl	%ecx
	leaq	(%rsi,%rbx,4), %rbx
	xorl	%r10d, %r10d
.LBB7_4:                                # %.lr.ph60
                                        # =>This Inner Loop Header: Depth=1
	movl	(%r11,%r10,4), %ebp
	movl	%ebp, (%rbx,%r10,4)
	incq	%r10
	cmpl	%r10d, %ecx
	jne	.LBB7_4
# %bb.5:                                # %.preheader50.loopexit
	addl	%r10d, %eax
.LBB7_6:                                # %.preheader50
	cmpl	%r9d, %r8d
	jg	.LBB7_9
# %bb.7:                                # %.lr.ph65.preheader
	movslq	%eax, %r10
	movslq	%r8d, %rax
	leaq	(%rdi,%rax,4), %rax
	movl	%r9d, %ecx
	subl	%r8d, %ecx
	incl	%ecx
	leaq	(%rsi,%r10,4), %r8
	xorl	%r10d, %r10d
.LBB7_8:                                # %.lr.ph65
                                        # =>This Inner Loop Header: Depth=1
	movl	(%rax,%r10,4), %r11d
	movl	%r11d, (%r8,%r10,4)
	incq	%r10
	cmpl	%r10d, %ecx
	jne	.LBB7_8
.LBB7_9:                                # %.preheader
	cmpl	%r9d, %edx
	jg	.LBB7_12
# %bb.10:                               # %.lr.ph68.preheader
	movslq	%edx, %rcx
	leaq	(%rsi,%rcx,4), %rax
	leaq	(%rdi,%rcx,4), %rcx
	subl	%edx, %r9d
	incl	%r9d
	xorl	%edx, %edx
.LBB7_11:                               # %.lr.ph68
                                        # =>This Inner Loop Header: Depth=1
	movl	(%rax,%rdx,4), %esi
	movl	%esi, (%rcx,%rdx,4)
	incq	%rdx
	cmpl	%edx, %r9d
	jne	.LBB7_11
.LBB7_12:                               # %._crit_edge
	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_end7:
	.size	_Z5MergePiS_iiii, .Lfunc_end7-_Z5MergePiS_iiii
	.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	.LBB8_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB8_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z9MergeSortPiS_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	$_Z11DoMergeOnlyPiS_i, %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_end8:
	.size	__hip_module_ctor, .Lfunc_end8-__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	.LBB9_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
.LBB9_2:
	retq
.Lfunc_end9:
	.size	__hip_module_dtor, .Lfunc_end9-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"Before MergeSort:"
	.size	.L.str, 18

	.type	_Z9MergeSortPiS_i,@object       # @_Z9MergeSortPiS_i
	.section	.rodata,"a",@progbits
	.globl	_Z9MergeSortPiS_i
	.p2align	3, 0x0
_Z9MergeSortPiS_i:
	.quad	_Z24__device_stub__MergeSortPiS_i
	.size	_Z9MergeSortPiS_i, 8

	.type	.L.str.2,@object                # @.str.2
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str.2:
	.asciz	"Intermediate:"
	.size	.L.str.2, 14

	.type	_Z11DoMergeOnlyPiS_i,@object    # @_Z11DoMergeOnlyPiS_i
	.section	.rodata,"a",@progbits
	.globl	_Z11DoMergeOnlyPiS_i
	.p2align	3, 0x0
_Z11DoMergeOnlyPiS_i:
	.quad	_Z26__device_stub__DoMergeOnlyPiS_i
	.size	_Z11DoMergeOnlyPiS_i, 8

	.type	.L.str.3,@object                # @.str.3
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str.3:
	.asciz	"After MergeSort:"
	.size	.L.str.3, 17

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

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z9MergeSortPiS_i"
	.size	.L__unnamed_1, 18

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"_Z11DoMergeOnlyPiS_i"
	.size	.L__unnamed_2, 21

	.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 _Z24__device_stub__MergeSortPiS_i
	.addrsig_sym _Z26__device_stub__DoMergeOnlyPiS_i
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z9MergeSortPiS_i
	.addrsig_sym _Z11DoMergeOnlyPiS_i
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z9MergeSortPiS_i       ; -- Begin function _Z9MergeSortPiS_i
	.globl	_Z9MergeSortPiS_i
	.p2align	8
	.type	_Z9MergeSortPiS_i,@function
_Z9MergeSortPiS_i:                      ; @_Z9MergeSortPiS_i
; %bb.0:
	s_clause 0x2
	s_load_b32 s3, s[0:1], 0x18
	s_load_b32 s4, s[0:1], 0x24
	s_load_b32 s5, s[0:1], 0x10
	s_mov_b32 s9, 0
	s_waitcnt lgkmcnt(0)
	v_cvt_f32_u32_e32 v1, s3
	s_sub_i32 s6, 0, s3
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_iflag_f32_e32 v1, v1
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v1, 0x4f7ffffe, v1
	v_cvt_u32_f32_e32 v1, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_readfirstlane_b32 s2, v1
	s_mul_i32 s6, s6, s2
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_mul_hi_u32 s6, s2, s6
	s_add_i32 s2, s2, s6
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_mul_hi_u32 s2, s5, s2
	s_mul_i32 s6, s2, s3
	s_add_i32 s7, s2, 1
	s_sub_i32 s6, s5, s6
	s_delay_alu instid0(SALU_CYCLE_1)
	s_sub_i32 s8, s6, s3
	s_cmp_ge_u32 s6, s3
	s_cselect_b32 s2, s7, s2
	s_cselect_b32 s6, s8, s6
	s_add_i32 s7, s2, 1
	s_cmp_ge_u32 s6, s3
	s_cselect_b32 s6, s7, s2
	s_add_i32 s3, s3, -1
	s_mul_i32 s2, s6, s15
	s_delay_alu instid0(SALU_CYCLE_1)
	s_add_i32 s7, s2, s6
	s_cmp_eq_u32 s15, s3
	s_cselect_b32 s11, s5, s7
	s_and_b32 s3, s4, 0xffff
	s_add_i32 s10, s11, -1
	v_cvt_f32_u32_e32 v1, s3
	s_sub_i32 s5, 0, s3
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_iflag_f32_e32 v1, v1
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v1, 0x4f7ffffe, v1
	v_cvt_u32_f32_e32 v1, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_readfirstlane_b32 s4, v1
	s_mul_i32 s5, s5, s4
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_mul_hi_u32 s5, s4, s5
	s_add_i32 s4, s4, s5
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_mul_hi_u32 s4, s6, s4
	s_mul_i32 s5, s4, s3
	s_delay_alu instid0(SALU_CYCLE_1)
	s_sub_i32 s5, s6, s5
	s_add_i32 s6, s4, 1
	s_sub_i32 s7, s5, s3
	s_cmp_ge_u32 s5, s3
	s_cselect_b32 s4, s6, s4
	s_cselect_b32 s5, s7, s5
	s_add_i32 s6, s4, 1
	s_cmp_ge_u32 s5, s3
	s_cselect_b32 s8, s6, s4
	s_load_b128 s[4:7], s[0:1], 0x0
	v_mad_u64_u32 v[1:2], null, s8, v0, s[2:3]
	s_add_i32 s0, s3, -1
	s_mov_b32 s1, exec_lo
	v_cmp_eq_u32_e32 vcc_lo, s0, v0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add3_u32 v2, s8, -1, v1
	v_cndmask_b32_e64 v8, v2, s10, vcc_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_le_i32_e64 v1, v8
	s_cbranch_execz .LBB0_8
; %bb.1:                                ; %.preheader.i.preheader
	v_ashrrev_i32_e32 v2, 31, v1
	v_mov_b32_e32 v9, v1
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[2:3], 2, v[1:2]
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v2, vcc_lo, s4, v2
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v3, vcc_lo, s5, v3, vcc_lo
.LBB0_2:                                ; %.preheader.i
                                        ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB0_4 Depth 2
	s_mov_b32 s0, exec_lo
	v_cmpx_gt_i32_e64 v9, v1
	s_cbranch_execz .LBB0_7
; %bb.3:                                ; %.lr.ph.i.preheader
                                        ;   in Loop: Header=BB0_2 Depth=1
	s_delay_alu instid0(VALU_DEP_2)
	v_dual_mov_b32 v5, v3 :: v_dual_mov_b32 v4, v2
	v_mov_b32_e32 v10, v9
	s_mov_b32 s12, 0
                                        ; implicit-def: $sgpr13
.LBB0_4:                                ; %.lr.ph.i
                                        ;   Parent Loop BB0_2 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	global_load_b64 v[6:7], v[4:5], off offset:-4
	s_or_b32 s13, s13, exec_lo
	s_mov_b32 s14, exec_lo
	s_waitcnt vmcnt(0)
	v_cmpx_lt_i32_e64 v7, v6
	s_cbranch_execz .LBB0_6
; %bb.5:                                ;   in Loop: Header=BB0_4 Depth=2
	v_dual_mov_b32 v11, v7 :: v_dual_mov_b32 v12, v6
	v_add_co_u32 v6, vcc_lo, v4, -4
	v_add_co_ci_u32_e32 v7, vcc_lo, -1, v5, vcc_lo
	v_add_nc_u32_e32 v10, -1, v10
	global_store_b64 v[4:5], v[11:12], off offset:-4
	v_dual_mov_b32 v4, v6 :: v_dual_mov_b32 v5, v7
	s_and_not1_b32 s13, s13, exec_lo
	v_cmp_le_i32_e32 vcc_lo, v10, v1
	s_and_b32 s15, vcc_lo, exec_lo
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 s13, s13, s15
.LBB0_6:                                ; %Flow184
                                        ;   in Loop: Header=BB0_4 Depth=2
	s_or_b32 exec_lo, exec_lo, s14
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b32 s14, exec_lo, s13
	s_or_b32 s12, s14, s12
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s12
	s_cbranch_execnz .LBB0_4
.LBB0_7:                                ; %Flow186
                                        ;   in Loop: Header=BB0_2 Depth=1
	s_or_b32 exec_lo, exec_lo, s0
	v_add_nc_u32_e32 v4, 1, v9
	v_cmp_ge_i32_e32 vcc_lo, v9, v8
	v_add_co_u32 v2, s0, v2, 4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_co_ci_u32_e64 v3, s0, 0, v3, s0
	v_mov_b32_e32 v9, v4
	s_or_b32 s9, vcc_lo, s9
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s9
	s_cbranch_execnz .LBB0_2
.LBB0_8:                                ; %Flow188
	s_or_b32 exec_lo, exec_lo, s1
	v_mov_b32_e32 v11, s8
	s_cmp_lt_u32 s3, 3
	s_waitcnt lgkmcnt(0)
	s_waitcnt_vscnt null, 0x0
	s_barrier
	buffer_gl0_inv
	s_cbranch_scc1 .LBB0_30
; %bb.9:                                ; %.lr.ph
	v_lshlrev_b32_e32 v12, 1, v0
	v_mov_b32_e32 v2, s8
	s_add_i32 s8, s2, -1
.LBB0_10:                               ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB0_15 Depth 2
                                        ;     Child Loop BB0_19 Depth 2
                                        ;     Child Loop BB0_23 Depth 2
                                        ;     Child Loop BB0_26 Depth 2
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	v_lshlrev_b32_e32 v11, 1, v2
	s_add_i32 s0, s3, 1
	s_lshr_b32 s9, s0, 1
	s_mov_b32 s0, exec_lo
	v_cmpx_le_u32_e64 s9, v0
	s_xor_b32 s0, exec_lo, s0
; %bb.11:                               ; %._Z5MergePiS_iiii.exit_crit_edge
                                        ;   in Loop: Header=BB0_10 Depth=1
	v_lshlrev_b32_e32 v11, 1, v2
                                        ; implicit-def: $vgpr2
; %bb.12:                               ; %Flow178
                                        ;   in Loop: Header=BB0_10 Depth=1
	s_and_not1_saveexec_b32 s12, s0
	s_cbranch_execz .LBB0_28
; %bb.13:                               ;   in Loop: Header=BB0_10 Depth=1
	v_mul_lo_u32 v14, v12, v2
	s_add_i32 s0, s9, -1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cmp_eq_u32_e32 vcc_lo, s0, v0
	v_add_nc_u32_e32 v1, s2, v14
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_4)
	v_add3_u32 v3, v11, v1, -1
	v_add_nc_u32_e32 v15, v1, v2
	v_mov_b32_e32 v5, v1
	v_mov_b32_e32 v7, v1
	v_cndmask_b32_e64 v13, v3, s10, vcc_lo
	v_cmp_lt_i32_e32 vcc_lo, 0, v2
	v_ashrrev_i32_e32 v2, 31, v1
	v_mov_b32_e32 v3, v15
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cmp_le_i32_e64 s0, v15, v13
	s_and_b32 s0, vcc_lo, s0
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_saveexec_b32 s13, s0
	s_cbranch_execz .LBB0_17
; %bb.14:                               ; %.lr.ph.i46.preheader
                                        ;   in Loop: Header=BB0_10 Depth=1
	v_lshlrev_b64 v[6:7], 2, v[1:2]
	v_mov_b32_e32 v5, v1
	v_mov_b32_e32 v3, v15
	s_mov_b32 s14, 0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_co_u32 v9, vcc_lo, s6, v6
	v_add_co_ci_u32_e32 v10, vcc_lo, s7, v7, vcc_lo
	v_mov_b32_e32 v7, v1
.LBB0_15:                               ; %.lr.ph.i46
                                        ;   Parent Loop BB0_10 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_ashrrev_i32_e32 v8, 31, v7
	v_ashrrev_i32_e32 v4, 31, v3
	v_add_nc_u32_e32 v5, 1, v5
	v_lshlrev_b64 v[16:17], 2, v[7:8]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_lshlrev_b64 v[18:19], 2, v[3:4]
	v_add_co_u32 v16, vcc_lo, s4, v16
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_ci_u32_e32 v17, vcc_lo, s5, v17, vcc_lo
	v_add_co_u32 v18, vcc_lo, s4, v18
	s_delay_alu instid0(VALU_DEP_4)
	v_add_co_ci_u32_e32 v19, vcc_lo, s5, v19, vcc_lo
	s_clause 0x1
	global_load_b32 v4, v[16:17], off
	global_load_b32 v6, v[18:19], off
	s_waitcnt vmcnt(0)
	v_cmp_lt_i32_e32 vcc_lo, v4, v6
	v_add_co_ci_u32_e32 v7, vcc_lo, 0, v7, vcc_lo
	v_cmp_ge_i32_e32 vcc_lo, v4, v6
	v_min_i32_e32 v4, v4, v6
	v_add_co_ci_u32_e32 v3, vcc_lo, 0, v3, vcc_lo
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_3) | instid1(VALU_DEP_1)
	v_cmp_ge_i32_e32 vcc_lo, v7, v15
	global_store_b32 v[9:10], v4, off
	v_add_co_u32 v9, s1, v9, 4
	v_cmp_gt_i32_e64 s0, v3, v13
	s_or_b32 s0, vcc_lo, s0
	v_add_co_ci_u32_e64 v10, vcc_lo, 0, v10, s1
	s_and_b32 s0, exec_lo, s0
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_or_b32 s14, s0, s14
	s_and_not1_b32 exec_lo, exec_lo, s14
	s_cbranch_execnz .LBB0_15
; %bb.16:                               ; %Flow175
                                        ;   in Loop: Header=BB0_10 Depth=1
	s_or_b32 exec_lo, exec_lo, s14
.LBB0_17:                               ; %Flow176
                                        ;   in Loop: Header=BB0_10 Depth=1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_or_b32 exec_lo, exec_lo, s13
	s_mov_b32 s1, exec_lo
	v_cmpx_lt_i32_e64 v7, v15
	s_cbranch_execz .LBB0_21
; %bb.18:                               ; %.lr.ph60.i.preheader
                                        ;   in Loop: Header=BB0_10 Depth=1
	v_ashrrev_i32_e32 v6, 31, v5
	v_ashrrev_i32_e32 v8, 31, v7
	v_add_nc_u32_e32 v10, -1, v7
	v_add_nc_u32_e32 v4, -1, v15
	s_mov_b32 s13, 0
	v_lshlrev_b64 v[16:17], 2, v[5:6]
	v_lshlrev_b64 v[8:9], 2, v[7:8]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_u32 v6, vcc_lo, s6, v16
	v_add_co_ci_u32_e32 v7, vcc_lo, s7, v17, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_co_u32 v8, vcc_lo, s4, v8
	v_add_co_ci_u32_e32 v9, vcc_lo, s5, v9, vcc_lo
.LBB0_19:                               ; %.lr.ph60.i
                                        ;   Parent Loop BB0_10 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	global_load_b32 v15, v[8:9], off
	v_add_nc_u32_e32 v10, 1, v10
	v_add_co_u32 v8, vcc_lo, v8, 4
	v_add_co_ci_u32_e32 v9, vcc_lo, 0, v9, vcc_lo
	s_delay_alu instid0(VALU_DEP_3)
	v_cmp_ge_i32_e32 vcc_lo, v10, v4
	v_add_nc_u32_e32 v5, 1, v5
	s_or_b32 s13, vcc_lo, s13
	s_waitcnt vmcnt(0)
	global_store_b32 v[6:7], v15, off
	v_add_co_u32 v6, s0, v6, 4
	s_delay_alu instid0(VALU_DEP_1)
	v_add_co_ci_u32_e64 v7, s0, 0, v7, s0
	s_and_not1_b32 exec_lo, exec_lo, s13
	s_cbranch_execnz .LBB0_19
; %bb.20:                               ; %Flow172
                                        ;   in Loop: Header=BB0_10 Depth=1
	s_or_b32 exec_lo, exec_lo, s13
.LBB0_21:                               ; %Flow173
                                        ;   in Loop: Header=BB0_10 Depth=1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_or_b32 exec_lo, exec_lo, s1
	s_mov_b32 s1, exec_lo
	v_cmpx_le_i32_e64 v3, v13
	s_cbranch_execz .LBB0_24
; %bb.22:                               ; %.lr.ph65.i.preheader
                                        ;   in Loop: Header=BB0_10 Depth=1
	v_ashrrev_i32_e32 v6, 31, v5
	v_ashrrev_i32_e32 v4, 31, v3
	v_add_nc_u32_e32 v7, -1, v3
	s_mov_b32 s13, 0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_lshlrev_b64 v[5:6], 2, v[5:6]
	v_lshlrev_b64 v[8:9], 2, v[3:4]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_u32 v3, vcc_lo, s6, v5
	v_add_co_ci_u32_e32 v4, vcc_lo, s7, v6, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_co_u32 v5, vcc_lo, s4, v8
	v_add_co_ci_u32_e32 v6, vcc_lo, s5, v9, vcc_lo
.LBB0_23:                               ; %.lr.ph65.i
                                        ;   Parent Loop BB0_10 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	global_load_b32 v8, v[5:6], off
	v_add_nc_u32_e32 v7, 1, v7
	v_add_co_u32 v5, vcc_lo, v5, 4
	v_add_co_ci_u32_e32 v6, vcc_lo, 0, v6, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_4) | instid1(VALU_DEP_1)
	v_cmp_ge_i32_e32 vcc_lo, v7, v13
	s_or_b32 s13, vcc_lo, s13
	s_waitcnt vmcnt(0)
	global_store_b32 v[3:4], v8, off
	v_add_co_u32 v3, s0, v3, 4
	v_add_co_ci_u32_e64 v4, s0, 0, v4, s0
	s_and_not1_b32 exec_lo, exec_lo, s13
	s_cbranch_execnz .LBB0_23
.LBB0_24:                               ; %Flow171
                                        ;   in Loop: Header=BB0_10 Depth=1
	s_or_b32 exec_lo, exec_lo, s1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_mov_b32 s1, exec_lo
	v_cmpx_le_i32_e64 v1, v13
	s_cbranch_execz .LBB0_27
; %bb.25:                               ; %.lr.ph68.i.preheader
                                        ;   in Loop: Header=BB0_10 Depth=1
	v_lshlrev_b64 v[3:4], 2, v[1:2]
	v_add_nc_u32_e32 v5, s8, v14
	s_mov_b32 s13, 0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_u32 v1, vcc_lo, s6, v3
	v_add_co_ci_u32_e32 v2, vcc_lo, s7, v4, vcc_lo
	v_add_co_u32 v3, vcc_lo, s4, v3
	v_add_co_ci_u32_e32 v4, vcc_lo, s5, v4, vcc_lo
.LBB0_26:                               ; %.lr.ph68.i
                                        ;   Parent Loop BB0_10 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	global_load_b32 v6, v[1:2], off
	v_add_nc_u32_e32 v5, 1, v5
	v_add_co_u32 v1, vcc_lo, v1, 4
	v_add_co_ci_u32_e32 v2, vcc_lo, 0, v2, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_4) | instid1(VALU_DEP_1)
	v_cmp_ge_i32_e32 vcc_lo, v5, v13
	s_or_b32 s13, vcc_lo, s13
	s_waitcnt vmcnt(0)
	global_store_b32 v[3:4], v6, off
	v_add_co_u32 v3, s0, v3, 4
	v_add_co_ci_u32_e64 v4, s0, 0, v4, s0
	s_and_not1_b32 exec_lo, exec_lo, s13
	s_cbranch_execnz .LBB0_26
.LBB0_27:                               ; %Flow169
                                        ;   in Loop: Header=BB0_10 Depth=1
	s_or_b32 exec_lo, exec_lo, s1
.LBB0_28:                               ; %Flow179
                                        ;   in Loop: Header=BB0_10 Depth=1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 exec_lo, exec_lo, s12
	s_cmp_gt_u32 s3, 4
	s_waitcnt_vscnt null, 0x0
	s_barrier
	buffer_gl0_inv
	s_cbranch_scc0 .LBB0_30
; %bb.29:                               ;   in Loop: Header=BB0_10 Depth=1
	v_mov_b32_e32 v2, v11
	s_mov_b32 s3, s9
	s_branch .LBB0_10
.LBB0_30:                               ; %._crit_edge
	s_mov_b32 s0, exec_lo
	v_cmpx_eq_u32_e32 0, v0
	s_cbranch_execz .LBB0_45
; %bb.31:
	v_add_nc_u32_e32 v6, s2, v11
	v_cmp_lt_i32_e32 vcc_lo, 0, v11
	v_mov_b32_e32 v2, s2
	v_mov_b32_e32 v4, s2
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_cmp_gt_i32_e64 s0, s11, v6
	v_mov_b32_e32 v0, v6
	s_and_b32 s0, vcc_lo, s0
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_saveexec_b32 s1, s0
	s_cbranch_execz .LBB0_35
; %bb.32:                               ; %.lr.ph.i68.preheader
	s_ashr_i32 s3, s2, 31
	v_dual_mov_b32 v4, s2 :: v_dual_mov_b32 v3, 0
	s_lshl_b64 s[8:9], s[2:3], 2
	v_mov_b32_e32 v0, v6
	s_add_u32 s8, s6, s8
	s_addc_u32 s9, s7, s9
	s_mov_b32 s3, 0
	s_mov_b32 s12, s2
.LBB0_33:                               ; %.lr.ph.i68
                                        ; =>This Inner Loop Header: Depth=1
	v_ashrrev_i32_e32 v5, 31, v4
	v_ashrrev_i32_e32 v1, 31, v0
	s_add_i32 s12, s12, 1
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_lshlrev_b64 v[7:8], 2, v[4:5]
	v_lshlrev_b64 v[1:2], 2, v[0:1]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_u32 v7, vcc_lo, s4, v7
	v_add_co_ci_u32_e32 v8, vcc_lo, s5, v8, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_co_u32 v1, vcc_lo, s4, v1
	v_add_co_ci_u32_e32 v2, vcc_lo, s5, v2, vcc_lo
	s_clause 0x1
	global_load_b32 v5, v[7:8], off
	global_load_b32 v1, v[1:2], off
	v_mov_b32_e32 v2, s12
	s_waitcnt vmcnt(0)
	v_cmp_lt_i32_e32 vcc_lo, v5, v1
	v_add_co_ci_u32_e32 v4, vcc_lo, 0, v4, vcc_lo
	v_cmp_ge_i32_e32 vcc_lo, v5, v1
	v_min_i32_e32 v1, v5, v1
	v_add_co_ci_u32_e32 v0, vcc_lo, 0, v0, vcc_lo
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_cmp_ge_i32_e32 vcc_lo, v4, v6
	global_store_b32 v3, v1, s[8:9]
	v_cmp_le_i32_e64 s0, s11, v0
	s_or_b32 s0, vcc_lo, s0
	s_add_u32 s8, s8, 4
	s_addc_u32 s9, s9, 0
	s_and_b32 s0, exec_lo, s0
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_or_b32 s3, s0, s3
	s_and_not1_b32 exec_lo, exec_lo, s3
	s_cbranch_execnz .LBB0_33
; %bb.34:                               ; %Flow165
	s_or_b32 exec_lo, exec_lo, s3
.LBB0_35:                               ; %Flow166
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_or_b32 exec_lo, exec_lo, s1
	s_mov_b32 s1, exec_lo
	v_cmpx_lt_i32_e64 v4, v6
	s_cbranch_execz .LBB0_39
; %bb.36:                               ; %.lr.ph60.i52.preheader
	v_ashrrev_i32_e32 v3, 31, v2
	v_ashrrev_i32_e32 v5, 31, v4
	v_add_nc_u32_e32 v1, -1, v6
	v_add_nc_u32_e32 v7, -1, v4
	s_mov_b32 s3, 0
	v_lshlrev_b64 v[8:9], 2, v[2:3]
	v_lshlrev_b64 v[5:6], 2, v[4:5]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_u32 v3, vcc_lo, s6, v8
	v_add_co_ci_u32_e32 v4, vcc_lo, s7, v9, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_co_u32 v5, vcc_lo, s4, v5
	v_add_co_ci_u32_e32 v6, vcc_lo, s5, v6, vcc_lo
.LBB0_37:                               ; %.lr.ph60.i52
                                        ; =>This Inner Loop Header: Depth=1
	global_load_b32 v8, v[5:6], off
	v_add_nc_u32_e32 v7, 1, v7
	v_add_co_u32 v5, vcc_lo, v5, 4
	v_add_co_ci_u32_e32 v6, vcc_lo, 0, v6, vcc_lo
	s_delay_alu instid0(VALU_DEP_3)
	v_cmp_ge_i32_e32 vcc_lo, v7, v1
	v_add_nc_u32_e32 v2, 1, v2
	s_or_b32 s3, vcc_lo, s3
	s_waitcnt vmcnt(0)
	global_store_b32 v[3:4], v8, off
	v_add_co_u32 v3, s0, v3, 4
	s_delay_alu instid0(VALU_DEP_1)
	v_add_co_ci_u32_e64 v4, s0, 0, v4, s0
	s_and_not1_b32 exec_lo, exec_lo, s3
	s_cbranch_execnz .LBB0_37
; %bb.38:                               ; %Flow163
	s_or_b32 exec_lo, exec_lo, s3
.LBB0_39:                               ; %Flow164
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_or_b32 exec_lo, exec_lo, s1
	s_mov_b32 s1, exec_lo
	v_cmpx_gt_i32_e64 s11, v0
	s_cbranch_execz .LBB0_42
; %bb.40:                               ; %.lr.ph65.i59.preheader
	v_ashrrev_i32_e32 v3, 31, v2
	v_ashrrev_i32_e32 v1, 31, v0
	v_add_nc_u32_e32 v4, -1, v0
	s_mov_b32 s3, 0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_lshlrev_b64 v[2:3], 2, v[2:3]
	v_lshlrev_b64 v[5:6], 2, v[0:1]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_u32 v0, vcc_lo, s6, v2
	v_add_co_ci_u32_e32 v1, vcc_lo, s7, v3, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_co_u32 v2, vcc_lo, s4, v5
	v_add_co_ci_u32_e32 v3, vcc_lo, s5, v6, vcc_lo
.LBB0_41:                               ; %.lr.ph65.i59
                                        ; =>This Inner Loop Header: Depth=1
	global_load_b32 v5, v[2:3], off
	v_add_nc_u32_e32 v4, 1, v4
	v_add_co_u32 v2, vcc_lo, v2, 4
	v_add_co_ci_u32_e32 v3, vcc_lo, 0, v3, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_4) | instid1(VALU_DEP_1)
	v_cmp_le_i32_e32 vcc_lo, s10, v4
	s_or_b32 s3, vcc_lo, s3
	s_waitcnt vmcnt(0)
	global_store_b32 v[0:1], v5, off
	v_add_co_u32 v0, s0, v0, 4
	v_add_co_ci_u32_e64 v1, s0, 0, v1, s0
	s_and_not1_b32 exec_lo, exec_lo, s3
	s_cbranch_execnz .LBB0_41
.LBB0_42:                               ; %Flow162
	s_or_b32 exec_lo, exec_lo, s1
	s_cmp_ge_i32 s2, s11
	s_cbranch_scc1 .LBB0_45
; %bb.43:                               ; %.lr.ph68.i65.preheader
	s_ashr_i32 s3, s2, 31
	s_add_i32 s8, s2, -1
	s_lshl_b64 s[12:13], s[2:3], 2
	v_mov_b32_e32 v0, 0
	s_add_u32 s0, s6, s12
	s_addc_u32 s1, s7, s13
	s_add_u32 s2, s4, s12
	s_addc_u32 s3, s5, s13
.LBB0_44:                               ; %.lr.ph68.i65
                                        ; =>This Inner Loop Header: Depth=1
	global_load_b32 v1, v0, s[0:1]
	s_add_i32 s8, s8, 1
	s_add_u32 s0, s0, 4
	s_addc_u32 s1, s1, 0
	s_waitcnt vmcnt(0)
	global_store_b32 v0, v1, s[2:3]
	s_add_u32 s2, s2, 4
	s_addc_u32 s3, s3, 0
	s_cmp_lt_i32 s8, s10
	s_cbranch_scc1 .LBB0_44
.LBB0_45:                               ; %_Z5MergePiS_iiii.exit77
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z9MergeSortPiS_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 20
		.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	_Z9MergeSortPiS_i, .Lfunc_end0-_Z9MergeSortPiS_i
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 2248
; NumSgprs: 18
; NumVgprs: 20
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 2
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 20
; 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	_Z11DoMergeOnlyPiS_i    ; -- Begin function _Z11DoMergeOnlyPiS_i
	.globl	_Z11DoMergeOnlyPiS_i
	.p2align	8
	.type	_Z11DoMergeOnlyPiS_i,@function
_Z11DoMergeOnlyPiS_i:                   ; @_Z11DoMergeOnlyPiS_i
; %bb.0:
	s_clause 0x1
	s_load_b32 s2, s[0:1], 0x24
	s_load_b32 s12, s[0:1], 0x10
	s_waitcnt lgkmcnt(0)
	s_and_b32 s2, s2, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_cvt_f32_u32_e32 v1, s2
	s_sub_i32 s4, 0, s2
	v_rcp_iflag_f32_e32 v1, v1
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v1, 0x4f7ffffe, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cvt_u32_f32_e32 v1, v1
	v_readfirstlane_b32 s3, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_mul_i32 s4, s4, s3
	s_mul_hi_u32 s8, s3, s4
	s_load_b128 s[4:7], s[0:1], 0x0
	s_add_i32 s3, s3, s8
	s_add_i32 s13, s12, -1
	s_mul_hi_u32 s0, s12, s3
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_2) | instid1(SALU_CYCLE_1)
	s_mul_i32 s1, s0, s2
	s_add_i32 s3, s0, 1
	s_sub_i32 s1, s12, s1
	s_sub_i32 s8, s1, s2
	s_cmp_ge_u32 s1, s2
	s_cselect_b32 s0, s3, s0
	s_cselect_b32 s1, s8, s1
	s_add_i32 s3, s0, 1
	s_cmp_ge_u32 s1, s2
	s_cselect_b32 s14, s3, s0
	s_cmp_lt_u32 s2, 3
	s_cbranch_scc1 .LBB1_19
; %bb.1:                                ; %.lr.ph
	v_lshlrev_b32_e32 v11, 1, v0
.LBB1_2:                                ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB1_5 Depth 2
                                        ;     Child Loop BB1_9 Depth 2
                                        ;     Child Loop BB1_13 Depth 2
                                        ;     Child Loop BB1_16 Depth 2
	s_add_i32 s0, s2, 1
	s_mov_b32 s8, exec_lo
	s_lshr_b32 s3, s0, 1
	s_delay_alu instid0(SALU_CYCLE_1)
	v_cmpx_gt_u32_e64 s3, v0
	s_cbranch_execz .LBB1_17
; %bb.3:                                ;   in Loop: Header=BB1_2 Depth=1
	v_mul_lo_u32 v1, s14, v11
	s_lshl_b32 s0, s14, 1
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_add3_u32 v2, v1, s0, -1
	s_add_i32 s0, s3, -1
	v_add_nc_u32_e32 v13, s14, v1
	v_cmp_eq_u32_e32 vcc_lo, s0, v0
	s_cmp_gt_i32 s14, 0
	v_mov_b32_e32 v5, v1
	v_mov_b32_e32 v7, v1
	v_mov_b32_e32 v3, v13
	v_cndmask_b32_e64 v12, v2, s13, vcc_lo
	v_ashrrev_i32_e32 v2, 31, v1
	s_cselect_b32 s0, -1, 0
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	v_cmp_le_i32_e32 vcc_lo, v13, v12
	s_and_b32 s0, s0, vcc_lo
	s_and_saveexec_b32 s9, s0
	s_cbranch_execz .LBB1_7
; %bb.4:                                ; %.lr.ph.i.preheader
                                        ;   in Loop: Header=BB1_2 Depth=1
	v_lshlrev_b64 v[6:7], 2, v[1:2]
	v_mov_b32_e32 v5, v1
	v_mov_b32_e32 v3, v13
	s_mov_b32 s10, 0
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_3)
	v_add_co_u32 v9, vcc_lo, s6, v6
	v_add_co_ci_u32_e32 v10, vcc_lo, s7, v7, vcc_lo
	v_mov_b32_e32 v7, v1
.LBB1_5:                                ; %.lr.ph.i
                                        ;   Parent Loop BB1_2 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_ashrrev_i32_e32 v8, 31, v7
	v_ashrrev_i32_e32 v4, 31, v3
	v_add_nc_u32_e32 v5, 1, v5
	v_lshlrev_b64 v[14:15], 2, v[7:8]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_lshlrev_b64 v[16:17], 2, v[3:4]
	v_add_co_u32 v14, vcc_lo, s4, v14
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_ci_u32_e32 v15, vcc_lo, s5, v15, vcc_lo
	v_add_co_u32 v16, vcc_lo, s4, v16
	s_delay_alu instid0(VALU_DEP_4)
	v_add_co_ci_u32_e32 v17, vcc_lo, s5, v17, vcc_lo
	s_clause 0x1
	global_load_b32 v4, v[14:15], off
	global_load_b32 v6, v[16:17], off
	s_waitcnt vmcnt(0)
	v_cmp_lt_i32_e32 vcc_lo, v4, v6
	v_add_co_ci_u32_e32 v7, vcc_lo, 0, v7, vcc_lo
	v_cmp_ge_i32_e32 vcc_lo, v4, v6
	v_min_i32_e32 v4, v4, v6
	v_add_co_ci_u32_e32 v3, vcc_lo, 0, v3, vcc_lo
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_3) | instid1(VALU_DEP_1)
	v_cmp_ge_i32_e32 vcc_lo, v7, v13
	global_store_b32 v[9:10], v4, off
	v_add_co_u32 v9, s1, v9, 4
	v_cmp_gt_i32_e64 s0, v3, v12
	s_or_b32 s0, vcc_lo, s0
	v_add_co_ci_u32_e64 v10, vcc_lo, 0, v10, s1
	s_and_b32 s0, exec_lo, s0
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_or_b32 s10, s0, s10
	s_and_not1_b32 exec_lo, exec_lo, s10
	s_cbranch_execnz .LBB1_5
; %bb.6:                                ; %Flow151
                                        ;   in Loop: Header=BB1_2 Depth=1
	s_or_b32 exec_lo, exec_lo, s10
.LBB1_7:                                ; %Flow152
                                        ;   in Loop: Header=BB1_2 Depth=1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_or_b32 exec_lo, exec_lo, s9
	s_mov_b32 s1, exec_lo
	v_cmpx_lt_i32_e64 v7, v13
	s_cbranch_execz .LBB1_11
; %bb.8:                                ; %.lr.ph60.i.preheader
                                        ;   in Loop: Header=BB1_2 Depth=1
	v_ashrrev_i32_e32 v6, 31, v5
	v_ashrrev_i32_e32 v8, 31, v7
	v_add_nc_u32_e32 v10, -1, v7
	v_add_nc_u32_e32 v4, -1, v13
	s_mov_b32 s9, 0
	v_lshlrev_b64 v[14:15], 2, v[5:6]
	v_lshlrev_b64 v[8:9], 2, v[7:8]
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_u32 v6, vcc_lo, s6, v14
	v_add_co_ci_u32_e32 v7, vcc_lo, s7, v15, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_co_u32 v8, vcc_lo, s4, v8
	v_add_co_ci_u32_e32 v9, vcc_lo, s5, v9, vcc_lo
.LBB1_9:                                ; %.lr.ph60.i
                                        ;   Parent Loop BB1_2 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	global_load_b32 v13, v[8:9], off
	v_add_nc_u32_e32 v10, 1, v10
	v_add_co_u32 v8, vcc_lo, v8, 4
	v_add_co_ci_u32_e32 v9, vcc_lo, 0, v9, vcc_lo
	s_delay_alu instid0(VALU_DEP_3)
	v_cmp_ge_i32_e32 vcc_lo, v10, v4
	v_add_nc_u32_e32 v5, 1, v5
	s_or_b32 s9, vcc_lo, s9
	s_waitcnt vmcnt(0)
	global_store_b32 v[6:7], v13, off
	v_add_co_u32 v6, s0, v6, 4
	s_delay_alu instid0(VALU_DEP_1)
	v_add_co_ci_u32_e64 v7, s0, 0, v7, s0
	s_and_not1_b32 exec_lo, exec_lo, s9
	s_cbranch_execnz .LBB1_9
; %bb.10:                               ; %Flow148
                                        ;   in Loop: Header=BB1_2 Depth=1
	s_or_b32 exec_lo, exec_lo, s9
.LBB1_11:                               ; %Flow149
                                        ;   in Loop: Header=BB1_2 Depth=1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_or_b32 exec_lo, exec_lo, s1
	s_mov_b32 s1, exec_lo
	v_cmpx_le_i32_e64 v3, v12
	s_cbranch_execz .LBB1_14
; %bb.12:                               ; %.lr.ph65.i.preheader
                                        ;   in Loop: Header=BB1_2 Depth=1
	v_ashrrev_i32_e32 v6, 31, v5
	v_ashrrev_i32_e32 v4, 31, v3
	v_add_nc_u32_e32 v7, -1, v3
	s_mov_b32 s9, 0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_lshlrev_b64 v[5:6], 2, v[5:6]
	v_lshlrev_b64 v[8:9], 2, v[3:4]
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_u32 v3, vcc_lo, s6, v5
	v_add_co_ci_u32_e32 v4, vcc_lo, s7, v6, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_co_u32 v5, vcc_lo, s4, v8
	v_add_co_ci_u32_e32 v6, vcc_lo, s5, v9, vcc_lo
.LBB1_13:                               ; %.lr.ph65.i
                                        ;   Parent Loop BB1_2 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	global_load_b32 v8, v[5:6], off
	v_add_nc_u32_e32 v7, 1, v7
	v_add_co_u32 v5, vcc_lo, v5, 4
	v_add_co_ci_u32_e32 v6, vcc_lo, 0, v6, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_4) | instid1(VALU_DEP_1)
	v_cmp_ge_i32_e32 vcc_lo, v7, v12
	s_or_b32 s9, vcc_lo, s9
	s_waitcnt vmcnt(0)
	global_store_b32 v[3:4], v8, off
	v_add_co_u32 v3, s0, v3, 4
	v_add_co_ci_u32_e64 v4, s0, 0, v4, s0
	s_and_not1_b32 exec_lo, exec_lo, s9
	s_cbranch_execnz .LBB1_13
.LBB1_14:                               ; %Flow147
                                        ;   in Loop: Header=BB1_2 Depth=1
	s_or_b32 exec_lo, exec_lo, s1
	v_cmp_le_i32_e32 vcc_lo, v1, v12
	s_and_b32 exec_lo, exec_lo, vcc_lo
	s_cbranch_execz .LBB1_17
; %bb.15:                               ; %.lr.ph68.i.preheader
                                        ;   in Loop: Header=BB1_2 Depth=1
	v_lshlrev_b64 v[3:4], 2, v[1:2]
	v_add_nc_u32_e32 v5, -1, v1
	s_mov_b32 s1, 0
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_u32 v1, vcc_lo, s6, v3
	v_add_co_ci_u32_e32 v2, vcc_lo, s7, v4, vcc_lo
	v_add_co_u32 v3, vcc_lo, s4, v3
	v_add_co_ci_u32_e32 v4, vcc_lo, s5, v4, vcc_lo
.LBB1_16:                               ; %.lr.ph68.i
                                        ;   Parent Loop BB1_2 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	global_load_b32 v6, v[1:2], off
	v_add_nc_u32_e32 v5, 1, v5
	v_add_co_u32 v1, vcc_lo, v1, 4
	v_add_co_ci_u32_e32 v2, vcc_lo, 0, v2, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_4) | instid1(VALU_DEP_1)
	v_cmp_ge_i32_e32 vcc_lo, v5, v12
	s_or_b32 s1, vcc_lo, s1
	s_waitcnt vmcnt(0)
	global_store_b32 v[3:4], v6, off
	v_add_co_u32 v3, s0, v3, 4
	v_add_co_ci_u32_e64 v4, s0, 0, v4, s0
	s_and_not1_b32 exec_lo, exec_lo, s1
	s_cbranch_execnz .LBB1_16
.LBB1_17:                               ; %Flow153
                                        ;   in Loop: Header=BB1_2 Depth=1
	s_or_b32 exec_lo, exec_lo, s8
	s_lshl_b32 s14, s14, 1
	s_cmp_gt_u32 s2, 4
	s_waitcnt lgkmcnt(0)
	s_waitcnt_vscnt null, 0x0
	s_barrier
	buffer_gl0_inv
	s_cbranch_scc0 .LBB1_19
; %bb.18:                               ;   in Loop: Header=BB1_2 Depth=1
	s_mov_b32 s2, s3
	s_branch .LBB1_2
.LBB1_19:                               ; %._crit_edge
	s_mov_b32 s0, 0
	s_mov_b32 s1, exec_lo
	v_cmpx_eq_u32_e32 0, v0
	s_cbranch_execz .LBB1_32
; %bb.20:
	s_cmp_lt_i32 s14, 1
	s_mov_b32 s8, 0
	s_cselect_b32 s1, -1, 0
	s_cmp_ge_i32 s14, s12
	s_cselect_b32 s2, -1, 0
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 s1, s1, s2
	s_mov_b32 s2, s14
	s_and_b32 vcc_lo, exec_lo, s1
	s_cbranch_vccnz .LBB1_23
; %bb.21:                               ; %.lr.ph.i54.preheader
	v_mov_b32_e32 v0, 0
	s_mov_b32 s9, 0
	s_waitcnt lgkmcnt(0)
	s_mov_b64 s[10:11], s[6:7]
	s_mov_b32 s2, s14
	s_mov_b32 s8, s9
.LBB1_22:                               ; %.lr.ph.i54
                                        ; =>This Inner Loop Header: Depth=1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_lshl_b64 s[16:17], s[8:9], 2
	s_add_u32 s16, s4, s16
	s_addc_u32 s17, s5, s17
	s_ashr_i32 s3, s2, 31
	global_load_b32 v1, v0, s[16:17]
	s_lshl_b64 s[18:19], s[2:3], 2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_add_u32 s16, s4, s18
	s_addc_u32 s17, s5, s19
	global_load_b32 v2, v0, s[16:17]
	s_waitcnt vmcnt(1)
	v_readfirstlane_b32 s1, v1
	s_waitcnt vmcnt(0)
	v_readfirstlane_b32 s3, v2
	v_min_i32_e32 v1, v1, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_4) | instid1(SALU_CYCLE_1)
	s_cmp_ge_i32 s1, s3
	global_store_b32 v0, v1, s[10:11]
	s_cselect_b32 s15, -1, 0
	s_cmp_lt_i32 s1, s3
	s_cselect_b32 s1, -1, 0
	s_cmp_lg_u32 s1, 0
	s_addc_u32 s8, s8, 0
	s_cmp_lg_u32 s15, 0
	s_addc_u32 s2, s2, 0
	s_add_i32 s0, s0, 1
	s_cmp_lt_i32 s8, s14
	s_cselect_b32 s1, -1, 0
	s_cmp_lt_i32 s2, s12
	s_cselect_b32 s3, -1, 0
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_b32 s1, s1, s3
	s_add_u32 s10, s10, 4
	s_addc_u32 s11, s11, 0
	s_and_not1_b32 vcc_lo, exec_lo, s1
	s_cbranch_vccz .LBB1_22
.LBB1_23:                               ; %.preheader51.i33
	s_cmp_ge_i32 s8, s14
	s_cbranch_scc1 .LBB1_26
; %bb.24:                               ; %.lr.ph60.i38.preheader
	s_ashr_i32 s1, s0, 31
	s_add_i32 s14, s14, -1
	s_lshl_b64 s[10:11], s[0:1], 2
	s_mov_b32 s9, 0
	s_waitcnt lgkmcnt(0)
	s_add_u32 s10, s6, s10
	s_addc_u32 s11, s7, s11
	s_lshl_b64 s[16:17], s[8:9], 2
	v_mov_b32_e32 v0, 0
	s_add_i32 s1, s8, -1
	s_add_u32 s8, s4, s16
	s_addc_u32 s9, s5, s17
.LBB1_25:                               ; %.lr.ph60.i38
                                        ; =>This Inner Loop Header: Depth=1
	global_load_b32 v1, v0, s[8:9]
	s_add_i32 s0, s0, 1
	s_waitcnt vmcnt(0)
	global_store_b32 v0, v1, s[10:11]
	s_add_u32 s10, s10, 4
	s_addc_u32 s11, s11, 0
	s_add_i32 s1, s1, 1
	s_add_u32 s8, s8, 4
	s_addc_u32 s9, s9, 0
	s_cmp_ge_i32 s1, s14
	s_cbranch_scc0 .LBB1_25
.LBB1_26:                               ; %.preheader50.i42
	s_cmp_ge_i32 s2, s12
	s_cbranch_scc1 .LBB1_29
; %bb.27:                               ; %.lr.ph65.i45.preheader
	s_ashr_i32 s1, s0, 31
	v_mov_b32_e32 v0, 0
	s_lshl_b64 s[0:1], s[0:1], 2
	s_waitcnt lgkmcnt(0)
	s_add_u32 s0, s6, s0
	s_addc_u32 s1, s7, s1
	s_ashr_i32 s3, s2, 31
	s_add_i32 s8, s2, -1
	s_lshl_b64 s[10:11], s[2:3], 2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_add_u32 s2, s4, s10
	s_addc_u32 s3, s5, s11
.LBB1_28:                               ; %.lr.ph65.i45
                                        ; =>This Inner Loop Header: Depth=1
	global_load_b32 v1, v0, s[2:3]
	s_waitcnt vmcnt(0)
	global_store_b32 v0, v1, s[0:1]
	s_add_u32 s0, s0, 4
	s_addc_u32 s1, s1, 0
	s_add_i32 s8, s8, 1
	s_add_u32 s2, s2, 4
	s_addc_u32 s3, s3, 0
	s_cmp_ge_i32 s8, s13
	s_cbranch_scc0 .LBB1_28
.LBB1_29:                               ; %.preheader.i49
	s_cmp_lt_i32 s12, 1
	s_cbranch_scc1 .LBB1_32
; %bb.30:
	v_mov_b32_e32 v0, 0
.LBB1_31:                               ; %.lr.ph68.i51
                                        ; =>This Inner Loop Header: Depth=1
	s_waitcnt lgkmcnt(0)
	global_load_b32 v1, v0, s[6:7]
	s_waitcnt vmcnt(0)
	global_store_b32 v0, v1, s[4:5]
	s_add_u32 s4, s4, 4
	s_addc_u32 s5, s5, 0
	s_add_u32 s6, s6, 4
	s_addc_u32 s7, s7, 0
	s_add_i32 s12, s12, -1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmp_lg_u32 s12, 0
	s_cbranch_scc1 .LBB1_31
.LBB1_32:                               ; %_Z5MergePiS_iiii.exit63
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z11DoMergeOnlyPiS_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 18
		.amdhsa_next_free_sgpr 20
		.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	_Z11DoMergeOnlyPiS_i, .Lfunc_end1-_Z11DoMergeOnlyPiS_i
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 1560
; NumSgprs: 22
; NumVgprs: 18
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 2
; NumSGPRsForWavesPerEU: 22
; NumVGPRsForWavesPerEU: 18
; 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:           _Z9MergeSortPiS_i
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z9MergeSortPiS_i.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     20
    .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
      - .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:           _Z11DoMergeOnlyPiS_i
    .private_segment_fixed_size: 0
    .sgpr_count:     22
    .sgpr_spill_count: 0
    .symbol:         _Z11DoMergeOnlyPiS_i.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     18
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
