// quicksort.cu
// CUDA implementation of Quick Sort for a graph with 64 nodes
// Multi-GPU version optimized for compute throughput on Ampere architecture

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

#define NUM_NODES 64
#define BLOCK_SIZE 32

// Kernel to perform partitioning step of Quick Sort
__global__ void partition(int* d_data, int* d_pivot_indices, int pivot) {
    extern __shared__ int s_data[];

    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < NUM_NODES) {
        s_data[threadIdx.x] = d_data[idx];
    }
    __syncthreads();

    if (idx < NUM_NODES) {
        if (s_data[threadIdx.x] < pivot) {
            atomicAdd(d_pivot_indices, 1);
            d_data[idx] = s_data[threadIdx.x];
        }
    }
}

// Kernel to perform the swap step of Quick Sort
__global__ void swap(int* d_data, int* d_pivot_indices, int pivot) {
    extern __shared__ int s_data[];

    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < NUM_NODES) {
        s_data[threadIdx.x] = d_data[idx];
    }
    __syncthreads();

    if (idx < NUM_NODES) {
        if (s_data[threadIdx.x] >= pivot) {
            int pivot_index = atomicAdd(d_pivot_indices, 1);
            d_data[pivot_index] = s_data[threadIdx.x];
        }
    }
}

// Function to perform Quick Sort on a single GPU
void quickSortGPU(int* d_data, int* d_pivot_indices, int num_nodes) {
    int pivot = d_data[num_nodes / 2];
    cudaMemset(d_pivot_indices, 0, sizeof(int));

    // Perform partitioning
    partition<<<(num_nodes + BLOCK_SIZE - 1) / BLOCK_SIZE, BLOCK_SIZE, BLOCK_SIZE * sizeof(int)>>>(d_data, d_pivot_indices, pivot);
    cudaDeviceSynchronize();

    // Perform swapping
    swap<<<(num_nodes + BLOCK_SIZE - 1) / BLOCK_SIZE, BLOCK_SIZE, BLOCK_SIZE * sizeof(int)>>>(d_data, d_pivot_indices, pivot);
    cudaDeviceSynchronize();
}

