// bucket_sort_graph.cu
// CUDA program to perform bucket sort on a graph with 32 nodes
// Optimized for bank conflict avoidance on Ada Lovelace architecture

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

#define NUM_NODES 32

// Kernel to perform bucket sort
__global__ void bucketSortKernel(int *d_values, int *d_buckets, int *d_output) {
    extern __shared__ int s_values[];

    // Load data into shared memory
    int idx = threadIdx.x;
    s_values[idx] = d_values[idx];
    __syncthreads();

    // Initialize bucket counters
    __shared__ int s_bucket_counts[NUM_NODES];
    if (idx < NUM_NODES) {
        s_bucket_counts[idx] = 0;
    }
    __syncthreads();

    // Count elements in each bucket
    atomicAdd(&s_bucket_counts[s_values[idx]], 1);
    __syncthreads();

    // Compute prefix sums for bucket offsets
    if (idx < NUM_NODES) {
        int offset = 0;
        for (int j = 0; j < idx; ++j) {
            offset += s_bucket_counts[j];
        }
        s_bucket_counts[idx] = offset;
    }
    __syncthreads();

    // Place elements into their respective buckets
    int bucket = s_values[idx];
    int pos = atomicAdd(&s_bucket_counts[bucket], 1);
    d_output[pos] = s_values[idx];
}

int main() {
    // Host data
    int h_values[NUM_NODES] = {27, 25, 48, 37, 12, 9, 82, 10, 58, 72, 21, 39, 29, 47, 63, 79, 81, 33, 17, 59, 73, 15, 23, 41, 57, 71, 83, 35, 19, 51, 67, 85};
    int h_buckets[NUM_NODES];
    int h_output[NUM_NODES];

    // Device data
    int *d_values, *d_buckets, *d_output;

    // Allocate device memory
    cudaError_t err = cudaMalloc((void **)&d_values, NUM_NODES * sizeof(int));
    if (err != cudaSuccess) {
        std::cerr << "Failed to allocate device memory for d_values" << std::endl;
        return -1;
    }
    err = cudaMalloc((void **)&d_buckets, NUM_NODES * sizeof(int));
    if (err != cudaSuccess) {
        std::cerr << "Failed to allocate device memory for d_buckets" << std::endl;
        cudaFree(d_values);
        return -1;
    }
    err = cudaMalloc((void **)&d_output, NUM_NODES * sizeof(int));
    if (err != cudaSuccess) {
        std::cerr << "Failed to allocate device memory for d_output" << std::endl;
        cudaFree(d_values);
        cudaFree(d_buckets);
        return -1;
    }

    // Copy data from host to device
    err = cudaMemcpy(d_values, h_values, NUM_NODES * sizeof(int), cudaMemcpyHostToDevice);
    if (err != cudaSuccess) {
        std::cerr << "Failed to copy data from host to device" << std::endl;
        cudaFree(d_values);
        cudaFree(d_buckets);
        cudaFree(d_output);
        return -1;
    }

    // Launch kernel
    bucketSortKernel<<<1, NUM_NODES, NUM_NODES * sizeof(int)>>>(d_values, d_buckets, d_output);
    err = cudaGetLastError();
    if (err != cudaSuccess) {
        std::cerr << "Kernel launch failed: " << cudaGetErrorString(err) << std::endl;
        cudaFree(d_values);
        cudaFree(d_buckets);
        cudaFree(d_output);
        return -1;
    }

    // Copy result back to host
    err = cudaMemcpy(h_output, d_output, NUM_NODES * sizeof(int), cudaMemcpyDeviceToHost);
    if (err != cudaSuccess) {
        std::cerr << "Failed to copy data from device to host" << std::endl;
        cudaFree(d_values);
        cudaFree(d_buckets);
        cudaFree(d_output);
        return -1;
    }

    // Free device memory
    cudaFree(d_values);
    cudaFree(d_buckets);
    cudaFree(d_output);

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

    return 0;
}
	.file	"tmpxft_0014f515_00000000-6_cuda_code_004933.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
	.align 2
	.type	_ZNSolsEPFRSoS_E.isra.0, @function
