#include <iostream>
#include <cuda_runtime.h>

// Function to perform bottom-up merge sort
__global__ void mergeSortKernel(int *arr, int n) {
    extern __shared__ int sharedArr[];

    // Load data into shared memory
    int tid = threadIdx.x;
    sharedArr[tid] = arr[tid];
    __syncthreads();

    // Perform merge sort in shared memory
    for (int size = 1; size < blockDim.x; size *= 2) {
        int stride = size * 2;
        int l = tid * stride;
        int r = l + size;
        int rlim = min(r + size, n);

        // Load elements into registers
        int a = l + threadIdx.x;
        int b = r + threadIdx.x;

        int x = (a < n) ? sharedArr[a] : INT_MAX;
        int y = (b < rlim) ? sharedArr[b] : INT_MAX;

        // Merge step
        for (int i = 0; i < stride && l + i < n; i++) {
            if (x <= y) {
                sharedArr[l + i] = x;
                x = (a + size < n) ? sharedArr[a + size] : INT_MAX;
                a++;
            } else {
                sharedArr[l + i] = y;
                y = (b + size < rlim) ? sharedArr[b + size] : INT_MAX;
                b++;
            }
        }
        __syncthreads();
    }

    // Write sorted data back to global memory
    arr[tid] = sharedArr[tid];
}

int main() {
    const int n = 128;
    int *arr;
    int h_arr[n] = {12, 11, 13, 5, 6, 7, 8, 9, 10, 1, 2, 3, 4, 15, 14, 16, 17, 18, 19, 20,
                    21, 22, 23, 24, 25, 26, 27, 28, 29, 30, 31, 32, 33, 34, 35, 36, 37, 38, 39, 40,
                    41, 42, 43, 44, 45, 46, 47, 48, 49, 50, 51, 52, 53, 54, 55, 56, 57, 58, 59, 60,
                    61, 62, 63, 64, 65, 66, 67, 68, 69, 70, 71, 72, 73, 74, 75, 76, 77, 78, 79, 80,
                    81, 82, 83, 84, 85, 86, 87, 88, 89, 90, 91, 92, 93, 94, 95, 96, 97, 98, 99, 100,
                    101, 102, 103, 104, 105, 106, 107, 108, 109, 110, 111, 112, 113, 114, 115, 116, 117, 118, 119, 120,
                    121, 122, 123, 124, 125, 126, 127, 128};

    // Allocate Unified Memory -- accessible from CPU or GPU
    cudaError_t err = cudaMallocManaged(&arr, n * sizeof(int));
    if (err != cudaSuccess) {
        std::cerr << "Failed to allocate Unified Memory" << std::endl;
        return -1;
    }

    // Copy data to device
    cudaMemcpy(arr, h_arr, n * sizeof(int), cudaMemcpyHostToDevice);

    // Launch kernel on 1D grid of 1D blocks
    mergeSortKernel<<<1, n, n * sizeof(int)>>>(arr, n);

    // Check for any errors launching the kernel
    err = cudaGetLastError();
    if (err != cudaSuccess) {
        std::cerr << "Failed to launch mergeSortKernel (error code " << err << ")" << std::endl;
        goto FreeMem;
    }

    // Wait for GPU to finish before accessing on host
    err = cudaDeviceSynchronize();
    if (err != cudaSuccess) {
        std::cerr << "Failed to synchronize (error code " << err << ")" << std::endl;
        goto FreeMem;
    }

    // Print sorted array
    std::cout << "Sorted array: ";
    for (int i = 0; i < n; i++) {
        std::cout << arr[i] << " ";
    }
    std::cout << std::endl;

FreeMem:
    cudaFree(arr);

    return 0;
}
	.file	"tmpxft_000be969_00000000-6_cuda_code_067853.cudafe1.cpp"
	.text
#APP
	.globl _ZSt21ios_base_library_initv

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB3638:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE3638:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z36__device_stub__Z15mergeSortKernelPiiPii
	.type	_Z36__device_stub__Z15mergeSortKernelPiiPii, @function
