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

#define NUM_NODES 32

// Kernel function for breadth-first search using persistent threads
__global__ void bfsKernel(int *rowPtr, int *colInd, int *visited, int *frontier, int *newFrontier, int *newFrontierSize) {
    extern __shared__ int sharedFrontier[];

    int tid = threadIdx.x;
    int blockId = blockIdx.x;
    int globalId = tid + blockId * blockDim.x;

    // Load frontier into shared memory
    if (globalId < *newFrontierSize) {
        sharedFrontier[tid] = frontier[globalId];
    }
    __syncthreads();

    // Each thread processes a node in the frontier
    if (tid < *newFrontierSize) {
        int currentNode = sharedFrontier[tid];
        for (int i = rowPtr[currentNode]; i < rowPtr[currentNode + 1]; ++i) {
            int neighbor = colInd[i];
            if (!atomicExch(&visited[neighbor], 1)) {
                int index = atomicAdd(newFrontierSize, 1);
                newFrontier[index] = neighbor;
            }
        }
    }
}

int main() {
    // Example graph in COO format
    int rowPtr[NUM_NODES + 1] = {0, 2, 5, 8, 10, 12, 14, 16, 18, 20, 22, 24, 26, 28, 30, 32, 34, 36, 38, 40, 42, 44, 46, 48, 50, 52, 54, 56, 58, 60, 62, 64};
    int colInd[] = {1, 2, 0, 3, 4, 1, 5, 6, 2, 7, 3, 8, 4, 9, 5, 10, 6, 11, 7, 12, 8, 13, 9, 14, 10, 15, 11, 16, 12, 17, 13, 18, 14, 19, 15, 20, 16, 21, 17, 22, 18, 23, 19, 24, 20, 25, 21, 26, 22, 27, 23, 28, 24, 29, 25, 30, 26, 31, 27, 28, 29, 30, 31};
    int numEdges = sizeof(colInd) / sizeof(colInd[0]);

    // Host arrays
    int *h_visited = new int[NUM_NODES];
    int *h_frontier = new int[NUM_NODES];
    int *h_newFrontier = new int[NUM_NODES];
    int h_newFrontierSize = 0;

    // Initialize visited array and frontier
    for (int i = 0; i < NUM_NODES; ++i) {
        h_visited[i] = 0;
    }
    h_frontier[0] = 0; // Start BFS from node 0
    h_visited[0] = 1;
    h_newFrontierSize = 1;

    // Device arrays
    int *d_rowPtr, *d_colInd, *d_visited, *d_frontier, *d_newFrontier, *d_newFrontierSize;

    // Allocate device memory
    cudaMalloc(&d_rowPtr, (NUM_NODES + 1) * sizeof(int));
    cudaMalloc(&d_colInd, numEdges * sizeof(int));
    cudaMalloc(&d_visited, NUM_NODES * sizeof(int));
    cudaMalloc(&d_frontier, NUM_NODES * sizeof(int));
    cudaMalloc(&d_newFrontier, NUM_NODES * sizeof(int));
    cudaMalloc(&d_newFrontierSize, sizeof(int));

    // Copy data to device
    cudaMemcpy(d_rowPtr, rowPtr, (NUM_NODES + 1) * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(d_colInd, colInd, numEdges * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(d_visited, h_visited, NUM_NODES * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(d_frontier, h_frontier, NUM_NODES * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(d_newFrontierSize, &h_newFrontierSize, sizeof(int), cudaMemcpyHostToDevice);

    // Kernel configuration
    int blockSize = 256;
    int numBlocks = (NUM_NODES + blockSize - 1) / blockSize;

    // Perform BFS
    while (h_newFrontierSize > 0) {
        // Reset new frontier size on device
        h_newFrontierSize = 0;
        cudaMemcpy(d_newFrontierSize, &h_newFrontierSize, sizeof(int), cudaMemcpyHostToDevice);

        // Launch kernel
        bfsKernel<<<numBlocks, blockSize, blockSize * sizeof(int)>>>(d_rowPtr, d_colInd, d_visited, d_frontier, d_newFrontier, d_newFrontierSize);

        // Check for errors
        cudaError_t err = cudaGetLastError();
        if (err != cudaSuccess) {
            std::cerr << "CUDA error: " << cudaGetErrorString(err) << std::endl;
            return -1;
        }

        // Copy new frontier size back to host
        cudaMemcpy(&h_newFrontierSize, d_newFrontierSize, sizeof(int), cudaMemcpyDeviceToHost);

        // Copy new frontier back to host
        if (h_newFrontierSize > 0) {
            cudaMemcpy(h_frontier, d_newFrontier, h_newFrontierSize * sizeof(int), cudaMemcpyDeviceToHost);
        }
    }

    // Copy visited array back to host
    cudaMemcpy(h_visited, d_visited, NUM_NODES * sizeof(int), cudaMemcpyDeviceToHost);

    // Print visited nodes
    std::cout << "Visited nodes: ";
    for (int i = 0; i < NUM_NODES; ++i) {
        if (h_visited[i]) {
            std::cout << i << " ";
        }
    }
    std::cout << std::endl;

    // Free device memory
    cudaFree(d_rowPtr);
    cudaFree(d_colInd);
    cudaFree(d_visited);
    cudaFree(d_frontier);
    cudaFree(d_newFrontier);
    cudaFree(d_newFrontierSize);

    // Free host memory
    delete[] h_visited;
    delete[] h_frontier;
    delete[] h_newFrontier;

    return 0;
}
	.file	"tmpxft_000144ba_00000000-6_cuda_code_066968.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	_Z38__device_stub__Z9bfsKernelPiS_S_S_S_S_PiS_S_S_S_S_
	.type	_Z38__device_stub__Z9bfsKernelPiS_S_S_S_S_PiS_S_S_S_S_, @function
_Z38__device_stub__Z9bfsKernelPiS_S_S_S_S_PiS_S_S_S_S_:
.LFB3660:
	.cfi_startproc
	endbr64
	subq	$184, %rsp
	.cfi_def_cfa_offset 192
	movq	%rdi, 40(%rsp)
	leaq	72(%rsp), %rdi
	movq	%rsi, 32(%rsp)
	leaq	84(%rsp), %rsi
	movq	%rdx, 24(%rsp)
	leaq	56(%rsp), %rdx
	movq	%rcx, 16(%rsp)
	leaq	64(%rsp), %rcx
	movq	%r8, 8(%rsp)
	movq	%r9, (%rsp)
	movq	%fs:40, %rax
	movq	%rax, 168(%rsp)
	xorl	%eax, %eax
	leaq	40(%rsp), %rax
	movl	$1, 80(%rsp)
	movq	%rax, 120(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 152(%rsp)
	movq	%rsp, %rax
	movq	%rax, 160(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 72(%rsp)
	movq	%rax, 84(%rsp)
	movl	$1, 92(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L2
	pushq	64(%rsp)
	.cfi_def_cfa_offset 200
	leaq	_Z9bfsKernelPiS_S_S_S_S_(%rip), %rdi
	pushq	64(%rsp)
	.cfi_def_cfa_offset 208
	movq	100(%rsp), %rcx
	movl	108(%rsp), %r8d
	movq	88(%rsp), %rsi
	movl	96(%rsp), %edx
	leaq	136(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 200
	popq	%rdx
	.cfi_def_cfa_offset 192
.L2:
	movq	168(%rsp), %rax
	subq	%fs:40, %rax
	je	.L4
	call	__stack_chk_fail@PLT
.L4:
	addq	$184, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3660:
	.size	_Z38__device_stub__Z9bfsKernelPiS_S_S_S_S_PiS_S_S_S_S_, .-_Z38__device_stub__Z9bfsKernelPiS_S_S_S_S_PiS_S_S_S_S_
	.globl	_Z9bfsKernelPiS_S_S_S_S_
	.type	_Z9bfsKernelPiS_S_S_S_S_, @function
_Z9bfsKernelPiS_S_S_S_S_:
.LFB3661:
	.cfi_startproc
	endbr64
	jmp	_Z38__device_stub__Z9bfsKernelPiS_S_S_S_S_PiS_S_S_S_S_
	.cfi_endproc
.LFE3661:
	.size	_Z9bfsKernelPiS_S_S_S_S_, .-_Z9bfsKernelPiS_S_S_S_S_
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC1:
	.string	"CUDA error: "
.LC2:
	.string	"Visited nodes: "
.LC3:
	.string	" "
	.section	.rodata
	.align 32
.LC0:
	.long	1
	.long	2
	.long	0
	.long	3
	.long	4
	.long	1
	.long	5
	.long	6
	.long	2
	.long	7
	.long	3
	.long	8
	.long	4
	.long	9
	.long	5
	.long	10
	.long	6
	.long	11
	.long	7
	.long	12
	.long	8
	.long	13
	.long	9
	.long	14
	.long	10
	.long	15
	.long	11
	.long	16
	.long	12
	.long	17
	.long	13
	.long	18
	.long	14
	.long	19
	.long	15
	.long	20
	.long	16
	.long	21
	.long	17
	.long	22
	.long	18
	.long	23
	.long	19
	.long	24
	.long	20
	.long	25
	.long	21
	.long	26
	.long	22
	.long	27
	.long	23
	.long	28
	.long	24
	.long	29
	.long	25
	.long	30
	.long	26
	.long	31
	.long	27
	.long	28
	.long	29
	.long	30
	.long	31
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB3635:
	.cfi_startproc
	endbr64
	pushq	%r14
	.cfi_def_cfa_offset 16
	.cfi_offset 14, -16
	xorl	%edx, %edx
	leaq	.LC0(%rip), %rsi
	movl	$63, %ecx
	pushq	%r13
	.cfi_def_cfa_offset 24
	.cfi_offset 13, -24
	movabsq	$4294967552, %r13
	pushq	%r12
	.cfi_def_cfa_offset 32
	.cfi_offset 12, -32
	pushq	%rbp
	.cfi_def_cfa_offset 40
	.cfi_offset 6, -40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset 3, -48
	subq	$480, %rsp
	.cfi_def_cfa_offset 528
	movq	%fs:40, %rax
	movq	%rax, 472(%rsp)
	movabsq	$8589934592, %rax
	leaq	220(%rsp), %rdi
	movl	%edx, 216(%rsp)
	leaq	12(%rsp), %r12
	movq	%rax, 88(%rsp)
	movabsq	$34359738373, %rax
	movq	%rax, 96(%rsp)
	movabsq	$51539607562, %rax
	movq	%rax, 104(%rsp)
	movabsq	$68719476750, %rax
	movq	%rax, 112(%rsp)
	movabsq	$85899345938, %rax
	movq	%rax, 120(%rsp)
	movabsq	$103079215126, %rax
	movq	%rax, 128(%rsp)
	movabsq	$120259084314, %rax
	movq	%rax, 136(%rsp)
	movabsq	$137438953502, %rax
	movq	%rax, 144(%rsp)
	movabsq	$154618822690, %rax
	movq	%rax, 152(%rsp)
	movabsq	$171798691878, %rax
	movq	%rax, 160(%rsp)
	movabsq	$188978561066, %rax
	movq	%rax, 168(%rsp)
	movabsq	$206158430254, %rax
	movq	%rax, 176(%rsp)
	movabsq	$223338299442, %rax
	movq	%rax, 184(%rsp)
	movabsq	$240518168630, %rax
	movq	%rax, 192(%rsp)
	movabsq	$257698037818, %rax
	movq	%rax, 200(%rsp)
	movabsq	$274877907006, %rax
	movq	%rax, 208(%rsp)
	rep movsl
	movl	$128, %edi
	call	_Znam@PLT
	movl	$128, %edi
	movq	%rax, %rbx
	call	_Znam@PLT
	leaq	4(%rbx), %rdx
	movl	$31, %ecx
	movl	$132, %esi
	movq	%rdx, %rdi
	movq	%rax, %rbp
	xorl	%eax, %eax
	movl	$1, 12(%rsp)
	rep stosl
	xorl	%ecx, %ecx
	movl	$1, (%rbx)
	leaq	16(%rsp), %rdi
	movl	%ecx, 0(%rbp)
	call	cudaMalloc@PLT
	leaq	24(%rsp), %rdi
	movl	$252, %esi
	call	cudaMalloc@PLT
	leaq	32(%rsp), %rdi
	movl	$128, %esi
	call	cudaMalloc@PLT
	leaq	40(%rsp), %rdi
	movl	$128, %esi
	call	cudaMalloc@PLT
	leaq	48(%rsp), %rdi
	movl	$128, %esi
	call	cudaMalloc@PLT
	leaq	56(%rsp), %rdi
	movl	$4, %esi
	call	cudaMalloc@PLT
	movq	16(%rsp), %rdi
	leaq	88(%rsp), %rsi
	movl	$1, %ecx
	movl	$132, %edx
	call	cudaMemcpy@PLT
	movq	24(%rsp), %rdi
	movl	$1, %ecx
	leaq	220(%rsp), %rsi
	movl	$252, %edx
	call	cudaMemcpy@PLT
	movq	32(%rsp), %rdi
	movl	$1, %ecx
	movq	%rbx, %rsi
	movl	$128, %edx
	call	cudaMemcpy@PLT
	movq	40(%rsp), %rdi
	movl	$1, %ecx
	movq	%rbp, %rsi
	movl	$128, %edx
	call	cudaMemcpy@PLT
	movq	56(%rsp), %rdi
	movl	$1, %ecx
	movq	%r12, %rsi
	movl	$4, %edx
	call	cudaMemcpy@PLT
.L13:
	cmpl	$0, 12(%rsp)
	jle	.L21
	movq	56(%rsp), %rdi
	xorl	%eax, %eax
	movl	$1, %ecx
	movq	%r12, %rsi
	movl	$4, %edx
	movl	%eax, 12(%rsp)
	call	cudaMemcpy@PLT
	xorl	%r9d, %r9d
	movq	%r13, %rdx
	movl	$1, %ecx
	movl	$1024, %r8d
	movl	$1, %esi
	movabsq	$4294967297, %rdi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L10
	movq	56(%rsp), %r9
	movq	48(%rsp), %r8
	movq	40(%rsp), %rcx
	movq	32(%rsp), %rdx
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	call	_Z38__device_stub__Z9bfsKernelPiS_S_S_S_S_PiS_S_S_S_S_
.L10:
	call	cudaGetLastError@PLT
	movl	%eax, %r14d
	testl	%eax, %eax
	je	.L11
	leaq	.LC1(%rip), %rsi
	leaq	_ZSt4cerr(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	%r14d, %edi
	movq	%rax, %rbx
	call	cudaGetErrorString@PLT
	movq	%rbx, %rdi
	movq	%rax, %rsi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%rax, %rdi
	call	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@PLT
	orl	$-1, %eax
	jmp	.L8
.L11:
	movq	56(%rsp), %rsi
	movl	$4, %edx
	movl	$2, %ecx
	movq	%r12, %rdi
	call	cudaMemcpy@PLT
	movslq	12(%rsp), %rdx
	testl	%edx, %edx
	jle	.L13
	movq	48(%rsp), %rsi
	salq	$2, %rdx
	movl	$2, %ecx
	movq	%rbp, %rdi
	call	cudaMemcpy@PLT
	jmp	.L13
.L21:
	movq	32(%rsp), %rsi
	movl	$2, %ecx
	movq	%rbx, %rdi
	xorl	%r12d, %r12d
	movl	$128, %edx
	leaq	_ZSt4cout(%rip), %r13
	leaq	.LC3(%rip), %r14
	call	cudaMemcpy@PLT
	leaq	.LC2(%rip), %rsi
	movq	%r13, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
.L16:
	cmpl	$0, (%rbx,%r12,4)
	je	.L15
	movl	%r12d, %esi
	movq	%r13, %rdi
	call	_ZNSolsEi@PLT
	movq	%r14, %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
.L15:
	incq	%r12
	cmpq	$32, %r12
	jne	.L16
	movq	%r13, %rdi
	call	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@PLT
	movq	16(%rsp), %rdi
	call	cudaFree@PLT
	movq	24(%rsp), %rdi
	call	cudaFree@PLT
	movq	32(%rsp), %rdi
	call	cudaFree@PLT
	movq	40(%rsp), %rdi
	call	cudaFree@PLT
	movq	48(%rsp), %rdi
	call	cudaFree@PLT
	movq	56(%rsp), %rdi
	call	cudaFree@PLT
	movq	%rbx, %rdi
	call	_ZdaPv@PLT
	movq	%rbp, %rdi
	call	_ZdaPv@PLT
	xorl	%eax, %eax
.L8:
	movq	472(%rsp), %rdx
	subq	%fs:40, %rdx
	je	.L17
	call	__stack_chk_fail@PLT
.L17:
	addq	$480, %rsp
	.cfi_def_cfa_offset 48
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%rbp
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r13
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3635:
	.size	main, .-main
	.section	.rodata.str1.1
.LC4:
	.string	"_Z9bfsKernelPiS_S_S_S_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	.LC4(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z9bfsKernelPiS_S_S_S_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:

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

#define NUM_NODES 32

// Kernel function for breadth-first search using persistent threads
__global__ void bfsKernel(int *rowPtr, int *colInd, int *visited, int *frontier, int *newFrontier, int *newFrontierSize) {
    extern __shared__ int sharedFrontier[];

    int tid = threadIdx.x;
    int blockId = blockIdx.x;
    int globalId = tid + blockId * blockDim.x;

    // Load frontier into shared memory
    if (globalId < *newFrontierSize) {
        sharedFrontier[tid] = frontier[globalId];
    }
    __syncthreads();

    // Each thread processes a node in the frontier
    if (tid < *newFrontierSize) {
        int currentNode = sharedFrontier[tid];
        for (int i = rowPtr[currentNode]; i < rowPtr[currentNode + 1]; ++i) {
            int neighbor = colInd[i];
            if (!atomicExch(&visited[neighbor], 1)) {
                int index = atomicAdd(newFrontierSize, 1);
                newFrontier[index] = neighbor;
            }
        }
    }
}

int main() {
    // Example graph in COO format
    int rowPtr[NUM_NODES + 1] = {0, 2, 5, 8, 10, 12, 14, 16, 18, 20, 22, 24, 26, 28, 30, 32, 34, 36, 38, 40, 42, 44, 46, 48, 50, 52, 54, 56, 58, 60, 62, 64};
    int colInd[] = {1, 2, 0, 3, 4, 1, 5, 6, 2, 7, 3, 8, 4, 9, 5, 10, 6, 11, 7, 12, 8, 13, 9, 14, 10, 15, 11, 16, 12, 17, 13, 18, 14, 19, 15, 20, 16, 21, 17, 22, 18, 23, 19, 24, 20, 25, 21, 26, 22, 27, 23, 28, 24, 29, 25, 30, 26, 31, 27, 28, 29, 30, 31};
    int numEdges = sizeof(colInd) / sizeof(colInd[0]);

    // Host arrays
    int *h_visited = new int[NUM_NODES];
    int *h_frontier = new int[NUM_NODES];
    int *h_newFrontier = new int[NUM_NODES];
    int h_newFrontierSize = 0;

    // Initialize visited array and frontier
    for (int i = 0; i < NUM_NODES; ++i) {
        h_visited[i] = 0;
    }
    h_frontier[0] = 0; // Start BFS from node 0
    h_visited[0] = 1;
    h_newFrontierSize = 1;

    // Device arrays
    int *d_rowPtr, *d_colInd, *d_visited, *d_frontier, *d_newFrontier, *d_newFrontierSize;

    // Allocate device memory
    hipMalloc(&d_rowPtr, (NUM_NODES + 1) * sizeof(int));
    hipMalloc(&d_colInd, numEdges * sizeof(int));
    hipMalloc(&d_visited, NUM_NODES * sizeof(int));
    hipMalloc(&d_frontier, NUM_NODES * sizeof(int));
    hipMalloc(&d_newFrontier, NUM_NODES * sizeof(int));
    hipMalloc(&d_newFrontierSize, sizeof(int));

    // Copy data to device
    hipMemcpy(d_rowPtr, rowPtr, (NUM_NODES + 1) * sizeof(int), hipMemcpyHostToDevice);
    hipMemcpy(d_colInd, colInd, numEdges * sizeof(int), hipMemcpyHostToDevice);
    hipMemcpy(d_visited, h_visited, NUM_NODES * sizeof(int), hipMemcpyHostToDevice);
    hipMemcpy(d_frontier, h_frontier, NUM_NODES * sizeof(int), hipMemcpyHostToDevice);
    hipMemcpy(d_newFrontierSize, &h_newFrontierSize, sizeof(int), hipMemcpyHostToDevice);

    // Kernel configuration
    int blockSize = 256;
    int numBlocks = (NUM_NODES + blockSize - 1) / blockSize;

    // Perform BFS
    while (h_newFrontierSize > 0) {
        // Reset new frontier size on device
        h_newFrontierSize = 0;
        hipMemcpy(d_newFrontierSize, &h_newFrontierSize, sizeof(int), hipMemcpyHostToDevice);

        // Launch kernel
        bfsKernel<<<numBlocks, blockSize, blockSize * sizeof(int)>>>(d_rowPtr, d_colInd, d_visited, d_frontier, d_newFrontier, d_newFrontierSize);

        // Check for errors
        hipError_t err = hipGetLastError();
        if (err != hipSuccess) {
            std::cerr << "CUDA error: " << hipGetErrorString(err) << std::endl;
            return -1;
        }

        // Copy new frontier size back to host
        hipMemcpy(&h_newFrontierSize, d_newFrontierSize, sizeof(int), hipMemcpyDeviceToHost);

        // Copy new frontier back to host
        if (h_newFrontierSize > 0) {
            hipMemcpy(h_frontier, d_newFrontier, h_newFrontierSize * sizeof(int), hipMemcpyDeviceToHost);
        }
    }

    // Copy visited array back to host
    hipMemcpy(h_visited, d_visited, NUM_NODES * sizeof(int), hipMemcpyDeviceToHost);

    // Print visited nodes
    std::cout << "Visited nodes: ";
    for (int i = 0; i < NUM_NODES; ++i) {
        if (h_visited[i]) {
            std::cout << i << " ";
        }
    }
    std::cout << std::endl;

    // Free device memory
    hipFree(d_rowPtr);
    hipFree(d_colInd);
    hipFree(d_visited);
    hipFree(d_frontier);
    hipFree(d_newFrontier);
    hipFree(d_newFrontierSize);

    // Free host memory
    delete[] h_visited;
    delete[] h_frontier;
    delete[] h_newFrontier;

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

                                        # End of file scope inline assembly
	.globl	_Z24__device_stub__bfsKernelPiS_S_S_S_S_ # -- Begin function _Z24__device_stub__bfsKernelPiS_S_S_S_S_
	.type	_Z24__device_stub__bfsKernelPiS_S_S_S_S_,@function
_Z24__device_stub__bfsKernelPiS_S_S_S_S_: # @_Z24__device_stub__bfsKernelPiS_S_S_S_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	$144, %rsp
	.cfi_def_cfa_offset 192
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	56(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	48(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	40(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	32(%rsp), %rdx
	movq	%rcx, (%rdx)
	leaq	24(%rsp), %rcx
	movq	%r8, (%rcx)
	leaq	16(%rsp), %r8
	movq	%r9, (%r8)
	leaq	96(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%rbx)
	movq	%r8, 40(%rbx)
	leaq	80(%rsp), %r14
	leaq	64(%rsp), %r15
	leaq	8(%rsp), %r12
	movq	%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	$_Z9bfsKernelPiS_S_S_S_S_, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$160, %rsp
	.cfi_adjust_cfa_offset -160
	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__bfsKernelPiS_S_S_S_S_, .Lfunc_end0-_Z24__device_stub__bfsKernelPiS_S_S_S_S_
	.cfi_endproc
                                        # -- End function
	.globl	main                            # -- Begin function main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	subq	$472, %rsp                      # imm = 0x1D8
	.cfi_def_cfa_offset 528
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	leaq	64(%rsp), %rdi
	movl	$.L__const.main.rowPtr, %esi
	movl	$132, %edx
	callq	memcpy@PLT
	leaq	208(%rsp), %rdi
	movl	$.L__const.main.colInd, %esi
	movl	$252, %edx
	callq	memcpy@PLT
	movl	$128, %edi
	callq	_Znam
	movq	%rax, %rbx
	movl	$128, %edi
	callq	_Znam
	movq	%rax, %r14
	xorps	%xmm0, %xmm0
	movups	%xmm0, 112(%rbx)
	movups	%xmm0, 100(%rbx)
	movups	%xmm0, 84(%rbx)
	movups	%xmm0, 68(%rbx)
	movups	%xmm0, 52(%rbx)
	movups	%xmm0, 36(%rbx)
	movups	%xmm0, 20(%rbx)
	movups	%xmm0, 4(%rbx)
	movl	$0, (%rax)
	movl	$1, %eax
	movl	%eax, (%rbx)
	leaq	12(%rsp), %rcx
	movl	%eax, (%rcx)
	leaq	56(%rsp), %rbp
	movl	$132, %esi
	movq	%rbp, %rdi
	callq	hipMalloc
	leaq	48(%rsp), %r15
	movl	$252, %esi
	movq	%r15, %rdi
	callq	hipMalloc
	leaq	40(%rsp), %r12
	movl	$128, %esi
	movq	%r12, %rdi
	callq	hipMalloc
	leaq	32(%rsp), %rdi
	movl	$128, %esi
	callq	hipMalloc
	leaq	24(%rsp), %rdi
	movl	$128, %esi
	callq	hipMalloc
	leaq	16(%rsp), %r13
	movl	$4, %esi
	movq	%r13, %rdi
	callq	hipMalloc
	movq	(%rbp), %rdi
	movl	$132, %edx
	leaq	64(%rsp), %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movq	(%r15), %rdi
	movl	$252, %edx
	leaq	208(%rsp), %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movq	(%r12), %rdi
	movl	$128, %edx
	movq	%rbx, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	32(%rsp), %rax
	movq	(%rax), %rdi
	movl	$128, %edx
	movq	%r14, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movq	(%r13), %rdi
	movl	$4, %edx
	leaq	12(%rsp), %r15
	movq	%r15, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	cmpl	$0, (%r15)
	jle	.LBB1_11
# %bb.1:                                # %.lr.ph.preheader
	movabsq	$4294967297, %r15               # imm = 0x100000001
	leaq	12(%rsp), %r12
	leaq	255(%r15), %r13
.LBB1_2:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movl	$0, 12(%rsp)
	movq	16(%rsp), %rdi
	movl	$4, %edx
	movq	%r12, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movl	$1024, %r8d                     # imm = 0x400
	movq	%r15, %rdi
	movl	$1, %esi
	movq	%r13, %rdx
	movl	$1, %ecx
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB1_4
# %bb.3:                                #   in Loop: Header=BB1_2 Depth=1
	movq	56(%rsp), %rdi
	movq	48(%rsp), %rsi
	movq	40(%rsp), %rdx
	movq	32(%rsp), %rcx
	movq	24(%rsp), %r8
	movq	16(%rsp), %r9
	callq	_Z24__device_stub__bfsKernelPiS_S_S_S_S_
.LBB1_4:                                #   in Loop: Header=BB1_2 Depth=1
	callq	hipGetLastError
	testl	%eax, %eax
	jne	.LBB1_5
# %bb.9:                                #   in Loop: Header=BB1_2 Depth=1
	movq	16(%rsp), %rsi
	movl	$4, %edx
	movq	%r12, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movl	12(%rsp), %edx
	testl	%edx, %edx
	jle	.LBB1_11
# %bb.10:                               #   in Loop: Header=BB1_2 Depth=1
	movq	24(%rsp), %rsi
	shlq	$2, %rdx
	movq	%r14, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	cmpl	$0, 12(%rsp)
	jg	.LBB1_2
.LBB1_11:                               # %._crit_edge
	movq	40(%rsp), %rsi
	movl	$128, %edx
	movq	%rbx, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movl	$_ZSt4cout, %edi
	movl	$.L.str.1, %esi
	movl	$15, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	xorl	%r15d, %r15d
.LBB1_12:                               # =>This Inner Loop Header: Depth=1
	cmpl	$0, (%rbx,%r15,4)
	je	.LBB1_14
# %bb.13:                               #   in Loop: Header=BB1_12 Depth=1
	movl	$_ZSt4cout, %edi
	movl	%r15d, %esi
	callq	_ZNSolsEi
	movl	$.L.str.2, %esi
	movl	$1, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
.LBB1_14:                               #   in Loop: Header=BB1_12 Depth=1
	incq	%r15
	cmpq	$32, %r15
	jne	.LBB1_12
# %bb.15:
	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	56(%rsp), %rdi
	callq	hipFree
	movq	48(%rsp), %rdi
	callq	hipFree
	movq	40(%rsp), %rdi
	callq	hipFree
	movq	32(%rsp), %rdi
	callq	hipFree
	movq	24(%rsp), %rdi
	callq	hipFree
	movq	16(%rsp), %rdi
	callq	hipFree
	movq	%rbx, %rdi
	callq	_ZdaPv
	movq	%r14, %rdi
	callq	_ZdaPv
	xorl	%eax, %eax
	jmp	.LBB1_16
.LBB1_5:
	movl	%eax, %ebp
	movl	$_ZSt4cerr, %edi
	movl	$.L.str, %esi
	movl	$12, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	%ebp, %edi
	callq	hipGetErrorString
	testq	%rax, %rax
	je	.LBB1_6
# %bb.7:
	movq	%rax, %rbx
	movq	%rax, %rdi
	callq	strlen
	movl	$_ZSt4cerr, %edi
	movq	%rbx, %rsi
	movq	%rax, %rdx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	jmp	.LBB1_8
.LBB1_6:
	movq	_ZSt4cerr(%rip), %rax
	movq	-24(%rax), %rax
	leaq	_ZSt4cerr(%rax), %rdi
	movl	32(%rdi), %esi
	orl	$1, %esi
	callq	_ZNSt9basic_iosIcSt11char_traitsIcEE5clearESt12_Ios_Iostate
.LBB1_8:                                # %.thread
	movq	_ZSt4cerr(%rip), %rax
	movl	$_ZSt4cerr, %ecx
	movq	-24(%rax), %rdi
	addq	%rcx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cerr, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	$-1, %eax
.LBB1_16:
	addq	$472, %rsp                      # imm = 0x1D8
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_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	$_Z9bfsKernelPiS_S_S_S_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	_Z9bfsKernelPiS_S_S_S_S_,@object # @_Z9bfsKernelPiS_S_S_S_S_
	.section	.rodata,"a",@progbits
	.globl	_Z9bfsKernelPiS_S_S_S_S_
	.p2align	3, 0x0
_Z9bfsKernelPiS_S_S_S_S_:
	.quad	_Z24__device_stub__bfsKernelPiS_S_S_S_S_
	.size	_Z9bfsKernelPiS_S_S_S_S_, 8

	.type	.L__const.main.rowPtr,@object   # @__const.main.rowPtr
	.p2align	4, 0x0
.L__const.main.rowPtr:
	.long	0                               # 0x0
	.long	2                               # 0x2
	.long	5                               # 0x5
	.long	8                               # 0x8
	.long	10                              # 0xa
	.long	12                              # 0xc
	.long	14                              # 0xe
	.long	16                              # 0x10
	.long	18                              # 0x12
	.long	20                              # 0x14
	.long	22                              # 0x16
	.long	24                              # 0x18
	.long	26                              # 0x1a
	.long	28                              # 0x1c
	.long	30                              # 0x1e
	.long	32                              # 0x20
	.long	34                              # 0x22
	.long	36                              # 0x24
	.long	38                              # 0x26
	.long	40                              # 0x28
	.long	42                              # 0x2a
	.long	44                              # 0x2c
	.long	46                              # 0x2e
	.long	48                              # 0x30
	.long	50                              # 0x32
	.long	52                              # 0x34
	.long	54                              # 0x36
	.long	56                              # 0x38
	.long	58                              # 0x3a
	.long	60                              # 0x3c
	.long	62                              # 0x3e
	.long	64                              # 0x40
	.long	0                               # 0x0
	.size	.L__const.main.rowPtr, 132

	.type	.L__const.main.colInd,@object   # @__const.main.colInd
	.p2align	4, 0x0
.L__const.main.colInd:
	.long	1                               # 0x1
	.long	2                               # 0x2
	.long	0                               # 0x0
	.long	3                               # 0x3
	.long	4                               # 0x4
	.long	1                               # 0x1
	.long	5                               # 0x5
	.long	6                               # 0x6
	.long	2                               # 0x2
	.long	7                               # 0x7
	.long	3                               # 0x3
	.long	8                               # 0x8
	.long	4                               # 0x4
	.long	9                               # 0x9
	.long	5                               # 0x5
	.long	10                              # 0xa
	.long	6                               # 0x6
	.long	11                              # 0xb
	.long	7                               # 0x7
	.long	12                              # 0xc
	.long	8                               # 0x8
	.long	13                              # 0xd
	.long	9                               # 0x9
	.long	14                              # 0xe
	.long	10                              # 0xa
	.long	15                              # 0xf
	.long	11                              # 0xb
	.long	16                              # 0x10
	.long	12                              # 0xc
	.long	17                              # 0x11
	.long	13                              # 0xd
	.long	18                              # 0x12
	.long	14                              # 0xe
	.long	19                              # 0x13
	.long	15                              # 0xf
	.long	20                              # 0x14
	.long	16                              # 0x10
	.long	21                              # 0x15
	.long	17                              # 0x11
	.long	22                              # 0x16
	.long	18                              # 0x12
	.long	23                              # 0x17
	.long	19                              # 0x13
	.long	24                              # 0x18
	.long	20                              # 0x14
	.long	25                              # 0x19
	.long	21                              # 0x15
	.long	26                              # 0x1a
	.long	22                              # 0x16
	.long	27                              # 0x1b
	.long	23                              # 0x17
	.long	28                              # 0x1c
	.long	24                              # 0x18
	.long	29                              # 0x1d
	.long	25                              # 0x19
	.long	30                              # 0x1e
	.long	26                              # 0x1a
	.long	31                              # 0x1f
	.long	27                              # 0x1b
	.long	28                              # 0x1c
	.long	29                              # 0x1d
	.long	30                              # 0x1e
	.long	31                              # 0x1f
	.size	.L__const.main.colInd, 252

	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"CUDA error: "
	.size	.L.str, 13

	.type	.L.str.1,@object                # @.str.1
.L.str.1:
	.asciz	"Visited nodes: "
	.size	.L.str.1, 16

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

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

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

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

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym _Z24__device_stub__bfsKernelPiS_S_S_S_S_
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z9bfsKernelPiS_S_S_S_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	_Z9bfsKernelPiS_S_S_S_S_ ; -- Begin function _Z9bfsKernelPiS_S_S_S_S_
	.globl	_Z9bfsKernelPiS_S_S_S_S_
	.p2align	8
	.type	_Z9bfsKernelPiS_S_S_S_S_,@function
_Z9bfsKernelPiS_S_S_S_S_:               ; @_Z9bfsKernelPiS_S_S_S_S_
; %bb.0:
	s_clause 0x1
	s_load_b128 s[8:11], s[0:1], 0x20
	s_load_b32 s12, s[0:1], 0x3c
	v_lshl_add_u32 v3, v0, 2, 0
	s_waitcnt lgkmcnt(0)
	s_load_b32 s13, s[10:11], 0x0
	s_load_b256 s[0:7], s[0:1], 0x0
	s_and_b32 s12, s12, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[1:2], null, s15, s12, v[0:1]
	s_mov_b32 s12, exec_lo
	s_waitcnt lgkmcnt(0)
	v_cmpx_gt_i32_e64 s13, v1
	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[1:2], 2, v[1:2]
	v_add_co_u32 v1, vcc_lo, s6, v1
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v2, vcc_lo, s7, v2, vcc_lo
	global_load_b32 v1, v[1:2], off
	s_waitcnt vmcnt(0)
	ds_store_b32 v3, v1
.LBB0_2:
	s_or_b32 exec_lo, exec_lo, s12
	s_waitcnt lgkmcnt(0)
	s_barrier
	buffer_gl0_inv
	s_load_b32 s6, s[10:11], 0x0
	s_waitcnt lgkmcnt(0)
	v_cmp_gt_i32_e32 vcc_lo, s6, v0
	s_and_saveexec_b32 s6, vcc_lo
	s_cbranch_execz .LBB0_10
; %bb.3:
	ds_load_b32 v0, v3
	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
	s_mov_b32 s1, 0
	global_load_b64 v[2:3], v[0:1], off
	s_waitcnt vmcnt(0)
	v_cmp_lt_i32_e32 vcc_lo, v2, v3
	s_and_b32 exec_lo, exec_lo, vcc_lo
	s_cbranch_execz .LBB0_10
; %bb.4:                                ; %.lr.ph.preheader
	v_ashrrev_i32_e32 v3, 31, v2
	v_dual_mov_b32 v7, 1 :: v_dual_mov_b32 v8, 0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[3:4], 2, v[2:3]
	v_add_co_u32 v3, vcc_lo, s2, v3
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v4, vcc_lo, s3, v4, vcc_lo
.LBB0_5:                                ; %.lr.ph
                                        ; =>This Inner Loop Header: Depth=1
	global_load_b32 v5, v[3:4], off
	s_mov_b32 s0, exec_lo
	s_waitcnt vmcnt(0)
	v_ashrrev_i32_e32 v6, 31, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[9:10], 2, v[5:6]
	v_add_co_u32 v9, vcc_lo, s4, v9
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v10, vcc_lo, s5, v10, vcc_lo
	global_atomic_swap_b32 v6, v[9:10], v7, off glc
	s_waitcnt vmcnt(0)
	v_cmpx_eq_u32_e32 0, v6
	s_cbranch_execz .LBB0_9
; %bb.6:                                ;   in Loop: Header=BB0_5 Depth=1
	s_mov_b32 s3, exec_lo
	s_mov_b32 s2, exec_lo
	v_mbcnt_lo_u32_b32 v6, s3, 0
                                        ; implicit-def: $vgpr9
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_eq_u32_e32 0, v6
	s_cbranch_execz .LBB0_8
; %bb.7:                                ;   in Loop: Header=BB0_5 Depth=1
	s_bcnt1_i32_b32 s3, s3
	s_delay_alu instid0(SALU_CYCLE_1)
	v_mov_b32_e32 v9, s3
	global_atomic_add_u32 v9, v8, v9, s[10:11] glc
.LBB0_8:                                ;   in Loop: Header=BB0_5 Depth=1
	s_or_b32 exec_lo, exec_lo, s2
	s_waitcnt vmcnt(0)
	v_readfirstlane_b32 s2, v9
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v9, s2, v6
	v_ashrrev_i32_e32 v10, 31, v9
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[9:10], 2, v[9:10]
	v_add_co_u32 v9, vcc_lo, s8, v9
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v10, vcc_lo, s9, v10, vcc_lo
	global_store_b32 v[9:10], v5, off
.LBB0_9:                                ; %Flow
                                        ;   in Loop: Header=BB0_5 Depth=1
	s_or_b32 exec_lo, exec_lo, s0
	global_load_b32 v5, v[0:1], off offset:4
	v_add_nc_u32_e32 v2, 1, v2
	v_add_co_u32 v3, s0, v3, 4
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_add_co_ci_u32_e64 v4, s0, 0, v4, s0
	s_waitcnt vmcnt(0)
	v_cmp_ge_i32_e32 vcc_lo, v2, v5
	s_or_b32 s1, vcc_lo, s1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s1
	s_cbranch_execnz .LBB0_5
.LBB0_10:                               ; %.loopexit
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z9bfsKernelPiS_S_S_S_S_
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 304
		.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 11
		.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	_Z9bfsKernelPiS_S_S_S_S_, .Lfunc_end0-_Z9bfsKernelPiS_S_S_S_S_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 544
; NumSgprs: 18
; NumVgprs: 11
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 11
; Occupancy: 16
; WaveLimiterHint : 1
; 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
      - .address_space:  global
        .offset:         24
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         32
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         40
        .size:           8
        .value_kind:     global_buffer
      - .offset:         48
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         52
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         56
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         60
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         62
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         64
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         66
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         68
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         70
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         88
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         96
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         104
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         112
        .size:           2
        .value_kind:     hidden_grid_dims
      - .offset:         168
        .size:           4
        .value_kind:     hidden_dynamic_lds_size
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 304
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z9bfsKernelPiS_S_S_S_S_
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z9bfsKernelPiS_S_S_S_S_.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     11
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