_ZNSolsEPFRSoS_E.isra.0:
.LFB4292:
	.cfi_startproc
	jmp	*%rsi
	.cfi_endproc
.LFE4292:
	.size	_ZNSolsEPFRSoS_E.isra.0, .-_ZNSolsEPFRSoS_E.isra.0
	.globl	_Z40__device_stub__Z16bucketSortKernelPiS_S_PiS_S_
	.type	_Z40__device_stub__Z16bucketSortKernelPiS_S_PiS_S_, @function
_Z40__device_stub__Z16bucketSortKernelPiS_S_PiS_S_:
.LFB3660:
	.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
	movq	%rdx, 8(%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	8(%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	.L3
	pushq	40(%rsp)
	.cfi_def_cfa_offset 152
	leaq	_Z16bucketSortKernelPiS_S_(%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
.L3:
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L5
	call	__stack_chk_fail@PLT
.L5:
	addq	$136, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3660:
	.size	_Z40__device_stub__Z16bucketSortKernelPiS_S_PiS_S_, .-_Z40__device_stub__Z16bucketSortKernelPiS_S_PiS_S_
	.globl	_Z16bucketSortKernelPiS_S_
	.type	_Z16bucketSortKernelPiS_S_, @function
_Z16bucketSortKernelPiS_S_:
.LFB3661:
	.cfi_startproc
	endbr64
	jmp	_Z40__device_stub__Z16bucketSortKernelPiS_S_PiS_S_
	.cfi_endproc
.LFE3661:
	.size	_Z16bucketSortKernelPiS_S_, .-_Z16bucketSortKernelPiS_S_
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC1:
	.string	"Failed to allocate device memory for d_values"
.LC2:
	.string	"Failed to allocate device memory for d_buckets"
.LC3:
	.string	"Failed to allocate device memory for d_output"
.LC4:
	.string	"Failed to copy data from host to device"
.LC5:
	.string	"Kernel launch failed: "
.LC6:
	.string	"Failed to copy data from device to host"
.LC7:
	.string	"Sorted array: "
.LC8:
	.string	" "
	.section	.rodata
	.align 32
.LC0:
	.long	27
	.long	25
	.long	48
	.long	37
	.long	12
	.long	9
	.long	82
	.long	10
	.long	58
	.long	72
	.long	21
	.long	39
	.long	29
	.long	47
	.long	63
	.long	79
	.long	81
	.long	33
	.long	17
	.long	59
	.long	73
	.long	15
	.long	23
	.long	41
	.long	57
	.long	71
	.long	83
	.long	35
	.long	19
	.long	51
	.long	67
	.long	85
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB3635:
	.cfi_startproc
	endbr64
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	leaq	.LC0(%rip), %rsi
	movl	$32, %ecx
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	subq	$328, %rsp
	.cfi_def_cfa_offset 368
	movq	%fs:40, %rax
	movq	%rax, 312(%rsp)
	xorl	%eax, %eax
	leaq	56(%rsp), %rdi
	rep movsl
	leaq	8(%rsp), %rdi
	movl	$128, %esi
	call	cudaMalloc@PLT
	testl	%eax, %eax
	je	.L10
	leaq	.LC1(%rip), %rsi
	leaq	_ZSt4cerr(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	jmp	.L11
.L10:
	leaq	16(%rsp), %rdi
	movl	$128, %esi
	call	cudaMalloc@PLT
	testl	%eax, %eax
	je	.L12
	leaq	.LC2(%rip), %rsi
	leaq	_ZSt4cerr(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	movq	8(%rsp), %rdi
.L23:
	call	cudaFree@PLT
.L11:
	orl	$-1, %eax
	jmp	.L9
.L12:
	leaq	24(%rsp), %rdi
	movl	$128, %esi
	call	cudaMalloc@PLT
	testl	%eax, %eax
	je	.L14
	leaq	.LC3(%rip), %rsi
	leaq	_ZSt4cerr(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	movq	8(%rsp), %rdi
	call	cudaFree@PLT
	movq	16(%rsp), %rdi
	jmp	.L23
.L14:
	movq	8(%rsp), %rdi
	leaq	56(%rsp), %rsi
	movl	$1, %ecx
	movl	$128, %edx
	call	cudaMemcpy@PLT
	leaq	.LC4(%rip), %rsi
	testl	%eax, %eax
	jne	.L25
	movl	$134217729, %edx
	xorl	%r9d, %r9d
	movl	$1, %ecx
	movabsq	$4294967297, %rdi
	salq	$5, %rdx
	movl	$128, %r8d
	movl	$1, %esi
	movl	$1, 52(%rsp)
	movq	%rdx, 44(%rsp)
	movq	%rdi, 32(%rsp)
	movl	$1, 40(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L16
	movq	24(%rsp), %rdx
	movq	16(%rsp), %rsi
	movq	8(%rsp), %rdi
	call	_Z40__device_stub__Z16bucketSortKernelPiS_S_PiS_S_
.L16:
	call	cudaGetLastError@PLT
	movl	%eax, %ebx
	testl	%eax, %eax
	je	.L17
	leaq	.LC5(%rip), %rsi
	leaq	_ZSt4cerr(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	%ebx, %edi
	movq	%rax, %rbp
	call	cudaGetErrorString@PLT
	movq	%rbp, %rdi
	movq	%rax, %rsi
	jmp	.L24
.L17:
	movq	24(%rsp), %rsi
	movl	$2, %ecx
	movl	$128, %edx
	leaq	184(%rsp), %r12
	movq	%r12, %rdi
	call	cudaMemcpy@PLT
	testl	%eax, %eax
	je	.L18
	leaq	.LC6(%rip), %rsi
.L25:
	leaq	_ZSt4cerr(%rip), %rdi
.L24:
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	movq	8(%rsp), %rdi
	call	cudaFree@PLT
	movq	16(%rsp), %rdi
	call	cudaFree@PLT
	movq	24(%rsp), %rdi
	jmp	.L23
.L18:
	movq	8(%rsp), %rdi
	leaq	_ZSt4cout(%rip), %rbp
	xorl	%ebx, %ebx
	leaq	.LC8(%rip), %r13
	call	cudaFree@PLT
	movq	16(%rsp), %rdi
	call	cudaFree@PLT
	movq	24(%rsp), %rdi
	call	cudaFree@PLT
	leaq	.LC7(%rip), %rsi
	movq	%rbp, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
.L19:
	movl	(%r12,%rbx,4), %esi
	movq	%rbp, %rdi
	incq	%rbx
	call	_ZNSolsEi@PLT
	movq	%r13, %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	cmpq	$32, %rbx
	jne	.L19
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rbp, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	xorl	%eax, %eax
.L9:
	movq	312(%rsp), %rdx
	subq	%fs:40, %rdx
	je	.L20
	call	__stack_chk_fail@PLT
.L20:
	addq	$328, %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
	ret
	.cfi_endproc
.LFE3635:
	.size	main, .-main
	.section	.rodata.str1.1
.LC9:
	.string	"_Z16bucketSortKernelPiS_S_"
	.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	.LC9(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z16bucketSortKernelPiS_S_(%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:

// bucket_sort_graph.cu
// CUDA program to perform bucket sort on a graph with 32 nodes
// Optimized for bank conflict avoidance on Ada Lovelace architecture

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

#define NUM_NODES 32

// Kernel to perform bucket sort
__global__ void bucketSortKernel(int *d_values, int *d_buckets, int *d_output) {
    extern __shared__ int s_values[];

    // Load data into shared memory
    int idx = threadIdx.x;
    s_values[idx] = d_values[idx];
    __syncthreads();

    // Initialize bucket counters
    __shared__ int s_bucket_counts[NUM_NODES];
    if (idx < NUM_NODES) {
        s_bucket_counts[idx] = 0;
    }
    __syncthreads();

    // Count elements in each bucket
    atomicAdd(&s_bucket_counts[s_values[idx]], 1);
    __syncthreads();

    // Compute prefix sums for bucket offsets
    if (idx < NUM_NODES) {
        int offset = 0;
        for (int j = 0; j < idx; ++j) {
            offset += s_bucket_counts[j];
        }
        s_bucket_counts[idx] = offset;
    }
    __syncthreads();

    // Place elements into their respective buckets
    int bucket = s_values[idx];
    int pos = atomicAdd(&s_bucket_counts[bucket], 1);
    d_output[pos] = s_values[idx];
}

int main() {
    // Host data
    int h_values[NUM_NODES] = {27, 25, 48, 37, 12, 9, 82, 10, 58, 72, 21, 39, 29, 47, 63, 79, 81, 33, 17, 59, 73, 15, 23, 41, 57, 71, 83, 35, 19, 51, 67, 85};
    int h_buckets[NUM_NODES];
    int h_output[NUM_NODES];

    // Device data
    int *d_values, *d_buckets, *d_output;

    // Allocate device memory
    hipError_t err = hipMalloc((void **)&d_values, NUM_NODES * sizeof(int));
    if (err != hipSuccess) {
        std::cerr << "Failed to allocate device memory for d_values" << std::endl;
        return -1;
    }
    err = hipMalloc((void **)&d_buckets, NUM_NODES * sizeof(int));
    if (err != hipSuccess) {
        std::cerr << "Failed to allocate device memory for d_buckets" << std::endl;
        hipFree(d_values);
        return -1;
    }
    err = hipMalloc((void **)&d_output, NUM_NODES * sizeof(int));
    if (err != hipSuccess) {
        std::cerr << "Failed to allocate device memory for d_output" << std::endl;
        hipFree(d_values);
        hipFree(d_buckets);
        return -1;
    }

    // Copy data from host to device
    err = hipMemcpy(d_values, h_values, NUM_NODES * sizeof(int), hipMemcpyHostToDevice);
    if (err != hipSuccess) {
        std::cerr << "Failed to copy data from host to device" << std::endl;
        hipFree(d_values);
        hipFree(d_buckets);
        hipFree(d_output);
        return -1;
    }

    // Launch kernel
    bucketSortKernel<<<1, NUM_NODES, NUM_NODES * sizeof(int)>>>(d_values, d_buckets, d_output);
    err = hipGetLastError();
    if (err != hipSuccess) {
        std::cerr << "Kernel launch failed: " << hipGetErrorString(err) << std::endl;
        hipFree(d_values);
        hipFree(d_buckets);
        hipFree(d_output);
        return -1;
    }

    // Copy result back to host
    err = hipMemcpy(h_output, d_output, NUM_NODES * sizeof(int), hipMemcpyDeviceToHost);
    if (err != hipSuccess) {
        std::cerr << "Failed to copy data from device to host" << std::endl;
        hipFree(d_values);
        hipFree(d_buckets);
        hipFree(d_output);
        return -1;
    }

    // Free device memory
    hipFree(d_values);
    hipFree(d_buckets);
    hipFree(d_output);

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

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

                                        # End of file scope inline assembly
	.globl	_Z31__device_stub__bucketSortKernelPiS_S_ # -- Begin function _Z31__device_stub__bucketSortKernelPiS_S_
	.type	_Z31__device_stub__bucketSortKernelPiS_S_,@function
_Z31__device_stub__bucketSortKernelPiS_S_: # @_Z31__device_stub__bucketSortKernelPiS_S_
	.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	24(%rsp), %rsi
	movq	%rdx, (%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	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	$_Z16bucketSortKernelPiS_S_, %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	_Z31__device_stub__bucketSortKernelPiS_S_, .Lfunc_end0-_Z31__device_stub__bucketSortKernelPiS_S_
	.cfi_endproc
                                        # -- End function
	.globl	main                            # -- Begin function main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	subq	$288, %rsp                      # imm = 0x120
	.cfi_def_cfa_offset 304
	.cfi_offset %rbx, -16
	leaq	160(%rsp), %rdi
	movl	$16, %ecx
	movl	$.L__const.main.h_values, %esi
	rep;movsq (%rsi), %es:(%rdi)
	leaq	8(%rsp), %rdi
	movl	$128, %esi
	callq	hipMalloc
	testl	%eax, %eax
	je	.LBB1_2
# %bb.1:
	movl	$_ZSt4cerr, %ebx
	movl	$_ZSt4cerr, %edi
	movl	$.L.str, %esi
	movl	$45, %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
	jmp	.LBB1_17
.LBB1_2:
	leaq	24(%rsp), %rdi
	movl	$128, %esi
	callq	hipMalloc
	testl	%eax, %eax
	je	.LBB1_4
# %bb.3:
	movl	$_ZSt4cerr, %ebx
	movl	$_ZSt4cerr, %edi
	movl	$.L.str.1, %esi
	movl	$46, %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
	movq	8(%rsp), %rdi
	jmp	.LBB1_16
.LBB1_4:
	leaq	16(%rsp), %rdi
	movl	$128, %esi
	callq	hipMalloc
	testl	%eax, %eax
	je	.LBB1_6
# %bb.5:
	movl	$_ZSt4cerr, %ebx
	movl	$_ZSt4cerr, %edi
	movl	$.L.str.2, %esi
	movl	$45, %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
	movq	8(%rsp), %rdi
	callq	hipFree
	movq	24(%rsp), %rdi
	jmp	.LBB1_16
.LBB1_6:
	movq	8(%rsp), %rdi
	leaq	160(%rsp), %rsi
	movl	$128, %edx
	movl	$1, %ecx
	callq	hipMemcpy
	testl	%eax, %eax
	je	.LBB1_8
# %bb.7:
	movl	$_ZSt4cerr, %ebx
	movl	$_ZSt4cerr, %edi
	movl	$.L.str.3, %esi
	movl	$39, %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
	jmp	.LBB1_15
.LBB1_8:
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	leaq	31(%rdi), %rdx
	movl	$128, %r8d
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB1_10
# %bb.9:
	movq	8(%rsp), %rdi
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdx
	callq	_Z31__device_stub__bucketSortKernelPiS_S_
.LBB1_10:
	callq	hipGetLastError
	testl	%eax, %eax
	je	.LBB1_12
# %bb.11:
	movl	%eax, %ebx
	movl	$_ZSt4cerr, %edi
	movl	$.L.str.4, %esi
	movl	$22, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	%ebx, %edi
	callq	hipGetErrorString
	movl	$_ZSt4cerr, %edi
	movq	%rax, %rsi
	callq	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc
	movq	%rax, %rdi
	jmp	.LBB1_14
.LBB1_12:
	movq	16(%rsp), %rsi
	leaq	32(%rsp), %rdi
	movl	$128, %edx
	movl	$2, %ecx
	callq	hipMemcpy
	testl	%eax, %eax
	je	.LBB1_19
# %bb.13:
	movl	$_ZSt4cerr, %edi
	movl	$.L.str.5, %esi
	movl	$39, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	$_ZSt4cerr, %edi
.LBB1_14:
	callq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_
.LBB1_15:
	movq	8(%rsp), %rdi
	callq	hipFree
	movq	24(%rsp), %rdi
	callq	hipFree
	movq	16(%rsp), %rdi
.LBB1_16:
	callq	hipFree
.LBB1_17:
	movl	$-1, %eax
.LBB1_18:
	addq	$288, %rsp                      # imm = 0x120
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	retq
.LBB1_19:
	.cfi_def_cfa_offset 304
	movq	8(%rsp), %rdi
	callq	hipFree
	movq	24(%rsp), %rdi
	callq	hipFree
	movq	16(%rsp), %rdi
	callq	hipFree
	movl	$_ZSt4cout, %edi
	movl	$.L.str.6, %esi
	movl	$14, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	xorl	%ebx, %ebx
.LBB1_20:                               # =>This Inner Loop Header: Depth=1
	movl	32(%rsp,%rbx,4), %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSolsEi
	movl	$.L.str.7, %esi
	movl	$1, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	incq	%rbx
	cmpq	$32, %rbx
	jne	.LBB1_20
# %bb.21:
	movl	$_ZSt4cout, %edi
	callq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_
	xorl	%eax, %eax
	jmp	.LBB1_18
.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	$_Z16bucketSortKernelPiS_S_, %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	_Z16bucketSortKernelPiS_S_,@object # @_Z16bucketSortKernelPiS_S_
	.section	.rodata,"a",@progbits
	.globl	_Z16bucketSortKernelPiS_S_
	.p2align	3, 0x0
_Z16bucketSortKernelPiS_S_:
	.quad	_Z31__device_stub__bucketSortKernelPiS_S_
	.size	_Z16bucketSortKernelPiS_S_, 8

	.type	.L__const.main.h_values,@object # @__const.main.h_values
	.p2align	4, 0x0
.L__const.main.h_values:
	.long	27                              # 0x1b
	.long	25                              # 0x19
	.long	48                              # 0x30
	.long	37                              # 0x25
	.long	12                              # 0xc
	.long	9                               # 0x9
	.long	82                              # 0x52
	.long	10                              # 0xa
	.long	58                              # 0x3a
	.long	72                              # 0x48
	.long	21                              # 0x15
	.long	39                              # 0x27
	.long	29                              # 0x1d
	.long	47                              # 0x2f
	.long	63                              # 0x3f
	.long	79                              # 0x4f
	.long	81                              # 0x51
	.long	33                              # 0x21
	.long	17                              # 0x11
	.long	59                              # 0x3b
	.long	73                              # 0x49
	.long	15                              # 0xf
	.long	23                              # 0x17
	.long	41                              # 0x29
	.long	57                              # 0x39
	.long	71                              # 0x47
	.long	83                              # 0x53
	.long	35                              # 0x23
	.long	19                              # 0x13
	.long	51                              # 0x33
	.long	67                              # 0x43
	.long	85                              # 0x55
	.size	.L__const.main.h_values, 128

	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"Failed to allocate device memory for d_values"
	.size	.L.str, 46

	.type	.L.str.1,@object                # @.str.1
.L.str.1:
	.asciz	"Failed to allocate device memory for d_buckets"
	.size	.L.str.1, 47

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	"Failed to allocate device memory for d_output"
	.size	.L.str.2, 46

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	"Failed to copy data from host to device"
	.size	.L.str.3, 40

	.type	.L.str.4,@object                # @.str.4
.L.str.4:
	.asciz	"Kernel launch failed: "
	.size	.L.str.4, 23

	.type	.L.str.5,@object                # @.str.5
.L.str.5:
	.asciz	"Failed to copy data from device to host"
	.size	.L.str.5, 40

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

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

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z16bucketSortKernelPiS_S_"
	.size	.L__unnamed_1, 27

	.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 _Z31__device_stub__bucketSortKernelPiS_S_
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z16bucketSortKernelPiS_S_
	.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	_Z16bucketSortKernelPiS_S_ ; -- Begin function _Z16bucketSortKernelPiS_S_
	.globl	_Z16bucketSortKernelPiS_S_
	.p2align	8
	.type	_Z16bucketSortKernelPiS_S_,@function
_Z16bucketSortKernelPiS_S_:             ; @_Z16bucketSortKernelPiS_S_
; %bb.0:
	s_load_b64 s[2:3], s[0:1], 0x0
	v_lshlrev_b32_e32 v2, 2, v0
	v_cmp_gt_u32_e32 vcc_lo, 32, v0
	s_delay_alu instid0(VALU_DEP_2)
	v_add_nc_u32_e32 v1, 0x80, v2
	s_waitcnt lgkmcnt(0)
	global_load_b32 v3, v2, s[2:3]
	s_waitcnt vmcnt(0)
	ds_store_b32 v1, v3
	s_waitcnt lgkmcnt(0)
	s_barrier
	buffer_gl0_inv
	s_and_saveexec_b32 s2, vcc_lo
	s_cbranch_execz .LBB0_2
; %bb.1:
	v_mov_b32_e32 v3, 0
	ds_store_b32 v2, v3
.LBB0_2:
	s_or_b32 exec_lo, exec_lo, s2
	s_waitcnt lgkmcnt(0)
	s_barrier
	buffer_gl0_inv
	ds_load_b32 v3, v1
	s_waitcnt lgkmcnt(0)
	v_dual_mov_b32 v4, 1 :: v_dual_lshlrev_b32 v3, 2, v3
	ds_add_u32 v3, v4
	s_waitcnt lgkmcnt(0)
	s_barrier
	buffer_gl0_inv
	s_and_saveexec_b32 s2, vcc_lo
	s_cbranch_execz .LBB0_8
; %bb.3:                                ; %.preheader
	v_mov_b32_e32 v3, 0
	s_mov_b32 s4, 0
	s_mov_b32 s3, exec_lo
	v_cmpx_ne_u32_e32 0, v0
	s_cbranch_execz .LBB0_7
; %bb.4:                                ; %.lr.ph.preheader
	v_mov_b32_e32 v3, 0
	s_mov_b32 s5, 0
.LBB0_5:                                ; %.lr.ph
                                        ; =>This Inner Loop Header: Depth=1
	v_mov_b32_e32 v4, s4
	v_add_nc_u32_e32 v0, -1, v0
	s_add_i32 s4, s4, 4
	ds_load_b32 v4, v4
	v_cmp_eq_u32_e32 vcc_lo, 0, v0
	s_or_b32 s5, vcc_lo, s5
	s_waitcnt lgkmcnt(0)
	v_add_nc_u32_e32 v3, v4, v3
	s_and_not1_b32 exec_lo, exec_lo, s5
	s_cbranch_execnz .LBB0_5
; %bb.6:                                ; %Flow
	s_or_b32 exec_lo, exec_lo, s5
.LBB0_7:                                ; %Flow29
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 exec_lo, exec_lo, s3
	ds_store_b32 v2, v3
.LBB0_8:                                ; %Flow30
	s_or_b32 exec_lo, exec_lo, s2
	s_waitcnt lgkmcnt(0)
	s_barrier
	buffer_gl0_inv
	ds_load_b32 v2, v1
	v_mov_b32_e32 v0, 1
	s_load_b64 s[0:1], s[0:1], 0x10
	s_waitcnt lgkmcnt(0)
	v_lshlrev_b32_e32 v1, 2, v2
	ds_add_rtn_u32 v0, v1, v0
	s_waitcnt lgkmcnt(0)
	v_ashrrev_i32_e32 v1, 31, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[0:1], 2, v[0:1]
	v_add_co_u32 v0, vcc_lo, s0, v0
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v1, vcc_lo, s1, v1, vcc_lo
	global_store_b32 v[0:1], v2, off
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z16bucketSortKernelPiS_S_
		.amdhsa_group_segment_fixed_size 128
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 24
		.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 5
		.amdhsa_next_free_sgpr 6
		.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	_Z16bucketSortKernelPiS_S_, .Lfunc_end0-_Z16bucketSortKernelPiS_S_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 364
; NumSgprs: 8
; NumVgprs: 5
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 128 bytes/workgroup (compile time only)
; SGPRBlocks: 0
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 8
; NumVGPRsForWavesPerEU: 5
; 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
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
    .group_segment_fixed_size: 128
    .kernarg_segment_align: 8
    .kernarg_segment_size: 24
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z16bucketSortKernelPiS_S_
    .private_segment_fixed_size: 0
    .sgpr_count:     8
    .sgpr_spill_count: 0
    .symbol:         _Z16bucketSortKernelPiS_S_.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     5
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