// Main function to demonstrate the usage of the Quick Sort kernel
int main() {
    int h_data[NUM_NODES] = {54, 26, 93, 17, 77, 31, 44, 55, 20, 65, 81, 56, 89, 40, 30, 73,
                             24, 32, 23, 50, 80, 90, 33, 45, 60, 70, 85, 95, 10, 25, 35, 46,
                             57, 67, 78, 88, 98, 11, 21, 36, 47, 58, 68, 79, 87, 97, 12, 22,
                             37, 48, 59, 69, 82, 91, 13, 27, 38, 49, 61, 71, 83, 92, 14, 28};

    int* d_data;
    int* d_pivot_indices;

    // Allocate memory on the device
    cudaMalloc((void**)&d_data, NUM_NODES * sizeof(int));
    cudaMalloc((void**)&d_pivot_indices, sizeof(int));

    // Copy data from host to device
    cudaMemcpy(d_data, h_data, NUM_NODES * sizeof(int), cudaMemcpyHostToDevice);

    // Perform Quick Sort on the GPU
    quickSortGPU(d_data, d_pivot_indices, NUM_NODES);

    // Copy sorted data back to host
    cudaMemcpy(h_data, d_data, NUM_NODES * sizeof(int), cudaMemcpyDeviceToHost);

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

    // Free device memory
    cudaFree(d_data);
    cudaFree(d_pivot_indices);

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

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB3639:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE3639:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z31__device_stub__Z9partitionPiS_iPiS_i
	.type	_Z31__device_stub__Z9partitionPiS_iPiS_i, @function
_Z31__device_stub__Z9partitionPiS_iPiS_i:
.LFB3661:
	.cfi_startproc
	endbr64
	subq	$136, %rsp
	.cfi_def_cfa_offset 144
	movq	%rdi, 24(%rsp)
	leaq	40(%rsp), %rcx
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movl	%edx, 12(%rsp)
	leaq	32(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 120(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 96(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 112(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L2
	pushq	40(%rsp)
	.cfi_def_cfa_offset 152
	leaq	_Z9partitionPiS_i(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 160
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	112(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 152
	popq	%rdx
	.cfi_def_cfa_offset 144
.L2:
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L4
	call	__stack_chk_fail@PLT
.L4:
	addq	$136, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3661:
	.size	_Z31__device_stub__Z9partitionPiS_iPiS_i, .-_Z31__device_stub__Z9partitionPiS_iPiS_i
	.globl	_Z9partitionPiS_i
	.type	_Z9partitionPiS_i, @function
_Z9partitionPiS_i:
.LFB3662:
	.cfi_startproc
	endbr64
	jmp	_Z31__device_stub__Z9partitionPiS_iPiS_i
	.cfi_endproc
.LFE3662:
	.size	_Z9partitionPiS_i, .-_Z9partitionPiS_i
	.globl	_Z26__device_stub__Z4swapPiS_iPiS_i
	.type	_Z26__device_stub__Z4swapPiS_iPiS_i, @function
_Z26__device_stub__Z4swapPiS_iPiS_i:
.LFB3663:
	.cfi_startproc
	endbr64
	subq	$136, %rsp
	.cfi_def_cfa_offset 144
	movq	%rdi, 24(%rsp)
	leaq	40(%rsp), %rcx
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movl	%edx, 12(%rsp)
	leaq	32(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 120(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 96(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 112(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L8
	pushq	40(%rsp)
	.cfi_def_cfa_offset 152
	leaq	_Z4swapPiS_i(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 160
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	112(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 152
	popq	%rdx
	.cfi_def_cfa_offset 144
.L8:
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L10
	call	__stack_chk_fail@PLT
.L10:
	addq	$136, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3663:
	.size	_Z26__device_stub__Z4swapPiS_iPiS_i, .-_Z26__device_stub__Z4swapPiS_iPiS_i
	.globl	_Z4swapPiS_i
	.type	_Z4swapPiS_i, @function
_Z4swapPiS_i:
.LFB3664:
	.cfi_startproc
	endbr64
	jmp	_Z26__device_stub__Z4swapPiS_iPiS_i
	.cfi_endproc
.LFE3664:
	.size	_Z4swapPiS_i, .-_Z4swapPiS_i
	.globl	_Z12quickSortGPUPiS_i
	.type	_Z12quickSortGPUPiS_i, @function
_Z12quickSortGPUPiS_i:
.LFB3635:
	.cfi_startproc
	endbr64
	movl	%edx, %eax
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	movl	$2, %ecx
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	movq	%rsi, %r12
	xorl	%esi, %esi
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	movq	%rdi, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	movl	%edx, %ebx
	cltd
	idivl	%ecx
	subq	$40, %rsp
	.cfi_def_cfa_offset 80
	addl	$31, %ebx
	movl	$4, %edx
	cltq
	movl	(%rdi,%rax,4), %r13d
	movq	%r12, %rdi
	call	cudaMemset@PLT
	movl	%ebx, %eax
	movl	$32, %esi
	xorl	%r9d, %r9d
	cltd
	movl	$134217729, %ecx
	movl	$128, %r8d
	movl	$1, 28(%rsp)
	idivl	%esi
	salq	$5, %rcx
	movq	%rcx, 20(%rsp)
	movq	%rcx, %rdx
	movl	$1, %ecx
	movl	%eax, 8(%rsp)
	movl	%eax, %ebx
	movabsq	$4294967297, %rax
	movq	%rax, 12(%rsp)
	movq	8(%rsp), %rdi
	movl	16(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L14
	movl	%r13d, %edx
	movq	%r12, %rsi
	movq	%rbp, %rdi
	call	_Z31__device_stub__Z9partitionPiS_iPiS_i
.L14:
	call	cudaDeviceSynchronize@PLT
	movl	%ebx, 8(%rsp)
	movl	$134217729, %edx
	movabsq	$4294967297, %rax
	movq	%rax, 12(%rsp)
	movq	8(%rsp), %rdi
	salq	$5, %rdx
	xorl	%r9d, %r9d
	movl	16(%rsp), %esi
	movl	$128, %r8d
	movl	$1, %ecx
	movq	%rdx, 20(%rsp)
	movl	$1, 28(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L15
	movl	%r13d, %edx
	movq	%r12, %rsi
	movq	%rbp, %rdi
	call	_Z26__device_stub__Z4swapPiS_iPiS_i
.L15:
	addq	$40, %rsp
	.cfi_def_cfa_offset 40
	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
	jmp	cudaDeviceSynchronize@PLT
	.cfi_endproc
.LFE3635:
	.size	_Z12quickSortGPUPiS_i, .-_Z12quickSortGPUPiS_i
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC1:
	.string	"Sorted array: "
.LC2:
	.string	" "
	.section	.rodata
	.align 32
.LC0:
	.long	54
	.long	26
	.long	93
	.long	17
	.long	77
	.long	31
	.long	44
	.long	55
	.long	20
	.long	65
	.long	81
	.long	56
	.long	89
	.long	40
	.long	30
	.long	73
	.long	24
	.long	32
	.long	23
	.long	50
	.long	80
	.long	90
	.long	33
	.long	45
	.long	60
	.long	70
	.long	85
	.long	95
	.long	10
	.long	25
	.long	35
	.long	46
	.long	57
	.long	67
	.long	78
	.long	88
	.long	98
	.long	11
	.long	21
	.long	36
	.long	47
	.long	58
	.long	68
	.long	79
	.long	87
	.long	97
	.long	12
	.long	22
	.long	37
	.long	48
	.long	59
	.long	69
	.long	82
	.long	91
	.long	13
	.long	27
	.long	38
	.long	49
	.long	61
	.long	71
	.long	83
	.long	92
	.long	14
	.long	28
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB3636:
	.cfi_startproc
	endbr64
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	leaq	.LC0(%rip), %rsi
	movl	$64, %ecx
	leaq	.LC2(%rip), %r13
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	leaq	_ZSt4cout(%rip), %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	xorl	%ebx, %ebx
	subq	$296, %rsp
	.cfi_def_cfa_offset 336
	movq	%fs:40, %rax
	movq	%rax, 280(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rdi
	leaq	24(%rsp), %r12
	rep movsl
	leaq	8(%rsp), %rdi
	movl	$256, %esi
	call	cudaMalloc@PLT
	leaq	16(%rsp), %rdi
	movl	$4, %esi
	call	cudaMalloc@PLT
	movq	8(%rsp), %rdi
	movl	$1, %ecx
	movq	%r12, %rsi
	movl	$256, %edx
	call	cudaMemcpy@PLT
	movq	16(%rsp), %rsi
	movq	8(%rsp), %rdi
	movl	$64, %edx
	call	_Z12quickSortGPUPiS_i
	movq	8(%rsp), %rsi
	movl	$2, %ecx
	movq	%r12, %rdi
	movl	$256, %edx
	call	cudaMemcpy@PLT
	movl	$14, %edx
	leaq	.LC1(%rip), %rsi
	movq	%rbp, %rdi
	call	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l@PLT
.L18:
	movl	(%r12,%rbx,4), %esi
	movq	%rbp, %rdi
	incq	%rbx
	call	_ZNSolsEi@PLT
	movl	$1, %edx
	movq	%r13, %rsi
	movq	%rax, %rdi
	call	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l@PLT
	cmpq	$64, %rbx
	jne	.L18
	movq	%rbp, %rdi
	call	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@PLT
	movq	8(%rsp), %rdi
	call	cudaFree@PLT
	movq	16(%rsp), %rdi
	call	cudaFree@PLT
	movq	280(%rsp), %rax
	subq	%fs:40, %rax
	je	.L19
	call	__stack_chk_fail@PLT
.L19:
	addq	$296, %rsp
	.cfi_def_cfa_offset 40
	xorl	%eax, %eax
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%rbp
	.cfi_def_cfa_offset 24
	popq	%r12
	.cfi_def_cfa_offset 16
	popq	%r13
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3636:
	.size	main, .-main
	.section	.rodata.str1.1
.LC3:
	.string	"_Z4swapPiS_i"
.LC4:
	.string	"_Z9partitionPiS_i"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB3666:
	.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	.LC3(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z4swapPiS_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	.LC4(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z9partitionPiS_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
.LFE3666:
	.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:

// quicksort.cu
// CUDA implementation of Quick Sort for a graph with 64 nodes
// Multi-GPU version optimized for compute throughput on Ampere architecture

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


#define NUM_NODES 64
#define BLOCK_SIZE 32

// Kernel to perform partitioning step of Quick Sort
__global__ void partition(int* d_data, int* d_pivot_indices, int pivot) {
    extern __shared__ int s_data[];

    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < NUM_NODES) {
        s_data[threadIdx.x] = d_data[idx];
    }
    __syncthreads();

    if (idx < NUM_NODES) {
        if (s_data[threadIdx.x] < pivot) {
            atomicAdd(d_pivot_indices, 1);
            d_data[idx] = s_data[threadIdx.x];
        }
    }
}

// Kernel to perform the swap step of Quick Sort
__global__ void swap(int* d_data, int* d_pivot_indices, int pivot) {
    extern __shared__ int s_data[];

    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < NUM_NODES) {
        s_data[threadIdx.x] = d_data[idx];
    }
    __syncthreads();

    if (idx < NUM_NODES) {
        if (s_data[threadIdx.x] >= pivot) {
            int pivot_index = atomicAdd(d_pivot_indices, 1);
            d_data[pivot_index] = s_data[threadIdx.x];
        }
    }
}

// Function to perform Quick Sort on a single GPU
void quickSortGPU(int* d_data, int* d_pivot_indices, int num_nodes) {
    int pivot = d_data[num_nodes / 2];
    hipMemset(d_pivot_indices, 0, sizeof(int));

    // Perform partitioning
    partition<<<(num_nodes + BLOCK_SIZE - 1) / BLOCK_SIZE, BLOCK_SIZE, BLOCK_SIZE * sizeof(int)>>>(d_data, d_pivot_indices, pivot);
    hipDeviceSynchronize();

    // Perform swapping
    swap<<<(num_nodes + BLOCK_SIZE - 1) / BLOCK_SIZE, BLOCK_SIZE, BLOCK_SIZE * sizeof(int)>>>(d_data, d_pivot_indices, pivot);
    hipDeviceSynchronize();
}

// Main function to demonstrate the usage of the Quick Sort kernel
int main() {
    int h_data[NUM_NODES] = {54, 26, 93, 17, 77, 31, 44, 55, 20, 65, 81, 56, 89, 40, 30, 73,
                             24, 32, 23, 50, 80, 90, 33, 45, 60, 70, 85, 95, 10, 25, 35, 46,
                             57, 67, 78, 88, 98, 11, 21, 36, 47, 58, 68, 79, 87, 97, 12, 22,
                             37, 48, 59, 69, 82, 91, 13, 27, 38, 49, 61, 71, 83, 92, 14, 28};

    int* d_data;
    int* d_pivot_indices;

    // Allocate memory on the device
    hipMalloc((void**)&d_data, NUM_NODES * sizeof(int));
    hipMalloc((void**)&d_pivot_indices, sizeof(int));

    // Copy data from host to device
    hipMemcpy(d_data, h_data, NUM_NODES * sizeof(int), hipMemcpyHostToDevice);

    // Perform Quick Sort on the GPU
    quickSortGPU(d_data, d_pivot_indices, NUM_NODES);

    // Copy sorted data back to host
    hipMemcpy(h_data, d_data, NUM_NODES * sizeof(int), hipMemcpyDeviceToHost);

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

    // Free device memory
    hipFree(d_data);
    hipFree(d_pivot_indices);

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

                                        # End of file scope inline assembly
	.globl	_Z24__device_stub__partitionPiS_i # -- Begin function _Z24__device_stub__partitionPiS_i
	.type	_Z24__device_stub__partitionPiS_i,@function
_Z24__device_stub__partitionPiS_i:      # @_Z24__device_stub__partitionPiS_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	$_Z9partitionPiS_i, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$128, %rsp
	.cfi_adjust_cfa_offset -128
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end0:
	.size	_Z24__device_stub__partitionPiS_i, .Lfunc_end0-_Z24__device_stub__partitionPiS_i
	.cfi_endproc
                                        # -- End function
	.globl	_Z19__device_stub__swapPiS_i    # -- Begin function _Z19__device_stub__swapPiS_i
	.type	_Z19__device_stub__swapPiS_i,@function
_Z19__device_stub__swapPiS_i:           # @_Z19__device_stub__swapPiS_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	$_Z4swapPiS_i, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$128, %rsp
	.cfi_adjust_cfa_offset -128
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end1:
	.size	_Z19__device_stub__swapPiS_i, .Lfunc_end1-_Z19__device_stub__swapPiS_i
	.cfi_endproc
                                        # -- End function
	.globl	_Z12quickSortGPUPiS_i           # -- Begin function _Z12quickSortGPUPiS_i
	.type	_Z12quickSortGPUPiS_i,@function
_Z12quickSortGPUPiS_i:                  # @_Z12quickSortGPUPiS_i
	.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	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movl	%edx, %r12d
	movq	%rsi, %rbx
	movq	%rdi, %r14
	movabsq	$4294967328, %r15               # imm = 0x100000020
	movl	%edx, %eax
	shrl	$31, %eax
	addl	%edx, %eax
	sarl	%eax
	cltq
	movl	(%rdi,%rax,4), %ebp
	movl	$4, %edx
	movq	%rsi, %rdi
	xorl	%esi, %esi
	callq	hipMemset
	leal	31(%r12), %eax
	addl	$62, %r12d
	testl	%eax, %eax
	cmovnsl	%eax, %r12d
	sarl	$5, %r12d
	addq	%r15, %r12
	addq	$-32, %r12
	movl	$128, %r8d
	movq	%r12, %rdi
	movl	$1, %esi
	movq	%r15, %rdx
	movl	$1, %ecx
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB2_2
# %bb.1:
	movq	%r14, %rdi
	movq	%rbx, %rsi
	movl	%ebp, %edx
	callq	_Z24__device_stub__partitionPiS_i
.LBB2_2:
	callq	hipDeviceSynchronize
	movl	$128, %r8d
	movq	%r12, %rdi
	movl	$1, %esi
	movq	%r15, %rdx
	movl	$1, %ecx
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB2_4
# %bb.3:
	movq	%r14, %rdi
	movq	%rbx, %rsi
	movl	%ebp, %edx
	callq	_Z19__device_stub__swapPiS_i
.LBB2_4:
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.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	hipDeviceSynchronize            # TAILCALL
.Lfunc_end2:
	.size	_Z12quickSortGPUPiS_i, .Lfunc_end2-_Z12quickSortGPUPiS_i
	.cfi_endproc
                                        # -- End function
	.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	$272, %rsp                      # imm = 0x110
	.cfi_def_cfa_offset 304
	.cfi_offset %rbx, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	16(%rsp), %rbx
	movl	$.L__const.main.h_data, %esi
	movl	$256, %edx                      # imm = 0x100
	movq	%rbx, %rdi
	callq	memcpy@PLT
	leaq	8(%rsp), %r14
	movl	$256, %esi                      # imm = 0x100
	movq	%r14, %rdi
	callq	hipMalloc
	movq	%rsp, %r15
	movl	$4, %esi
	movq	%r15, %rdi
	callq	hipMalloc
	movq	(%r14), %rdi
	movl	$256, %edx                      # imm = 0x100
	movq	%rbx, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movq	(%r14), %rdi
	movq	(%r15), %rsi
	movl	$64, %edx
	callq	_Z12quickSortGPUPiS_i
	movq	(%r14), %rsi
	movl	$256, %edx                      # imm = 0x100
	movq	%rbx, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movl	$_ZSt4cout, %edi
	movl	$.L.str, %esi
	movl	$14, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	xorl	%ebx, %ebx
.LBB3_1:                                # =>This Inner Loop Header: Depth=1
	movl	16(%rsp,%rbx,4), %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSolsEi
	movl	$.L.str.1, %esi
	movl	$1, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	incq	%rbx
	cmpq	$64, %rbx
	jne	.LBB3_1
# %bb.2:
	movq	_ZSt4cout(%rip), %rax
	movl	$_ZSt4cout, %ecx
	movq	-24(%rax), %rdi
	addq	%rcx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movq	8(%rsp), %rdi
	callq	hipFree
	movq	(%rsp), %rdi
	callq	hipFree
	xorl	%eax, %eax
	addq	$272, %rsp                      # imm = 0x110
	.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_end3:
	.size	main, .Lfunc_end3-main
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_ctor
	.type	__hip_module_ctor,@function
__hip_module_ctor:                      # @__hip_module_ctor
	.cfi_startproc
# %bb.0:
	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	.LBB4_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB4_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z9partitionPiS_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	$_Z4swapPiS_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_end4:
	.size	__hip_module_ctor, .Lfunc_end4-__hip_module_ctor
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_dtor
	.type	__hip_module_dtor,@function
__hip_module_dtor:                      # @__hip_module_dtor
	.cfi_startproc
# %bb.0:
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	je	.LBB5_2
# %bb.1:
	pushq	%rax
	.cfi_def_cfa_offset 16
	callq	__hipUnregisterFatBinary
	movq	$0, __hip_gpubin_handle(%rip)
	addq	$8, %rsp
	.cfi_def_cfa_offset 8
.LBB5_2:
	retq
.Lfunc_end5:
	.size	__hip_module_dtor, .Lfunc_end5-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_Z9partitionPiS_i,@object       # @_Z9partitionPiS_i
	.section	.rodata,"a",@progbits
	.globl	_Z9partitionPiS_i
	.p2align	3, 0x0
_Z9partitionPiS_i:
	.quad	_Z24__device_stub__partitionPiS_i
	.size	_Z9partitionPiS_i, 8

	.type	_Z4swapPiS_i,@object            # @_Z4swapPiS_i
	.globl	_Z4swapPiS_i
	.p2align	3, 0x0
_Z4swapPiS_i:
	.quad	_Z19__device_stub__swapPiS_i
	.size	_Z4swapPiS_i, 8

	.type	.L__const.main.h_data,@object   # @__const.main.h_data
	.p2align	4, 0x0
.L__const.main.h_data:
	.long	54                              # 0x36
	.long	26                              # 0x1a
	.long	93                              # 0x5d
	.long	17                              # 0x11
	.long	77                              # 0x4d
	.long	31                              # 0x1f
	.long	44                              # 0x2c
	.long	55                              # 0x37
	.long	20                              # 0x14
	.long	65                              # 0x41
	.long	81                              # 0x51
	.long	56                              # 0x38
	.long	89                              # 0x59
	.long	40                              # 0x28
	.long	30                              # 0x1e
	.long	73                              # 0x49
	.long	24                              # 0x18
	.long	32                              # 0x20
	.long	23                              # 0x17
	.long	50                              # 0x32
	.long	80                              # 0x50
	.long	90                              # 0x5a
	.long	33                              # 0x21
	.long	45                              # 0x2d
	.long	60                              # 0x3c
	.long	70                              # 0x46
	.long	85                              # 0x55
	.long	95                              # 0x5f
	.long	10                              # 0xa
	.long	25                              # 0x19
	.long	35                              # 0x23
	.long	46                              # 0x2e
	.long	57                              # 0x39
	.long	67                              # 0x43
	.long	78                              # 0x4e
	.long	88                              # 0x58
	.long	98                              # 0x62
	.long	11                              # 0xb
	.long	21                              # 0x15
	.long	36                              # 0x24
	.long	47                              # 0x2f
	.long	58                              # 0x3a
	.long	68                              # 0x44
	.long	79                              # 0x4f
	.long	87                              # 0x57
	.long	97                              # 0x61
	.long	12                              # 0xc
	.long	22                              # 0x16
	.long	37                              # 0x25
	.long	48                              # 0x30
	.long	59                              # 0x3b
	.long	69                              # 0x45
	.long	82                              # 0x52
	.long	91                              # 0x5b
	.long	13                              # 0xd
	.long	27                              # 0x1b
	.long	38                              # 0x26
	.long	49                              # 0x31
	.long	61                              # 0x3d
	.long	71                              # 0x47
	.long	83                              # 0x53
	.long	92                              # 0x5c
	.long	14                              # 0xe
	.long	28                              # 0x1c
	.size	.L__const.main.h_data, 256

	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"Sorted array: "
	.size	.L.str, 15

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

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

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"_Z4swapPiS_i"
	.size	.L__unnamed_2, 13

	.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__partitionPiS_i
	.addrsig_sym _Z19__device_stub__swapPiS_i
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z9partitionPiS_i
	.addrsig_sym _Z4swapPiS_i
	.addrsig_sym _ZSt4cout
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z9partitionPiS_i       ; -- Begin function _Z9partitionPiS_i
	.globl	_Z9partitionPiS_i
	.p2align	8
	.type	_Z9partitionPiS_i,@function
_Z9partitionPiS_i:                      ; @_Z9partitionPiS_i
; %bb.0:
	s_clause 0x1
	s_load_b32 s2, s[0:1], 0x24
	s_load_b128 s[4:7], s[0:1], 0x0
	s_waitcnt lgkmcnt(0)
	s_and_b32 s2, s2, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_mad_u64_u32 v[1:2], null, s15, s2, v[0:1]
	v_lshl_add_u32 v0, v0, 2, 0
	v_cmp_gt_i32_e32 vcc_lo, 64, v1
	s_and_saveexec_b32 s3, vcc_lo
	s_cbranch_execz .LBB0_2
; %bb.1:
	v_ashrrev_i32_e32 v2, 31, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[2:3], 2, v[1:2]
	v_add_co_u32 v2, s2, s4, v2
	s_delay_alu instid0(VALU_DEP_1)
	v_add_co_ci_u32_e64 v3, s2, s5, v3, s2
	global_load_b32 v2, v[2:3], off
	s_waitcnt vmcnt(0)
	ds_store_b32 v0, v2
.LBB0_2:
	s_or_b32 exec_lo, exec_lo, s3
	s_waitcnt lgkmcnt(0)
	s_barrier
	buffer_gl0_inv
	s_and_saveexec_b32 s2, vcc_lo
	s_cbranch_execz .LBB0_7
; %bb.3:
	s_load_b32 s0, s[0:1], 0x10
	ds_load_b32 v0, v0
	s_waitcnt lgkmcnt(0)
	v_cmp_gt_i32_e32 vcc_lo, s0, v0
	s_and_b32 exec_lo, exec_lo, vcc_lo
	s_cbranch_execz .LBB0_7
; %bb.4:
	s_mov_b32 s1, exec_lo
	s_mov_b32 s0, exec_lo
	v_mbcnt_lo_u32_b32 v2, s1, 0
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_eq_u32_e32 0, v2
	s_cbranch_execz .LBB0_6
; %bb.5:
	s_bcnt1_i32_b32 s1, s1
	s_delay_alu instid0(SALU_CYCLE_1)
	v_dual_mov_b32 v2, 0 :: v_dual_mov_b32 v3, s1
	global_atomic_add_u32 v2, v3, s[6:7]
.LBB0_6:
	s_or_b32 exec_lo, exec_lo, s0
	v_ashrrev_i32_e32 v2, 31, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[1:2], 2, v[1:2]
	v_add_co_u32 v1, vcc_lo, s4, v1
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v2, vcc_lo, s5, v2, vcc_lo
	global_store_b32 v[1:2], v0, off
.LBB0_7:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z9partitionPiS_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 4
		.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	_Z9partitionPiS_i, .Lfunc_end0-_Z9partitionPiS_i
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 288
; NumSgprs: 18
; NumVgprs: 4
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 4
; 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	_Z4swapPiS_i            ; -- Begin function _Z4swapPiS_i
	.globl	_Z4swapPiS_i
	.p2align	8
	.type	_Z4swapPiS_i,@function
_Z4swapPiS_i:                           ; @_Z4swapPiS_i
; %bb.0:
	s_clause 0x1
	s_load_b32 s2, s[0:1], 0x24
	s_load_b128 s[4:7], s[0:1], 0x0
	s_waitcnt lgkmcnt(0)
	s_and_b32 s2, s2, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_mad_u64_u32 v[1:2], null, s15, s2, v[0:1]
	v_lshl_add_u32 v0, v0, 2, 0
	v_cmp_gt_i32_e32 vcc_lo, 64, v1
	s_and_saveexec_b32 s3, vcc_lo
	s_cbranch_execz .LBB1_2
; %bb.1:
	v_ashrrev_i32_e32 v2, 31, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[1:2], 2, v[1:2]
	v_add_co_u32 v1, s2, s4, v1
	s_delay_alu instid0(VALU_DEP_1)
	v_add_co_ci_u32_e64 v2, s2, s5, v2, s2
	global_load_b32 v1, v[1:2], off
	s_waitcnt vmcnt(0)
	ds_store_b32 v0, v1
.LBB1_2:
	s_or_b32 exec_lo, exec_lo, s3
	s_waitcnt lgkmcnt(0)
	s_barrier
	buffer_gl0_inv
	s_and_saveexec_b32 s2, vcc_lo
	s_cbranch_execz .LBB1_7
; %bb.3:
	s_load_b32 s0, s[0:1], 0x10
	ds_load_b32 v0, v0
	s_waitcnt lgkmcnt(0)
	v_cmp_le_i32_e32 vcc_lo, s0, v0
	s_and_b32 exec_lo, exec_lo, vcc_lo
	s_cbranch_execz .LBB1_7
; %bb.4:
	s_mov_b32 s1, exec_lo
	s_mov_b32 s0, exec_lo
	v_mbcnt_lo_u32_b32 v1, s1, 0
                                        ; implicit-def: $vgpr2
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_eq_u32_e32 0, v1
	s_cbranch_execz .LBB1_6
; %bb.5:
	s_bcnt1_i32_b32 s1, s1
	s_delay_alu instid0(SALU_CYCLE_1)
	v_dual_mov_b32 v2, 0 :: v_dual_mov_b32 v3, s1
	global_atomic_add_u32 v2, v2, v3, s[6:7] glc
.LBB1_6:
	s_or_b32 exec_lo, exec_lo, s0
	s_waitcnt vmcnt(0)
	v_readfirstlane_b32 s0, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v1, s0, v1
	v_ashrrev_i32_e32 v2, 31, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[1:2], 2, v[1:2]
	v_add_co_u32 v1, vcc_lo, s4, v1
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v2, vcc_lo, s5, v2, vcc_lo
	global_store_b32 v[1:2], v0, off
.LBB1_7:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z4swapPiS_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 4
		.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_end1:
	.size	_Z4swapPiS_i, .Lfunc_end1-_Z4swapPiS_i
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 304
; NumSgprs: 18
; NumVgprs: 4
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 4
; 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
      - .offset:         144
        .size:           4
        .value_kind:     hidden_dynamic_lds_size
    .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:           _Z9partitionPiS_i
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z9partitionPiS_i.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     4
    .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
      - .offset:         144
        .size:           4
        .value_kind:     hidden_dynamic_lds_size
    .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:           _Z4swapPiS_i
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z4swapPiS_i.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     4
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
