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

#define MAX_EDGES 128
#define MAX_NODES 32

__global__ void bfs_kernel(int *frontier, int *visited, int *edge_list, int *edge_offsets, int num_nodes, int *new_frontier, int *new_frontier_size) {
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    if (tid < num_nodes && frontier[tid] != -1) {
        int node = frontier[tid];
        for (int i = edge_offsets[node]; i < edge_offsets[node + 1]; i++) {
            int neighbor = edge_list[i];
            if (!visited[neighbor]) {
                atomicExch(&visited[neighbor], 1);
                int index = atomicAdd(new_frontier_size, 1);
                new_frontier[index] = neighbor;
            }
        }
        frontier[tid] = -1; // Mark as processed
    }
}

__global__ void initialize_frontier(int *frontier, int start_node) {
    frontier[start_node] = start_node;
}

__global__ void initialize_visited(int *visited) {
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    if (tid < MAX_NODES) {
        visited[tid] = 0;
    }
}

void bfs(int *edge_list, int *edge_offsets, int start_node, int num_nodes, int num_gpus) {
    int *d_frontier, *d_visited, *d_edge_list, *d_edge_offsets, *d_new_frontier, *d_new_frontier_size;
    int *h_frontier = new int[num_nodes];
    int *h_visited = new int[num_nodes];
    int *h_new_frontier = new int[num_nodes];
    int *h_new_frontier_size = new int[1];

    cudaSetDevice(0);
    cudaMalloc(&d_frontier, num_nodes * sizeof(int));
    cudaMalloc(&d_visited, num_nodes * sizeof(int));
    cudaMalloc(&d_edge_list, MAX_EDGES * sizeof(int));
    cudaMalloc(&d_edge_offsets, (num_nodes + 1) * sizeof(int));
    cudaMalloc(&d_new_frontier, num_nodes * sizeof(int));
    cudaMalloc(&d_new_frontier_size, sizeof(int));

    // Initialize host data
    for (int i = 0; i < num_nodes; i++) {
        h_frontier[i] = -1;
        h_visited[i] = 0;
    }
    h_frontier[start_node] = start_node;
    h_visited[start_node] = 1;
    *h_new_frontier_size = 0;

    // Copy data to device
    cudaMemcpy(d_frontier, h_frontier, num_nodes * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(d_visited, h_visited, num_nodes * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(d_edge_list, edge_list, MAX_EDGES * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(d_edge_offsets, edge_offsets, (num_nodes + 1) * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(d_new_frontier_size, h_new_frontier_size, sizeof(int), cudaMemcpyHostToDevice);

    dim3 threadsPerBlock(256);
    dim3 blocksPerGrid((num_nodes + threadsPerBlock.x - 1) / threadsPerBlock.x);

    while (true) {
        // Reset new frontier size
        *h_new_frontier_size = 0;
        cudaMemcpy(d_new_frontier_size, h_new_frontier_size, sizeof(int), cudaMemcpyHostToDevice);

        // Launch BFS kernel
        bfs_kernel<<<blocksPerGrid, threadsPerBlock>>>(d_frontier, d_visited, d_edge_list, d_edge_offsets, num_nodes, d_new_frontier, d_new_frontier_size);
        cudaDeviceSynchronize();

        // Copy new frontier size back to host
        cudaMemcpy(h_new_frontier_size, d_new_frontier_size, sizeof(int), cudaMemcpyDeviceToHost);

        if (*h_new_frontier_size == 0) break;

        // Copy new frontier to host
        cudaMemcpy(h_new_frontier, d_new_frontier, *h_new_frontier_size * sizeof(int), cudaMemcpyDeviceToHost);

        // Copy new frontier to old frontier
        cudaMemcpy(d_frontier, d_new_frontier, *h_new_frontier_size * sizeof(int), cudaMemcpyDeviceToDevice);
    }

    // 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_frontier);
    cudaFree(d_visited);
    cudaFree(d_edge_list);
    cudaFree(d_edge_offsets);
    cudaFree(d_new_frontier);
    cudaFree(d_new_frontier_size);

    // Free host memory
    delete[] h_frontier;
    delete[] h_visited;
    delete[] h_new_frontier;
    delete[] h_new_frontier_size;
}

int main() {
    // Example graph represented as an edge list
    int edge_list[MAX_EDGES] = {1, 2, 3, 0, 4, 5, 1, 6, 7, 2, 8, 9, 3, 10, 11, 4, 12, 13, 5, 14, 15, 6, 16, 17, 7, 18, 19, 8, 20, 21, 9, 22, 23, 10, 24, 25, 11, 26, 27, 12, 28, 29, 13, 30, 31};
    int edge_offsets[MAX_NODES + 1] = {0, 2, 4, 6, 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 start_node = 0;
    int num_nodes = MAX_NODES;
    int num_gpus = 1;

    bfs(edge_list, edge_offsets, start_node, num_nodes, num_gpus);

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

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB4376:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE4376:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z41__device_stub__Z10bfs_kernelPiS_S_S_iS_S_PiS_S_S_iS_S_
	.type	_Z41__device_stub__Z10bfs_kernelPiS_S_S_iS_S_PiS_S_S_iS_S_, @function
_Z41__device_stub__Z10bfs_kernelPiS_S_S_iS_S_PiS_S_S_iS_S_:
.LFB4398:
	.cfi_startproc
	endbr64
	subq	$200, %rsp
	.cfi_def_cfa_offset 208
	movq	208(%rsp), %rax
	movq	%rdi, 56(%rsp)
	leaq	80(%rsp), %rdi
	movq	%rsi, 48(%rsp)
	leaq	92(%rsp), %rsi
	movq	%rdx, 40(%rsp)
	leaq	64(%rsp), %rdx
	movq	%rcx, 32(%rsp)
	leaq	72(%rsp), %rcx
	movl	%r8d, 28(%rsp)
	movq	%r9, 16(%rsp)
	movq	%rax, 8(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 184(%rsp)
	xorl	%eax, %eax
	leaq	56(%rsp), %rax
	movl	$1, 88(%rsp)
	movq	%rax, 128(%rsp)
	leaq	48(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	28(%rsp), %rax
	movq	%rax, 160(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 168(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 176(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 80(%rsp)
	movq	%rax, 92(%rsp)
	movl	$1, 100(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L2
	pushq	72(%rsp)
	.cfi_def_cfa_offset 216
	leaq	_Z10bfs_kernelPiS_S_S_iS_S_(%rip), %rdi
	pushq	72(%rsp)
	.cfi_def_cfa_offset 224
	movq	108(%rsp), %rcx
	movl	116(%rsp), %r8d
	movq	96(%rsp), %rsi
	movl	104(%rsp), %edx
	leaq	144(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 216
	popq	%rdx
	.cfi_def_cfa_offset 208
.L2:
	movq	184(%rsp), %rax
	subq	%fs:40, %rax
	je	.L4
	call	__stack_chk_fail@PLT
.L4:
	addq	$200, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4398:
	.size	_Z41__device_stub__Z10bfs_kernelPiS_S_S_iS_S_PiS_S_S_iS_S_, .-_Z41__device_stub__Z10bfs_kernelPiS_S_S_iS_S_PiS_S_S_iS_S_
	.globl	_Z10bfs_kernelPiS_S_S_iS_S_
	.type	_Z10bfs_kernelPiS_S_S_iS_S_, @function
_Z10bfs_kernelPiS_S_S_iS_S_:
.LFB4399:
	.cfi_startproc
	endbr64
	jmp	_Z41__device_stub__Z10bfs_kernelPiS_S_S_iS_S_PiS_S_S_iS_S_
	.cfi_endproc
.LFE4399:
	.size	_Z10bfs_kernelPiS_S_S_iS_S_, .-_Z10bfs_kernelPiS_S_S_iS_S_
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC1:
	.string	"Visited nodes: "
.LC2:
	.string	" "
	.text
	.globl	_Z3bfsPiS_iii
	.type	_Z3bfsPiS_iii, @function
_Z3bfsPiS_iii:
.LFB4372:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movslq	%ecx, %rbx
	subq	$120, %rsp
	.cfi_def_cfa_offset 176
	movq	%rsi, 16(%rsp)
	movl	%edx, 28(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	movabsq	$2305843009213693950, %rax
	cmpq	%rbx, %rax
	jb	.L9
	movq	%rbx, %r13
	salq	$2, %rbx
	movq	%rdi, %r15
	movq	%rbx, %rdi
	call	_Znam@PLT
	movq	%rbx, %rdi
	movq	%rax, %r14
	call	_Znam@PLT
	movq	%rbx, %rdi
	movq	%rax, %r12
	call	_Znam@PLT
	movl	$4, %edi
	movq	%rax, 8(%rsp)
	call	_Znam@PLT
	xorl	%edi, %edi
	movq	%rax, %rbp
	call	cudaSetDevice@PLT
	leaq	32(%rsp), %rdi
	movq	%rbx, %rsi
	call	cudaMalloc@PLT
	leaq	40(%rsp), %rdi
	movq	%rbx, %rsi
	call	cudaMalloc@PLT
	leaq	48(%rsp), %rdi
	movl	$512, %esi
	call	cudaMalloc@PLT
	leaq	4(%rbx), %rax
	leaq	56(%rsp), %rdi
	movq	%rax, %rsi
	movq	%rax, (%rsp)
	call	cudaMalloc@PLT
	leaq	64(%rsp), %rdi
	movq	%rbx, %rsi
	call	cudaMalloc@PLT
	leaq	72(%rsp), %rdi
	movl	$4, %esi
	call	cudaMalloc@PLT
	movl	28(%rsp), %edx
	xorl	%eax, %eax
	jmp	.L10
.L9:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L21
	call	__cxa_throw_bad_array_new_length@PLT
.L10:
	cmpl	%eax, %r13d
	jle	.L22
	xorl	%r8d, %r8d
	movl	$-1, (%r14,%rax,4)
	movl	%r8d, (%r12,%rax,4)
	incq	%rax
	jmp	.L10
.L22:
	movslq	%edx, %rax
	xorl	%edi, %edi
	movl	$1, %ecx
	movq	%r14, %rsi
	movl	%edx, (%r14,%rax,4)
	movq	%rbx, %rdx
	movl	%edi, 0(%rbp)
	movq	32(%rsp), %rdi
	movl	$1, (%r12,%rax,4)
	call	cudaMemcpy@PLT
	movq	40(%rsp), %rdi
	movq	%rbx, %rdx
	movq	%r12, %rsi
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	movq	48(%rsp), %rdi
	movl	$1, %ecx
	movq	%r15, %rsi
	movl	$512, %edx
	call	cudaMemcpy@PLT
	movq	(%rsp), %rdx
	movq	16(%rsp), %rsi
	movl	$1, %ecx
	movq	56(%rsp), %rdi
	call	cudaMemcpy@PLT
	movq	72(%rsp), %rdi
	movl	$4, %edx
	movq	%rbp, %rsi
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	leal	255(%r13), %eax
	movabsq	$4294967297, %rdx
	shrl	$8, %eax
	movq	%rdx, 84(%rsp)
	movl	%eax, 92(%rsp)
	movq	%rdx, 96(%rsp)
.L15:
	xorl	%esi, %esi
	movq	72(%rsp), %rdi
	movl	$1, %ecx
	movl	$4, %edx
	movl	%esi, 0(%rbp)
	movq	%rbp, %rsi
	call	cudaMemcpy@PLT
	movl	88(%rsp), %ecx
	movl	100(%rsp), %esi
	xorl	%r9d, %r9d
	movl	$256, 80(%rsp)
	movq	92(%rsp), %rdi
	xorl	%r8d, %r8d
	movq	80(%rsp), %rdx
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L13
	pushq	%rax
	.cfi_def_cfa_offset 184
	movl	%r13d, %r8d
	pushq	80(%rsp)
	.cfi_def_cfa_offset 192
	movq	80(%rsp), %r9
	movq	72(%rsp), %rcx
	movq	64(%rsp), %rdx
	movq	56(%rsp), %rsi
	movq	48(%rsp), %rdi
	call	_Z41__device_stub__Z10bfs_kernelPiS_S_S_iS_S_PiS_S_S_iS_S_
	popq	%rdx
	.cfi_def_cfa_offset 184
	popq	%rcx
	.cfi_def_cfa_offset 176
.L13:
	call	cudaDeviceSynchronize@PLT
	movq	72(%rsp), %rsi
	movl	$4, %edx
	movq	%rbp, %rdi
	movl	$2, %ecx
	call	cudaMemcpy@PLT
	movslq	0(%rbp), %rdx
	testl	%edx, %edx
	je	.L14
	movq	64(%rsp), %rsi
	movq	8(%rsp), %rdi
	salq	$2, %rdx
	movl	$2, %ecx
	call	cudaMemcpy@PLT
	movslq	0(%rbp), %rdx
	movq	64(%rsp), %rsi
	movl	$3, %ecx
	movq	32(%rsp), %rdi
	salq	$2, %rdx
	call	cudaMemcpy@PLT
	jmp	.L15
.L14:
	movq	40(%rsp), %rsi
	movq	%rbx, %rdx
	movl	$2, %ecx
	movq	%r12, %rdi
	xorl	%ebx, %ebx
	leaq	.LC2(%rip), %r15
	call	cudaMemcpy@PLT
	movl	$15, %edx
	leaq	.LC1(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l@PLT
.L16:
	cmpl	%ebx, %r13d
	jle	.L23
	cmpl	$0, (%r12,%rbx,4)
	je	.L17
	movl	%ebx, %esi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZNSolsEi@PLT
	movl	$1, %edx
	movq	%r15, %rsi
	movq	%rax, %rdi
	call	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l@PLT
.L17:
	incq	%rbx
	jmp	.L16
.L23:
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@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	64(%rsp), %rdi
	call	cudaFree@PLT
	movq	72(%rsp), %rdi
	call	cudaFree@PLT
	movq	%r14, %rdi
	call	_ZdaPv@PLT
	movq	%r12, %rdi
	call	_ZdaPv@PLT
	movq	8(%rsp), %rdi
	call	_ZdaPv@PLT
	movq	%rbp, %rdi
	call	_ZdaPv@PLT
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L19
.L21:
	call	__stack_chk_fail@PLT
.L19:
	addq	$120, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.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
	ret
	.cfi_endproc
.LFE4372:
	.size	_Z3bfsPiS_iii, .-_Z3bfsPiS_iii
	.section	.rodata
	.align 32
.LC0:
	.long	0
	.long	2
	.long	4
	.long	6
	.long	8
	.long	10
	.long	12
	.long	14
	.long	16
	.long	18
	.long	20
	.long	22
	.long	24
	.long	26
	.long	28
	.long	30
	.long	32
	.long	34
	.long	36
	.long	38
	.long	40
	.long	42
	.long	44
	.long	46
	.long	48
	.long	50
	.long	52
	.long	54
	.long	56
	.long	58
	.long	60
	.long	62
	.long	64
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB4373:
	.cfi_startproc
	endbr64
	subq	$664, %rsp
	.cfi_def_cfa_offset 672
	movl	$128, %ecx
	leaq	.LC0(%rip), %rsi
	xorl	%edx, %edx
	movq	%fs:40, %rax
	movq	%rax, 648(%rsp)
	xorl	%eax, %eax
	leaq	136(%rsp), %rdi
	movl	$1, %r8d
	rep stosl
	leaq	4(%rsp), %rdi
	movl	$33, %ecx
	movabsq	$8589934593, %rax
	rep movsl
	leaq	4(%rsp), %rsi
	movq	%rax, 136(%rsp)
	movabsq	$21474836484, %rax
	movq	%rax, 152(%rsp)
	movl	$32, %ecx
	movabsq	$25769803777, %rax
	leaq	136(%rsp), %rdi
	movq	%rax, 160(%rsp)
	movabsq	$8589934599, %rax
	movq	%rax, 168(%rsp)
	movabsq	$38654705672, %rax
	movq	%rax, 176(%rsp)
	movabsq	$42949672963, %rax
	movq	%rax, 184(%rsp)
	movabsq	$17179869195, %rax
	movq	%rax, 192(%rsp)
	movabsq	$55834574860, %rax
	movq	%rax, 200(%rsp)
	movabsq	$60129542149, %rax
	movq	%rax, 208(%rsp)
	movabsq	$25769803791, %rax
	movq	%rax, 216(%rsp)
	movabsq	$73014444048, %rax
	movq	%rax, 224(%rsp)
	movabsq	$77309411335, %rax
	movq	%rax, 232(%rsp)
	movabsq	$34359738387, %rax
	movq	%rax, 240(%rsp)
	movabsq	$90194313236, %rax
	movq	%rax, 248(%rsp)
	movabsq	$94489280521, %rax
	movq	%rax, 256(%rsp)
	movabsq	$42949672983, %rax
	movq	%rax, 264(%rsp)
	movabsq	$107374182424, %rax
	movq	%rax, 272(%rsp)
	movabsq	$111669149707, %rax
	movq	%rax, 280(%rsp)
	movabsq	$51539607579, %rax
	movq	%rax, 288(%rsp)
	movabsq	$124554051612, %rax
	movq	%rax, 296(%rsp)
	movabsq	$128849018893, %rax
	movq	%rax, 304(%rsp)
	movl	$3, 144(%rsp)
	movl	$31, 312(%rsp)
	call	_Z3bfsPiS_iii
	movq	648(%rsp), %rax
	subq	%fs:40, %rax
	je	.L25
	call	__stack_chk_fail@PLT
.L25:
	xorl	%eax, %eax
	addq	$664, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4373:
	.size	main, .-main
	.text
	.globl	_Z40__device_stub__Z19initialize_frontierPiiPii
	.type	_Z40__device_stub__Z19initialize_frontierPiiPii, @function
_Z40__device_stub__Z19initialize_frontierPiiPii:
.LFB4400:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movq	%rdi, 8(%rsp)
	leaq	32(%rsp), %rcx
	leaq	24(%rsp), %rdx
	movl	%esi, 4(%rsp)
	leaq	40(%rsp), %rdi
	leaq	52(%rsp), %rsi
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rax
	movl	$1, 48(%rsp)
	movq	%rax, 88(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 40(%rsp)
	movq	%rax, 52(%rsp)
	movl	$1, 60(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L27
	pushq	32(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z19initialize_frontierPii(%rip), %rdi
	pushq	32(%rsp)
	.cfi_def_cfa_offset 144
	movq	68(%rsp), %rcx
	movl	76(%rsp), %r8d
	movq	56(%rsp), %rsi
	movl	64(%rsp), %edx
	leaq	104(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L27:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L29
	call	__stack_chk_fail@PLT
.L29:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4400:
	.size	_Z40__device_stub__Z19initialize_frontierPiiPii, .-_Z40__device_stub__Z19initialize_frontierPiiPii
	.globl	_Z19initialize_frontierPii
	.type	_Z19initialize_frontierPii, @function
_Z19initialize_frontierPii:
.LFB4401:
	.cfi_startproc
	endbr64
	jmp	_Z40__device_stub__Z19initialize_frontierPiiPii
	.cfi_endproc
.LFE4401:
	.size	_Z19initialize_frontierPii, .-_Z19initialize_frontierPii
	.globl	_Z38__device_stub__Z18initialize_visitedPiPi
	.type	_Z38__device_stub__Z18initialize_visitedPiPi, @function
_Z38__device_stub__Z18initialize_visitedPiPi:
.LFB4402:
	.cfi_startproc
	endbr64
	subq	$104, %rsp
	.cfi_def_cfa_offset 112
	movq	%rdi, 8(%rsp)
	leaq	24(%rsp), %rcx
	leaq	16(%rsp), %rdx
	leaq	44(%rsp), %rsi
	leaq	32(%rsp), %rdi
	movq	%fs:40, %rax
	movq	%rax, 88(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rax
	movl	$1, 40(%rsp)
	movq	%rax, 80(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 32(%rsp)
	movq	%rax, 44(%rsp)
	movl	$1, 52(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L32
	pushq	24(%rsp)
	.cfi_def_cfa_offset 120
	leaq	_Z18initialize_visitedPi(%rip), %rdi
	pushq	24(%rsp)
	.cfi_def_cfa_offset 128
	movq	60(%rsp), %rcx
	movl	68(%rsp), %r8d
	movq	48(%rsp), %rsi
	movl	56(%rsp), %edx
	leaq	96(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 120
	popq	%rdx
	.cfi_def_cfa_offset 112
.L32:
	movq	88(%rsp), %rax
	subq	%fs:40, %rax
	je	.L34
	call	__stack_chk_fail@PLT
.L34:
	addq	$104, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4402:
	.size	_Z38__device_stub__Z18initialize_visitedPiPi, .-_Z38__device_stub__Z18initialize_visitedPiPi
	.globl	_Z18initialize_visitedPi
	.type	_Z18initialize_visitedPi, @function
_Z18initialize_visitedPi:
.LFB4403:
	.cfi_startproc
	endbr64
	jmp	_Z38__device_stub__Z18initialize_visitedPiPi
	.cfi_endproc
.LFE4403:
	.size	_Z18initialize_visitedPi, .-_Z18initialize_visitedPi
	.section	.rodata.str1.1
.LC3:
	.string	"_Z18initialize_visitedPi"
.LC4:
	.string	"_Z19initialize_frontierPii"
.LC5:
	.string	"_Z10bfs_kernelPiS_S_S_iS_S_"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB4405:
	.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	_Z18initialize_visitedPi(%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	_Z19initialize_frontierPii(%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
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC5(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z10bfs_kernelPiS_S_S_iS_S_(%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
.LFE4405:
	.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 <vector>
#include <hip/hip_runtime.h>
#include <hip/hip_fp16.h>


#define MAX_EDGES 128
#define MAX_NODES 32

__global__ void bfs_kernel(int *frontier, int *visited, int *edge_list, int *edge_offsets, int num_nodes, int *new_frontier, int *new_frontier_size) {
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    if (tid < num_nodes && frontier[tid] != -1) {
        int node = frontier[tid];
        for (int i = edge_offsets[node]; i < edge_offsets[node + 1]; i++) {
            int neighbor = edge_list[i];
            if (!visited[neighbor]) {
                atomicExch(&visited[neighbor], 1);
                int index = atomicAdd(new_frontier_size, 1);
                new_frontier[index] = neighbor;
            }
        }
        frontier[tid] = -1; // Mark as processed
    }
}

__global__ void initialize_frontier(int *frontier, int start_node) {
    frontier[start_node] = start_node;
}

__global__ void initialize_visited(int *visited) {
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    if (tid < MAX_NODES) {
        visited[tid] = 0;
    }
}

void bfs(int *edge_list, int *edge_offsets, int start_node, int num_nodes, int num_gpus) {
    int *d_frontier, *d_visited, *d_edge_list, *d_edge_offsets, *d_new_frontier, *d_new_frontier_size;
    int *h_frontier = new int[num_nodes];
    int *h_visited = new int[num_nodes];
    int *h_new_frontier = new int[num_nodes];
    int *h_new_frontier_size = new int[1];

    hipSetDevice(0);
    hipMalloc(&d_frontier, num_nodes * sizeof(int));
    hipMalloc(&d_visited, num_nodes * sizeof(int));
    hipMalloc(&d_edge_list, MAX_EDGES * sizeof(int));
    hipMalloc(&d_edge_offsets, (num_nodes + 1) * sizeof(int));
    hipMalloc(&d_new_frontier, num_nodes * sizeof(int));
    hipMalloc(&d_new_frontier_size, sizeof(int));

    // Initialize host data
    for (int i = 0; i < num_nodes; i++) {
        h_frontier[i] = -1;
        h_visited[i] = 0;
    }
    h_frontier[start_node] = start_node;
    h_visited[start_node] = 1;
    *h_new_frontier_size = 0;

    // Copy data to device
    hipMemcpy(d_frontier, h_frontier, num_nodes * sizeof(int), hipMemcpyHostToDevice);
    hipMemcpy(d_visited, h_visited, num_nodes * sizeof(int), hipMemcpyHostToDevice);
    hipMemcpy(d_edge_list, edge_list, MAX_EDGES * sizeof(int), hipMemcpyHostToDevice);
    hipMemcpy(d_edge_offsets, edge_offsets, (num_nodes + 1) * sizeof(int), hipMemcpyHostToDevice);
    hipMemcpy(d_new_frontier_size, h_new_frontier_size, sizeof(int), hipMemcpyHostToDevice);

    dim3 threadsPerBlock(256);
    dim3 blocksPerGrid((num_nodes + threadsPerBlock.x - 1) / threadsPerBlock.x);

    while (true) {
        // Reset new frontier size
        *h_new_frontier_size = 0;
        hipMemcpy(d_new_frontier_size, h_new_frontier_size, sizeof(int), hipMemcpyHostToDevice);

        // Launch BFS kernel
        bfs_kernel<<<blocksPerGrid, threadsPerBlock>>>(d_frontier, d_visited, d_edge_list, d_edge_offsets, num_nodes, d_new_frontier, d_new_frontier_size);
        hipDeviceSynchronize();

        // Copy new frontier size back to host
        hipMemcpy(h_new_frontier_size, d_new_frontier_size, sizeof(int), hipMemcpyDeviceToHost);

        if (*h_new_frontier_size == 0) break;

        // Copy new frontier to host
        hipMemcpy(h_new_frontier, d_new_frontier, *h_new_frontier_size * sizeof(int), hipMemcpyDeviceToHost);

        // Copy new frontier to old frontier
        hipMemcpy(d_frontier, d_new_frontier, *h_new_frontier_size * sizeof(int), hipMemcpyDeviceToDevice);
    }

    // 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_frontier);
    hipFree(d_visited);
    hipFree(d_edge_list);
    hipFree(d_edge_offsets);
    hipFree(d_new_frontier);
    hipFree(d_new_frontier_size);

    // Free host memory
    delete[] h_frontier;
    delete[] h_visited;
    delete[] h_new_frontier;
    delete[] h_new_frontier_size;
}

int main() {
    // Example graph represented as an edge list
    int edge_list[MAX_EDGES] = {1, 2, 3, 0, 4, 5, 1, 6, 7, 2, 8, 9, 3, 10, 11, 4, 12, 13, 5, 14, 15, 6, 16, 17, 7, 18, 19, 8, 20, 21, 9, 22, 23, 10, 24, 25, 11, 26, 27, 12, 28, 29, 13, 30, 31};
    int edge_offsets[MAX_NODES + 1] = {0, 2, 4, 6, 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 start_node = 0;
    int num_nodes = MAX_NODES;
    int num_gpus = 1;

    bfs(edge_list, edge_offsets, start_node, num_nodes, num_gpus);

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

                                        # End of file scope inline assembly
	.globl	_Z25__device_stub__bfs_kernelPiS_S_S_iS_S_ # -- Begin function _Z25__device_stub__bfs_kernelPiS_S_S_iS_S_
	.type	_Z25__device_stub__bfs_kernelPiS_S_S_iS_S_,@function
_Z25__device_stub__bfs_kernelPiS_S_S_iS_S_: # @_Z25__device_stub__bfs_kernelPiS_S_S_iS_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	$160, %rsp
	.cfi_def_cfa_offset 208
	.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	4(%rsp), %rcx
	movl	%r8d, (%rcx)
	leaq	24(%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	208(%rsp), %rax
	movq	%rax, 48(%rbx)
	leaq	80(%rsp), %r14
	leaq	64(%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	$_Z10bfs_kernelPiS_S_S_iS_S_, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$176, %rsp
	.cfi_adjust_cfa_offset -176
	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	_Z25__device_stub__bfs_kernelPiS_S_S_iS_S_, .Lfunc_end0-_Z25__device_stub__bfs_kernelPiS_S_S_iS_S_
	.cfi_endproc
                                        # -- End function
	.globl	_Z34__device_stub__initialize_frontierPii # -- Begin function _Z34__device_stub__initialize_frontierPii
	.type	_Z34__device_stub__initialize_frontierPii,@function
_Z34__device_stub__initialize_frontierPii: # @_Z34__device_stub__initialize_frontierPii
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$80, %rsp
	.cfi_def_cfa_offset 128
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	24(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	4(%rsp), %rcx
	movl	%esi, (%rcx)
	leaq	64(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%rbx)
	leaq	48(%rsp), %r14
	leaq	32(%rsp), %r15
	leaq	16(%rsp), %r12
	leaq	8(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z19initialize_frontierPii, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$96, %rsp
	.cfi_adjust_cfa_offset -96
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end1:
	.size	_Z34__device_stub__initialize_frontierPii, .Lfunc_end1-_Z34__device_stub__initialize_frontierPii
	.cfi_endproc
                                        # -- End function
	.globl	_Z33__device_stub__initialize_visitedPi # -- Begin function _Z33__device_stub__initialize_visitedPi
	.type	_Z33__device_stub__initialize_visitedPi,@function
_Z33__device_stub__initialize_visitedPi: # @_Z33__device_stub__initialize_visitedPi
	.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	$64, %rsp
	.cfi_def_cfa_offset 112
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	24(%rsp), %rax
	movq	%rdi, (%rax)
	movq	%rsp, %rbx
	movq	%rax, (%rbx)
	leaq	48(%rsp), %r14
	leaq	32(%rsp), %r15
	leaq	16(%rsp), %r12
	leaq	8(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z18initialize_visitedPi, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$80, %rsp
	.cfi_adjust_cfa_offset -80
	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_end2:
	.size	_Z33__device_stub__initialize_visitedPi, .Lfunc_end2-_Z33__device_stub__initialize_visitedPi
	.cfi_endproc
                                        # -- End function
	.globl	_Z3bfsPiS_iii                   # -- Begin function _Z3bfsPiS_iii
	.type	_Z3bfsPiS_iii,@function
_Z3bfsPiS_iii:                          # @_Z3bfsPiS_iii
	.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	$120, %rsp
	.cfi_def_cfa_offset 176
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
                                        # kill: def $ecx killed $ecx def $rcx
	movq	%rcx, 24(%rsp)                  # 8-byte Spill
	movl	%edx, 76(%rsp)                  # 4-byte Spill
	movq	%rsi, 112(%rsp)                 # 8-byte Spill
	movq	%rdi, 104(%rsp)                 # 8-byte Spill
	movslq	%ecx, %rbp
	leaq	(,%rbp,4), %r13
	testl	%ebp, %ebp
	movq	$-1, %rbx
	cmovnsq	%r13, %rbx
	movq	%rbx, %rdi
	callq	_Znam
	movq	%rax, 88(%rsp)                  # 8-byte Spill
	movq	%rbx, %rdi
	callq	_Znam
	movq	%rax, %r14
	movq	%rbx, %rdi
	callq	_Znam
	movq	%rax, %r15
	movl	$4, %edi
	callq	_Znam
	movq	%rax, %r12
	xorl	%edi, %edi
	callq	hipSetDevice
	leaq	48(%rsp), %rdi
	movq	%r13, %rsi
	callq	hipMalloc
	leaq	40(%rsp), %rdi
	movq	%r13, %rsi
	callq	hipMalloc
	leaq	64(%rsp), %rdi
	movl	$512, %esi                      # imm = 0x200
	callq	hipMalloc
	leaq	4(,%rbp,4), %rsi
	leaq	56(%rsp), %rdi
	movq	%rsi, 96(%rsp)                  # 8-byte Spill
	callq	hipMalloc
	leaq	32(%rsp), %rdi
	movq	%r13, 80(%rsp)                  # 8-byte Spill
	movq	%r13, %rsi
	callq	hipMalloc
	leaq	16(%rsp), %rdi
	movl	$4, %esi
	callq	hipMalloc
	movl	24(%rsp), %r13d                 # 4-byte Reload
	testl	%ebp, %ebp
	movq	88(%rsp), %rbp                  # 8-byte Reload
	jle	.LBB3_2
# %bb.1:                                # %.lr.ph.preheader
	leaq	(,%r13,4), %rbx
	movq	%rbp, %rdi
	movl	$255, %esi
	movq	%rbx, %rdx
	callq	memset@PLT
	movq	%r14, %rdi
	xorl	%esi, %esi
	movq	%rbx, %rdx
	callq	memset@PLT
.LBB3_2:                                # %._crit_edge
	movabsq	$4294967296, %rbx               # imm = 0x100000000
	movslq	76(%rsp), %rax                  # 4-byte Folded Reload
	movl	%eax, (%rbp,%rax,4)
	movl	$1, (%r14,%rax,4)
	movl	$0, (%r12)
	movq	48(%rsp), %rdi
	movq	%rbp, %rsi
	movq	80(%rsp), %rbp                  # 8-byte Reload
	movq	%rbp, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	40(%rsp), %rdi
	movq	%r14, %rsi
	movq	%rbp, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	64(%rsp), %rdi
	movl	$512, %edx                      # imm = 0x200
	movq	104(%rsp), %rsi                 # 8-byte Reload
	movl	$1, %ecx
	callq	hipMemcpy
	movq	56(%rsp), %rdi
	movq	112(%rsp), %rsi                 # 8-byte Reload
	movq	96(%rsp), %rdx                  # 8-byte Reload
	movl	$1, %ecx
	callq	hipMemcpy
	movq	16(%rsp), %rdi
	movl	$4, %edx
	movq	%r12, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movq	24(%rsp), %rax                  # 8-byte Reload
	leal	255(%rax), %ebp
	shrl	$8, %ebp
	btsq	$32, %rbp
	addq	$256, %rbx                      # imm = 0x100
.LBB3_3:                                # =>This Inner Loop Header: Depth=1
	movl	$0, (%r12)
	movq	16(%rsp), %rdi
	movl	$4, %edx
	movq	%r12, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movq	%rbp, %rdi
	movl	$1, %esi
	movq	%rbx, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB3_5
# %bb.4:                                #   in Loop: Header=BB3_3 Depth=1
	movq	48(%rsp), %rdi
	movq	40(%rsp), %rsi
	movq	64(%rsp), %rdx
	movq	56(%rsp), %rcx
	movq	32(%rsp), %r9
	movq	16(%rsp), %rax
	movq	%rax, (%rsp)
	movq	24(%rsp), %r8                   # 8-byte Reload
                                        # kill: def $r8d killed $r8d killed $r8
	callq	_Z25__device_stub__bfs_kernelPiS_S_S_iS_S_
.LBB3_5:                                #   in Loop: Header=BB3_3 Depth=1
	callq	hipDeviceSynchronize
	movq	16(%rsp), %rsi
	movl	$4, %edx
	movq	%r12, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movslq	(%r12), %rdx
	testq	%rdx, %rdx
	je	.LBB3_7
# %bb.6:                                #   in Loop: Header=BB3_3 Depth=1
	movq	32(%rsp), %rsi
	shlq	$2, %rdx
	movq	%r15, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movq	48(%rsp), %rdi
	movq	32(%rsp), %rsi
	movslq	(%r12), %rdx
	shlq	$2, %rdx
	movl	$3, %ecx
	callq	hipMemcpy
	jmp	.LBB3_3
.LBB3_7:
	movq	40(%rsp), %rsi
	movq	%r14, %rdi
	movq	80(%rsp), %rdx                  # 8-byte Reload
	movl	$2, %ecx
	callq	hipMemcpy
	movl	$_ZSt4cout, %ebp
	movl	$_ZSt4cout, %edi
	movl	$.L.str, %esi
	movl	$15, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	cmpl	$0, 24(%rsp)                    # 4-byte Folded Reload
	jle	.LBB3_12
# %bb.8:                                # %.lr.ph63.preheader
	xorl	%ebx, %ebx
.LBB3_9:                                # %.lr.ph63
                                        # =>This Inner Loop Header: Depth=1
	cmpl	$0, (%r14,%rbx,4)
	je	.LBB3_11
# %bb.10:                               #   in Loop: Header=BB3_9 Depth=1
	movl	$_ZSt4cout, %edi
	movl	%ebx, %esi
	callq	_ZNSolsEi
	movl	$.L.str.1, %esi
	movl	$1, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
.LBB3_11:                               #   in Loop: Header=BB3_9 Depth=1
	incq	%rbx
	cmpq	%rbx, %r13
	jne	.LBB3_9
.LBB3_12:                               # %._crit_edge64
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%rbp, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movq	48(%rsp), %rdi
	callq	hipFree
	movq	40(%rsp), %rdi
	callq	hipFree
	movq	64(%rsp), %rdi
	callq	hipFree
	movq	56(%rsp), %rdi
	callq	hipFree
	movq	32(%rsp), %rdi
	callq	hipFree
	movq	16(%rsp), %rdi
	callq	hipFree
	movq	88(%rsp), %rdi                  # 8-byte Reload
	callq	_ZdaPv
	movq	%r14, %rdi
	callq	_ZdaPv
	movq	%r15, %rdi
	callq	_ZdaPv
	movq	%r12, %rdi
	callq	_ZdaPv
	addq	$120, %rsp
	.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_end3:
	.size	_Z3bfsPiS_iii, .Lfunc_end3-_Z3bfsPiS_iii
	.cfi_endproc
                                        # -- End function
	.globl	main                            # -- Begin function main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%r14
	.cfi_def_cfa_offset 16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	subq	$664, %rsp                      # imm = 0x298
	.cfi_def_cfa_offset 688
	.cfi_offset %rbx, -24
	.cfi_offset %r14, -16
	leaq	144(%rsp), %rbx
	movl	$.L__const.main.edge_list, %esi
	movl	$512, %edx                      # imm = 0x200
	movq	%rbx, %rdi
	callq	memcpy@PLT
	movq	%rsp, %r14
	movl	$.L__const.main.edge_offsets, %esi
	movl	$132, %edx
	movq	%r14, %rdi
	callq	memcpy@PLT
	movq	%rbx, %rdi
	movq	%r14, %rsi
	xorl	%edx, %edx
	movl	$32, %ecx
	callq	_Z3bfsPiS_iii
	xorl	%eax, %eax
	addq	$664, %rsp                      # imm = 0x298
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end4:
	.size	main, .Lfunc_end4-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	.LBB5_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB5_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z10bfs_kernelPiS_S_S_iS_S_, %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	$_Z19initialize_frontierPii, %esi
	movl	$.L__unnamed_2, %edx
	movl	$.L__unnamed_2, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z18initialize_visitedPi, %esi
	movl	$.L__unnamed_3, %edx
	movl	$.L__unnamed_3, %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_end5:
	.size	__hip_module_ctor, .Lfunc_end5-__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	.LBB6_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
.LBB6_2:
	retq
.Lfunc_end6:
	.size	__hip_module_dtor, .Lfunc_end6-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_Z10bfs_kernelPiS_S_S_iS_S_,@object # @_Z10bfs_kernelPiS_S_S_iS_S_
	.section	.rodata,"a",@progbits
	.globl	_Z10bfs_kernelPiS_S_S_iS_S_
	.p2align	3, 0x0
_Z10bfs_kernelPiS_S_S_iS_S_:
	.quad	_Z25__device_stub__bfs_kernelPiS_S_S_iS_S_
	.size	_Z10bfs_kernelPiS_S_S_iS_S_, 8

	.type	_Z19initialize_frontierPii,@object # @_Z19initialize_frontierPii
	.globl	_Z19initialize_frontierPii
	.p2align	3, 0x0
_Z19initialize_frontierPii:
	.quad	_Z34__device_stub__initialize_frontierPii
	.size	_Z19initialize_frontierPii, 8

	.type	_Z18initialize_visitedPi,@object # @_Z18initialize_visitedPi
	.globl	_Z18initialize_visitedPi
	.p2align	3, 0x0
_Z18initialize_visitedPi:
	.quad	_Z33__device_stub__initialize_visitedPi
	.size	_Z18initialize_visitedPi, 8

	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"Visited nodes: "
	.size	.L.str, 16

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

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

	.type	.L__const.main.edge_offsets,@object # @__const.main.edge_offsets
	.p2align	4, 0x0
.L__const.main.edge_offsets:
	.long	0                               # 0x0
	.long	2                               # 0x2
	.long	4                               # 0x4
	.long	6                               # 0x6
	.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
	.size	.L__const.main.edge_offsets, 132

	.type	.L__unnamed_1,@object           # @0
	.section	.rodata.str1.1,"aMS",@progbits,1
.L__unnamed_1:
	.asciz	"_Z10bfs_kernelPiS_S_S_iS_S_"
	.size	.L__unnamed_1, 28

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"_Z19initialize_frontierPii"
	.size	.L__unnamed_2, 27

	.type	.L__unnamed_3,@object           # @2
.L__unnamed_3:
	.asciz	"_Z18initialize_visitedPi"
	.size	.L__unnamed_3, 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 _Z25__device_stub__bfs_kernelPiS_S_S_iS_S_
	.addrsig_sym _Z34__device_stub__initialize_frontierPii
	.addrsig_sym _Z33__device_stub__initialize_visitedPi
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z10bfs_kernelPiS_S_S_iS_S_
	.addrsig_sym _Z19initialize_frontierPii
	.addrsig_sym _Z18initialize_visitedPi
	.addrsig_sym _ZSt4cout
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z10bfs_kernelPiS_S_S_iS_S_ ; -- Begin function _Z10bfs_kernelPiS_S_S_iS_S_
	.globl	_Z10bfs_kernelPiS_S_S_iS_S_
	.p2align	8
	.type	_Z10bfs_kernelPiS_S_S_iS_S_,@function
_Z10bfs_kernelPiS_S_S_iS_S_:            ; @_Z10bfs_kernelPiS_S_S_iS_S_
; %bb.0:
	s_clause 0x1
	s_load_b32 s2, s[0:1], 0x44
	s_load_b32 s3, s[0:1], 0x20
	s_waitcnt lgkmcnt(0)
	s_and_b32 s2, s2, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[1:2], null, s15, s2, v[0:1]
	s_mov_b32 s2, exec_lo
	v_cmpx_gt_i32_e64 s3, v1
	s_cbranch_execz .LBB0_10
; %bb.1:
	s_load_b256 s[4:11], s[0:1], 0x0
	v_ashrrev_i32_e32 v2, 31, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[0:1], 2, v[1:2]
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v0, vcc_lo, s4, v0
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v1, vcc_lo, s5, v1, vcc_lo
	global_load_b32 v2, v[0:1], off
	s_waitcnt vmcnt(0)
	v_cmp_ne_u32_e32 vcc_lo, -1, v2
	s_and_b32 exec_lo, exec_lo, vcc_lo
	s_cbranch_execz .LBB0_10
; %bb.2:
	v_ashrrev_i32_e32 v3, 31, v2
	s_mov_b32 s2, exec_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[2:3], 2, v[2:3]
	v_add_co_u32 v2, vcc_lo, s10, v2
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v3, vcc_lo, s11, v3, vcc_lo
	global_load_b64 v[4:5], v[2:3], off
	s_waitcnt vmcnt(0)
	v_cmpx_lt_i32_e64 v4, v5
	s_cbranch_execz .LBB0_9
; %bb.3:                                ; %.lr.ph.preheader
	s_load_b128 s[12:15], s[0:1], 0x28
	v_ashrrev_i32_e32 v7, 31, v4
	v_dual_mov_b32 v6, v4 :: v_dual_mov_b32 v11, 1
	v_mov_b32_e32 v12, 0
	s_mov_b32 s1, 0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[6:7], 2, v[6:7]
	v_add_co_u32 v6, vcc_lo, s8, v6
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v7, vcc_lo, s9, v7, vcc_lo
.LBB0_4:                                ; %.lr.ph
                                        ; =>This Inner Loop Header: Depth=1
	global_load_b32 v8, v[6:7], off
	s_mov_b32 s0, exec_lo
	s_waitcnt vmcnt(0)
	v_ashrrev_i32_e32 v9, 31, v8
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[9:10], 2, v[8:9]
	v_add_co_u32 v9, vcc_lo, s6, v9
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v10, vcc_lo, s7, v10, vcc_lo
	global_load_b32 v13, v[9:10], off
	s_waitcnt vmcnt(0)
	v_cmpx_eq_u32_e32 0, v13
	s_cbranch_execz .LBB0_8
; %bb.5:                                ;   in Loop: Header=BB0_4 Depth=1
	global_atomic_swap_b32 v[9:10], v11, off
	s_mov_b32 s4, exec_lo
	s_mov_b32 s3, exec_lo
	v_mbcnt_lo_u32_b32 v5, s4, 0
                                        ; implicit-def: $vgpr9
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_eq_u32_e32 0, v5
	s_cbranch_execz .LBB0_7
; %bb.6:                                ;   in Loop: Header=BB0_4 Depth=1
	s_bcnt1_i32_b32 s4, s4
	s_delay_alu instid0(SALU_CYCLE_1)
	v_mov_b32_e32 v9, s4
	s_waitcnt lgkmcnt(0)
	global_atomic_add_u32 v9, v12, v9, s[14:15] glc
.LBB0_7:                                ;   in Loop: Header=BB0_4 Depth=1
	s_or_b32 exec_lo, exec_lo, s3
	s_waitcnt vmcnt(0)
	v_readfirstlane_b32 s3, v9
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v9, s3, v5
	v_ashrrev_i32_e32 v10, 31, v9
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[9:10], 2, v[9:10]
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v9, vcc_lo, s12, v9
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v10, vcc_lo, s13, v10, vcc_lo
	global_store_b32 v[9:10], v8, off
	global_load_b32 v5, v[2:3], off offset:4
.LBB0_8:                                ; %Flow
                                        ;   in Loop: Header=BB0_4 Depth=1
	s_or_b32 exec_lo, exec_lo, s0
	v_add_nc_u32_e32 v4, 1, v4
	v_add_co_u32 v6, s0, v6, 4
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_add_co_ci_u32_e64 v7, s0, 0, v7, s0
	s_waitcnt vmcnt(0)
	v_cmp_ge_i32_e32 vcc_lo, v4, 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_4
.LBB0_9:                                ; %Flow54
	s_or_b32 exec_lo, exec_lo, s2
	v_mov_b32_e32 v2, -1
	global_store_b32 v[0:1], v2, off
.LBB0_10:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z10bfs_kernelPiS_S_S_iS_S_
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 312
		.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 14
		.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	_Z10bfs_kernelPiS_S_S_iS_S_, .Lfunc_end0-_Z10bfs_kernelPiS_S_S_iS_S_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 524
; NumSgprs: 18
; NumVgprs: 14
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 14
; 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
	.protected	_Z19initialize_frontierPii ; -- Begin function _Z19initialize_frontierPii
	.globl	_Z19initialize_frontierPii
	.p2align	8
	.type	_Z19initialize_frontierPii,@function
_Z19initialize_frontierPii:             ; @_Z19initialize_frontierPii
; %bb.0:
	s_clause 0x1
	s_load_b32 s2, s[0:1], 0x8
	s_load_b64 s[0:1], s[0:1], 0x0
	s_waitcnt lgkmcnt(0)
	v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v1, s2
	s_ashr_i32 s3, s2, 31
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_lshl_b64 s[4:5], s[2:3], 2
	s_add_u32 s0, s0, s4
	s_addc_u32 s1, s1, s5
	global_store_b32 v0, v1, s[0:1]
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z19initialize_frontierPii
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 12
		.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 2
		.amdhsa_next_free_sgpr 6
		.amdhsa_reserve_vcc 0
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end1:
	.size	_Z19initialize_frontierPii, .Lfunc_end1-_Z19initialize_frontierPii
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 72
; NumSgprs: 6
; NumVgprs: 2
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 0
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 6
; NumVGPRsForWavesPerEU: 2
; 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	_Z18initialize_visitedPi ; -- Begin function _Z18initialize_visitedPi
	.globl	_Z18initialize_visitedPi
	.p2align	8
	.type	_Z18initialize_visitedPi,@function
_Z18initialize_visitedPi:               ; @_Z18initialize_visitedPi
; %bb.0:
	s_load_b32 s2, s[0:1], 0x14
	s_waitcnt lgkmcnt(0)
	s_and_b32 s2, s2, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[1:2], null, s15, s2, v[0:1]
	s_mov_b32 s2, exec_lo
	v_cmpx_gt_i32_e32 32, v1
	s_cbranch_execz .LBB2_2
; %bb.1:
	s_load_b64 s[0:1], s[0:1], 0x0
	v_ashrrev_i32_e32 v2, 31, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_lshlrev_b64 v[0:1], 2, v[1:2]
	v_mov_b32_e32 v2, 0
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v0, vcc_lo, s0, v0
	s_delay_alu instid0(VALU_DEP_3)
	v_add_co_ci_u32_e32 v1, vcc_lo, s1, v1, vcc_lo
	global_store_b32 v[0:1], v2, off
.LBB2_2:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z18initialize_visitedPi
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 264
		.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 3
		.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_end2:
	.size	_Z18initialize_visitedPi, .Lfunc_end2-_Z18initialize_visitedPi
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 112
; NumSgprs: 18
; NumVgprs: 3
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 3
; 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
      - .address_space:  global
        .offset:         24
        .size:           8
        .value_kind:     global_buffer
      - .offset:         32
        .size:           4
        .value_kind:     by_value
      - .address_space:  global
        .offset:         40
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         48
        .size:           8
        .value_kind:     global_buffer
      - .offset:         56
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         60
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         64
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         68
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         70
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         72
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         74
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         76
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         78
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         96
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         104
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         112
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         120
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 312
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z10bfs_kernelPiS_S_S_iS_S_
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z10bfs_kernelPiS_S_S_iS_S_.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     14
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .offset:         8
        .size:           4
        .value_kind:     by_value
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 12
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z19initialize_frontierPii
    .private_segment_fixed_size: 0
    .sgpr_count:     6
    .sgpr_spill_count: 0
    .symbol:         _Z19initialize_frontierPii.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     2
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .offset:         8
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         12
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         16
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         20
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         22
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         24
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         26
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         28
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         30
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         48
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         56
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         64
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         72
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 264
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z18initialize_visitedPi
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z18initialize_visitedPi.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     3
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
