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

// Define constants
#define MAX_NODES 128
#define MAX_EDGES 512 // Assuming each node can have up to 4 edges on average

// Kernel to initialize the frontier
__global__ void initFrontier(bool* visited, int* frontier, int* frontierSize) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx == 0) {
        frontier[0] = 0; // Start BFS from node 0
        *frontierSize = 1;
        visited[0] = true;
    }
}

// Kernel to perform BFS
__global__ void bfsKernel(int* adjacencyList, int* edgeList, bool* visited, int* frontier, int* newFrontier, int* frontierSize, int* newFrontierSize, int numNodes) {
    extern __shared__ int sharedFrontier[];

    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    int numThreads = blockDim.x * gridDim.x;

    // Load frontier into shared memory
    int numFrontier = *frontierSize;
    int start = tid;
    while (start < numFrontier) {
        sharedFrontier[start] = frontier[start];
        start += numThreads;
    }
    __syncthreads();

    // Each thread processes a node in the frontier
    int node = sharedFrontier[tid];
    if (tid < numFrontier) {
        int edgeStart = adjacencyList[node];
        int edgeEnd = adjacencyList[node + 1];
        for (int e = edgeStart; e < edgeEnd; ++e) {
            int neighbor = edgeList[e];
            if (!visited[neighbor]) {
                int idx = atomicAdd(newFrontierSize, 1);
                newFrontier[idx] = neighbor;
                visited[neighbor] = true;
            }
        }
    }
}