_Z36__device_stub__Z15mergeSortKernelPiiPii:
.LFB3660:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movq	%rdi, 8(%rsp)
	leaq	32(%rsp), %rcx
	leaq	24(%rsp), %rdx
	movl	%esi, 4(%rsp)
	leaq	40(%rsp), %rdi
	leaq	52(%rsp), %rsi
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rax
	movl	$1, 48(%rsp)
	movq	%rax, 88(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 40(%rsp)
	movq	%rax, 52(%rsp)
	movl	$1, 60(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L2
	pushq	32(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z15mergeSortKernelPii(%rip), %rdi
	pushq	32(%rsp)
	.cfi_def_cfa_offset 144
	movq	68(%rsp), %rcx
	movl	76(%rsp), %r8d
	movq	56(%rsp), %rsi
	movl	64(%rsp), %edx
	leaq	104(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L2:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L4
	call	__stack_chk_fail@PLT
.L4:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3660:
	.size	_Z36__device_stub__Z15mergeSortKernelPiiPii, .-_Z36__device_stub__Z15mergeSortKernelPiiPii
	.globl	_Z15mergeSortKernelPii
	.type	_Z15mergeSortKernelPii, @function
_Z15mergeSortKernelPii:
.LFB3661:
	.cfi_startproc
	endbr64
	jmp	_Z36__device_stub__Z15mergeSortKernelPiiPii
	.cfi_endproc
.LFE3661:
	.size	_Z15mergeSortKernelPii, .-_Z15mergeSortKernelPii
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC1:
	.string	"Failed to allocate Unified Memory"
.LC2:
	.string	"Failed to launch mergeSortKernel (error code "
.LC3:
	.string	")"
.LC4:
	.string	"Failed to synchronize (error code "
.LC5:
	.string	"Sorted array: "
.LC6:
	.string	" "
	.section	.rodata
	.align 32
.LC0:
	.long	12
	.long	11
	.long	13
	.long	5
	.long	6
	.long	7
	.long	8
	.long	9
	.long	10
	.long	1
	.long	2
	.long	3
	.long	4
	.long	15
	.long	14
	.long	16
	.long	17
	.long	18
	.long	19
	.long	20
	.long	21
	.long	22
	.long	23
	.long	24
	.long	25
	.long	26
	.long	27
	.long	28
	.long	29
	.long	30
	.long	31
	.long	32
	.long	33
	.long	34
	.long	35
	.long	36
	.long	37
	.long	38
	.long	39
	.long	40
	.long	41
	.long	42
	.long	43
	.long	44
	.long	45
	.long	46
	.long	47
	.long	48
	.long	49
	.long	50
	.long	51
	.long	52
	.long	53
	.long	54
	.long	55
	.long	56
	.long	57
	.long	58
	.long	59
	.long	60
	.long	61
	.long	62
	.long	63
	.long	64
	.long	65
	.long	66
	.long	67
	.long	68
	.long	69
	.long	70
	.long	71
	.long	72
	.long	73
	.long	74
	.long	75
	.long	76
	.long	77
	.long	78
	.long	79
	.long	80
	.long	81
	.long	82
	.long	83
	.long	84
	.long	85
	.long	86
	.long	87
	.long	88
	.long	89
	.long	90
	.long	91
	.long	92
	.long	93
	.long	94
	.long	95
	.long	96
	.long	97
	.long	98
	.long	99
	.long	100
	.long	101
	.long	102
	.long	103
	.long	104
	.long	105
	.long	106
	.long	107
	.long	108
	.long	109
	.long	110
	.long	111
	.long	112
	.long	113
	.long	114
	.long	115
	.long	116
	.long	117
	.long	118
	.long	119
	.long	120
	.long	121
	.long	122
	.long	123
	.long	124
	.long	125
	.long	126
	.long	127
	.long	128
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB3635:
	.cfi_startproc
	endbr64
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	movl	$128, %ecx
	movl	$1, %edx
	leaq	.LC0(%rip), %rsi
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	subq	$560, %rsp
	.cfi_def_cfa_offset 592
	movq	%fs:40, %rax
	movq	%rax, 552(%rsp)
	xorl	%eax, %eax
	leaq	40(%rsp), %rdi
	rep movsl
	leaq	8(%rsp), %rdi
	movl	$512, %esi
	call	cudaMallocManaged@PLT
	testl	%eax, %eax
	je	.L9
	leaq	.LC1(%rip), %rsi
	leaq	_ZSt4cerr(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%rax, %rdi
	call	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@PLT
	orl	$-1, %eax
	jmp	.L8
.L9:
	movq	8(%rsp), %rdi
	movl	$1, %ecx
	movl	$512, %edx
	leaq	40(%rsp), %rsi
	call	cudaMemcpy@PLT
	movl	$33554433, %edx
	xorl	%r9d, %r9d
	movabsq	$4294967297, %rdi
	salq	$7, %rdx
	movl	$512, %r8d
	movl	$1, %ecx
	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	.L11
	movq	8(%rsp), %rdi
	movl	$128, %esi
	call	_Z36__device_stub__Z15mergeSortKernelPiiPii
.L11:
	call	cudaGetLastError@PLT
	leaq	.LC2(%rip), %rsi
	movl	%eax, %ebx
	testl	%eax, %eax
	jne	.L20
	call	cudaDeviceSynchronize@PLT
	movl	%eax, %ebx
	testl	%eax, %eax
	je	.L14
	leaq	.LC4(%rip), %rsi
.L20:
	leaq	_ZSt4cerr(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	%ebx, %esi
	movq	%rax, %rdi
	call	_ZNSolsEi@PLT
	leaq	.LC3(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%rax, %rdi
	jmp	.L19
.L14:
	leaq	_ZSt4cout(%rip), %rbp
	leaq	.LC5(%rip), %rsi
	xorl	%ebx, %ebx
	movq	%rbp, %rdi
	leaq	.LC6(%rip), %r12
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
.L15:
	movq	8(%rsp), %rax
	movq	%rbp, %rdi
	movl	(%rax,%rbx), %esi
	addq	$4, %rbx
	call	_ZNSolsEi@PLT
	movq	%r12, %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	cmpq	$512, %rbx
	jne	.L15
	movq	%rbp, %rdi
.L19:
.L13:
	call	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@PLT
	movq	8(%rsp), %rdi
	call	cudaFree@PLT
	xorl	%eax, %eax
.L8:
	movq	552(%rsp), %rdx
	subq	%fs:40, %rdx
	je	.L16
	call	__stack_chk_fail@PLT
.L16:
	addq	$560, %rsp
	.cfi_def_cfa_offset 32
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3635:
	.size	main, .-main
	.section	.rodata.str1.1
.LC7:
	.string	"_Z15mergeSortKernelPii"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB3663:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	pushq	$0
	.cfi_def_cfa_offset 24
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC7(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z15mergeSortKernelPii(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	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	%rdx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE3663:
	.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 <iostream>
#include <hip/hip_runtime.h>

// Function to perform bottom-up merge sort
__global__ void mergeSortKernel(int *arr, int n) {
    extern __shared__ int sharedArr[];

    // Load data into shared memory
    int tid = threadIdx.x;
    sharedArr[tid] = arr[tid];
    __syncthreads();

    // Perform merge sort in shared memory
    for (int size = 1; size < blockDim.x; size *= 2) {
        int stride = size * 2;
        int l = tid * stride;
        int r = l + size;
        int rlim = min(r + size, n);

        // Load elements into registers
        int a = l + threadIdx.x;
        int b = r + threadIdx.x;

        int x = (a < n) ? sharedArr[a] : INT_MAX;
        int y = (b < rlim) ? sharedArr[b] : INT_MAX;

        // Merge step
        for (int i = 0; i < stride && l + i < n; i++) {
            if (x <= y) {
                sharedArr[l + i] = x;
                x = (a + size < n) ? sharedArr[a + size] : INT_MAX;
                a++;
            } else {
                sharedArr[l + i] = y;
                y = (b + size < rlim) ? sharedArr[b + size] : INT_MAX;
                b++;
            }
        }
        __syncthreads();
    }

    // Write sorted data back to global memory
    arr[tid] = sharedArr[tid];
}

int main() {
    const int n = 128;
    int *arr;
    int h_arr[n] = {12, 11, 13, 5, 6, 7, 8, 9, 10, 1, 2, 3, 4, 15, 14, 16, 17, 18, 19, 20,
                    21, 22, 23, 24, 25, 26, 27, 28, 29, 30, 31, 32, 33, 34, 35, 36, 37, 38, 39, 40,
                    41, 42, 43, 44, 45, 46, 47, 48, 49, 50, 51, 52, 53, 54, 55, 56, 57, 58, 59, 60,
                    61, 62, 63, 64, 65, 66, 67, 68, 69, 70, 71, 72, 73, 74, 75, 76, 77, 78, 79, 80,
                    81, 82, 83, 84, 85, 86, 87, 88, 89, 90, 91, 92, 93, 94, 95, 96, 97, 98, 99, 100,
                    101, 102, 103, 104, 105, 106, 107, 108, 109, 110, 111, 112, 113, 114, 115, 116, 117, 118, 119, 120,
                    121, 122, 123, 124, 125, 126, 127, 128};

    // Allocate Unified Memory -- accessible from CPU or GPU
    hipError_t err = hipMallocManaged(&arr, n * sizeof(int));
    if (err != hipSuccess) {
        std::cerr << "Failed to allocate Unified Memory" << std::endl;
        return -1;
    }

    // Copy data to device
    hipMemcpy(arr, h_arr, n * sizeof(int), hipMemcpyHostToDevice);

    // Launch kernel on 1D grid of 1D blocks
    mergeSortKernel<<<1, n, n * sizeof(int)>>>(arr, n);

    // Check for any errors launching the kernel
    err = hipGetLastError();
    if (err != hipSuccess) {
        std::cerr << "Failed to launch mergeSortKernel (error code " << err << ")" << std::endl;
        goto FreeMem;
    }

    // Wait for GPU to finish before accessing on host
    err = hipDeviceSynchronize();
    if (err != hipSuccess) {
        std::cerr << "Failed to synchronize (error code " << err << ")" << std::endl;
        goto FreeMem;
    }

    // Print sorted array
    std::cout << "Sorted array: ";
    for (int i = 0; i < n; i++) {
        std::cout << arr[i] << " ";
    }
    std::cout << std::endl;

FreeMem:
    hipFree(arr);

    return 0;
}
	.text
	.file	"cuda_code_067853.hip"
                                        # Start of file scope inline assembly
	.globl	_ZSt21ios_base_library_initv

                                        # End of file scope inline assembly
	.globl	_Z30__device_stub__mergeSortKernelPii # -- Begin function _Z30__device_stub__mergeSortKernelPii
	.type	_Z30__device_stub__mergeSortKernelPii,@function
_Z30__device_stub__mergeSortKernelPii:  # @_Z30__device_stub__mergeSortKernelPii
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$80, %rsp
	.cfi_def_cfa_offset 128
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	24(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	4(%rsp), %rcx
	movl	%esi, (%rcx)
	leaq	64(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%rbx)
	leaq	48(%rsp), %r14
	leaq	32(%rsp), %r15
	leaq	16(%rsp), %r12
	leaq	8(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z15mergeSortKernelPii, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$96, %rsp
	.cfi_adjust_cfa_offset -96
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end0:
	.size	_Z30__device_stub__mergeSortKernelPii, .Lfunc_end0-_Z30__device_stub__mergeSortKernelPii
	.cfi_endproc
                                        # -- End function
	.globl	main                            # -- Begin function main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%r14
	.cfi_def_cfa_offset 16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	subq	$536, %rsp                      # imm = 0x218
	.cfi_def_cfa_offset 560
	.cfi_offset %rbx, -24
	.cfi_offset %r14, -16
	leaq	16(%rsp), %rdi
	movl	$.L__const.main.h_arr, %esi
	movl	$512, %edx                      # imm = 0x200
	callq	memcpy@PLT
	leaq	8(%rsp), %rdi
	movl	$512, %esi                      # imm = 0x200
	movl	$1, %edx
	callq	hipMallocManaged
	testl	%eax, %eax
	je	.LBB1_2
# %bb.1:
	movl	$_ZSt4cerr, %ebx
	movl	$_ZSt4cerr, %edi
	movl	$.L.str, %esi
	movl	$33, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	_ZSt4cerr(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cerr, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	$-1, %eax
	jmp	.LBB1_12
.LBB1_2:
	movq	8(%rsp), %rdi
	leaq	16(%rsp), %rsi
	movl	$512, %edx                      # imm = 0x200
	movl	$1, %ecx
	callq	hipMemcpy
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	leaq	127(%rdi), %rdx
	movl	$512, %r8d                      # imm = 0x200
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB1_4
# %bb.3:
	movq	8(%rsp), %rdi
	movl	$128, %esi
	callq	_Z30__device_stub__mergeSortKernelPii
.LBB1_4:
	callq	hipGetLastError
	testl	%eax, %eax
	je	.LBB1_7
# %bb.5:
	movl	%eax, %ebx
	movl	$_ZSt4cerr, %edi
	movl	$.L.str.1, %esi
	movl	$45, %edx
	jmp	.LBB1_6
.LBB1_7:
	callq	hipDeviceSynchronize
	testl	%eax, %eax
	je	.LBB1_9
# %bb.8:
	movl	%eax, %ebx
	movl	$_ZSt4cerr, %edi
	movl	$.L.str.3, %esi
	movl	$34, %edx
.LBB1_6:                                # %.sink.split
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	$_ZSt4cerr, %edi
	movl	%ebx, %esi
	callq	_ZNSolsEi
	movq	%rax, %rbx
	movl	$.L.str.2, %esi
	movl	$1, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
.LBB1_11:                               # %.loopexit
	movq	(%rbx), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%rbx, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movq	8(%rsp), %rdi
	callq	hipFree
	xorl	%eax, %eax
.LBB1_12:
	addq	$536, %rsp                      # imm = 0x218
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	retq
.LBB1_9:
	.cfi_def_cfa_offset 560
	movl	$_ZSt4cout, %ebx
	movl	$_ZSt4cout, %edi
	movl	$.L.str.4, %esi
	movl	$14, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	xorl	%r14d, %r14d
.LBB1_10:                               # =>This Inner Loop Header: Depth=1
	movq	8(%rsp), %rax
	movl	(%rax,%r14,4), %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSolsEi
	movl	$.L.str.5, %esi
	movl	$1, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	incq	%r14
	cmpq	$128, %r14
	jne	.LBB1_10
	jmp	.LBB1_11
.Lfunc_end1:
	.size	main, .Lfunc_end1-main
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_ctor
	.type	__hip_module_ctor,@function
__hip_module_ctor:                      # @__hip_module_ctor
	.cfi_startproc
# %bb.0:
	subq	$40, %rsp
	.cfi_def_cfa_offset 48
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	jne	.LBB2_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rdi
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB2_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z15mergeSortKernelPii, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$40, %rsp
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end2:
	.size	__hip_module_ctor, .Lfunc_end2-__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	.LBB3_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
.LBB3_2:
	retq
.Lfunc_end3:
	.size	__hip_module_dtor, .Lfunc_end3-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_Z15mergeSortKernelPii,@object  # @_Z15mergeSortKernelPii
	.section	.rodata,"a",@progbits
	.globl	_Z15mergeSortKernelPii
	.p2align	3, 0x0
_Z15mergeSortKernelPii:
	.quad	_Z30__device_stub__mergeSortKernelPii
	.size	_Z15mergeSortKernelPii, 8

	.type	.L__const.main.h_arr,@object    # @__const.main.h_arr
	.p2align	4, 0x0
.L__const.main.h_arr:
	.long	12                              # 0xc
	.long	11                              # 0xb
	.long	13                              # 0xd
	.long	5                               # 0x5
	.long	6                               # 0x6
	.long	7                               # 0x7
	.long	8                               # 0x8
	.long	9                               # 0x9
	.long	10                              # 0xa
	.long	1                               # 0x1
	.long	2                               # 0x2
	.long	3                               # 0x3
	.long	4                               # 0x4
	.long	15                              # 0xf
	.long	14                              # 0xe
	.long	16                              # 0x10
	.long	17                              # 0x11
	.long	18                              # 0x12
	.long	19                              # 0x13
	.long	20                              # 0x14
	.long	21                              # 0x15
	.long	22                              # 0x16
	.long	23                              # 0x17
	.long	24                              # 0x18
	.long	25                              # 0x19
	.long	26                              # 0x1a
	.long	27                              # 0x1b
	.long	28                              # 0x1c
	.long	29                              # 0x1d
	.long	30                              # 0x1e
	.long	31                              # 0x1f
	.long	32                              # 0x20
	.long	33                              # 0x21
	.long	34                              # 0x22
	.long	35                              # 0x23
	.long	36                              # 0x24
	.long	37                              # 0x25
	.long	38                              # 0x26
	.long	39                              # 0x27
	.long	40                              # 0x28
	.long	41                              # 0x29
	.long	42                              # 0x2a
	.long	43                              # 0x2b
	.long	44                              # 0x2c
	.long	45                              # 0x2d
	.long	46                              # 0x2e
	.long	47                              # 0x2f
	.long	48                              # 0x30
	.long	49                              # 0x31
	.long	50                              # 0x32
	.long	51                              # 0x33
	.long	52                              # 0x34
	.long	53                              # 0x35
	.long	54                              # 0x36
	.long	55                              # 0x37
	.long	56                              # 0x38
	.long	57                              # 0x39
	.long	58                              # 0x3a
	.long	59                              # 0x3b
	.long	60                              # 0x3c
	.long	61                              # 0x3d
	.long	62                              # 0x3e
	.long	63                              # 0x3f
	.long	64                              # 0x40
	.long	65                              # 0x41
	.long	66                              # 0x42
	.long	67                              # 0x43
	.long	68                              # 0x44
	.long	69                              # 0x45
	.long	70                              # 0x46
	.long	71                              # 0x47
	.long	72                              # 0x48
	.long	73                              # 0x49
	.long	74                              # 0x4a
	.long	75                              # 0x4b
	.long	76                              # 0x4c
	.long	77                              # 0x4d
	.long	78                              # 0x4e
	.long	79                              # 0x4f
	.long	80                              # 0x50
	.long	81                              # 0x51
	.long	82                              # 0x52
	.long	83                              # 0x53
	.long	84                              # 0x54
	.long	85                              # 0x55
	.long	86                              # 0x56
	.long	87                              # 0x57
	.long	88                              # 0x58
	.long	89                              # 0x59
	.long	90                              # 0x5a
	.long	91                              # 0x5b
	.long	92                              # 0x5c
	.long	93                              # 0x5d
	.long	94                              # 0x5e
	.long	95                              # 0x5f
	.long	96                              # 0x60
	.long	97                              # 0x61
	.long	98                              # 0x62
	.long	99                              # 0x63
	.long	100                             # 0x64
	.long	101                             # 0x65
	.long	102                             # 0x66
	.long	103                             # 0x67
	.long	104                             # 0x68
	.long	105                             # 0x69
	.long	106                             # 0x6a
	.long	107                             # 0x6b
	.long	108                             # 0x6c
	.long	109                             # 0x6d
	.long	110                             # 0x6e
	.long	111                             # 0x6f
	.long	112                             # 0x70
	.long	113                             # 0x71
	.long	114                             # 0x72
	.long	115                             # 0x73
	.long	116                             # 0x74
	.long	117                             # 0x75
	.long	118                             # 0x76
	.long	119                             # 0x77
	.long	120                             # 0x78
	.long	121                             # 0x79
	.long	122                             # 0x7a
	.long	123                             # 0x7b
	.long	124                             # 0x7c
	.long	125                             # 0x7d
	.long	126                             # 0x7e
	.long	127                             # 0x7f
	.long	128                             # 0x80
	.size	.L__const.main.h_arr, 512

	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"Failed to allocate Unified Memory"
	.size	.L.str, 34

	.type	.L.str.1,@object                # @.str.1
.L.str.1:
	.asciz	"Failed to launch mergeSortKernel (error code "
	.size	.L.str.1, 46

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

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	"Failed to synchronize (error code "
	.size	.L.str.3, 35

	.type	.L.str.4,@object                # @.str.4
.L.str.4:
	.asciz	"Sorted array: "
	.size	.L.str.4, 15

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

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z15mergeSortKernelPii"
	.size	.L__unnamed_1, 23

	.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 _Z30__device_stub__mergeSortKernelPii
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z15mergeSortKernelPii
	.addrsig_sym _ZSt4cerr
	.addrsig_sym _ZSt4cout
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z15mergeSortKernelPii  ; -- Begin function _Z15mergeSortKernelPii
	.globl	_Z15mergeSortKernelPii
	.p2align	8
	.type	_Z15mergeSortKernelPii,@function
_Z15mergeSortKernelPii:                 ; @_Z15mergeSortKernelPii
; %bb.0:
	s_clause 0x1
	s_load_b64 s[4:5], s[0:1], 0x0
	s_load_b32 s3, s[0:1], 0x1c
	v_lshlrev_b32_e32 v1, 2, v0
	s_mov_b32 s2, 0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v5, 0, v1
	s_waitcnt lgkmcnt(0)
	global_load_b32 v3, v1, s[4:5]
	v_add_co_u32 v1, s4, s4, v1
	v_add_co_ci_u32_e64 v2, null, s5, 0, s4
	s_and_b32 s3, s3, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmp_lt_u32 s3, 2
	s_waitcnt vmcnt(0)
	ds_store_b32 v5, v3
	s_waitcnt lgkmcnt(0)
	s_barrier
	buffer_gl0_inv
	s_cbranch_scc1 .LBB0_19
; %bb.1:                                ; %.lr.ph.preheader
	s_load_b32 s0, s[0:1], 0x8
	v_lshlrev_b32_e32 v6, 3, v0
	s_mov_b32 s1, 1
.LBB0_2:                                ; %.lr.ph
                                        ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB0_7 Depth 2
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_4) | instid1(VALU_DEP_2)
	s_lshl_b32 s4, s1, 1
	v_bfrev_b32_e32 v9, -2
	v_mul_lo_u32 v7, s4, v0
	v_bfrev_b32_e32 v10, -2
	s_mov_b32 s5, exec_lo
	v_add_nc_u32_e32 v8, v7, v0
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_gt_i32_e64 s0, v8
	s_cbranch_execz .LBB0_4
; %bb.3:                                ;   in Loop: Header=BB0_2 Depth=1
	v_lshl_add_u32 v3, v7, 2, v5
	ds_load_b32 v10, v3
.LBB0_4:                                ;   in Loop: Header=BB0_2 Depth=1
	s_or_b32 exec_lo, exec_lo, s5
	v_add_nc_u32_e32 v3, s1, v7
	s_mov_b32 s5, exec_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v4, s1, v3
	v_add_nc_u32_e32 v12, v3, v0
	v_min_i32_e32 v11, s0, v4
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_lt_i32_e64 v12, v11
	s_cbranch_execz .LBB0_6
; %bb.5:                                ;   in Loop: Header=BB0_2 Depth=1
	v_lshl_add_u32 v3, v3, 2, v5
	ds_load_b32 v9, v3
.LBB0_6:                                ;   in Loop: Header=BB0_2 Depth=1
	s_or_b32 exec_lo, exec_lo, s5
	v_mad_u64_u32 v[3:4], null, v6, s1, s[2:3]
	s_max_u32 s5, s4, 1
	s_mov_b32 s6, 0
                                        ; implicit-def: $sgpr7
.LBB0_7:                                ;   Parent Loop BB0_2 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_or_b32 s7, s7, exec_lo
	s_mov_b32 s8, exec_lo
	v_cmpx_gt_i32_e64 s0, v7
	s_cbranch_execz .LBB0_17
; %bb.8:                                ;   in Loop: Header=BB0_7 Depth=2
	s_mov_b32 s9, exec_lo
	s_waitcnt lgkmcnt(0)
	v_cmpx_le_i32_e64 v10, v9
	s_xor_b32 s9, exec_lo, s9
	s_cbranch_execz .LBB0_12
; %bb.9:                                ;   in Loop: Header=BB0_7 Depth=2
	v_add_nc_u32_e32 v4, s1, v8
	ds_store_b32 v3, v10
	v_bfrev_b32_e32 v10, -2
	s_mov_b32 s10, exec_lo
	v_cmpx_gt_i32_e64 s0, v4
	s_cbranch_execz .LBB0_11
; %bb.10:                               ;   in Loop: Header=BB0_7 Depth=2
	v_lshl_add_u32 v4, v4, 2, 0
	ds_load_b32 v10, v4
.LBB0_11:                               ;   in Loop: Header=BB0_7 Depth=2
	s_or_b32 exec_lo, exec_lo, s10
	v_add_nc_u32_e32 v8, 1, v8
.LBB0_12:                               ; %Flow
                                        ;   in Loop: Header=BB0_7 Depth=2
	s_and_not1_saveexec_b32 s9, s9
	s_cbranch_execz .LBB0_16
; %bb.13:                               ;   in Loop: Header=BB0_7 Depth=2
	v_add_nc_u32_e32 v4, s1, v12
	ds_store_b32 v3, v9
	v_bfrev_b32_e32 v9, -2
	s_mov_b32 s10, exec_lo
	v_cmpx_lt_i32_e64 v4, v11
	s_cbranch_execz .LBB0_15
; %bb.14:                               ;   in Loop: Header=BB0_7 Depth=2
	v_lshl_add_u32 v4, v4, 2, 0
	ds_load_b32 v9, v4
.LBB0_15:                               ;   in Loop: Header=BB0_7 Depth=2
	s_or_b32 exec_lo, exec_lo, s10
	v_add_nc_u32_e32 v12, 1, v12
.LBB0_16:                               ; %Flow77
                                        ;   in Loop: Header=BB0_7 Depth=2
	s_or_b32 exec_lo, exec_lo, s9
	s_add_i32 s5, s5, -1
	v_add_nc_u32_e32 v3, 4, v3
	s_cmp_eq_u32 s5, 0
	v_add_nc_u32_e32 v7, 1, v7
	s_cselect_b32 s9, -1, 0
	s_and_not1_b32 s7, s7, exec_lo
	s_and_b32 s9, s9, exec_lo
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 s7, s7, s9
.LBB0_17:                               ; %Flow78
                                        ;   in Loop: Header=BB0_7 Depth=2
	s_or_b32 exec_lo, exec_lo, s8
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b32 s8, exec_lo, s7
	s_or_b32 s6, s8, s6
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s6
	s_cbranch_execnz .LBB0_7
; %bb.18:                               ; %.critedge
                                        ;   in Loop: Header=BB0_2 Depth=1
	s_or_b32 exec_lo, exec_lo, s6
	s_cmp_ge_u32 s4, s3
	s_mov_b32 s1, s4
	s_waitcnt lgkmcnt(0)
	s_barrier
	buffer_gl0_inv
	s_cbranch_scc0 .LBB0_2
.LBB0_19:                               ; %._crit_edge
	ds_load_b32 v0, v5
	s_waitcnt lgkmcnt(0)
	global_store_b32 v[1:2], v0, off
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z15mergeSortKernelPii
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 272
		.amdhsa_user_sgpr_count 15
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 13
		.amdhsa_next_free_sgpr 11
		.amdhsa_reserve_vcc 0
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end0:
	.size	_Z15mergeSortKernelPii, .Lfunc_end0-_Z15mergeSortKernelPii
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 572
; NumSgprs: 11
; NumVgprs: 13
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 1
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 11
; NumVGPRsForWavesPerEU: 13
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_
	.amdgpu_metadata
---
amdhsa.kernels:
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .offset:         8
        .size:           4
        .value_kind:     by_value
      - .offset:         16
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         20
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         24
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         28
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         30
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         32
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         34
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         36
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         38
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         56
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         64
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         80
        .size:           2
        .value_kind:     hidden_grid_dims
      - .offset:         136
        .size:           4
        .value_kind:     hidden_dynamic_lds_size
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 272
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z15mergeSortKernelPii
    .private_segment_fixed_size: 0
    .sgpr_count:     11
    .sgpr_spill_count: 0
    .symbol:         _Z15mergeSortKernelPii.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     13
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