int main() {
    // Host data
    std::vector<int> adjacencyList = {0, 2, 4, 7, 9, 11, 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, 66, 68, 70, 72, 74, 76, 78, 80, 82, 84, 86, 88, 90, 92, 94, 96, 98, 100, 102, 104, 106, 108, 110, 112, 114, 116, 118, 120, 122, 124, 126, 128};
    std::vector<int> edgeList(MAX_EDGES);
    for (int i = 0; i < MAX_EDGES; ++i) {
        edgeList[i] = i % MAX_NODES; // Randomly assign edges for demonstration
    }

    // Device data
    int* d_adjacencyList;
    int* d_edgeList;
    bool* d_visited;
    int* d_frontier;
    int* d_newFrontier;
    int* d_frontierSize;
    int* d_newFrontierSize;

    // Allocate memory
    cudaMalloc(&d_adjacencyList, (MAX_NODES + 1) * sizeof(int));
    cudaMalloc(&d_edgeList, MAX_EDGES * sizeof(int));
    cudaMalloc(&d_visited, MAX_NODES * sizeof(bool));
    cudaMalloc(&d_frontier, MAX_NODES * sizeof(int));
    cudaMalloc(&d_newFrontier, MAX_NODES * sizeof(int));
    cudaMalloc(&d_frontierSize, sizeof(int));
    cudaMalloc(&d_newFrontierSize, sizeof(int));

    // Copy data to device
    cudaMemcpy(d_adjacencyList, adjacencyList.data(), (MAX_NODES + 1) * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(d_edgeList, edgeList.data(), MAX_EDGES * sizeof(int), cudaMemcpyHostToDevice);

    // Initialize visited array to false
    cudaMemset(d_visited, 0, MAX_NODES * sizeof(bool));

    // Initialize frontier
    initFrontier<<<1, 1>>>(d_visited, d_frontier, d_frontierSize);
    cudaDeviceSynchronize();

    // BFS parameters
    int blockSize = 256;
    int gridSize = (MAX_NODES + blockSize - 1) / blockSize;
    int maxIterations = 10; // Arbitrary number of iterations for demonstration

    // Perform BFS
    for (int i = 0; i < maxIterations; ++i) {
        cudaMemset(d_newFrontierSize, 0, sizeof(int));
        bfsKernel<<<gridSize, blockSize, blockSize * sizeof(int)>>>(d_adjacencyList, d_edgeList, d_visited, d_frontier, d_newFrontier, d_frontierSize, d_newFrontierSize, MAX_NODES);
        cudaDeviceSynchronize();

        // Swap frontiers
        int* temp = d_frontier;
        d_frontier = d_newFrontier;
        d_newFrontier = temp;
        temp = d_frontierSize;
        d_frontierSize = d_newFrontierSize;
        d_newFrontierSize = temp;
    }

    // Free memory
    cudaFree(d_adjacencyList);
    cudaFree(d_edgeList);
    cudaFree(d_visited);
    cudaFree(d_frontier);
    cudaFree(d_newFrontier);
    cudaFree(d_frontierSize);
    cudaFree(d_newFrontierSize);

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

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB4001:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE4001:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z37__device_stub__Z12initFrontierPbPiS0_PbPiS0_
	.type	_Z37__device_stub__Z12initFrontierPbPiS0_PbPiS0_, @function
_Z37__device_stub__Z12initFrontierPbPiS0_PbPiS0_:
.LFB4023:
	.cfi_startproc
	endbr64
	subq	$136, %rsp
	.cfi_def_cfa_offset 144
	movq	%rdi, 24(%rsp)
	leaq	40(%rsp), %rcx
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	32(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 120(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 96(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 112(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L2
	pushq	40(%rsp)
	.cfi_def_cfa_offset 152
	leaq	_Z12initFrontierPbPiS0_(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 160
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	112(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 152
	popq	%rdx
	.cfi_def_cfa_offset 144
.L2:
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L4
	call	__stack_chk_fail@PLT
.L4:
	addq	$136, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4023:
	.size	_Z37__device_stub__Z12initFrontierPbPiS0_PbPiS0_, .-_Z37__device_stub__Z12initFrontierPbPiS0_PbPiS0_
	.globl	_Z12initFrontierPbPiS0_
	.type	_Z12initFrontierPbPiS0_, @function
_Z12initFrontierPbPiS0_:
.LFB4024:
	.cfi_startproc
	endbr64
	jmp	_Z37__device_stub__Z12initFrontierPbPiS0_PbPiS0_
	.cfi_endproc
.LFE4024:
	.size	_Z12initFrontierPbPiS0_, .-_Z12initFrontierPbPiS0_
	.globl	_Z41__device_stub__Z9bfsKernelPiS_PbS_S_S_S_iPiS_PbS_S_S_S_i
	.type	_Z41__device_stub__Z9bfsKernelPiS_PbS_S_S_S_iPiS_PbS_S_S_S_i, @function
_Z41__device_stub__Z9bfsKernelPiS_PbS_S_S_S_iPiS_PbS_S_S_S_i:
.LFB4025:
	.cfi_startproc
	endbr64
	subq	$216, %rsp
	.cfi_def_cfa_offset 224
	movq	224(%rsp), %rax
	movq	%rdi, 56(%rsp)
	leaq	88(%rsp), %rdi
	movq	%rsi, 48(%rsp)
	leaq	100(%rsp), %rsi
	movq	%rdx, 40(%rsp)
	leaq	72(%rsp), %rdx
	movq	%rcx, 32(%rsp)
	leaq	80(%rsp), %rcx
	movq	%r8, 24(%rsp)
	movq	%r9, 16(%rsp)
	movq	%rax, 8(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 200(%rsp)
	xorl	%eax, %eax
	leaq	56(%rsp), %rax
	movl	$1, 96(%rsp)
	movq	%rax, 136(%rsp)
	leaq	48(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 160(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 168(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 176(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 184(%rsp)
	leaq	232(%rsp), %rax
	movq	%rax, 192(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 88(%rsp)
	movq	%rax, 100(%rsp)
	movl	$1, 108(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L8
	pushq	80(%rsp)
	.cfi_def_cfa_offset 232
	leaq	_Z9bfsKernelPiS_PbS_S_S_S_i(%rip), %rdi
	pushq	80(%rsp)
	.cfi_def_cfa_offset 240
	movq	116(%rsp), %rcx
	movl	124(%rsp), %r8d
	movq	104(%rsp), %rsi
	movl	112(%rsp), %edx
	leaq	152(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 232
	popq	%rdx
	.cfi_def_cfa_offset 224
.L8:
	movq	200(%rsp), %rax
	subq	%fs:40, %rax
	je	.L10
	call	__stack_chk_fail@PLT
.L10:
	addq	$216, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4025:
	.size	_Z41__device_stub__Z9bfsKernelPiS_PbS_S_S_S_iPiS_PbS_S_S_S_i, .-_Z41__device_stub__Z9bfsKernelPiS_PbS_S_S_S_iPiS_PbS_S_S_S_i
	.globl	_Z9bfsKernelPiS_PbS_S_S_S_i
	.type	_Z9bfsKernelPiS_PbS_S_S_S_i, @function
_Z9bfsKernelPiS_PbS_S_S_S_i:
.LFB4026:
	.cfi_startproc
	endbr64
	jmp	_Z41__device_stub__Z9bfsKernelPiS_PbS_S_S_S_iPiS_PbS_S_S_S_i
	.cfi_endproc
.LFE4026:
	.size	_Z9bfsKernelPiS_PbS_S_S_S_i, .-_Z9bfsKernelPiS_PbS_S_S_S_i
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC1:
	.string	"_Z9bfsKernelPiS_PbS_S_S_S_i"
.LC2:
	.string	"_Z12initFrontierPbPiS0_"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB4028:
	.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	.LC1(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z9bfsKernelPiS_PbS_S_S_S_i(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC2(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z12initFrontierPbPiS0_(%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
.LFE4028:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.text._ZNSt12_Vector_baseIiSaIiEED2Ev,"axG",@progbits,_ZNSt12_Vector_baseIiSaIiEED5Ev,comdat
	.align 2
	.weak	_ZNSt12_Vector_baseIiSaIiEED2Ev
	.type	_ZNSt12_Vector_baseIiSaIiEED2Ev, @function
_ZNSt12_Vector_baseIiSaIiEED2Ev:
.LFB4502:
	.cfi_startproc
	endbr64
	movq	(%rdi), %rax
	testq	%rax, %rax
	je	.L15
	movq	16(%rdi), %rsi
	movq	%rax, %rdi
	subq	%rax, %rsi
	jmp	_ZdlPvm@PLT
.L15:
	ret
	.cfi_endproc
.LFE4502:
	.size	_ZNSt12_Vector_baseIiSaIiEED2Ev, .-_ZNSt12_Vector_baseIiSaIiEED2Ev
	.weak	_ZNSt12_Vector_baseIiSaIiEED1Ev
	.set	_ZNSt12_Vector_baseIiSaIiEED1Ev,_ZNSt12_Vector_baseIiSaIiEED2Ev
	.section	.rodata
	.align 32
.LC0:
	.long	0
	.long	2
	.long	4
	.long	7
	.long	9
	.long	11
	.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
	.long	66
	.long	68
	.long	70
	.long	72
	.long	74
	.long	76
	.long	78
	.long	80
	.long	82
	.long	84
	.long	86
	.long	88
	.long	90
	.long	92
	.long	94
	.long	96
	.long	98
	.long	100
	.long	102
	.long	104
	.long	106
	.long	108
	.long	110
	.long	112
	.long	114
	.long	116
	.long	118
	.long	120
	.long	122
	.long	124
	.long	126
	.long	128
	.section	.text.startup
	.globl	main
	.type	main, @function
main:
.LFB3998:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA3998
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	xorl	%r8d, %r8d
	leaq	.LC0(%rip), %rsi
	movl	$64, %ecx
	movq	%rsp, %rbp
	.cfi_def_cfa_register 6
	pushq	%r14
	pushq	%rbx
	leaq	-280(%rbp), %rdi
	subq	$400, %rsp
	.cfi_offset 14, -24
	.cfi_offset 3, -32
	movq	%fs:40, %rax
	movq	%rax, -24(%rbp)
	xorl	%eax, %eax
	movq	%r8, -328(%rbp)
	rep movsl
	movl	$256, %edi
	movq	%r8, -320(%rbp)
	movq	%r8, -312(%rbp)
.LEHB0:
	call	_Znwm@PLT
.LEHE0:
	movq	%rax, %r14
	leaq	-280(%rbp), %rbx
	movl	$64, %ecx
	movq	%rax, -328(%rbp)
	movq	%r14, %rdi
	movq	%rbx, %rsi
	addq	$256, %rax
	rep movsl
	xorl	%edi, %edi
	movq	%rax, -312(%rbp)
	movq	%rdi, -296(%rbp)
	movq	%rdi, -288(%rbp)
	movl	$2048, %edi
	movq	%rax, -320(%rbp)
.LEHB1:
	call	_Znwm@PLT
.LEHE1:
	leaq	2048(%rax), %rdx
	xorl	%esi, %esi
	movq	%rax, -304(%rbp)
	movq	%rax, %rbx
	movl	%esi, (%rax)
	xorl	%eax, %eax
	movq	%rdx, -288(%rbp)
	jmp	.L21
.L37:
	endbr64
	jmp	.L45
.L21:
	addq	$4, %rax
	cmpq	$2048, %rax
	je	.L46
	xorl	%ecx, %ecx
	movl	%ecx, (%rbx,%rax)
	jmp	.L21
.L46:
	movq	%rdx, -296(%rbp)
	xorl	%eax, %eax
.L23:
	movl	%eax, %edx
	andl	$127, %edx
	movl	%edx, (%rbx,%rax,4)
	incq	%rax
	cmpq	$512, %rax
	jne	.L23
	leaq	-408(%rbp), %rdi
	movl	$516, %esi
.LEHB2:
	call	cudaMalloc@PLT
	leaq	-400(%rbp), %rdi
	movl	$2048, %esi
	call	cudaMalloc@PLT
	leaq	-392(%rbp), %rdi
	movl	$128, %esi
	call	cudaMalloc@PLT
	leaq	-384(%rbp), %rdi
	movl	$512, %esi
	call	cudaMalloc@PLT
	leaq	-376(%rbp), %rdi
	movl	$512, %esi
	call	cudaMalloc@PLT
	leaq	-368(%rbp), %rdi
	movl	$4, %esi
	call	cudaMalloc@PLT
	leaq	-360(%rbp), %rdi
	movl	$4, %esi
	call	cudaMalloc@PLT
	movq	-408(%rbp), %rdi
	movl	$1, %ecx
	movl	$516, %edx
	movq	%r14, %rsi
	call	cudaMemcpy@PLT
	movq	-400(%rbp), %rdi
	movl	$1, %ecx
	movl	$2048, %edx
	movq	%rbx, %rsi
	call	cudaMemcpy@PLT
	movq	-392(%rbp), %rdi
	movl	$128, %edx
	xorl	%esi, %esi
	call	cudaMemset@PLT
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$1, %ecx
	movl	$1, %esi
	movabsq	$4294967297, %rdi
	movl	$1, -332(%rbp)
	movq	%rdi, %rdx
	movq	%rdi, -340(%rbp)
	movq	%rdi, -352(%rbp)
	movl	$1, -344(%rbp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	je	.L24
.L26:
	call	cudaDeviceSynchronize@PLT
	movl	$10, %ebx
	jmp	.L25
.L24:
	movq	-368(%rbp), %rdx
	movq	-384(%rbp), %rsi
	movq	-392(%rbp), %rdi
	call	_Z37__device_stub__Z12initFrontierPbPiS0_PbPiS0_
	jmp	.L26
.L48:
	movl	$16777217, %eax
	xorl	%r9d, %r9d
	movl	$1, %ecx
	movabsq	$4294967297, %rdi
	salq	$8, %rax
	movl	$1024, %r8d
	movl	$1, %esi
	movl	$1, -332(%rbp)
	movq	%rax, %rdx
	movq	%rax, -340(%rbp)
	movq	%rdi, -352(%rbp)
	movl	$1, -344(%rbp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	je	.L27
.L29:
	call	cudaDeviceSynchronize@PLT
	movq	-384(%rbp), %rax
	movq	-376(%rbp), %rdx
	movq	%rax, -376(%rbp)
	movq	-368(%rbp), %rax
	movq	%rdx, -384(%rbp)
	movq	-360(%rbp), %rdx
	movq	%rax, -360(%rbp)
	movq	%rdx, -368(%rbp)
	decl	%ebx
	je	.L47
.L25:
	movq	-360(%rbp), %rdi
	movl	$4, %edx
	xorl	%esi, %esi
	call	cudaMemset@PLT
	jmp	.L48
.L27:
	pushq	$128
	movq	-392(%rbp), %rdx
	pushq	-360(%rbp)
	movq	-368(%rbp), %r9
	movq	-376(%rbp), %r8
	movq	-384(%rbp), %rcx
	movq	-400(%rbp), %rsi
	movq	-408(%rbp), %rdi
	.cfi_escape 0x2e,0x10
	call	_Z41__device_stub__Z9bfsKernelPiS_PbS_S_S_S_iPiS_PbS_S_S_S_i
	popq	%rax
	popq	%rdx
	jmp	.L29
.L47:
	movq	-408(%rbp), %rdi
	.cfi_escape 0x2e,0
	call	cudaFree@PLT
	movq	-400(%rbp), %rdi
	call	cudaFree@PLT
	movq	-392(%rbp), %rdi
	call	cudaFree@PLT
	movq	-384(%rbp), %rdi
	call	cudaFree@PLT
	movq	-376(%rbp), %rdi
	call	cudaFree@PLT
	movq	-368(%rbp), %rdi
	call	cudaFree@PLT
	movq	-360(%rbp), %rdi
	call	cudaFree@PLT
.LEHE2:
	leaq	-304(%rbp), %rdi
	call	_ZNSt12_Vector_baseIiSaIiEED2Ev
	leaq	-328(%rbp), %rdi
	call	_ZNSt12_Vector_baseIiSaIiEED2Ev
	movq	-24(%rbp), %rax
	subq	%fs:40, %rax
	je	.L33
	jmp	.L42
.L36:
	endbr64
	movq	%rax, %rbx
.L30:
	leaq	-304(%rbp), %rdi
	call	_ZNSt12_Vector_baseIiSaIiEED2Ev
	jmp	.L31
.L35:
	endbr64
.L45:
	movq	%rax, %rbx
.L31:
	leaq	-328(%rbp), %rdi
	call	_ZNSt12_Vector_baseIiSaIiEED2Ev
	movq	-24(%rbp), %rax
	subq	%fs:40, %rax
	jne	.L42
	movq	%rbx, %rdi
.LEHB3:
	call	_Unwind_Resume@PLT
.LEHE3:
.L42:
	call	__stack_chk_fail@PLT
.L33:
	leaq	-16(%rbp), %rsp
	xorl	%eax, %eax
	popq	%rbx
	popq	%r14
	popq	%rbp
	.cfi_def_cfa 7, 8
	ret
	.cfi_endproc
.LFE3998:
	.globl	__gxx_personality_v0
	.section	.gcc_except_table,"a",@progbits
.LLSDA3998:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE3998-.LLSDACSB3998
.LLSDACSB3998:
	.uleb128 .LEHB0-.LFB3998
	.uleb128 .LEHE0-.LEHB0
	.uleb128 .L37-.LFB3998
	.uleb128 0
	.uleb128 .LEHB1-.LFB3998
	.uleb128 .LEHE1-.LEHB1
	.uleb128 .L35-.LFB3998
	.uleb128 0
	.uleb128 .LEHB2-.LFB3998
	.uleb128 .LEHE2-.LEHB2
	.uleb128 .L36-.LFB3998
	.uleb128 0
	.uleb128 .LEHB3-.LFB3998
	.uleb128 .LEHE3-.LEHB3
	.uleb128 0
	.uleb128 0
.LLSDACSE3998:
	.section	.text.startup
	.size	main, .-main
	.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
	.hidden	DW.ref.__gxx_personality_v0
	.weak	DW.ref.__gxx_personality_v0
	.section	.data.rel.local.DW.ref.__gxx_personality_v0,"awG",@progbits,DW.ref.__gxx_personality_v0,comdat
	.align 8
	.type	DW.ref.__gxx_personality_v0, @object
	.size	DW.ref.__gxx_personality_v0, 8
DW.ref.__gxx_personality_v0:
	.quad	__gxx_personality_v0
	.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>


// Define constants
#define MAX_NODES 128
#define MAX_EDGES 512 // Assuming each node can have up to 4 edges on average

// Kernel to initialize the frontier
__global__ void initFrontier(bool* visited, int* frontier, int* frontierSize) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx == 0) {
        frontier[0] = 0; // Start BFS from node 0
        *frontierSize = 1;
        visited[0] = true;
    }
}

// Kernel to perform BFS
__global__ void bfsKernel(int* adjacencyList, int* edgeList, bool* visited, int* frontier, int* newFrontier, int* frontierSize, int* newFrontierSize, int numNodes) {
    extern __shared__ int sharedFrontier[];

    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    int numThreads = blockDim.x * gridDim.x;

    // Load frontier into shared memory
    int numFrontier = *frontierSize;
    int start = tid;
    while (start < numFrontier) {
        sharedFrontier[start] = frontier[start];
        start += numThreads;
    }
    __syncthreads();

    // Each thread processes a node in the frontier
    int node = sharedFrontier[tid];
    if (tid < numFrontier) {
        int edgeStart = adjacencyList[node];
        int edgeEnd = adjacencyList[node + 1];
        for (int e = edgeStart; e < edgeEnd; ++e) {
            int neighbor = edgeList[e];
            if (!visited[neighbor]) {
                int idx = atomicAdd(newFrontierSize, 1);
                newFrontier[idx] = neighbor;
                visited[neighbor] = true;
            }
        }
    }
}

int main() {
    // Host data
    std::vector<int> adjacencyList = {0, 2, 4, 7, 9, 11, 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, 66, 68, 70, 72, 74, 76, 78, 80, 82, 84, 86, 88, 90, 92, 94, 96, 98, 100, 102, 104, 106, 108, 110, 112, 114, 116, 118, 120, 122, 124, 126, 128};
    std::vector<int> edgeList(MAX_EDGES);
    for (int i = 0; i < MAX_EDGES; ++i) {
        edgeList[i] = i % MAX_NODES; // Randomly assign edges for demonstration
    }

    // Device data
    int* d_adjacencyList;
    int* d_edgeList;
    bool* d_visited;
    int* d_frontier;
    int* d_newFrontier;
    int* d_frontierSize;
    int* d_newFrontierSize;

    // Allocate memory
    hipMalloc(&d_adjacencyList, (MAX_NODES + 1) * sizeof(int));
    hipMalloc(&d_edgeList, MAX_EDGES * sizeof(int));
    hipMalloc(&d_visited, MAX_NODES * sizeof(bool));
    hipMalloc(&d_frontier, MAX_NODES * sizeof(int));
    hipMalloc(&d_newFrontier, MAX_NODES * sizeof(int));
    hipMalloc(&d_frontierSize, sizeof(int));
    hipMalloc(&d_newFrontierSize, sizeof(int));

    // Copy data to device
    hipMemcpy(d_adjacencyList, adjacencyList.data(), (MAX_NODES + 1) * sizeof(int), hipMemcpyHostToDevice);
    hipMemcpy(d_edgeList, edgeList.data(), MAX_EDGES * sizeof(int), hipMemcpyHostToDevice);

    // Initialize visited array to false
    hipMemset(d_visited, 0, MAX_NODES * sizeof(bool));

    // Initialize frontier
    initFrontier<<<1, 1>>>(d_visited, d_frontier, d_frontierSize);
    hipDeviceSynchronize();

    // BFS parameters
    int blockSize = 256;
    int gridSize = (MAX_NODES + blockSize - 1) / blockSize;
    int maxIterations = 10; // Arbitrary number of iterations for demonstration

    // Perform BFS
    for (int i = 0; i < maxIterations; ++i) {
        hipMemset(d_newFrontierSize, 0, sizeof(int));
        bfsKernel<<<gridSize, blockSize, blockSize * sizeof(int)>>>(d_adjacencyList, d_edgeList, d_visited, d_frontier, d_newFrontier, d_frontierSize, d_newFrontierSize, MAX_NODES);
        hipDeviceSynchronize();

        // Swap frontiers
        int* temp = d_frontier;
        d_frontier = d_newFrontier;
        d_newFrontier = temp;
        temp = d_frontierSize;
        d_frontierSize = d_newFrontierSize;
        d_newFrontierSize = temp;
    }

    // Free memory
    hipFree(d_adjacencyList);
    hipFree(d_edgeList);
    hipFree(d_visited);
    hipFree(d_frontier);
    hipFree(d_newFrontier);
    hipFree(d_frontierSize);
    hipFree(d_newFrontierSize);

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

                                        # End of file scope inline assembly
	.globl	_Z27__device_stub__initFrontierPbPiS0_ # -- Begin function _Z27__device_stub__initFrontierPbPiS0_
	.type	_Z27__device_stub__initFrontierPbPiS0_,@function
_Z27__device_stub__initFrontierPbPiS0_: # @_Z27__device_stub__initFrontierPbPiS0_
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$112, %rsp
	.cfi_def_cfa_offset 160
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	40(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	32(%rsp), %rcx
	movq	%rsi, (%rcx)
	leaq	24(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%rbx)
	movq	%rsi, 16(%rbx)
	leaq	64(%rsp), %r14
	leaq	48(%rsp), %r15
	leaq	16(%rsp), %r12
	leaq	8(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z12initFrontierPbPiS0_, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$128, %rsp
	.cfi_adjust_cfa_offset -128
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end0:
	.size	_Z27__device_stub__initFrontierPbPiS0_, .Lfunc_end0-_Z27__device_stub__initFrontierPbPiS0_
	.cfi_endproc
                                        # -- End function
	.globl	_Z24__device_stub__bfsKernelPiS_PbS_S_S_S_i # -- Begin function _Z24__device_stub__bfsKernelPiS_PbS_S_S_S_i
	.type	_Z24__device_stub__bfsKernelPiS_PbS_S_S_S_i,@function
_Z24__device_stub__bfsKernelPiS_PbS_S_S_S_i: # @_Z24__device_stub__bfsKernelPiS_PbS_S_S_S_i
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$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	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	208(%rsp), %rax
	movq	%rax, 48(%rbx)
	leaq	216(%rsp), %rax
	movq	%rax, 56(%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_PbS_S_S_S_i, %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_end1:
	.size	_Z24__device_stub__bfsKernelPiS_PbS_S_S_S_i, .Lfunc_end1-_Z24__device_stub__bfsKernelPiS_PbS_S_S_S_i
	.cfi_endproc
                                        # -- End function
	.globl	main                            # -- Begin function main
	.type	main,@function
main:                                   # @main
.Lfunc_begin0:
	.cfi_startproc
	.cfi_personality 3, __gxx_personality_v0
	.cfi_lsda 3, .Lexception0
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	subq	$328, %rsp                      # imm = 0x148
	.cfi_def_cfa_offset 368
	.cfi_offset %rbx, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	.cfi_escape 0x2e, 0x00
	leaq	72(%rsp), %rbx
	movl	$.Lconstinit, %esi
	movl	$256, %edx                      # imm = 0x100
	movq	%rbx, %rdi
	callq	memcpy@PLT
	.cfi_escape 0x2e, 0x00
	leaq	48(%rsp), %rdi
	leaq	40(%rsp), %rcx
	movl	$64, %edx
	movq	%rbx, %rsi
	callq	_ZNSt6vectorIiSaIiEEC2ESt16initializer_listIiERKS0_
.Ltmp0:
	.cfi_escape 0x2e, 0x00
	movl	$2048, %edi                     # imm = 0x800
	callq	_Znwm
.Ltmp1:
# %bb.1:
	movq	%rax, %rbx
	.cfi_escape 0x2e, 0x00
	xorl	%r14d, %r14d
	movl	$2048, %edx                     # imm = 0x800
	movq	%rax, %rdi
	xorl	%esi, %esi
	callq	memset@PLT
.LBB2_2:                                # =>This Inner Loop Header: Depth=1
	movl	%r14d, %eax
	andl	$127, %eax
	movl	%eax, (%rbx,%r14,4)
	incq	%r14
	cmpq	$512, %r14                      # imm = 0x200
	jne	.LBB2_2
# %bb.3:
.Ltmp3:
	.cfi_escape 0x2e, 0x00
	leaq	72(%rsp), %rdi
	movl	$516, %esi                      # imm = 0x204
	callq	hipMalloc
.Ltmp4:
# %bb.4:                                # %_ZL9hipMallocIiE10hipError_tPPT_m.exit
.Ltmp5:
	.cfi_escape 0x2e, 0x00
	leaq	40(%rsp), %rdi
	movl	$2048, %esi                     # imm = 0x800
	callq	hipMalloc
.Ltmp6:
# %bb.5:                                # %_ZL9hipMallocIiE10hipError_tPPT_m.exit35
.Ltmp7:
	.cfi_escape 0x2e, 0x00
	leaq	32(%rsp), %rdi
	movl	$128, %esi
	callq	hipMalloc
.Ltmp8:
# %bb.6:                                # %_ZL9hipMallocIbE10hipError_tPPT_m.exit
.Ltmp9:
	.cfi_escape 0x2e, 0x00
	leaq	16(%rsp), %rdi
	movl	$512, %esi                      # imm = 0x200
	callq	hipMalloc
.Ltmp10:
# %bb.7:                                # %_ZL9hipMallocIiE10hipError_tPPT_m.exit36
.Ltmp11:
	.cfi_escape 0x2e, 0x00
	leaq	24(%rsp), %rdi
	movl	$512, %esi                      # imm = 0x200
	callq	hipMalloc
.Ltmp12:
# %bb.8:                                # %_ZL9hipMallocIiE10hipError_tPPT_m.exit37
.Ltmp13:
	.cfi_escape 0x2e, 0x00
	leaq	8(%rsp), %rdi
	movl	$4, %esi
	callq	hipMalloc
.Ltmp14:
# %bb.9:                                # %_ZL9hipMallocIiE10hipError_tPPT_m.exit38
.Ltmp15:
	.cfi_escape 0x2e, 0x00
	movq	%rsp, %rdi
	movl	$4, %esi
	callq	hipMalloc
.Ltmp16:
# %bb.10:                               # %_ZL9hipMallocIiE10hipError_tPPT_m.exit39
	movq	72(%rsp), %rdi
	movq	48(%rsp), %rsi
.Ltmp17:
	.cfi_escape 0x2e, 0x00
	movl	$516, %edx                      # imm = 0x204
	movl	$1, %ecx
	callq	hipMemcpy
.Ltmp18:
# %bb.11:
	movq	40(%rsp), %rdi
.Ltmp19:
	.cfi_escape 0x2e, 0x00
	movl	$2048, %edx                     # imm = 0x800
	movq	%rbx, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
.Ltmp20:
# %bb.12:
	movq	32(%rsp), %rdi
.Ltmp21:
	.cfi_escape 0x2e, 0x00
	movl	$128, %edx
	xorl	%esi, %esi
	callq	hipMemset
.Ltmp22:
# %bb.13:
.Ltmp23:
	.cfi_escape 0x2e, 0x00
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
.Ltmp24:
# %bb.14:
	testl	%eax, %eax
	jne	.LBB2_16
# %bb.15:
	movq	32(%rsp), %rdi
	movq	16(%rsp), %rsi
	movq	8(%rsp), %rdx
.Ltmp25:
	.cfi_escape 0x2e, 0x00
	callq	_Z27__device_stub__initFrontierPbPiS0_
.Ltmp26:
.LBB2_16:
.Ltmp27:
	.cfi_escape 0x2e, 0x00
	callq	hipDeviceSynchronize
.Ltmp28:
# %bb.17:                               # %.preheader.preheader
	movl	$10, %ebp
	movq	(%rsp), %rdi
	movabsq	$4294967297, %r14               # imm = 0x100000001
	movabsq	$4294967552, %r15               # imm = 0x100000100
.LBB2_18:                               # %.preheader
                                        # =>This Inner Loop Header: Depth=1
.Ltmp30:
	.cfi_escape 0x2e, 0x00
	movl	$4, %edx
	xorl	%esi, %esi
	callq	hipMemset
.Ltmp31:
# %bb.19:                               #   in Loop: Header=BB2_18 Depth=1
.Ltmp32:
	.cfi_escape 0x2e, 0x00
	movl	$1024, %r8d                     # imm = 0x400
	movq	%r14, %rdi
	movl	$1, %esi
	movq	%r15, %rdx
	movl	$1, %ecx
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
.Ltmp33:
# %bb.20:                               #   in Loop: Header=BB2_18 Depth=1
	testl	%eax, %eax
	jne	.LBB2_22
# %bb.21:                               #   in Loop: Header=BB2_18 Depth=1
	movq	72(%rsp), %rdi
	movq	40(%rsp), %rsi
	movq	32(%rsp), %rdx
	movq	16(%rsp), %rcx
	movq	24(%rsp), %r8
	movq	8(%rsp), %r9
.Ltmp34:
	.cfi_escape 0x2e, 0x10
	pushq	$128
	.cfi_adjust_cfa_offset 8
	pushq	8(%rsp)
	.cfi_adjust_cfa_offset 8
	callq	_Z24__device_stub__bfsKernelPiS_PbS_S_S_S_i
	addq	$16, %rsp
	.cfi_adjust_cfa_offset -16
.Ltmp35:
.LBB2_22:                               #   in Loop: Header=BB2_18 Depth=1
.Ltmp36:
	.cfi_escape 0x2e, 0x00
	callq	hipDeviceSynchronize
.Ltmp37:
# %bb.23:                               #   in Loop: Header=BB2_18 Depth=1
	movq	16(%rsp), %rax
	movq	24(%rsp), %rcx
	movq	%rcx, 16(%rsp)
	movq	%rax, 24(%rsp)
	movq	8(%rsp), %rdi
	movq	(%rsp), %rax
	movq	%rax, 8(%rsp)
	movq	%rdi, (%rsp)
	decl	%ebp
	jne	.LBB2_18
# %bb.24:
	movq	72(%rsp), %rdi
.Ltmp39:
	.cfi_escape 0x2e, 0x00
	callq	hipFree
.Ltmp40:
# %bb.25:
	movq	40(%rsp), %rdi
.Ltmp41:
	.cfi_escape 0x2e, 0x00
	callq	hipFree
.Ltmp42:
# %bb.26:
	movq	32(%rsp), %rdi
.Ltmp43:
	.cfi_escape 0x2e, 0x00
	callq	hipFree
.Ltmp44:
# %bb.27:
	movq	16(%rsp), %rdi
.Ltmp45:
	.cfi_escape 0x2e, 0x00
	callq	hipFree
.Ltmp46:
# %bb.28:
	movq	24(%rsp), %rdi
.Ltmp47:
	.cfi_escape 0x2e, 0x00
	callq	hipFree
.Ltmp48:
# %bb.29:
	movq	8(%rsp), %rdi
.Ltmp49:
	.cfi_escape 0x2e, 0x00
	callq	hipFree
.Ltmp50:
# %bb.30:
	movq	(%rsp), %rdi
.Ltmp51:
	.cfi_escape 0x2e, 0x00
	callq	hipFree
.Ltmp52:
# %bb.31:                               # %_ZNSt6vectorIiSaIiEED2Ev.exit
	.cfi_escape 0x2e, 0x00
	movq	%rbx, %rdi
	callq	_ZdlPv
	movq	48(%rsp), %rdi
	testq	%rdi, %rdi
	je	.LBB2_33
# %bb.32:
	.cfi_escape 0x2e, 0x00
	callq	_ZdlPv
.LBB2_33:                               # %_ZNSt6vectorIiSaIiEED2Ev.exit41
	xorl	%eax, %eax
	addq	$328, %rsp                      # imm = 0x148
	.cfi_def_cfa_offset 40
	popq	%rbx
	.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
.LBB2_39:
	.cfi_def_cfa_offset 368
.Ltmp2:
	movq	%rax, %r14
	jmp	.LBB2_36
.LBB2_34:
.Ltmp53:
	jmp	.LBB2_35
.LBB2_40:
.Ltmp29:
	jmp	.LBB2_35
.LBB2_41:
.Ltmp38:
.LBB2_35:                               # %_ZNSt6vectorIiSaIiEED2Ev.exit43
	movq	%rax, %r14
	.cfi_escape 0x2e, 0x00
	movq	%rbx, %rdi
	callq	_ZdlPv
.LBB2_36:
	movq	48(%rsp), %rdi
	testq	%rdi, %rdi
	je	.LBB2_38
# %bb.37:
	.cfi_escape 0x2e, 0x00
	callq	_ZdlPv
.LBB2_38:                               # %_ZNSt6vectorIiSaIiEED2Ev.exit45
	.cfi_escape 0x2e, 0x00
	movq	%r14, %rdi
	callq	_Unwind_Resume@PLT
.Lfunc_end2:
	.size	main, .Lfunc_end2-main
	.cfi_endproc
	.section	.gcc_except_table,"a",@progbits
	.p2align	2, 0x0
GCC_except_table2:
.Lexception0:
	.byte	255                             # @LPStart Encoding = omit
	.byte	255                             # @TType Encoding = omit
	.byte	1                               # Call site Encoding = uleb128
	.uleb128 .Lcst_end0-.Lcst_begin0
.Lcst_begin0:
	.uleb128 .Lfunc_begin0-.Lfunc_begin0    # >> Call Site 1 <<
	.uleb128 .Ltmp0-.Lfunc_begin0           #   Call between .Lfunc_begin0 and .Ltmp0
	.byte	0                               #     has no landing pad
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp0-.Lfunc_begin0           # >> Call Site 2 <<
	.uleb128 .Ltmp1-.Ltmp0                  #   Call between .Ltmp0 and .Ltmp1
	.uleb128 .Ltmp2-.Lfunc_begin0           #     jumps to .Ltmp2
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp1-.Lfunc_begin0           # >> Call Site 3 <<
	.uleb128 .Ltmp3-.Ltmp1                  #   Call between .Ltmp1 and .Ltmp3
	.byte	0                               #     has no landing pad
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp3-.Lfunc_begin0           # >> Call Site 4 <<
	.uleb128 .Ltmp28-.Ltmp3                 #   Call between .Ltmp3 and .Ltmp28
	.uleb128 .Ltmp29-.Lfunc_begin0          #     jumps to .Ltmp29
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp30-.Lfunc_begin0          # >> Call Site 5 <<
	.uleb128 .Ltmp37-.Ltmp30                #   Call between .Ltmp30 and .Ltmp37
	.uleb128 .Ltmp38-.Lfunc_begin0          #     jumps to .Ltmp38
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp39-.Lfunc_begin0          # >> Call Site 6 <<
	.uleb128 .Ltmp52-.Ltmp39                #   Call between .Ltmp39 and .Ltmp52
	.uleb128 .Ltmp53-.Lfunc_begin0          #     jumps to .Ltmp53
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp52-.Lfunc_begin0          # >> Call Site 7 <<
	.uleb128 .Lfunc_end2-.Ltmp52            #   Call between .Ltmp52 and .Lfunc_end2
	.byte	0                               #     has no landing pad
	.byte	0                               #   On action: cleanup
.Lcst_end0:
	.p2align	2, 0x0
                                        # -- End function
	.section	.text._ZNSt6vectorIiSaIiEEC2ESt16initializer_listIiERKS0_,"axG",@progbits,_ZNSt6vectorIiSaIiEEC2ESt16initializer_listIiERKS0_,comdat
	.weak	_ZNSt6vectorIiSaIiEEC2ESt16initializer_listIiERKS0_ # -- Begin function _ZNSt6vectorIiSaIiEEC2ESt16initializer_listIiERKS0_
	.p2align	1, 0x90
	.type	_ZNSt6vectorIiSaIiEEC2ESt16initializer_listIiERKS0_,@function
_ZNSt6vectorIiSaIiEEC2ESt16initializer_listIiERKS0_: # @_ZNSt6vectorIiSaIiEEC2ESt16initializer_listIiERKS0_
.Lfunc_begin1:
	.cfi_startproc
	.cfi_personality 3, __gxx_personality_v0
	.cfi_lsda 3, .Lexception1
# %bb.0:
	pushq	%r14
	.cfi_def_cfa_offset 16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	pushq	%rax
	.cfi_def_cfa_offset 32
	.cfi_offset %rbx, -24
	.cfi_offset %r14, -16
	movq	%rdi, %rbx
	xorps	%xmm0, %xmm0
	movups	%xmm0, (%rdi)
	movq	$0, 16(%rdi)
	leaq	(%rsi,%rdx,4), %rdx
.Ltmp54:
	callq	_ZNSt6vectorIiSaIiEE19_M_range_initializeIPKiEEvT_S5_St20forward_iterator_tag
.Ltmp55:
# %bb.1:
	addq	$8, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	retq
.LBB3_2:
	.cfi_def_cfa_offset 32
.Ltmp56:
	movq	%rax, %r14
	movq	(%rbx), %rdi
	testq	%rdi, %rdi
	je	.LBB3_4
# %bb.3:
	callq	_ZdlPv
.LBB3_4:                                # %_ZNSt12_Vector_baseIiSaIiEED2Ev.exit
	movq	%r14, %rdi
	callq	_Unwind_Resume@PLT
.Lfunc_end3:
	.size	_ZNSt6vectorIiSaIiEEC2ESt16initializer_listIiERKS0_, .Lfunc_end3-_ZNSt6vectorIiSaIiEEC2ESt16initializer_listIiERKS0_
	.cfi_endproc
	.section	.gcc_except_table._ZNSt6vectorIiSaIiEEC2ESt16initializer_listIiERKS0_,"aG",@progbits,_ZNSt6vectorIiSaIiEEC2ESt16initializer_listIiERKS0_,comdat
	.p2align	2, 0x0
GCC_except_table3:
.Lexception1:
	.byte	255                             # @LPStart Encoding = omit
	.byte	255                             # @TType Encoding = omit
	.byte	1                               # Call site Encoding = uleb128
	.uleb128 .Lcst_end1-.Lcst_begin1
.Lcst_begin1:
	.uleb128 .Ltmp54-.Lfunc_begin1          # >> Call Site 1 <<
	.uleb128 .Ltmp55-.Ltmp54                #   Call between .Ltmp54 and .Ltmp55
	.uleb128 .Ltmp56-.Lfunc_begin1          #     jumps to .Ltmp56
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp55-.Lfunc_begin1          # >> Call Site 2 <<
	.uleb128 .Lfunc_end3-.Ltmp55            #   Call between .Ltmp55 and .Lfunc_end3
	.byte	0                               #     has no landing pad
	.byte	0                               #   On action: cleanup
.Lcst_end1:
	.p2align	2, 0x0
                                        # -- End function
	.section	.text._ZNSt6vectorIiSaIiEE19_M_range_initializeIPKiEEvT_S5_St20forward_iterator_tag,"axG",@progbits,_ZNSt6vectorIiSaIiEE19_M_range_initializeIPKiEEvT_S5_St20forward_iterator_tag,comdat
	.weak	_ZNSt6vectorIiSaIiEE19_M_range_initializeIPKiEEvT_S5_St20forward_iterator_tag # -- Begin function _ZNSt6vectorIiSaIiEE19_M_range_initializeIPKiEEvT_S5_St20forward_iterator_tag
	.p2align	1, 0x90
	.type	_ZNSt6vectorIiSaIiEE19_M_range_initializeIPKiEEvT_S5_St20forward_iterator_tag,@function
_ZNSt6vectorIiSaIiEE19_M_range_initializeIPKiEEvT_S5_St20forward_iterator_tag: # @_ZNSt6vectorIiSaIiEE19_M_range_initializeIPKiEEvT_S5_St20forward_iterator_tag
	.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
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movq	%rdx, %r14
	subq	%rsi, %r14
	movq	%r14, %r13
	sarq	$2, %r13
	movq	%r13, %rax
	shrq	$61, %rax
	jne	.LBB4_9
# %bb.1:                                # %_ZNSt6vectorIiSaIiEE17_S_check_init_lenEmRKS0_.exit
	movq	%rsi, %r15
	movq	%rdi, %rbx
	cmpq	%rsi, %rdx
	je	.LBB4_2
# %bb.3:
	movq	%rbx, %rdi
	movq	%r13, %rsi
	xorl	%edx, %edx
	callq	_ZNSt15__new_allocatorIiE8allocateEmPKv
	movq	%rax, %r12
	jmp	.LBB4_4
.LBB4_2:
	xorl	%r12d, %r12d
.LBB4_4:                                # %_ZNSt12_Vector_baseIiSaIiEE11_M_allocateEm.exit
	movq	%r12, (%rbx)
	leaq	(%r12,%r13,4), %rax
	movq	%rax, 16(%rbx)
	cmpq	$5, %r14
	jl	.LBB4_6
# %bb.5:
	movq	%r12, %rdi
	movq	%r15, %rsi
	movq	%r14, %rdx
	callq	memmove@PLT
.LBB4_8:                                # %_ZSt22__uninitialized_copy_aIPKiPiiET0_T_S4_S3_RSaIT1_E.exit
	addq	%r14, %r12
	movq	%r12, 8(%rbx)
	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
.LBB4_6:
	.cfi_def_cfa_offset 48
	cmpq	$4, %r14
	jne	.LBB4_8
# %bb.7:
	movl	(%r15), %eax
	movl	%eax, (%r12)
	jmp	.LBB4_8
.LBB4_9:
	movl	$.L.str, %edi
	callq	_ZSt20__throw_length_errorPKc
.Lfunc_end4:
	.size	_ZNSt6vectorIiSaIiEE19_M_range_initializeIPKiEEvT_S5_St20forward_iterator_tag, .Lfunc_end4-_ZNSt6vectorIiSaIiEE19_M_range_initializeIPKiEEvT_S5_St20forward_iterator_tag
	.cfi_endproc
                                        # -- End function
	.section	.text._ZNSt15__new_allocatorIiE8allocateEmPKv,"axG",@progbits,_ZNSt15__new_allocatorIiE8allocateEmPKv,comdat
	.weak	_ZNSt15__new_allocatorIiE8allocateEmPKv # -- Begin function _ZNSt15__new_allocatorIiE8allocateEmPKv
	.p2align	1, 0x90
	.type	_ZNSt15__new_allocatorIiE8allocateEmPKv,@function
_ZNSt15__new_allocatorIiE8allocateEmPKv: # @_ZNSt15__new_allocatorIiE8allocateEmPKv
	.cfi_startproc
# %bb.0:
	pushq	%rax
	.cfi_def_cfa_offset 16
	movq	%rsi, %rax
	shrq	$61, %rax
	jne	.LBB5_1
# %bb.3:
	shlq	$2, %rsi
	movq	%rsi, %rdi
	popq	%rax
	.cfi_def_cfa_offset 8
	jmp	_Znwm                           # TAILCALL
.LBB5_1:
	.cfi_def_cfa_offset 16
	shrq	$62, %rsi
	je	.LBB5_2
# %bb.4:
	callq	_ZSt28__throw_bad_array_new_lengthv
.LBB5_2:
	callq	_ZSt17__throw_bad_allocv
.Lfunc_end5:
	.size	_ZNSt15__new_allocatorIiE8allocateEmPKv, .Lfunc_end5-_ZNSt15__new_allocatorIiE8allocateEmPKv
	.cfi_endproc
                                        # -- End function
	.text
	.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	.LBB6_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB6_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z12initFrontierPbPiS0_, %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	$_Z9bfsKernelPiS_PbS_S_S_S_i, %esi
	movl	$.L__unnamed_2, %edx
	movl	$.L__unnamed_2, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end6:
	.size	__hip_module_ctor, .Lfunc_end6-__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	.LBB7_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
.LBB7_2:
	retq
.Lfunc_end7:
	.size	__hip_module_dtor, .Lfunc_end7-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_Z12initFrontierPbPiS0_,@object # @_Z12initFrontierPbPiS0_
	.section	.rodata,"a",@progbits
	.globl	_Z12initFrontierPbPiS0_
	.p2align	3, 0x0
_Z12initFrontierPbPiS0_:
	.quad	_Z27__device_stub__initFrontierPbPiS0_
	.size	_Z12initFrontierPbPiS0_, 8

	.type	_Z9bfsKernelPiS_PbS_S_S_S_i,@object # @_Z9bfsKernelPiS_PbS_S_S_S_i
	.globl	_Z9bfsKernelPiS_PbS_S_S_S_i
	.p2align	3, 0x0
_Z9bfsKernelPiS_PbS_S_S_S_i:
	.quad	_Z24__device_stub__bfsKernelPiS_PbS_S_S_S_i
	.size	_Z9bfsKernelPiS_PbS_S_S_S_i, 8

	.type	.Lconstinit,@object             # @constinit
	.p2align	2, 0x0
.Lconstinit:
	.long	0                               # 0x0
	.long	2                               # 0x2
	.long	4                               # 0x4
	.long	7                               # 0x7
	.long	9                               # 0x9
	.long	11                              # 0xb
	.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	66                              # 0x42
	.long	68                              # 0x44
	.long	70                              # 0x46
	.long	72                              # 0x48
	.long	74                              # 0x4a
	.long	76                              # 0x4c
	.long	78                              # 0x4e
	.long	80                              # 0x50
	.long	82                              # 0x52
	.long	84                              # 0x54
	.long	86                              # 0x56
	.long	88                              # 0x58
	.long	90                              # 0x5a
	.long	92                              # 0x5c
	.long	94                              # 0x5e
	.long	96                              # 0x60
	.long	98                              # 0x62
	.long	100                             # 0x64
	.long	102                             # 0x66
	.long	104                             # 0x68
	.long	106                             # 0x6a
	.long	108                             # 0x6c
	.long	110                             # 0x6e
	.long	112                             # 0x70
	.long	114                             # 0x72
	.long	116                             # 0x74
	.long	118                             # 0x76
	.long	120                             # 0x78
	.long	122                             # 0x7a
	.long	124                             # 0x7c
	.long	126                             # 0x7e
	.long	128                             # 0x80
	.size	.Lconstinit, 256

	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"cannot create std::vector larger than max_size()"
	.size	.L.str, 49

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z12initFrontierPbPiS0_"
	.size	.L__unnamed_1, 24

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"_Z9bfsKernelPiS_PbS_S_S_S_i"
	.size	.L__unnamed_2, 28

	.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 _Z27__device_stub__initFrontierPbPiS0_
	.addrsig_sym _Z24__device_stub__bfsKernelPiS_PbS_S_S_S_i
	.addrsig_sym __gxx_personality_v0
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Unwind_Resume
	.addrsig_sym _Z12initFrontierPbPiS0_
	.addrsig_sym _Z9bfsKernelPiS_PbS_S_S_S_i
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z12initFrontierPbPiS0_ ; -- Begin function _Z12initFrontierPbPiS0_
	.globl	_Z12initFrontierPbPiS0_
	.p2align	8
	.type	_Z12initFrontierPbPiS0_,@function
_Z12initFrontierPbPiS0_:                ; @_Z12initFrontierPbPiS0_
; %bb.0:
	s_load_b32 s2, s[0:1], 0x24
	s_waitcnt lgkmcnt(0)
	s_and_b32 s2, s2, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_mul_i32 s15, s15, s2
	s_sub_i32 s2, 0, s15
	s_delay_alu instid0(SALU_CYCLE_1)
	v_cmp_eq_u32_e32 vcc_lo, s2, v0
	s_and_saveexec_b32 s2, vcc_lo
	s_cbranch_execz .LBB0_2
; %bb.1:
	s_clause 0x1
	s_load_b128 s[4:7], s[0:1], 0x0
	s_load_b64 s[0:1], s[0:1], 0x10
	v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v1, 1
	s_waitcnt lgkmcnt(0)
	s_clause 0x2
	global_store_b32 v0, v0, s[6:7]
	global_store_b32 v0, v1, s[0:1]
	global_store_b8 v0, v1, s[4:5]
.LBB0_2:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z12initFrontierPbPiS0_
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 280
		.amdhsa_user_sgpr_count 15
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 2
		.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	_Z12initFrontierPbPiS0_, .Lfunc_end0-_Z12initFrontierPbPiS0_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 120
; NumSgprs: 18
; NumVgprs: 2
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 18
; 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	_Z9bfsKernelPiS_PbS_S_S_S_i ; -- Begin function _Z9bfsKernelPiS_PbS_S_S_S_i
	.globl	_Z9bfsKernelPiS_PbS_S_S_S_i
	.p2align	8
	.type	_Z9bfsKernelPiS_PbS_S_S_S_i,@function
_Z9bfsKernelPiS_PbS_S_S_S_i:            ; @_Z9bfsKernelPiS_PbS_S_S_S_i
; %bb.0:
	s_clause 0x1
	s_load_b256 s[4:11], s[0:1], 0x20
	s_load_b32 s2, s[0:1], 0x4c
	s_waitcnt lgkmcnt(0)
	s_load_b32 s10, s[6:7], 0x0
	s_load_b256 s[16:23], s[0:1], 0x0
	s_add_u32 s0, s0, 64
	s_addc_u32 s1, s1, 0
	s_and_b32 s2, s2, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[1:2], null, s15, s2, v[0:1]
	v_lshl_add_u32 v0, v1, 2, 0
	s_waitcnt lgkmcnt(0)
	v_cmp_gt_i32_e32 vcc_lo, s10, v1
	s_and_saveexec_b32 s11, vcc_lo
	s_cbranch_execz .LBB1_3
; %bb.1:                                ; %.lr.ph.preheader
	s_load_b32 s0, s[0:1], 0x0
	v_ashrrev_i32_e32 v2, 31, v1
	v_mov_b32_e32 v4, v0
	s_mov_b32 s12, 0
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[2:3], 2, v[1:2]
	s_waitcnt lgkmcnt(0)
	s_mul_i32 s2, s0, s2
	v_add_co_u32 v2, s0, s22, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	v_add_co_ci_u32_e64 v3, s0, s23, v3, s0
	s_ashr_i32 s3, s2, 31
	s_lshl_b64 s[6:7], s[2:3], 2
	s_lshl_b32 s3, s2, 2
.LBB1_2:                                ; %.lr.ph
                                        ; =>This Inner Loop Header: Depth=1
	global_load_b32 v5, v[2:3], off
	v_add_nc_u32_e32 v1, s2, v1
	v_add_co_u32 v2, s0, v2, s6
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_ci_u32_e64 v3, s0, s7, v3, s0
	v_cmp_le_i32_e64 s1, s10, v1
	s_delay_alu instid0(VALU_DEP_1)
	s_or_b32 s12, s1, s12
	s_waitcnt vmcnt(0)
	ds_store_b32 v4, v5
	v_add_nc_u32_e32 v4, s3, v4
	s_and_not1_b32 exec_lo, exec_lo, s12
	s_cbranch_execnz .LBB1_2
.LBB1_3:                                ; %Flow72
	s_or_b32 exec_lo, exec_lo, s11
	s_waitcnt lgkmcnt(0)
	s_barrier
	buffer_gl0_inv
	s_and_saveexec_b32 s0, vcc_lo
	s_cbranch_execz .LBB1_11
; %bb.4:
	ds_load_b32 v0, v0
	s_mov_b32 s1, 0
	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, s16, v0
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v1, vcc_lo, s17, v1, vcc_lo
	global_load_b64 v[0:1], v[0:1], off
	s_waitcnt vmcnt(0)
	v_cmp_lt_i32_e32 vcc_lo, v0, v1
	s_and_b32 exec_lo, exec_lo, vcc_lo
	s_cbranch_execz .LBB1_11
; %bb.5:                                ; %.lr.ph41.preheader
	v_ashrrev_i32_e32 v3, 31, v0
	v_dual_mov_b32 v2, v0 :: v_dual_mov_b32 v7, 1
	v_mov_b32_e32 v6, 0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[2:3], 2, v[2:3]
	v_add_co_u32 v2, vcc_lo, s18, v2
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v3, vcc_lo, s19, v3, vcc_lo
.LBB1_6:                                ; %.lr.ph41
                                        ; =>This Inner Loop Header: Depth=1
	global_load_b32 v8, v[2:3], off
	s_mov_b32 s0, exec_lo
	s_waitcnt vmcnt(0)
	v_ashrrev_i32_e32 v5, 31, v8
	v_add_co_u32 v4, vcc_lo, s20, v8
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v5, vcc_lo, s21, v5, vcc_lo
	global_load_u8 v9, v[4:5], off
	s_waitcnt vmcnt(0)
	v_cmpx_eq_u16_e32 0, v9
	s_cbranch_execz .LBB1_10
; %bb.7:                                ;   in Loop: Header=BB1_6 Depth=1
	s_mov_b32 s3, exec_lo
	s_mov_b32 s2, exec_lo
	v_mbcnt_lo_u32_b32 v9, s3, 0
                                        ; implicit-def: $vgpr10
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_eq_u32_e32 0, v9
	s_cbranch_execz .LBB1_9
; %bb.8:                                ;   in Loop: Header=BB1_6 Depth=1
	s_bcnt1_i32_b32 s3, s3
	s_delay_alu instid0(SALU_CYCLE_1)
	v_mov_b32_e32 v10, s3
	global_atomic_add_u32 v10, v6, v10, s[8:9] glc
.LBB1_9:                                ;   in Loop: Header=BB1_6 Depth=1
	s_or_b32 exec_lo, exec_lo, s2
	s_waitcnt vmcnt(0)
	v_readfirstlane_b32 s2, v10
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v9, s2, v9
	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, s4, v9
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v10, vcc_lo, s5, v10, vcc_lo
	global_store_b32 v[9:10], v8, off
	global_store_b8 v[4:5], v7, off
.LBB1_10:                               ; %Flow
                                        ;   in Loop: Header=BB1_6 Depth=1
	s_or_b32 exec_lo, exec_lo, s0
	v_add_nc_u32_e32 v0, 1, v0
	v_add_co_u32 v2, s0, v2, 4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_ci_u32_e64 v3, s0, 0, v3, s0
	v_cmp_ge_i32_e32 vcc_lo, v0, v1
	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 .LBB1_6
.LBB1_11:                               ; %.loopexit
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z9bfsKernelPiS_PbS_S_S_S_i
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 320
		.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 24
		.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	_Z9bfsKernelPiS_PbS_S_S_S_i, .Lfunc_end1-_Z9bfsKernelPiS_PbS_S_S_S_i
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 612
; NumSgprs: 26
; NumVgprs: 11
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 3
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 26
; 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
      - .offset:         24
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         28
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         32
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         36
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         38
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         40
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         42
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         44
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         46
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         64
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         88
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 280
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z12initFrontierPbPiS0_
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z12initFrontierPbPiS0_.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
      - .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
      - .address_space:  global
        .offset:         48
        .size:           8
        .value_kind:     global_buffer
      - .offset:         56
        .size:           4
        .value_kind:     by_value
      - .offset:         64
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         68
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         72
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         76
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         78
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         80
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         82
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         84
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         86
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         104
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         112
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         120
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         128
        .size:           2
        .value_kind:     hidden_grid_dims
      - .offset:         184
        .size:           4
        .value_kind:     hidden_dynamic_lds_size
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 320
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z9bfsKernelPiS_PbS_S_S_S_i
    .private_segment_fixed_size: 0
    .sgpr_count:     26
    .sgpr_spill_count: 0
    .symbol:         _Z9bfsKernelPiS_PbS_S_S_S_i.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
