#include<iostream>
#include<stdio.h>
#include<stdlib.h>
#include<math.h>
#include<fstream>
#include<time.h>
#include<sys/time.h>

using namespace std;

#define num_threads 1000
// #define num_edges 700000
// #define num_vertices1 10000
// #define num_vertices2 10000

// #define num_edges 1000000
// #define num_vertices1 1000
// #define num_vertices2 1000


#define lli long long int

// #define num_edges 2998468
// #define num_vertices1 100000
// #define num_vertices2 100000



const lli num_edges = 1000000;
const lli num_vertices1 = 1000;
const lli num_vertices2 = 1000;

// const lli num_edges = 291;
// const lli num_vertices1 = 100;
// const lli num_vertices2 = 100;

__device__ int d_flat_adj_list[2*num_edges];
__device__ int d_degree[num_vertices1+num_vertices2+1]={0};      //store degree of each vertex
__device__ int d_list_ptr[num_vertices1+num_vertices2+2];        //1-indexed and extra element at the end for easy size access  // Pointer to the start of adjacency list
__device__ int d_list_ptr_copy[num_vertices1+num_vertices2+2];    // Temporrary stuff, gotta sleep

__device__ bool d_is_matched_edge[(num_vertices1+ num_vertices2 + 1)*(num_vertices1 + num_vertices2+1)] = {0} ;     // Adjacency matrix (1-indexed)
__device__ bool d_is_matched_vertex[num_vertices1 + num_vertices2 + 1] = {0};	//is the vertex matched
__device__ int d_partner_vertex[num_vertices1 + num_vertices2 + 1];
__device__ int d_visited[num_vertices1 + num_vertices2 + 1] = {0};
__device__ int d_bfs_parent[num_vertices1 +  num_vertices2 + 1];
__device__ bool d_is_parent_change[num_vertices1 + num_vertices2 + 1] = {0};

__device__ int d_frontier[num_vertices1 + num_vertices2+1] = {0};
__device__ int d_next_frontier[num_vertices1+num_vertices2+1] = {0};

__device__ int num_aug_paths = 0;



int *h_flat_adj_list;
int *h_degree;
int * h_list_ptr;
int *h_list_ptr_copy;

bool *h_is_matched_edge;
bool *h_is_matched_vertex;
int *h_partner_vertex;
int *h_visited;
int *h_bfs_parent;
bool *h_is_parent_change;

int fc = num_vertices1;
// int num_aug_paths = 0;

int *h_frontier;
int *h_next_frontier;


__device__ 
int get_is_matched_edge(int i, int j){
	return d_is_matched_edge[i*(num_vertices1 + num_vertices2+1) + j ];
}

void print_matchings(){
	cout << "Matchings: " << endl;
    for(int i=1;i<=num_vertices1+num_vertices2; i++){
    	cout<< i << " " << h_partner_vertex[i] << endl;
    }
}

int get_is_matched_edge_h(int i, int j){
	return h_is_matched_edge[i*(num_vertices1 + num_vertices2+1) + j ];
}

__device__ 
void set_is_matched_edge(int i, int j, int value){
	d_is_matched_edge[i*(num_vertices1 + num_vertices2+1) + j ] = value;
}

__device__
void match_edges(int u, int v){
	set_is_matched_edge(u,v,1);
	set_is_matched_edge(v,u,1);
	d_is_matched_vertex[u] = 1;
	d_is_matched_vertex[v] = 1;
	d_partner_vertex[u] = v;
	d_partner_vertex[v] = u;

}

// Unmatching edges also unmatches the vertices since the graph is a matching
__device__
void unmatch_edges(int u, int v){
	set_is_matched_edge(u,v,0);
	set_is_matched_edge(v,u,0);
	if(d_partner_vertex[u]==v){
		d_is_matched_vertex[u] = 0;
		d_partner_vertex[u] = -1;
	}
	if(d_partner_vertex[v]==u){
		d_is_matched_vertex[v] = 0;
		d_partner_vertex[v] = -1;
	}
}

// Make this parallel
__device__
void update_matchings(){
	for(int i=1; i<=num_vertices1+num_vertices2; i++){
		int vertex = i;
		if(d_is_parent_change[vertex] == true){
			
			// cout << "Found aug. path till " << vertex << endl;
			// There should always be odd number of vertices in aug. path
			int path_length = 1;
			int parent = d_bfs_parent[vertex];
			while(parent!=vertex){
				// cout << vertex << " " <<parent << endl;
				if(path_length%2==1){
					match_edges(vertex, parent);
					// printf("Matching %d and %d \n", vertex, parent);
				}
				else{
					unmatch_edges(vertex, parent);
					// printf("Unmatching %d and %d \n", vertex, parent);
				}
				vertex =  d_bfs_parent[vertex];
				parent = d_bfs_parent[vertex];
				path_length++;
			}
		}
	}
}



__device__
void clear_visited(){
	int tid = blockIdx.x*1024 + threadIdx.x;
	int vertex1 = tid + 1;

	if(vertex1<=num_vertices1 + num_vertices2){
		d_visited[vertex1] = 0;
	}
}


__device__ 
void clear_bfs_parent(){
	int tid = blockIdx.x*1024 + threadIdx.x;
	int vertex1 = tid + 1;

	if(vertex1<=num_vertices1 + num_vertices2){
		d_bfs_parent[vertex1] = vertex1;
	}
}

__device__ 
void initialise_partner_vertex(){
	int tid = blockIdx.x*1024 + threadIdx.x;
	int vertex1 = tid + 1;

	if(vertex1<=num_vertices1 + num_vertices2){
		d_partner_vertex[vertex1] = -1;
	}
}

__device__ 
void clear_is_parent_change(){
	int tid = blockIdx.x*1024 + threadIdx.x;
	int vertex1 = tid + 1;

	if(vertex1<=num_vertices1 + num_vertices2){
		d_is_parent_change[vertex1] = -1;
	}
}

__device__ 
void copy_frontier(int *my_frontier, int *my_next_frontier){
	
	for (int i=1;i<=num_vertices1+num_vertices2;i++){
		my_frontier[i] = my_next_frontier[i];
	}
}

__device__ 
void clear_frontier(int *my_frontier, int *my_next_frontier ){
	for (int i=1;i<=num_vertices1+num_vertices2;i++){
			my_frontier[i] = 0;
			my_next_frontier[i] = 0;
	}
}
__device__
void vertex_disjoint_bfs(int binary_level, int vertex, int tid){
	int frontier_element = vertex;
	// printf("Frontier element: %d \n", frontier_element );
	if(!d_frontier[vertex]){
		return;
	}

	// my_frontier[vertex] = 1;
	d_frontier[vertex] = 1;


	// Iterate all frontier elements
	if(frontier_element!=-1){
	
		int vertex = frontier_element;
		
		// Make this atomic
		d_visited[vertex] = true;
		
		// cout << "Frontier: " << frontier_element << endl;
		// cout << "Continuining for vertex: " << vertex << endl;
		
		bool found_path = false;
		int start_edge = d_list_ptr[vertex];
		int end_edge = d_list_ptr[vertex + 1]; 
		
		// cout << "Start-End edge " <<  start_edge << " " << end_edge  << endl;
		// printf ("Start-End edge %d %d \n", start_edge, end_edge);
		for(int j=start_edge;j<end_edge;j++){
			if(found_path)
				break;


			int neighbor = d_flat_adj_list[j];

			// printf("[%d]Processing %d %d \n", tid, vertex, neighbor);

			int visited = atomicExch(&d_visited[neighbor], 1);

			if(!visited){
				// We want to alternate between unmatched and matched edges, otherwise we ignore
				d_visited[neighbor] = true;
				// cout << "Processing: " << vertex << " " << neighbor << endl;
				// printf("[%d]Processing %d %d \n", tid, vertex, neighbor);
				// exit(0);
				d_bfs_parent[neighbor] = vertex;

				if( binary_level==0 && get_is_matched_edge(vertex, neighbor)==0 && d_is_matched_vertex[neighbor]==1 ){
					// next_frontier.push_back(neighbor);
					d_next_frontier[neighbor] = 1;
				}

				// is_matched_vertex is implicitly true since the edge is matched
				// In level 1, we are only interested in matched edges
				else if( binary_level==1 && get_is_matched_edge(vertex, neighbor)==1 ){
					// next_frontier.push_back(neighbor);
					d_next_frontier[neighbor] = 1;
					// If I have found a path to the next level; I have to break
					// found_path = 1;
					return;
				}

				// Changing parent change only for this node
				else if(binary_level==0 && get_is_matched_edge(vertex, neighbor)==0 && d_is_matched_vertex[neighbor]==0){
					// printf("Found a aug. path with %d with parent: %d \n", neighbor, vertex);
					d_is_parent_change[neighbor] = 1;
					num_aug_paths++ ;
					// remove this return so that multiple paths can be found 
					return;
				}
			}
		}

		// frontier_element = get_frontier_element(vertex);
		// Getting next frontier element
		for(int x = vertex+1; x <=num_vertices1 + num_vertices2; x++){
			if(d_frontier[x]==1){
				frontier_element = x;
				break;
			}
		}
	}
	copy_frontier(d_frontier, d_next_frontier);
	// Only working for one level for now
	// bfs(binary_level = !binary_level);

}


__global__
void vertex_disjoint_bfs_util(){

	// parallelise these functions
	clear_visited();
	clear_bfs_parent();
	clear_is_parent_change();

	// clear_frontier(my_frontier, my_next_frontier );
	initialise_partner_vertex();


	// int num_aug_paths = 1000;

	int tid = blockIdx.x*1024 + threadIdx.x;


	int vertex = tid+1;
	if(vertex > num_vertices1)
		return;

	if(!d_visited[vertex] && !d_is_matched_vertex[vertex]){
		d_frontier[vertex] = 1;
		vertex_disjoint_bfs(0, vertex, tid);
		__syncthreads();
	}


	if(num_aug_paths > 0){
		update_matchings();
	}

	// return num_aug_paths;
}

int check_matching(){
	int total_matched = 0;
	for(int i=1;i<=num_vertices1+num_vertices2;i++){
		int vertex = i;
		int num_matched = 0;


		for(int j=h_list_ptr[i];j<h_list_ptr[i+1];j++){
			int neighbor = h_flat_adj_list[j];
			// cout << vertex << " " << neighbor << endl;
			if(get_is_matched_edge_h(vertex, neighbor)){
				num_matched++;
			}
		}


		if(num_matched==1){
			// cout << "Hi" << endl;
			total_matched++;
		}
		if(num_matched>1){
			cout << vertex << endl;
			cout << "Error! Not a matching!";
			exit(0);
		}
	}
	return total_matched/2;
}


int main(){
	struct timespec start, end;

	h_is_matched_edge = (bool *)calloc( (num_vertices1+ num_vertices2 + 1)*(num_vertices1 + num_vertices2+1), sizeof(bool));

	h_flat_adj_list = (int *)malloc(2*num_edges*sizeof(int));
	h_degree = (int *)malloc((num_vertices1+num_vertices2+1)*sizeof(int));
	h_list_ptr = (int *)malloc((num_vertices1+num_vertices2+2)*sizeof(int));
	h_list_ptr_copy = (int *)malloc((num_vertices1+num_vertices2+2)*sizeof(int));
	h_is_matched_vertex = (bool *)malloc((num_vertices1+num_vertices2+1)*sizeof(bool));
	h_partner_vertex = (int *)malloc((num_vertices1+num_vertices2+1)*sizeof(int));
	h_visited = (int *)malloc((num_vertices1+num_vertices2+1)*sizeof(int));
	h_bfs_parent = (int *)malloc((num_vertices1+num_vertices2+1)*sizeof(int));
	h_is_parent_change = (bool *)malloc((num_vertices1+num_vertices2+1)*sizeof(bool));
	h_frontier = (int *)malloc((num_vertices1+num_vertices2+1)*sizeof(int));
	h_next_frontier = (int *)malloc((num_vertices1+num_vertices2+1)*sizeof(int));

	// Add a check for null memory

	memset(h_degree, 0, num_vertices1 + num_vertices2 +1);
	// memset(h_is_matched_edge, 0, (num_vertices1 + num_vertices2 +1)*(num_vertices1+num_vertices2+1));
	memset(h_is_matched_vertex, 0, num_vertices1 + num_vertices2 +1);
	memset(h_visited, 0, num_vertices1 + num_vertices2 +1);
	memset(h_is_parent_change, 0, num_vertices1 + num_vertices2 +1);
	memset(h_frontier, 0, num_vertices1 + num_vertices2 +1);
	memset(h_next_frontier, 0, num_vertices1 + num_vertices2 +1);



	// to and from of edges
	// int h_edges_u[num_edges], h_edges_v[num_edges];			// Make this dynamic memory and free it once we have our 2 pass initialisation phase
	int *h_edges_u, *h_edges_v;
	h_edges_u = (int *)malloc((num_edges)*sizeof(int));
	h_edges_v = (int *)malloc((num_edges)*sizeof(int));


	ifstream fin;
    fin.open("FC_" + to_string(fc) + "_" + to_string(fc) + ".txt", ios::in);
    // fin.open("random_" + to_string(num_vertices1) + "_" + to_string(num_vertices2) + ".txt", ios::in);
    // cout << "random_" + to_string(num_vertices1) + "_" + to_string(num_vertices2) + ".txt" <<endl;
    int u, v;

    // cout << "Printing all the edges: \n";

    // Vertices with 0 edges are implicitly ignored while reading the file itself
    for(int i=0;i<num_edges;i++){
    		// cout << i << endl;
            fin >> u >> v;
            h_edges_u[i] = u;
            h_edges_v[i] = v;
            h_degree[u]++;
            h_degree[v]++;
    }

    cout << "Done reading edges" << endl;

    // Get pointer to adjacency list using prefix sum (no opti here since other parts are more complex anyway)
    // Index 0 will never be used.... the last elem
    h_list_ptr[1] = 0;
    h_list_ptr_copy[1] = h_list_ptr[1];
    for(int i=2;i<=num_vertices1+num_vertices2;i++){
    	h_list_ptr[i] = h_list_ptr[i-1] + h_degree[i-1];
    	h_list_ptr_copy[i] = h_list_ptr[i];
    }
    h_list_ptr[num_vertices1+num_vertices2+1] = 2*num_edges;       //For easy coding
    h_list_ptr_copy[num_vertices1+num_vertices2+1] = 2*num_edges;  // list_ptr has the start of the adj list ; list_ptr_copy has the current position

    
    for(int i=0;i<num_edges;i++){
    	h_flat_adj_list[h_list_ptr_copy[h_edges_u[i]]] = h_edges_v[i];
    	h_flat_adj_list[h_list_ptr_copy[h_edges_v[i]]] = h_edges_u[i];
    	h_list_ptr_copy[h_edges_u[i]]++;
    	h_list_ptr_copy[h_edges_v[i]]++;
    }
    

    clock_gettime( CLOCK_REALTIME,&start);

	cudaMemcpyToSymbol(d_is_matched_edge, h_is_matched_edge, (num_vertices1+ num_vertices2 + 1)*(num_vertices1 + num_vertices2+1)*sizeof(int),0,cudaMemcpyHostToDevice);
	cudaMemcpyToSymbol(d_flat_adj_list, h_flat_adj_list, 2*num_edges*sizeof(int),0,cudaMemcpyHostToDevice);
	cudaMemcpyToSymbol(d_degree, h_degree, (num_vertices1+num_vertices2+1)*sizeof(int),0,cudaMemcpyHostToDevice);
	cudaMemcpyToSymbol(d_list_ptr, h_list_ptr, (num_vertices1+num_vertices2+2)*sizeof(int),0,cudaMemcpyHostToDevice);
	cudaMemcpyToSymbol(d_is_matched_vertex, h_is_matched_vertex, (num_vertices1+num_vertices2+1)*sizeof(int),0,cudaMemcpyHostToDevice);
	cudaMemcpyToSymbol(d_visited, h_visited, (num_vertices1+num_vertices2+1)*sizeof(int),0,cudaMemcpyHostToDevice);
	cudaMemcpyToSymbol(d_frontier, h_frontier, (num_vertices1+num_vertices2+2)*sizeof(int),0,cudaMemcpyHostToDevice);
	
	// cudaMemcpy(d_is_matched_edge, h_is_matched_edge, (num_vertices1+ num_vertices2 + 1)*(num_vertices1 + num_vertices2+1)*sizeof(int),cudaMemcpyHostToDevice);
	// cudaMemcpy(d_flat_adj_list, h_flat_adj_list, 2*num_edges*sizeof(int),cudaMemcpyHostToDevice);
	// cudaMemcpy(d_degree, h_degree, (num_vertices1+num_vertices2+1)*sizeof(int),cudaMemcpyHostToDevice);
	// cudaMemcpy(d_list_ptr, h_list_ptr, (num_vertices1+num_vertices2+2)*sizeof(int),cudaMemcpyHostToDevice);
	// cudaMemcpy(d_is_matched_vertex, h_is_matched_vertex, (num_vertices1+num_vertices2+1)*sizeof(int),cudaMemcpyHostToDevice);
	// cudaMemcpy(d_visited, h_visited, (num_vertices1+num_vertices2+1)*sizeof(int),cudaMemcpyHostToDevice);
	// cudaMemcpy(d_frontier, h_frontier, (num_vertices1+num_vertices2+2)*sizeof(int),cudaMemcpyHostToDevice);


  	cudaDeviceSynchronize();
  	vertex_disjoint_bfs_util<<<1, num_threads>>>();
// 1219611
  	cudaDeviceSynchronize();
  	cudaError_t bla = cudaMemcpyFromSymbol(h_is_matched_edge, d_is_matched_edge, sizeof(d_is_matched_edge),0,cudaMemcpyDeviceToHost);
  	cudaMemcpyFromSymbol(h_partner_vertex, d_partner_vertex, sizeof(d_partner_vertex),0,cudaMemcpyDeviceToHost);
  	
  	clock_gettime( CLOCK_REALTIME,&end);
  	cudaDeviceSynchronize();
  

  	int num_matches = check_matching();
  	
  	// print_matchings();

  	printf("Number of matchings: %d \n", num_matches);

  	double elapsed = (end.tv_sec-start.tv_sec)*1000000000 + end.tv_nsec-start.tv_nsec;
  	printf("Time elapsed %lf\n", elapsed/1e6);
	

  	cudaDeviceSynchronize();
  


}
	.file	"tmpxft_0037170e_00000000-6_list_parallel_matching.cudafe1.cpp"
	.text
#APP
	.globl _ZSt21ios_base_library_initv

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB3784:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE3784:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.section	.rodata._ZNSt7__cxx119to_stringEi.str1.1,"aMS",@progbits,1
.LC0:
	.string	"00010203040506070809101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899"
	.section	.text._ZNSt7__cxx119to_stringEi,"axG",@progbits,_ZNSt7__cxx119to_stringEi,comdat
	.weak	_ZNSt7__cxx119to_stringEi
	.type	_ZNSt7__cxx119to_stringEi, @function
_ZNSt7__cxx119to_stringEi:
.LFB2826:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA2826
	endbr64
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	movl	%esi, %r13d
	movl	$10000, %ecx
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	shrl	$31, %r13d
	movq	%rdi, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	movl	$1, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	movl	%esi, %ebx
	subq	$232, %rsp
	.cfi_def_cfa_offset 272
	movq	%fs:40, %rax
	movq	%rax, 216(%rsp)
	xorl	%eax, %eax
	negl	%ebx
	cmovs	%esi, %ebx
	movl	%ebx, %eax
.L8:
	cmpl	$9, %eax
	jbe	.L4
	cmpl	$99, %eax
	ja	.L5
	incl	%ebp
	jmp	.L4
.L5:
	cmpl	$999, %eax
	ja	.L6
	addl	$2, %ebp
	jmp	.L4
.L6:
	cmpl	$9999, %eax
	ja	.L7
	addl	$3, %ebp
	jmp	.L4
.L7:
	movl	%eax, %eax
	xorl	%edx, %edx
	addl	$4, %ebp
	divq	%rcx
	jmp	.L8
.L4:
	leaq	16(%r12), %rax
	leal	0(%r13,%rbp), %esi
	movq	%r12, %rdi
	decl	%ebp
	movq	%rax, (%r12)
	movl	$45, %edx
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE12_M_constructEmc@PLT
	movzbl	%r13b, %r8d
	leaq	15(%rsp), %rdi
	leaq	.LC0(%rip), %rsi
	movl	$201, %ecx
	addq	(%r12), %r8
	rep movsb
	movl	$100, %ecx
.L9:
	cmpl	$99, %ebx
	jbe	.L16
	movl	%ebx, %eax
	xorl	%edx, %edx
	movl	%ebp, %esi
	divl	%ecx
	addl	%edx, %edx
	movl	%eax, %ebx
	leal	1(%rdx), %eax
	movb	15(%rsp,%rdx), %dl
	movb	15(%rsp,%rax), %al
	movb	%al, (%r8,%rsi)
	leal	-1(%rbp), %eax
	subl	$2, %ebp
	movb	%dl, (%r8,%rax)
	jmp	.L9
.L16:
	leal	48(%rbx), %eax
	cmpl	$9, %ebx
	jbe	.L12
	addl	%ebx, %ebx
	leal	1(%rbx), %eax
	movb	15(%rsp,%rax), %al
	movb	%al, 1(%r8)
	movb	15(%rsp,%rbx), %al
.L12:
	movb	%al, (%r8)
	movq	216(%rsp), %rax
	subq	%fs:40, %rax
	je	.L13
	call	__stack_chk_fail@PLT
.L13:
	addq	$232, %rsp
	.cfi_def_cfa_offset 40
	movq	%r12, %rax
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%rbp
	.cfi_def_cfa_offset 24
	popq	%r12
	.cfi_def_cfa_offset 16
	popq	%r13
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2826:
	.globl	__gxx_personality_v0
	.section	.gcc_except_table._ZNSt7__cxx119to_stringEi,"aG",@progbits,_ZNSt7__cxx119to_stringEi,comdat
.LLSDA2826:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE2826-.LLSDACSB2826
.LLSDACSB2826:
.LLSDACSE2826:
	.section	.text._ZNSt7__cxx119to_stringEi,"axG",@progbits,_ZNSt7__cxx119to_stringEi,comdat
	.size	_ZNSt7__cxx119to_stringEi, .-_ZNSt7__cxx119to_stringEi
	.text
	.globl	_Z19get_is_matched_edgeii
	.type	_Z19get_is_matched_edgeii, @function
_Z19get_is_matched_edgeii:
.LFB3766:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE3766:
	.size	_Z19get_is_matched_edgeii, .-_Z19get_is_matched_edgeii
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC1:
	.string	"Matchings: "
.LC2:
	.string	" "
	.text
	.globl	_Z15print_matchingsv
	.type	_Z15print_matchingsv, @function
_Z15print_matchingsv:
.LFB3767:
	.cfi_startproc
	endbr64
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	leaq	.LC1(%rip), %rsi
	leaq	.LC2(%rip), %r13
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	leaq	_ZSt4cout(%rip), %r12
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	movq	%r12, %rdi
	movl	$4, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	movl	$1, %ebx
	pushq	%rdx
	.cfi_def_cfa_offset 48
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%rax, %rdi
	call	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@PLT
.L20:
	movl	%ebx, %esi
	movq	%r12, %rdi
	incl	%ebx
	call	_ZNSolsEi@PLT
	movq	%r13, %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%rax, %rdi
	movq	h_partner_vertex(%rip), %rax
	movl	(%rax,%rbp), %esi
	addq	$4, %rbp
	call	_ZNSolsEi@PLT
	movq	%rax, %rdi
	call	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@PLT
	cmpl	$2001, %ebx
	jne	.L20
	popq	%rax
	.cfi_def_cfa_offset 40
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%rbp
	.cfi_def_cfa_offset 24
	popq	%r12
	.cfi_def_cfa_offset 16
	popq	%r13
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3767:
	.size	_Z15print_matchingsv, .-_Z15print_matchingsv
	.globl	_Z21get_is_matched_edge_hii
	.type	_Z21get_is_matched_edge_hii, @function
_Z21get_is_matched_edge_hii:
.LFB3768:
	.cfi_startproc
	endbr64
	movslq	%edi, %rdi
	movslq	%esi, %rsi
	imulq	$2001, %rdi, %rdi
	addq	h_is_matched_edge(%rip), %rdi
	movzbl	(%rdi,%rsi), %eax
	ret
	.cfi_endproc
.LFE3768:
	.size	_Z21get_is_matched_edge_hii, .-_Z21get_is_matched_edge_hii
	.globl	_Z19set_is_matched_edgeiii
	.type	_Z19set_is_matched_edgeiii, @function
_Z19set_is_matched_edgeiii:
.LFB3769:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE3769:
	.size	_Z19set_is_matched_edgeiii, .-_Z19set_is_matched_edgeiii
	.globl	_Z11match_edgesii
	.type	_Z11match_edgesii, @function
_Z11match_edgesii:
.LFB3770:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE3770:
	.size	_Z11match_edgesii, .-_Z11match_edgesii
	.globl	_Z13unmatch_edgesii
	.type	_Z13unmatch_edgesii, @function
_Z13unmatch_edgesii:
.LFB4509:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	pushq	%rax
	.cfi_def_cfa_offset 16
	call	_Z11match_edgesii
	.cfi_endproc
.LFE4509:
	.size	_Z13unmatch_edgesii, .-_Z13unmatch_edgesii
	.globl	_Z16update_matchingsv
	.type	_Z16update_matchingsv, @function
_Z16update_matchingsv:
.LFB3772:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE3772:
	.size	_Z16update_matchingsv, .-_Z16update_matchingsv
	.globl	_Z13clear_visitedv
	.type	_Z13clear_visitedv, @function
_Z13clear_visitedv:
.LFB4515:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	pushq	%rax
	.cfi_def_cfa_offset 16
	call	_Z16update_matchingsv
	.cfi_endproc
.LFE4515:
	.size	_Z13clear_visitedv, .-_Z13clear_visitedv
	.globl	_Z16clear_bfs_parentv
	.type	_Z16clear_bfs_parentv, @function
_Z16clear_bfs_parentv:
.LFB4517:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	pushq	%rax
	.cfi_def_cfa_offset 16
	call	_Z16update_matchingsv
	.cfi_endproc
.LFE4517:
	.size	_Z16clear_bfs_parentv, .-_Z16clear_bfs_parentv
	.globl	_Z25initialise_partner_vertexv
	.type	_Z25initialise_partner_vertexv, @function
_Z25initialise_partner_vertexv:
.LFB4519:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	pushq	%rax
	.cfi_def_cfa_offset 16
	call	_Z16update_matchingsv
	.cfi_endproc
.LFE4519:
	.size	_Z25initialise_partner_vertexv, .-_Z25initialise_partner_vertexv
	.globl	_Z22clear_is_parent_changev
	.type	_Z22clear_is_parent_changev, @function
_Z22clear_is_parent_changev:
.LFB4521:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	pushq	%rax
	.cfi_def_cfa_offset 16
	call	_Z16update_matchingsv
	.cfi_endproc
.LFE4521:
	.size	_Z22clear_is_parent_changev, .-_Z22clear_is_parent_changev
	.globl	_Z13copy_frontierPiS_
	.type	_Z13copy_frontierPiS_, @function
_Z13copy_frontierPiS_:
.LFB3777:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE3777:
	.size	_Z13copy_frontierPiS_, .-_Z13copy_frontierPiS_
	.globl	_Z14clear_frontierPiS_
	.type	_Z14clear_frontierPiS_, @function
_Z14clear_frontierPiS_:
.LFB4511:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	pushq	%rax
	.cfi_def_cfa_offset 16
	call	_Z13copy_frontierPiS_
	.cfi_endproc
.LFE4511:
	.size	_Z14clear_frontierPiS_, .-_Z14clear_frontierPiS_
	.globl	_Z19vertex_disjoint_bfsiii
	.type	_Z19vertex_disjoint_bfsiii, @function
_Z19vertex_disjoint_bfsiii:
.LFB4513:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE4513:
	.size	_Z19vertex_disjoint_bfsiii, .-_Z19vertex_disjoint_bfsiii
	.section	.rodata.str1.1
.LC3:
	.string	"Error! Not a matching!"
	.text
	.globl	_Z14check_matchingv
	.type	_Z14check_matchingv, @function
_Z14check_matchingv:
.LFB3780:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	xorl	%edx, %edx
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	pushq	%rcx
	.cfi_def_cfa_offset 32
	movq	h_list_ptr(%rip), %r10
	movl	$1, %ecx
	movq	h_flat_adj_list(%rip), %rbx
.L52:
	movl	4(%r10,%rcx,4), %ebp
	movslq	(%r10,%rcx,4), %r8
	movl	%ecx, %r11d
	xorl	%r9d, %r9d
.L47:
	cmpl	%r8d, %ebp
	jle	.L57
	movl	(%rbx,%r8,4), %esi
	movl	%r11d, %edi
	call	_Z21get_is_matched_edge_hii
	cmpl	$1, %eax
	sbbl	$-1, %r9d
	incq	%r8
	jmp	.L47
.L57:
	decl	%r9d
	jne	.L50
	incl	%edx
	jmp	.L51
.L50:
	jle	.L51
	leaq	_ZSt4cout(%rip), %rbx
	movl	%r11d, %esi
	movq	%rbx, %rdi
	call	_ZNSolsEi@PLT
	movq	%rax, %rdi
	call	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@PLT
	movq	%rbx, %rdi
	leaq	.LC3(%rip), %rsi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	xorl	%edi, %edi
	call	exit@PLT
.L51:
	incq	%rcx
	cmpq	$2001, %rcx
	jne	.L52
	movl	%edx, %eax
	movl	$2, %ecx
	cltd
	idivl	%ecx
	popq	%rdx
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3780:
	.size	_Z14check_matchingv, .-_Z14check_matchingv
	.globl	_Z43__device_stub__Z24vertex_disjoint_bfs_utilvv
	.type	_Z43__device_stub__Z24vertex_disjoint_bfs_utilvv, @function
_Z43__device_stub__Z24vertex_disjoint_bfs_utilvv:
.LFB3806:
	.cfi_startproc
	endbr64
	subq	$88, %rsp
	.cfi_def_cfa_offset 96
	movq	%fs:40, %rax
	movq	%rax, 72(%rsp)
	movabsq	$4294967297, %rax
	leaq	8(%rsp), %rcx
	movq	%rsp, %rdx
	leaq	28(%rsp), %rsi
	leaq	16(%rsp), %rdi
	movl	$1, 24(%rsp)
	movl	$1, 36(%rsp)
	movq	%rax, 16(%rsp)
	movq	%rax, 28(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L58
	pushq	8(%rsp)
	.cfi_def_cfa_offset 104
	leaq	_Z24vertex_disjoint_bfs_utilv(%rip), %rdi
	pushq	8(%rsp)
	.cfi_def_cfa_offset 112
	movq	44(%rsp), %rcx
	movl	52(%rsp), %r8d
	movq	32(%rsp), %rsi
	movl	40(%rsp), %edx
	leaq	80(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 104
	popq	%rdx
	.cfi_def_cfa_offset 96
.L58:
	movq	72(%rsp), %rax
	subq	%fs:40, %rax
	je	.L60
	call	__stack_chk_fail@PLT
.L60:
	addq	$88, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3806:
	.size	_Z43__device_stub__Z24vertex_disjoint_bfs_utilvv, .-_Z43__device_stub__Z24vertex_disjoint_bfs_utilvv
	.globl	_Z24vertex_disjoint_bfs_utilv
	.type	_Z24vertex_disjoint_bfs_utilv, @function
_Z24vertex_disjoint_bfs_utilv:
.LFB3807:
	.cfi_startproc
	endbr64
	jmp	_Z43__device_stub__Z24vertex_disjoint_bfs_utilvv
	.cfi_endproc
.LFE3807:
	.size	_Z24vertex_disjoint_bfs_utilv, .-_Z24vertex_disjoint_bfs_utilv
	.section	.rodata.str1.1
.LC4:
	.string	"_Z24vertex_disjoint_bfs_utilv"
.LC5:
	.string	"d_flat_adj_list"
.LC6:
	.string	"d_degree"
.LC7:
	.string	"d_list_ptr"
.LC8:
	.string	"d_list_ptr_copy"
.LC9:
	.string	"d_is_matched_edge"
.LC10:
	.string	"d_is_matched_vertex"
.LC11:
	.string	"d_partner_vertex"
.LC12:
	.string	"d_visited"
.LC13:
	.string	"d_bfs_parent"
.LC14:
	.string	"d_is_parent_change"
.LC15:
	.string	"d_frontier"
.LC16:
	.string	"d_next_frontier"
.LC17:
	.string	"num_aug_paths"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB3809:
	.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	.LC4(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z24vertex_disjoint_bfs_utilv(%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	%r8d, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC5(%rip), %rdx
	movl	$8000000, %r9d
	leaq	_ZL15d_flat_adj_list(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	call	__cudaRegisterVar@PLT
	popq	%rax
	.cfi_def_cfa_offset 24
	popq	%rdx
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC6(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$8004, %r9d
	leaq	_ZL8d_degree(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%rcx
	.cfi_def_cfa_offset 24
	popq	%rsi
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC7(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$8008, %r9d
	leaq	_ZL10d_list_ptr(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%rdi
	.cfi_def_cfa_offset 24
	popq	%r8
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC8(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$8008, %r9d
	leaq	_ZL15d_list_ptr_copy(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%r9
	.cfi_def_cfa_offset 24
	popq	%r10
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC9(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$4004001, %r9d
	leaq	_ZL17d_is_matched_edge(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%r11
	.cfi_def_cfa_offset 24
	popq	%rax
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC10(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$2001, %r9d
	leaq	_ZL19d_is_matched_vertex(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%rax
	.cfi_def_cfa_offset 24
	popq	%rdx
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC11(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$8004, %r9d
	leaq	_ZL16d_partner_vertex(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%rcx
	.cfi_def_cfa_offset 24
	popq	%rsi
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC12(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$8004, %r9d
	leaq	_ZL9d_visited(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%rdi
	.cfi_def_cfa_offset 24
	popq	%r8
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC13(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$8004, %r9d
	leaq	_ZL12d_bfs_parent(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%r9
	.cfi_def_cfa_offset 24
	popq	%r10
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC14(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$2001, %r9d
	leaq	_ZL18d_is_parent_change(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%r11
	.cfi_def_cfa_offset 24
	popq	%rax
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC15(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$8004, %r9d
	leaq	_ZL10d_frontier(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%rax
	.cfi_def_cfa_offset 24
	popq	%rdx
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC16(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$8004, %r9d
	leaq	_ZL15d_next_frontier(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%rcx
	.cfi_def_cfa_offset 24
	popq	%rsi
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC17(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movl	$4, %r9d
	movq	%rdx, %rcx
	leaq	_ZL13num_aug_paths(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%rdi
	.cfi_def_cfa_offset 24
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	popq	%r8
	.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
.LFE3809:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.rodata.str1.1
.LC18:
	.string	"FC_"
.LC19:
	.string	"_"
.LC20:
	.string	".txt"
.LC21:
	.string	"Done reading edges"
.LC22:
	.string	"Number of matchings: %d \n"
.LC24:
	.string	"Time elapsed %lf\n"
	.section	.text.startup
	.globl	main
	.type	main, @function
main:
.LFB3781:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA3781
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movl	$1, %esi
	movl	$4004001, %edi
	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
	subq	$824, %rsp
	.cfi_def_cfa_offset 880
	movq	%fs:40, %rax
	movq	%rax, 808(%rsp)
	xorl	%eax, %eax
	call	calloc@PLT
	movl	$8000000, %edi
	movq	%rax, h_is_matched_edge(%rip)
	call	malloc@PLT
	movl	$8004, %edi
	movq	%rax, h_flat_adj_list(%rip)
	call	malloc@PLT
	movl	$8008, %edi
	movq	%rax, %r14
	movq	%rax, h_degree(%rip)
	call	malloc@PLT
	movl	$8008, %edi
	movq	%rax, h_list_ptr(%rip)
	call	malloc@PLT
	movl	$2001, %edi
	movq	%rax, h_list_ptr_copy(%rip)
	call	malloc@PLT
	movl	$8004, %edi
	movq	%rax, %r13
	movq	%rax, h_is_matched_vertex(%rip)
	call	malloc@PLT
	movl	$8004, %edi
	movq	%rax, h_partner_vertex(%rip)
	call	malloc@PLT
	movl	$8004, %edi
	movq	%rax, %r12
	movq	%rax, h_visited(%rip)
	call	malloc@PLT
	movl	$2001, %edi
	movq	%rax, h_bfs_parent(%rip)
	call	malloc@PLT
	movl	$8004, %edi
	movq	%rax, %rbp
	movq	%rax, h_is_parent_change(%rip)
	call	malloc@PLT
	movl	$8004, %edi
	movq	%rax, %rbx
	movq	%rax, h_frontier(%rip)
	call	malloc@PLT
	movq	%r14, %rdi
	movl	$2001, %ecx
	movq	%rax, %rdx
	movq	%rax, h_next_frontier(%rip)
	xorl	%eax, %eax
	rep stosb
	movq	%r13, %rdi
	movl	$2001, %ecx
	rep stosb
	movq	%r12, %rdi
	movl	$2001, %ecx
	leaq	288(%rsp), %r12
	rep stosb
	movl	$2001, %ecx
	movq	%rbp, %rdi
	rep stosb
	movl	$2001, %ecx
	movq	%rbx, %rdi
	rep stosb
	movq	%rdx, %rdi
	movl	$2001, %ecx
	rep stosb
	movl	$4000000, %edi
	call	malloc@PLT
	movl	$4000000, %edi
	movq	%rax, %r14
	call	malloc@PLT
	movq	%r12, %rdi
	movq	%rax, %r13
.LEHB0:
	call	_ZNSt14basic_ifstreamIcSt11char_traitsIcEEC1Ev@PLT
.LEHE0:
	leaq	192(%rsp), %rax
	movl	fc(%rip), %esi
	movq	%rax, %rdi
	movq	%rax, 8(%rsp)
	call	_ZNSt7__cxx119to_stringEi
	leaq	96(%rsp), %rax
	movl	fc(%rip), %esi
	movq	%rax, %rdi
	movq	%rax, 16(%rsp)
	call	_ZNSt7__cxx119to_stringEi
	movq	16(%rsp), %rdi
	leaq	.LC18(%rip), %rdx
	xorl	%esi, %esi
.LEHB1:
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE6insertEmPKc@PLT
.LEHE1:
	movq	%rax, %rsi
	leaq	128(%rsp), %rax
	movq	%rax, %rdi
	movq	%rax, 24(%rsp)
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC1EOS4_@PLT
	movq	24(%rsp), %rdi
	leaq	.LC19(%rip), %rsi
.LEHB2:
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE6appendEPKc@PLT
.LEHE2:
	leaq	160(%rsp), %rbx
	movq	%rax, %rsi
	leaq	224(%rsp), %rbp
	movq	%rbx, %rdi
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC1EOS4_@PLT
	movq	%rbx, %rdi
	movq	200(%rsp), %r15
	addq	168(%rsp), %r15
	call	_ZNKSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE8capacityEv@PLT
	cmpq	%r15, %rax
	jnb	.L66
	movq	8(%rsp), %rdi
	call	_ZNKSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE8capacityEv@PLT
	cmpq	%r15, %rax
	jb	.L66
	movq	8(%rsp), %rdi
	movq	%rbx, %rdx
	xorl	%esi, %esi
.LEHB3:
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE6insertEmRKS4_@PLT
	jmp	.L93
.L66:
	movq	8(%rsp), %rsi
	movq	%rbx, %rdi
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE6appendERKS4_@PLT
.LEHE3:
.L93:
	movq	%rax, %rsi
	movq	%rbp, %rdi
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC1EOS4_@PLT
	leaq	.LC20(%rip), %rsi
	movq	%rbp, %rdi
.LEHB4:
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE6appendEPKc@PLT
.LEHE4:
	leaq	256(%rsp), %r15
	movq	%rax, %rsi
	movq	%r15, %rdi
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC1EOS4_@PLT
	movq	256(%rsp), %rsi
	movl	$8, %edx
	leaq	304(%rsp), %rdi
.LEHB5:
	call	_ZNSt13basic_filebufIcSt11char_traitsIcEE4openEPKcSt13_Ios_Openmode@PLT
	movq	288(%rsp), %rdx
	movq	-24(%rdx), %rdi
	addq	%r12, %rdi
	testq	%rax, %rax
	jne	.L68
	movl	$4, %esi
	call	_ZNSt9basic_iosIcSt11char_traitsIcEE8setstateESt12_Ios_Iostate@PLT
	jmp	.L69
.L68:
	xorl	%esi, %esi
	call	_ZNSt9basic_iosIcSt11char_traitsIcEE5clearESt12_Ios_Iostate@PLT
.LEHE5:
.L69:
	movq	%r15, %rdi
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE10_M_disposeEv@PLT
	movq	%rbp, %rdi
	leaq	44(%rsp), %rbp
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE10_M_disposeEv@PLT
	movq	%rbx, %rdi
	xorl	%ebx, %ebx
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE10_M_disposeEv@PLT
	movq	24(%rsp), %rdi
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE10_M_disposeEv@PLT
	movq	16(%rsp), %rdi
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE10_M_disposeEv@PLT
	movq	8(%rsp), %rdi
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE10_M_disposeEv@PLT
.L70:
	movq	%rbp, %rsi
	movq	%r12, %rdi
.LEHB6:
	call	_ZNSirsERi@PLT
	movq	%rax, %rdi
	leaq	48(%rsp), %rsi
	call	_ZNSirsERi@PLT
	movl	48(%rsp), %edx
	movslq	44(%rsp), %rax
	movl	%edx, 0(%r13,%rbx)
	movq	h_degree(%rip), %rdx
	movl	%eax, (%r14,%rbx)
	addq	$4, %rbx
	incl	(%rdx,%rax,4)
	movslq	48(%rsp), %rax
	incl	(%rdx,%rax,4)
	cmpq	$4000000, %rbx
	jne	.L70
	leaq	.LC21(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%rax, %rdi
	call	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@PLT
	movq	h_list_ptr(%rip), %rcx
	movq	h_list_ptr_copy(%rip), %rdx
	xorl	%eax, %eax
	xorl	%esi, %esi
	movq	h_degree(%rip), %rdi
	movl	%eax, 4(%rcx)
	movl	$4, %eax
	movl	%esi, 4(%rdx)
.L71:
	movl	(%rdi,%rax), %esi
	addl	(%rcx,%rax), %esi
	movl	%esi, 4(%rcx,%rax)
	movl	%esi, 4(%rdx,%rax)
	addq	$4, %rax
	cmpq	$8000, %rax
	jne	.L71
	movl	$2000000, 8004(%rcx)
	movq	h_flat_adj_list(%rip), %rdi
	xorl	%eax, %eax
	movl	$2000000, 8004(%rdx)
.L72:
	movslq	(%r14,%rax), %rcx
	leaq	(%rdx,%rcx,4), %rsi
	movq	%rcx, %r8
	movslq	0(%r13,%rax), %rcx
	addq	$4, %rax
	movslq	(%rsi), %r9
	movl	%ecx, (%rdi,%r9,4)
	leaq	(%rdx,%rcx,4), %rcx
	movslq	(%rcx), %r9
	movl	%r8d, (%rdi,%r9,4)
	incl	(%rsi)
	incl	(%rcx)
	cmpq	$4000000, %rax
	jne	.L72
	leaq	64(%rsp), %rsi
	xorl	%edi, %edi
	leaq	_ZL17d_is_matched_edge(%rip), %rbx
	call	clock_gettime@PLT
	movl	$1, %r8d
	xorl	%ecx, %ecx
	movq	%rbx, %rdi
	movq	h_is_matched_edge(%rip), %rsi
	movl	$16016004, %edx
	call	cudaMemcpyToSymbol@PLT
	movq	h_flat_adj_list(%rip), %rsi
	movl	$1, %r8d
	xorl	%ecx, %ecx
	movl	$8000000, %edx
	leaq	_ZL15d_flat_adj_list(%rip), %rdi
	call	cudaMemcpyToSymbol@PLT
	movq	h_degree(%rip), %rsi
	movl	$1, %r8d
	xorl	%ecx, %ecx
	movl	$8004, %edx
	leaq	_ZL8d_degree(%rip), %rdi
	call	cudaMemcpyToSymbol@PLT
	movq	h_list_ptr(%rip), %rsi
	movl	$1, %r8d
	xorl	%ecx, %ecx
	movl	$8008, %edx
	leaq	_ZL10d_list_ptr(%rip), %rdi
	call	cudaMemcpyToSymbol@PLT
	movq	h_is_matched_vertex(%rip), %rsi
	movl	$1, %r8d
	xorl	%ecx, %ecx
	movl	$8004, %edx
	leaq	_ZL19d_is_matched_vertex(%rip), %rdi
	call	cudaMemcpyToSymbol@PLT
	movq	h_visited(%rip), %rsi
	movl	$1, %r8d
	xorl	%ecx, %ecx
	movl	$8004, %edx
	leaq	_ZL9d_visited(%rip), %rdi
	call	cudaMemcpyToSymbol@PLT
	movq	h_frontier(%rip), %rsi
	movl	$1, %r8d
	xorl	%ecx, %ecx
	movl	$8008, %edx
	leaq	_ZL10d_frontier(%rip), %rdi
	call	cudaMemcpyToSymbol@PLT
	call	cudaDeviceSynchronize@PLT
	movl	$536871037, %edx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$1, %ecx
	salq	$3, %rdx
	movl	$1, %esi
	movabsq	$4294967297, %rdi
	movl	$1, 88(%rsp)
	movq	%rdx, 80(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L73
	call	_Z43__device_stub__Z24vertex_disjoint_bfs_utilvv
.L73:
	call	cudaDeviceSynchronize@PLT
	movl	$2, %r8d
	xorl	%ecx, %ecx
	movl	$4004001, %edx
	movq	%rbx, %rsi
	movq	h_is_matched_edge(%rip), %rdi
	call	cudaMemcpyFromSymbol@PLT
	movq	h_partner_vertex(%rip), %rdi
	movl	$2, %r8d
	xorl	%ecx, %ecx
	movl	$8004, %edx
	leaq	_ZL16d_partner_vertex(%rip), %rsi
	call	cudaMemcpyFromSymbol@PLT
	leaq	80(%rsp), %rsi
	xorl	%edi, %edi
	call	clock_gettime@PLT
	call	cudaDeviceSynchronize@PLT
	call	_Z14check_matchingv
	movl	%eax, %edx
	leaq	.LC22(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	80(%rsp), %rax
	subq	64(%rsp), %rax
	movl	$2, %edi
	leaq	.LC24(%rip), %rsi
	imulq	$1000000000, %rax, %rax
	addq	88(%rsp), %rax
	subq	72(%rsp), %rax
	cvtsi2sdq	%rax, %xmm0
	movb	$1, %al
	divsd	.LC23(%rip), %xmm0
	call	__printf_chk@PLT
	call	cudaDeviceSynchronize@PLT
.LEHE6:
	leaq	288(%rsp), %rdi
	call	_ZNSt14basic_ifstreamIcSt11char_traitsIcEED1Ev@PLT
	movq	808(%rsp), %rax
	subq	%fs:40, %rax
	je	.L81
	jmp	.L91
.L86:
	endbr64
	movq	%rax, %r13
.L74:
	movq	%r15, %rdi
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE10_M_disposeEv@PLT
	jmp	.L75
.L85:
	endbr64
	movq	%rax, %r13
.L75:
	movq	%rbp, %rdi
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE10_M_disposeEv@PLT
	jmp	.L76
.L84:
	endbr64
	movq	%rax, %r13
.L76:
	movq	%rbx, %rdi
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE10_M_disposeEv@PLT
	jmp	.L77
.L83:
	endbr64
	movq	%rax, %r13
.L77:
	movq	24(%rsp), %rdi
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE10_M_disposeEv@PLT
	jmp	.L78
.L82:
	endbr64
	movq	%rax, %r13
.L78:
	movq	16(%rsp), %rdi
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE10_M_disposeEv@PLT
	movq	8(%rsp), %rdi
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE10_M_disposeEv@PLT
	jmp	.L79
.L87:
	endbr64
	movq	%rax, %r13
.L79:
	movq	%r12, %rdi
	call	_ZNSt14basic_ifstreamIcSt11char_traitsIcEED1Ev@PLT
	movq	808(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L91
	movq	%r13, %rdi
.LEHB7:
	call	_Unwind_Resume@PLT
.LEHE7:
.L91:
	call	__stack_chk_fail@PLT
.L81:
	addq	$824, %rsp
	.cfi_def_cfa_offset 56
	xorl	%eax, %eax
	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
.LFE3781:
	.section	.gcc_except_table,"a",@progbits
.LLSDA3781:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE3781-.LLSDACSB3781
.LLSDACSB3781:
	.uleb128 .LEHB0-.LFB3781
	.uleb128 .LEHE0-.LEHB0
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB1-.LFB3781
	.uleb128 .LEHE1-.LEHB1
	.uleb128 .L82-.LFB3781
	.uleb128 0
	.uleb128 .LEHB2-.LFB3781
	.uleb128 .LEHE2-.LEHB2
	.uleb128 .L83-.LFB3781
	.uleb128 0
	.uleb128 .LEHB3-.LFB3781
	.uleb128 .LEHE3-.LEHB3
	.uleb128 .L84-.LFB3781
	.uleb128 0
	.uleb128 .LEHB4-.LFB3781
	.uleb128 .LEHE4-.LEHB4
	.uleb128 .L85-.LFB3781
	.uleb128 0
	.uleb128 .LEHB5-.LFB3781
	.uleb128 .LEHE5-.LEHB5
	.uleb128 .L86-.LFB3781
	.uleb128 0
	.uleb128 .LEHB6-.LFB3781
	.uleb128 .LEHE6-.LEHB6
	.uleb128 .L87-.LFB3781
	.uleb128 0
	.uleb128 .LEHB7-.LFB3781
	.uleb128 .LEHE7-.LEHB7
	.uleb128 0
	.uleb128 0
.LLSDACSE3781:
	.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
	.globl	h_next_frontier
	.bss
	.align 8
	.type	h_next_frontier, @object
	.size	h_next_frontier, 8
h_next_frontier:
	.zero	8
	.globl	h_frontier
	.align 8
	.type	h_frontier, @object
	.size	h_frontier, 8
h_frontier:
	.zero	8
	.globl	fc
	.data
	.align 4
	.type	fc, @object
	.size	fc, 4
fc:
	.long	1000
	.globl	h_is_parent_change
	.bss
	.align 8
	.type	h_is_parent_change, @object
	.size	h_is_parent_change, 8
h_is_parent_change:
	.zero	8
	.globl	h_bfs_parent
	.align 8
	.type	h_bfs_parent, @object
	.size	h_bfs_parent, 8
h_bfs_parent:
	.zero	8
	.globl	h_visited
	.align 8
	.type	h_visited, @object
	.size	h_visited, 8
h_visited:
	.zero	8
	.globl	h_partner_vertex
	.align 8
	.type	h_partner_vertex, @object
	.size	h_partner_vertex, 8
h_partner_vertex:
	.zero	8
	.globl	h_is_matched_vertex
	.align 8
	.type	h_is_matched_vertex, @object
	.size	h_is_matched_vertex, 8
h_is_matched_vertex:
	.zero	8
	.globl	h_is_matched_edge
	.align 8
	.type	h_is_matched_edge, @object
	.size	h_is_matched_edge, 8
h_is_matched_edge:
	.zero	8
	.globl	h_list_ptr_copy
	.align 8
	.type	h_list_ptr_copy, @object
	.size	h_list_ptr_copy, 8
h_list_ptr_copy:
	.zero	8
	.globl	h_list_ptr
	.align 8
	.type	h_list_ptr, @object
	.size	h_list_ptr, 8
h_list_ptr:
	.zero	8
	.globl	h_degree
	.align 8
	.type	h_degree, @object
	.size	h_degree, 8
h_degree:
	.zero	8
	.globl	h_flat_adj_list
	.align 8
	.type	h_flat_adj_list, @object
	.size	h_flat_adj_list, 8
h_flat_adj_list:
	.zero	8
	.local	_ZL13num_aug_paths
	.comm	_ZL13num_aug_paths,4,4
	.local	_ZL15d_next_frontier
	.comm	_ZL15d_next_frontier,8004,32
	.local	_ZL10d_frontier
	.comm	_ZL10d_frontier,8004,32
	.local	_ZL18d_is_parent_change
	.comm	_ZL18d_is_parent_change,2001,32
	.local	_ZL12d_bfs_parent
	.comm	_ZL12d_bfs_parent,8004,32
	.local	_ZL9d_visited
	.comm	_ZL9d_visited,8004,32
	.local	_ZL16d_partner_vertex
	.comm	_ZL16d_partner_vertex,8004,32
	.local	_ZL19d_is_matched_vertex
	.comm	_ZL19d_is_matched_vertex,2001,32
	.local	_ZL17d_is_matched_edge
	.comm	_ZL17d_is_matched_edge,4004001,32
	.local	_ZL15d_list_ptr_copy
	.comm	_ZL15d_list_ptr_copy,8008,32
	.local	_ZL10d_list_ptr
	.comm	_ZL10d_list_ptr,8008,32
	.local	_ZL8d_degree
	.comm	_ZL8d_degree,8004,32
	.local	_ZL15d_flat_adj_list
	.comm	_ZL15d_flat_adj_list,8000000,32
	.section	.rodata.cst8,"aM",@progbits,8
	.align 8
.LC23:
	.long	0
	.long	1093567616
	.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 <hip/hip_runtime.h>
#include<iostream>
#include<stdio.h>
#include<stdlib.h>
#include<math.h>
#include<fstream>
#include<time.h>
#include<sys/time.h>

using namespace std;

#define num_threads 1000
// #define num_edges 700000
// #define num_vertices1 10000
// #define num_vertices2 10000

// #define num_edges 1000000
// #define num_vertices1 1000
// #define num_vertices2 1000


#define lli long long int

// #define num_edges 2998468
// #define num_vertices1 100000
// #define num_vertices2 100000



const lli num_edges = 1000000;
const lli num_vertices1 = 1000;
const lli num_vertices2 = 1000;

// const lli num_edges = 291;
// const lli num_vertices1 = 100;
// const lli num_vertices2 = 100;

__device__ int d_flat_adj_list[2*num_edges];
__device__ int d_degree[num_vertices1+num_vertices2+1]={0};      //store degree of each vertex
__device__ int d_list_ptr[num_vertices1+num_vertices2+2];        //1-indexed and extra element at the end for easy size access  // Pointer to the start of adjacency list
__device__ int d_list_ptr_copy[num_vertices1+num_vertices2+2];    // Temporrary stuff, gotta sleep

__device__ bool d_is_matched_edge[(num_vertices1+ num_vertices2 + 1)*(num_vertices1 + num_vertices2+1)] = {0} ;     // Adjacency matrix (1-indexed)
__device__ bool d_is_matched_vertex[num_vertices1 + num_vertices2 + 1] = {0};	//is the vertex matched
__device__ int d_partner_vertex[num_vertices1 + num_vertices2 + 1];
__device__ int d_visited[num_vertices1 + num_vertices2 + 1] = {0};
__device__ int d_bfs_parent[num_vertices1 +  num_vertices2 + 1];
__device__ bool d_is_parent_change[num_vertices1 + num_vertices2 + 1] = {0};

__device__ int d_frontier[num_vertices1 + num_vertices2+1] = {0};
__device__ int d_next_frontier[num_vertices1+num_vertices2+1] = {0};

__device__ int num_aug_paths = 0;



int *h_flat_adj_list;
int *h_degree;
int * h_list_ptr;
int *h_list_ptr_copy;

bool *h_is_matched_edge;
bool *h_is_matched_vertex;
int *h_partner_vertex;
int *h_visited;
int *h_bfs_parent;
bool *h_is_parent_change;

int fc = num_vertices1;
// int num_aug_paths = 0;

int *h_frontier;
int *h_next_frontier;


__device__ 
int get_is_matched_edge(int i, int j){
	return d_is_matched_edge[i*(num_vertices1 + num_vertices2+1) + j ];
}

void print_matchings(){
	cout << "Matchings: " << endl;
    for(int i=1;i<=num_vertices1+num_vertices2; i++){
    	cout<< i << " " << h_partner_vertex[i] << endl;
    }
}

int get_is_matched_edge_h(int i, int j){
	return h_is_matched_edge[i*(num_vertices1 + num_vertices2+1) + j ];
}

__device__ 
void set_is_matched_edge(int i, int j, int value){
	d_is_matched_edge[i*(num_vertices1 + num_vertices2+1) + j ] = value;
}

__device__
void match_edges(int u, int v){
	set_is_matched_edge(u,v,1);
	set_is_matched_edge(v,u,1);
	d_is_matched_vertex[u] = 1;
	d_is_matched_vertex[v] = 1;
	d_partner_vertex[u] = v;
	d_partner_vertex[v] = u;

}

// Unmatching edges also unmatches the vertices since the graph is a matching
__device__
void unmatch_edges(int u, int v){
	set_is_matched_edge(u,v,0);
	set_is_matched_edge(v,u,0);
	if(d_partner_vertex[u]==v){
		d_is_matched_vertex[u] = 0;
		d_partner_vertex[u] = -1;
	}
	if(d_partner_vertex[v]==u){
		d_is_matched_vertex[v] = 0;
		d_partner_vertex[v] = -1;
	}
}

// Make this parallel
__device__
void update_matchings(){
	for(int i=1; i<=num_vertices1+num_vertices2; i++){
		int vertex = i;
		if(d_is_parent_change[vertex] == true){
			
			// cout << "Found aug. path till " << vertex << endl;
			// There should always be odd number of vertices in aug. path
			int path_length = 1;
			int parent = d_bfs_parent[vertex];
			while(parent!=vertex){
				// cout << vertex << " " <<parent << endl;
				if(path_length%2==1){
					match_edges(vertex, parent);
					// printf("Matching %d and %d \n", vertex, parent);
				}
				else{
					unmatch_edges(vertex, parent);
					// printf("Unmatching %d and %d \n", vertex, parent);
				}
				vertex =  d_bfs_parent[vertex];
				parent = d_bfs_parent[vertex];
				path_length++;
			}
		}
	}
}



__device__
void clear_visited(){
	int tid = blockIdx.x*1024 + threadIdx.x;
	int vertex1 = tid + 1;

	if(vertex1<=num_vertices1 + num_vertices2){
		d_visited[vertex1] = 0;
	}
}


__device__ 
void clear_bfs_parent(){
	int tid = blockIdx.x*1024 + threadIdx.x;
	int vertex1 = tid + 1;

	if(vertex1<=num_vertices1 + num_vertices2){
		d_bfs_parent[vertex1] = vertex1;
	}
}

__device__ 
void initialise_partner_vertex(){
	int tid = blockIdx.x*1024 + threadIdx.x;
	int vertex1 = tid + 1;

	if(vertex1<=num_vertices1 + num_vertices2){
		d_partner_vertex[vertex1] = -1;
	}
}

__device__ 
void clear_is_parent_change(){
	int tid = blockIdx.x*1024 + threadIdx.x;
	int vertex1 = tid + 1;

	if(vertex1<=num_vertices1 + num_vertices2){
		d_is_parent_change[vertex1] = -1;
	}
}

__device__ 
void copy_frontier(int *my_frontier, int *my_next_frontier){
	
	for (int i=1;i<=num_vertices1+num_vertices2;i++){
		my_frontier[i] = my_next_frontier[i];
	}
}

__device__ 
void clear_frontier(int *my_frontier, int *my_next_frontier ){
	for (int i=1;i<=num_vertices1+num_vertices2;i++){
			my_frontier[i] = 0;
			my_next_frontier[i] = 0;
	}
}
__device__
void vertex_disjoint_bfs(int binary_level, int vertex, int tid){
	int frontier_element = vertex;
	// printf("Frontier element: %d \n", frontier_element );
	if(!d_frontier[vertex]){
		return;
	}

	// my_frontier[vertex] = 1;
	d_frontier[vertex] = 1;


	// Iterate all frontier elements
	if(frontier_element!=-1){
	
		int vertex = frontier_element;
		
		// Make this atomic
		d_visited[vertex] = true;
		
		// cout << "Frontier: " << frontier_element << endl;
		// cout << "Continuining for vertex: " << vertex << endl;
		
		bool found_path = false;
		int start_edge = d_list_ptr[vertex];
		int end_edge = d_list_ptr[vertex + 1]; 
		
		// cout << "Start-End edge " <<  start_edge << " " << end_edge  << endl;
		// printf ("Start-End edge %d %d \n", start_edge, end_edge);
		for(int j=start_edge;j<end_edge;j++){
			if(found_path)
				break;


			int neighbor = d_flat_adj_list[j];

			// printf("[%d]Processing %d %d \n", tid, vertex, neighbor);

			int visited = atomicExch(&d_visited[neighbor], 1);

			if(!visited){
				// We want to alternate between unmatched and matched edges, otherwise we ignore
				d_visited[neighbor] = true;
				// cout << "Processing: " << vertex << " " << neighbor << endl;
				// printf("[%d]Processing %d %d \n", tid, vertex, neighbor);
				// exit(0);
				d_bfs_parent[neighbor] = vertex;

				if( binary_level==0 && get_is_matched_edge(vertex, neighbor)==0 && d_is_matched_vertex[neighbor]==1 ){
					// next_frontier.push_back(neighbor);
					d_next_frontier[neighbor] = 1;
				}

				// is_matched_vertex is implicitly true since the edge is matched
				// In level 1, we are only interested in matched edges
				else if( binary_level==1 && get_is_matched_edge(vertex, neighbor)==1 ){
					// next_frontier.push_back(neighbor);
					d_next_frontier[neighbor] = 1;
					// If I have found a path to the next level; I have to break
					// found_path = 1;
					return;
				}

				// Changing parent change only for this node
				else if(binary_level==0 && get_is_matched_edge(vertex, neighbor)==0 && d_is_matched_vertex[neighbor]==0){
					// printf("Found a aug. path with %d with parent: %d \n", neighbor, vertex);
					d_is_parent_change[neighbor] = 1;
					num_aug_paths++ ;
					// remove this return so that multiple paths can be found 
					return;
				}
			}
		}

		// frontier_element = get_frontier_element(vertex);
		// Getting next frontier element
		for(int x = vertex+1; x <=num_vertices1 + num_vertices2; x++){
			if(d_frontier[x]==1){
				frontier_element = x;
				break;
			}
		}
	}
	copy_frontier(d_frontier, d_next_frontier);
	// Only working for one level for now
	// bfs(binary_level = !binary_level);

}


__global__
void vertex_disjoint_bfs_util(){

	// parallelise these functions
	clear_visited();
	clear_bfs_parent();
	clear_is_parent_change();

	// clear_frontier(my_frontier, my_next_frontier );
	initialise_partner_vertex();


	// int num_aug_paths = 1000;

	int tid = blockIdx.x*1024 + threadIdx.x;


	int vertex = tid+1;
	if(vertex > num_vertices1)
		return;

	if(!d_visited[vertex] && !d_is_matched_vertex[vertex]){
		d_frontier[vertex] = 1;
		vertex_disjoint_bfs(0, vertex, tid);
		__syncthreads();
	}


	if(num_aug_paths > 0){
		update_matchings();
	}

	// return num_aug_paths;
}

int check_matching(){
	int total_matched = 0;
	for(int i=1;i<=num_vertices1+num_vertices2;i++){
		int vertex = i;
		int num_matched = 0;


		for(int j=h_list_ptr[i];j<h_list_ptr[i+1];j++){
			int neighbor = h_flat_adj_list[j];
			// cout << vertex << " " << neighbor << endl;
			if(get_is_matched_edge_h(vertex, neighbor)){
				num_matched++;
			}
		}


		if(num_matched==1){
			// cout << "Hi" << endl;
			total_matched++;
		}
		if(num_matched>1){
			cout << vertex << endl;
			cout << "Error! Not a matching!";
			exit(0);
		}
	}
	return total_matched/2;
}


int main(){
	struct timespec start, end;

	h_is_matched_edge = (bool *)calloc( (num_vertices1+ num_vertices2 + 1)*(num_vertices1 + num_vertices2+1), sizeof(bool));

	h_flat_adj_list = (int *)malloc(2*num_edges*sizeof(int));
	h_degree = (int *)malloc((num_vertices1+num_vertices2+1)*sizeof(int));
	h_list_ptr = (int *)malloc((num_vertices1+num_vertices2+2)*sizeof(int));
	h_list_ptr_copy = (int *)malloc((num_vertices1+num_vertices2+2)*sizeof(int));
	h_is_matched_vertex = (bool *)malloc((num_vertices1+num_vertices2+1)*sizeof(bool));
	h_partner_vertex = (int *)malloc((num_vertices1+num_vertices2+1)*sizeof(int));
	h_visited = (int *)malloc((num_vertices1+num_vertices2+1)*sizeof(int));
	h_bfs_parent = (int *)malloc((num_vertices1+num_vertices2+1)*sizeof(int));
	h_is_parent_change = (bool *)malloc((num_vertices1+num_vertices2+1)*sizeof(bool));
	h_frontier = (int *)malloc((num_vertices1+num_vertices2+1)*sizeof(int));
	h_next_frontier = (int *)malloc((num_vertices1+num_vertices2+1)*sizeof(int));

	// Add a check for null memory

	memset(h_degree, 0, num_vertices1 + num_vertices2 +1);
	// memset(h_is_matched_edge, 0, (num_vertices1 + num_vertices2 +1)*(num_vertices1+num_vertices2+1));
	memset(h_is_matched_vertex, 0, num_vertices1 + num_vertices2 +1);
	memset(h_visited, 0, num_vertices1 + num_vertices2 +1);
	memset(h_is_parent_change, 0, num_vertices1 + num_vertices2 +1);
	memset(h_frontier, 0, num_vertices1 + num_vertices2 +1);
	memset(h_next_frontier, 0, num_vertices1 + num_vertices2 +1);



	// to and from of edges
	// int h_edges_u[num_edges], h_edges_v[num_edges];			// Make this dynamic memory and free it once we have our 2 pass initialisation phase
	int *h_edges_u, *h_edges_v;
	h_edges_u = (int *)malloc((num_edges)*sizeof(int));
	h_edges_v = (int *)malloc((num_edges)*sizeof(int));


	ifstream fin;
    fin.open("FC_" + to_string(fc) + "_" + to_string(fc) + ".txt", ios::in);
    // fin.open("random_" + to_string(num_vertices1) + "_" + to_string(num_vertices2) + ".txt", ios::in);
    // cout << "random_" + to_string(num_vertices1) + "_" + to_string(num_vertices2) + ".txt" <<endl;
    int u, v;

    // cout << "Printing all the edges: \n";

    // Vertices with 0 edges are implicitly ignored while reading the file itself
    for(int i=0;i<num_edges;i++){
    		// cout << i << endl;
            fin >> u >> v;
            h_edges_u[i] = u;
            h_edges_v[i] = v;
            h_degree[u]++;
            h_degree[v]++;
    }

    cout << "Done reading edges" << endl;

    // Get pointer to adjacency list using prefix sum (no opti here since other parts are more complex anyway)
    // Index 0 will never be used.... the last elem
    h_list_ptr[1] = 0;
    h_list_ptr_copy[1] = h_list_ptr[1];
    for(int i=2;i<=num_vertices1+num_vertices2;i++){
    	h_list_ptr[i] = h_list_ptr[i-1] + h_degree[i-1];
    	h_list_ptr_copy[i] = h_list_ptr[i];
    }
    h_list_ptr[num_vertices1+num_vertices2+1] = 2*num_edges;       //For easy coding
    h_list_ptr_copy[num_vertices1+num_vertices2+1] = 2*num_edges;  // list_ptr has the start of the adj list ; list_ptr_copy has the current position

    
    for(int i=0;i<num_edges;i++){
    	h_flat_adj_list[h_list_ptr_copy[h_edges_u[i]]] = h_edges_v[i];
    	h_flat_adj_list[h_list_ptr_copy[h_edges_v[i]]] = h_edges_u[i];
    	h_list_ptr_copy[h_edges_u[i]]++;
    	h_list_ptr_copy[h_edges_v[i]]++;
    }
    

    clock_gettime( CLOCK_REALTIME,&start);

	hipMemcpyToSymbol(HIP_SYMBOL(d_is_matched_edge), h_is_matched_edge, (num_vertices1+ num_vertices2 + 1)*(num_vertices1 + num_vertices2+1)*sizeof(int),0,hipMemcpyHostToDevice);
	hipMemcpyToSymbol(HIP_SYMBOL(d_flat_adj_list), h_flat_adj_list, 2*num_edges*sizeof(int),0,hipMemcpyHostToDevice);
	hipMemcpyToSymbol(HIP_SYMBOL(d_degree), h_degree, (num_vertices1+num_vertices2+1)*sizeof(int),0,hipMemcpyHostToDevice);
	hipMemcpyToSymbol(HIP_SYMBOL(d_list_ptr), h_list_ptr, (num_vertices1+num_vertices2+2)*sizeof(int),0,hipMemcpyHostToDevice);
	hipMemcpyToSymbol(HIP_SYMBOL(d_is_matched_vertex), h_is_matched_vertex, (num_vertices1+num_vertices2+1)*sizeof(int),0,hipMemcpyHostToDevice);
	hipMemcpyToSymbol(HIP_SYMBOL(d_visited), h_visited, (num_vertices1+num_vertices2+1)*sizeof(int),0,hipMemcpyHostToDevice);
	hipMemcpyToSymbol(HIP_SYMBOL(d_frontier), h_frontier, (num_vertices1+num_vertices2+2)*sizeof(int),0,hipMemcpyHostToDevice);
	
	// cudaMemcpy(d_is_matched_edge, h_is_matched_edge, (num_vertices1+ num_vertices2 + 1)*(num_vertices1 + num_vertices2+1)*sizeof(int),cudaMemcpyHostToDevice);
	// cudaMemcpy(d_flat_adj_list, h_flat_adj_list, 2*num_edges*sizeof(int),cudaMemcpyHostToDevice);
	// cudaMemcpy(d_degree, h_degree, (num_vertices1+num_vertices2+1)*sizeof(int),cudaMemcpyHostToDevice);
	// cudaMemcpy(d_list_ptr, h_list_ptr, (num_vertices1+num_vertices2+2)*sizeof(int),cudaMemcpyHostToDevice);
	// cudaMemcpy(d_is_matched_vertex, h_is_matched_vertex, (num_vertices1+num_vertices2+1)*sizeof(int),cudaMemcpyHostToDevice);
	// cudaMemcpy(d_visited, h_visited, (num_vertices1+num_vertices2+1)*sizeof(int),cudaMemcpyHostToDevice);
	// cudaMemcpy(d_frontier, h_frontier, (num_vertices1+num_vertices2+2)*sizeof(int),cudaMemcpyHostToDevice);


  	hipDeviceSynchronize();
  	vertex_disjoint_bfs_util<<<1, num_threads>>>();
// 1219611
  	hipDeviceSynchronize();
  	hipError_t bla = hipMemcpyFromSymbol(h_is_matched_edge, HIP_SYMBOL(d_is_matched_edge), sizeof(d_is_matched_edge),0,hipMemcpyDeviceToHost);
  	hipMemcpyFromSymbol(h_partner_vertex, HIP_SYMBOL(d_partner_vertex), sizeof(d_partner_vertex),0,hipMemcpyDeviceToHost);
  	
  	clock_gettime( CLOCK_REALTIME,&end);
  	hipDeviceSynchronize();
  

  	int num_matches = check_matching();
  	
  	// print_matchings();

  	printf("Number of matchings: %d \n", num_matches);

  	double elapsed = (end.tv_sec-start.tv_sec)*1000000000 + end.tv_nsec-start.tv_nsec;
  	printf("Time elapsed %lf\n", elapsed/1e6);
	

  	hipDeviceSynchronize();
  


}
	.text
	.file	"list_parallel_matching.hip"
                                        # Start of file scope inline assembly
	.globl	_ZSt21ios_base_library_initv

                                        # End of file scope inline assembly
	.globl	_Z15print_matchingsv            # -- Begin function _Z15print_matchingsv
	.type	_Z15print_matchingsv,@function
_Z15print_matchingsv:                   # @_Z15print_matchingsv
	.cfi_startproc
# %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
	movl	$_ZSt4cout, %ebx
	movl	$_ZSt4cout, %edi
	movl	$.L.str, %esi
	movl	$11, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	$1, %ebx
.LBB0_1:                                # =>This Inner Loop Header: Depth=1
	movl	$_ZSt4cout, %edi
	movl	%ebx, %esi
	callq	_ZNSolsEi
	movq	%rax, %r14
	movl	$.L.str.1, %esi
	movl	$1, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	h_partner_vertex(%rip), %rax
	movl	(%rax,%rbx,4), %esi
	movq	%r14, %rdi
	callq	_ZNSolsEi
	movq	%rax, %r14
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%r14, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%r14, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	incq	%rbx
	cmpq	$2001, %rbx                     # imm = 0x7D1
	jne	.LBB0_1
# %bb.2:
	addq	$8, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end0:
	.size	_Z15print_matchingsv, .Lfunc_end0-_Z15print_matchingsv
	.cfi_endproc
                                        # -- End function
	.globl	_Z21get_is_matched_edge_hii     # -- Begin function _Z21get_is_matched_edge_hii
	.type	_Z21get_is_matched_edge_hii,@function
_Z21get_is_matched_edge_hii:            # @_Z21get_is_matched_edge_hii
	.cfi_startproc
# %bb.0:
	movslq	%edi, %rax
	imulq	$2001, %rax, %rax               # imm = 0x7D1
	movslq	%esi, %rcx
	addq	h_is_matched_edge(%rip), %rax
	movzbl	(%rcx,%rax), %eax
	retq
.Lfunc_end1:
	.size	_Z21get_is_matched_edge_hii, .Lfunc_end1-_Z21get_is_matched_edge_hii
	.cfi_endproc
                                        # -- End function
	.globl	_Z39__device_stub__vertex_disjoint_bfs_utilv # -- Begin function _Z39__device_stub__vertex_disjoint_bfs_utilv
	.type	_Z39__device_stub__vertex_disjoint_bfs_utilv,@function
_Z39__device_stub__vertex_disjoint_bfs_utilv: # @_Z39__device_stub__vertex_disjoint_bfs_utilv
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r12
	.cfi_def_cfa_offset 32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	subq	$56, %rsp
	.cfi_def_cfa_offset 96
	.cfi_offset %rbx, -40
	.cfi_offset %r12, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	40(%rsp), %rbx
	leaq	24(%rsp), %r14
	leaq	16(%rsp), %r15
	leaq	8(%rsp), %r12
	movq	%rbx, %rdi
	movq	%r14, %rsi
	movq	%r15, %rdx
	movq	%r12, %rcx
	callq	__hipPopCallConfiguration
	movq	(%rbx), %rsi
	movl	8(%rbx), %edx
	movq	(%r14), %rcx
	movl	8(%r14), %r8d
	movq	%rsp, %r9
	movl	$_Z24vertex_disjoint_bfs_utilv, %edi
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	pushq	(%r15)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$72, %rsp
	.cfi_adjust_cfa_offset -72
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end2:
	.size	_Z39__device_stub__vertex_disjoint_bfs_utilv, .Lfunc_end2-_Z39__device_stub__vertex_disjoint_bfs_utilv
	.cfi_endproc
                                        # -- End function
	.globl	_Z14check_matchingv             # -- Begin function _Z14check_matchingv
	.type	_Z14check_matchingv,@function
_Z14check_matchingv:                    # @_Z14check_matchingv
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset %rbx, -16
	movq	h_list_ptr(%rip), %rcx
	movq	h_flat_adj_list(%rip), %rdx
	xorl	%eax, %eax
	movl	$1, %r9d
	movq	h_is_matched_edge(%rip), %rdi
.LBB3_1:                                # =>This Loop Header: Depth=1
                                        #     Child Loop BB3_4 Depth 2
	cmpq	$2001, %r9                      # imm = 0x7D1
	je	.LBB3_7
# %bb.2:                                #   in Loop: Header=BB3_1 Depth=1
	movq	%r9, %rsi
	movslq	(%rcx,%r9,4), %r9
	movslq	4(%rcx,%rsi,4), %r10
	xorl	%r8d, %r8d
	cmpl	%r10d, %r9d
	jge	.LBB3_5
# %bb.3:                                # %.lr.ph
                                        #   in Loop: Header=BB3_1 Depth=1
	imulq	$2001, %rsi, %r11               # imm = 0x7D1
	addq	%rdi, %r11
	xorl	%r8d, %r8d
.LBB3_4:                                #   Parent Loop BB3_1 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movslq	(%rdx,%r9,4), %rbx
	movzbl	(%r11,%rbx), %ebx
	addl	%ebx, %r8d
	incq	%r9
	cmpq	%r9, %r10
	jne	.LBB3_4
.LBB3_5:                                # %._crit_edge
                                        #   in Loop: Header=BB3_1 Depth=1
	leaq	1(%rsi), %r9
	xorl	%r10d, %r10d
	cmpl	$1, %r8d
	sete	%r10b
	addl	%r10d, %eax
	cmpl	$1, %r8d
	jle	.LBB3_1
# %bb.6:
	movl	$_ZSt4cout, %edi
                                        # kill: def $esi killed $esi killed $rsi
	callq	_ZNSolsEi
	movq	%rax, %rdi
	callq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_
	movl	$_ZSt4cout, %edi
	movl	$.L.str.2, %esi
	callq	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc
	xorl	%edi, %edi
	callq	exit
.LBB3_7:
	shrl	%eax
	popq	%rbx
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end3:
	.size	_Z14check_matchingv, .Lfunc_end3-_Z14check_matchingv
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function main
.LCPI4_0:
	.quad	0x412e848000000000              # double 1.0E+6
	.text
	.globl	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	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	subq	$744, %rsp                      # imm = 0x2E8
	.cfi_def_cfa_offset 800
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movl	$4004001, %edi                  # imm = 0x3D18A1
	movl	$1, %esi
	callq	calloc
	movq	%rax, h_is_matched_edge(%rip)
	movl	$8000000, %edi                  # imm = 0x7A1200
	callq	malloc
	movq	%rax, h_flat_adj_list(%rip)
	movl	$8004, %edi                     # imm = 0x1F44
	callq	malloc
	movq	%rax, %rbx
	movq	%rax, h_degree(%rip)
	movl	$8008, %edi                     # imm = 0x1F48
	callq	malloc
	movq	%rax, h_list_ptr(%rip)
	movl	$8008, %edi                     # imm = 0x1F48
	callq	malloc
	movq	%rax, h_list_ptr_copy(%rip)
	movl	$1, %edi
	movl	$2001, %esi                     # imm = 0x7D1
	callq	calloc
	movq	%rax, h_is_matched_vertex(%rip)
	movl	$8004, %edi                     # imm = 0x1F44
	callq	malloc
	movq	%rax, h_partner_vertex(%rip)
	movl	$8004, %edi                     # imm = 0x1F44
	callq	malloc
	movq	%rax, %r14
	movq	%rax, h_visited(%rip)
	movl	$8004, %edi                     # imm = 0x1F44
	callq	malloc
	movq	%rax, h_bfs_parent(%rip)
	movl	$1, %edi
	movl	$2001, %esi                     # imm = 0x7D1
	callq	calloc
	movq	%rax, h_is_parent_change(%rip)
	movl	$8004, %edi                     # imm = 0x1F44
	callq	malloc
	movq	%rax, %r12
	movq	%rax, h_frontier(%rip)
	movl	$8004, %edi                     # imm = 0x1F44
	callq	malloc
	movq	%rax, %r15
	movq	%rax, h_next_frontier(%rip)
	movl	$2001, %edx                     # imm = 0x7D1
	movq	%rbx, %rdi
	xorl	%esi, %esi
	callq	memset@PLT
	movl	$2001, %edx                     # imm = 0x7D1
	movq	%r14, %rdi
	xorl	%esi, %esi
	callq	memset@PLT
	movl	$2001, %edx                     # imm = 0x7D1
	movq	%r12, %rdi
	xorl	%esi, %esi
	callq	memset@PLT
	movl	$2001, %edx                     # imm = 0x7D1
	movq	%r15, %rdi
	xorl	%esi, %esi
	callq	memset@PLT
	movl	$4000000, %edi                  # imm = 0x3D0900
	callq	malloc
	movq	%rax, %rbx
	movl	$4000000, %edi                  # imm = 0x3D0900
	callq	malloc
	movq	%rax, %r14
	leaq	224(%rsp), %rdi
	callq	_ZNSt14basic_ifstreamIcSt11char_traitsIcEEC1Ev
	movl	fc(%rip), %r15d
	movl	%r15d, %ebp
	negl	%ebp
	cmovsl	%r15d, %ebp
	movl	$1, %r12d
	cmpl	$10, %ebp
	jb	.LBB4_9
# %bb.1:                                # %.lr.ph.i.i.preheader
	movl	$4, %r12d
	movl	$3518437209, %eax               # imm = 0xD1B71759
	movl	%ebp, %ecx
.LBB4_2:                                # %.lr.ph.i.i
                                        # =>This Inner Loop Header: Depth=1
	cmpl	$99, %ecx
	jbe	.LBB4_7
# %bb.3:                                #   in Loop: Header=BB4_2 Depth=1
	cmpl	$999, %ecx                      # imm = 0x3E7
	jbe	.LBB4_8
# %bb.4:                                #   in Loop: Header=BB4_2 Depth=1
	cmpl	$10000, %ecx                    # imm = 0x2710
	jb	.LBB4_9
# %bb.5:                                #   in Loop: Header=BB4_2 Depth=1
	movl	%ecx, %edx
	imulq	%rax, %rdx
	shrq	$45, %rdx
	addl	$4, %r12d
	cmpl	$99999, %ecx                    # imm = 0x1869F
	movl	%edx, %ecx
	ja	.LBB4_2
# %bb.6:                                # %_ZNSt8__detail14__to_chars_lenIjEEjT_i.exit.i.loopexit
	addl	$-3, %r12d
	jmp	.LBB4_9
.LBB4_7:
	addl	$-2, %r12d
	jmp	.LBB4_9
.LBB4_8:
	decl	%r12d
.LBB4_9:                                # %_ZNSt8__detail14__to_chars_lenIjEEjT_i.exit.i
	shrl	$31, %r15d
	leal	(%r12,%r15), %esi
	leaq	176(%rsp), %rax
	movq	%rax, -16(%rax)
.Ltmp0:
	leaq	160(%rsp), %rdi
	movl	$45, %edx
	callq	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE12_M_constructEmc
.Ltmp1:
# %bb.10:                               # %_ZNSt7__cxx119to_stringEi.exit
	leaq	160(%rsp), %r13
	addq	(%r13), %r15
	movq	%r15, %rdi
	movl	%r12d, %esi
	movl	%ebp, %edx
	callq	_ZNSt8__detail18__to_chars_10_implIjEEvPcjT_
.Ltmp3:
	movl	$.L.str.3, %ecx
	movl	$3, %r8d
	movq	%r13, %rdi
	xorl	%esi, %esi
	xorl	%edx, %edx
	callq	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE7replaceEmmPKcm
.Ltmp4:
# %bb.11:                               # %.noexc
	movq	%rax, %r15
	leaq	112(%rsp), %rdi
	movq	%rdi, -16(%rdi)
	movq	(%rax), %rsi
	movq	%rax, %r12
	addq	$16, %r12
	cmpq	%r12, %rsi
	je	.LBB4_13
# %bb.12:                               # %_ZNKSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE11_M_is_localEv.exit.i.i
	movq	%rsi, 96(%rsp)
	movq	16(%r15), %rax
	movq	%rax, 112(%rsp)
	movq	8(%r15), %r13
	jmp	.LBB4_14
.LBB4_13:
	movq	8(%r15), %r13
	leaq	1(%r13), %rdx
	callq	memcpy@PLT
.LBB4_14:
	leaq	96(%rsp), %rdi
	movq	%r13, 8(%rdi)
	movq	%r12, (%r15)
	movq	$0, 8(%r15)
	movb	$0, 16(%r15)
.Ltmp6:
	movl	$.L.str.4, %esi
	callq	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE6appendEPKc
.Ltmp7:
# %bb.15:                               # %.noexc60
	movq	%rax, %r15
	leaq	48(%rsp), %rdi
	movq	%rdi, -16(%rdi)
	movq	(%rax), %rsi
	movq	%rax, %r12
	addq	$16, %r12
	cmpq	%r12, %rsi
	je	.LBB4_17
# %bb.16:                               # %_ZNKSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE11_M_is_localEv.exit.i.i57
	movq	%rsi, 32(%rsp)
	movq	16(%r15), %rax
	movq	%rax, 48(%rsp)
	movq	8(%r15), %r13
	jmp	.LBB4_18
.LBB4_17:
	movq	8(%r15), %r13
	leaq	1(%r13), %rdx
	callq	memcpy@PLT
.LBB4_18:
	movq	%r13, 40(%rsp)
	movq	%r12, (%r15)
	movq	$0, 8(%r15)
	movb	$0, 16(%r15)
	movl	fc(%rip), %r13d
	movl	%r13d, %ebp
	negl	%ebp
	cmovsl	%r13d, %ebp
	movl	$1, %r15d
	cmpl	$10, %ebp
	jb	.LBB4_27
# %bb.19:                               # %.lr.ph.i.i61.preheader
	movl	$4, %r15d
	movl	$3518437209, %eax               # imm = 0xD1B71759
	movl	%ebp, %ecx
.LBB4_20:                               # %.lr.ph.i.i61
                                        # =>This Inner Loop Header: Depth=1
	cmpl	$99, %ecx
	jbe	.LBB4_25
# %bb.21:                               #   in Loop: Header=BB4_20 Depth=1
	cmpl	$999, %ecx                      # imm = 0x3E7
	jbe	.LBB4_26
# %bb.22:                               #   in Loop: Header=BB4_20 Depth=1
	cmpl	$10000, %ecx                    # imm = 0x2710
	jb	.LBB4_27
# %bb.23:                               #   in Loop: Header=BB4_20 Depth=1
	movl	%ecx, %edx
	imulq	%rax, %rdx
	shrq	$45, %rdx
	addl	$4, %r15d
	cmpl	$99999, %ecx                    # imm = 0x1869F
	movl	%edx, %ecx
	ja	.LBB4_20
# %bb.24:                               # %_ZNSt8__detail14__to_chars_lenIjEEjT_i.exit.i64.loopexit
	addl	$-3, %r15d
	jmp	.LBB4_27
.LBB4_25:
	addl	$-2, %r15d
	jmp	.LBB4_27
.LBB4_26:
	decl	%r15d
.LBB4_27:                               # %_ZNSt8__detail14__to_chars_lenIjEEjT_i.exit.i64
	shrl	$31, %r13d
	leal	(%r15,%r13), %esi
	leaq	80(%rsp), %r12
	movq	%r12, -16(%r12)
.Ltmp9:
	leaq	64(%rsp), %rdi
	movl	$45, %edx
	callq	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE12_M_constructEmc
.Ltmp10:
# %bb.28:                               # %_ZNSt7__cxx119to_stringEi.exit67
	addq	64(%rsp), %r13
	movq	%r13, %rdi
	movl	%r15d, %esi
	movl	%ebp, %edx
	callq	_ZNSt8__detail18__to_chars_10_implIjEEvPcjT_
	movq	32(%rsp), %rcx
	movq	40(%rsp), %r8
	movq	72(%rsp), %rdx
	leaq	(%rdx,%r8), %rax
	movl	$15, %edi
	leaq	48(%rsp), %rsi
	cmpq	%rsi, %rcx
	je	.LBB4_30
# %bb.29:                               # %_ZNSt7__cxx119to_stringEi.exit67
	movq	48(%rsp), %rdi
.LBB4_30:                               # %_ZNSt7__cxx119to_stringEi.exit67
	movq	64(%rsp), %rsi
	cmpq	%rdi, %rax
	jbe	.LBB4_34
# %bb.31:                               # %_ZNKSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE8capacityEv.exit12.i
	movl	$15, %edi
	cmpq	%r12, %rsi
	je	.LBB4_33
# %bb.32:                               # %_ZNKSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE8capacityEv.exit12.i
	movq	80(%rsp), %rdi
.LBB4_33:                               # %_ZNKSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE8capacityEv.exit12.i
	cmpq	%rdi, %rax
	jbe	.LBB4_35
.LBB4_34:
.Ltmp14:
	leaq	32(%rsp), %rdi
	callq	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE6appendEPKcm
.Ltmp15:
	jmp	.LBB4_36
.LBB4_35:                               # %.critedge.i
.Ltmp12:
	leaq	64(%rsp), %rdi
	xorl	%esi, %esi
	xorl	%edx, %edx
	callq	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE7replaceEmmPKcm
.Ltmp13:
.LBB4_36:                               # %.noexc70
	movq	%rax, %r15
	leaq	16(%rsp), %rdi
	movq	%rdi, -16(%rdi)
	movq	(%rax), %rsi
	leaq	16(%rax), %r13
	cmpq	%r13, %rsi
	je	.LBB4_38
# %bb.37:                               # %_ZNKSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE11_M_is_localEv.exit.i13.i
	movq	%rsi, (%rsp)
	movq	(%r13), %rax
	movq	%rax, 16(%rsp)
	jmp	.LBB4_39
.LBB4_38:
	movq	8(%r15), %rdx
	incq	%rdx
	callq	memcpy@PLT
.LBB4_39:                               # %_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2EOS4_.exit14.i
	movq	%r15, %rax
	addq	$8, %rax
	movq	8(%r15), %rcx
	movq	%rcx, 8(%rsp)
	movq	%r13, (%r15)
	movq	$0, (%rax)
	movb	$0, (%r13)
.Ltmp17:
	movq	%rsp, %rdi
	movl	$.L.str.5, %esi
	callq	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE6appendEPKc
.Ltmp18:
# %bb.40:                               # %.noexc76
	movq	%rax, %r13
	leaq	144(%rsp), %r15
	movq	%r15, -16(%r15)
	movq	(%rax), %rsi
	movq	%rax, %rbp
	addq	$16, %rbp
	cmpq	%rbp, %rsi
	je	.LBB4_42
# %bb.41:                               # %_ZNKSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE11_M_is_localEv.exit.i.i72
	movq	%rsi, 128(%rsp)
	movq	16(%r13), %rax
	movq	%rax, 144(%rsp)
	movq	8(%r13), %r12
	jmp	.LBB4_43
.LBB4_42:
	movq	8(%r13), %r12
	leaq	1(%r12), %rdx
	movq	%r15, %rdi
	callq	memcpy@PLT
.LBB4_43:
	leaq	128(%rsp), %rsi
	movq	%r12, 8(%rsi)
	movq	%rbp, (%r13)
	movq	$0, 8(%r13)
	movb	$0, 16(%r13)
.Ltmp20:
	leaq	224(%rsp), %rdi
	movl	$8, %edx
	callq	_ZNSt14basic_ifstreamIcSt11char_traitsIcEE4openERKNSt7__cxx1112basic_stringIcS1_SaIcEEESt13_Ios_Openmode
.Ltmp21:
# %bb.44:
	movq	128(%rsp), %rdi
	cmpq	%r15, %rdi
	leaq	80(%rsp), %r15
	je	.LBB4_46
# %bb.45:                               # %_ZNKSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE11_M_is_localEv.exit.i.i78
	callq	_ZdlPv
.LBB4_46:                               # %_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev.exit
	leaq	16(%rsp), %rax
	movq	-16(%rax), %rdi
	cmpq	%rax, %rdi
	je	.LBB4_48
# %bb.47:                               # %_ZNKSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE11_M_is_localEv.exit.i.i79
	callq	_ZdlPv
.LBB4_48:                               # %_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev.exit81
	movq	64(%rsp), %rdi
	cmpq	%r15, %rdi
	je	.LBB4_50
# %bb.49:                               # %_ZNKSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE11_M_is_localEv.exit.i.i82
	callq	_ZdlPv
.LBB4_50:                               # %_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev.exit84
	movq	32(%rsp), %rdi
	leaq	48(%rsp), %rax
	cmpq	%rax, %rdi
	je	.LBB4_52
# %bb.51:                               # %_ZNKSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE11_M_is_localEv.exit.i.i85
	callq	_ZdlPv
.LBB4_52:                               # %_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev.exit87
	movq	96(%rsp), %rdi
	leaq	112(%rsp), %rax
	cmpq	%rax, %rdi
	je	.LBB4_54
# %bb.53:                               # %_ZNKSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE11_M_is_localEv.exit.i.i88
	callq	_ZdlPv
.LBB4_54:                               # %_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev.exit90
	movq	160(%rsp), %rdi
	leaq	176(%rsp), %rax
	cmpq	%rax, %rdi
	je	.LBB4_56
# %bb.55:                               # %_ZNKSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE11_M_is_localEv.exit.i.i91
	callq	_ZdlPv
.LBB4_56:                               # %_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev.exit93
	xorl	%ebp, %ebp
	leaq	224(%rsp), %r15
	leaq	128(%rsp), %r12
	movq	%rsp, %r13
.LBB4_57:                               # =>This Inner Loop Header: Depth=1
.Ltmp23:
	movq	%r15, %rdi
	movq	%r12, %rsi
	callq	_ZNSirsERi
.Ltmp24:
# %bb.58:                               #   in Loop: Header=BB4_57 Depth=1
.Ltmp25:
	movq	%rax, %rdi
	movq	%r13, %rsi
	callq	_ZNSirsERi
.Ltmp26:
# %bb.59:                               #   in Loop: Header=BB4_57 Depth=1
	movslq	128(%rsp), %rax
	movl	%eax, (%rbx,%rbp,4)
	movl	(%rsp), %ecx
	movl	%ecx, (%r14,%rbp,4)
	movq	h_degree(%rip), %rcx
	incl	(%rcx,%rax,4)
	movslq	(%rsp), %rax
	incl	(%rcx,%rax,4)
	incq	%rbp
	cmpq	$1000000, %rbp                  # imm = 0xF4240
	jne	.LBB4_57
# %bb.60:
.Ltmp28:
	movl	$_ZSt4cout, %r15d
	movl	$_ZSt4cout, %edi
	movl	$.L.str.6, %esi
	movl	$18, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
.Ltmp29:
# %bb.61:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc.exit
	movq	_ZSt4cout(%rip), %rax
	addq	-24(%rax), %r15
.Ltmp30:
	movq	%r15, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
.Ltmp31:
# %bb.62:                               # %.noexc125
.Ltmp32:
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
.Ltmp33:
# %bb.63:                               # %.noexc126
.Ltmp34:
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
.Ltmp35:
# %bb.64:                               # %_ZNSolsEPFRSoS_E.exit
	movq	h_list_ptr(%rip), %rcx
	xorl	%edx, %edx
	movl	%edx, 4(%rcx)
	movq	h_list_ptr_copy(%rip), %rax
	movl	%edx, 4(%rax)
	movq	h_degree(%rip), %rsi
.LBB4_65:                               # =>This Inner Loop Header: Depth=1
	movl	4(%rsi,%rdx,4), %edi
	addl	4(%rcx,%rdx,4), %edi
	movl	%edi, 8(%rcx,%rdx,4)
	movl	%edi, 8(%rax,%rdx,4)
	incq	%rdx
	cmpq	$1999, %rdx                     # imm = 0x7CF
	jne	.LBB4_65
# %bb.66:
	movl	$2000000, %edx                  # imm = 0x1E8480
	movl	%edx, 8004(%rcx)
	movl	%edx, 8004(%rax)
	xorl	%ecx, %ecx
	movq	h_flat_adj_list(%rip), %rdx
.LBB4_67:                               # =>This Inner Loop Header: Depth=1
	movslq	(%r14,%rcx,4), %rsi
	movslq	(%rbx,%rcx,4), %rdi
	movslq	(%rax,%rdi,4), %r8
	movl	%esi, (%rdx,%r8,4)
	movslq	(%rax,%rsi,4), %r8
	movl	%edi, (%rdx,%r8,4)
	incl	(%rax,%rdi,4)
	incl	(%rax,%rsi,4)
	incq	%rcx
	cmpq	$1000000, %rcx                  # imm = 0xF4240
	jne	.LBB4_67
# %bb.68:
	leaq	208(%rsp), %rsi
	xorl	%edi, %edi
	callq	clock_gettime
	movq	h_is_matched_edge(%rip), %rsi
.Ltmp36:
	movl	$d_is_matched_edge, %edi
	movl	$16016004, %edx                 # imm = 0xF46284
	xorl	%ecx, %ecx
	movl	$1, %r8d
	callq	hipMemcpyToSymbol
.Ltmp37:
# %bb.69:                               # %_Z17hipMemcpyToSymbolIA4004001_bE10hipError_tRKT_PKvmm13hipMemcpyKind.exit
	movq	h_flat_adj_list(%rip), %rsi
.Ltmp38:
	movl	$d_flat_adj_list, %edi
	movl	$8000000, %edx                  # imm = 0x7A1200
	xorl	%ecx, %ecx
	movl	$1, %r8d
	callq	hipMemcpyToSymbol
.Ltmp39:
# %bb.70:                               # %_Z17hipMemcpyToSymbolIA2000000_iE10hipError_tRKT_PKvmm13hipMemcpyKind.exit
	movq	h_degree(%rip), %rsi
.Ltmp40:
	movl	$d_degree, %edi
	movl	$8004, %edx                     # imm = 0x1F44
	xorl	%ecx, %ecx
	movl	$1, %r8d
	callq	hipMemcpyToSymbol
.Ltmp41:
# %bb.71:                               # %_Z17hipMemcpyToSymbolIA2001_iE10hipError_tRKT_PKvmm13hipMemcpyKind.exit
	movq	h_list_ptr(%rip), %rsi
.Ltmp42:
	movl	$d_list_ptr, %edi
	movl	$8008, %edx                     # imm = 0x1F48
	xorl	%ecx, %ecx
	movl	$1, %r8d
	callq	hipMemcpyToSymbol
.Ltmp43:
# %bb.72:                               # %_Z17hipMemcpyToSymbolIA2002_iE10hipError_tRKT_PKvmm13hipMemcpyKind.exit
	movq	h_is_matched_vertex(%rip), %rsi
.Ltmp44:
	movl	$d_is_matched_vertex, %edi
	movl	$8004, %edx                     # imm = 0x1F44
	xorl	%ecx, %ecx
	movl	$1, %r8d
	callq	hipMemcpyToSymbol
.Ltmp45:
# %bb.73:                               # %_Z17hipMemcpyToSymbolIA2001_bE10hipError_tRKT_PKvmm13hipMemcpyKind.exit
	movq	h_visited(%rip), %rsi
.Ltmp46:
	movl	$d_visited, %edi
	movl	$8004, %edx                     # imm = 0x1F44
	xorl	%ecx, %ecx
	movl	$1, %r8d
	callq	hipMemcpyToSymbol
.Ltmp47:
# %bb.74:                               # %_Z17hipMemcpyToSymbolIA2001_iE10hipError_tRKT_PKvmm13hipMemcpyKind.exit120
	movq	h_frontier(%rip), %rsi
.Ltmp48:
	movl	$d_frontier, %edi
	movl	$8008, %edx                     # imm = 0x1F48
	xorl	%ecx, %ecx
	movl	$1, %r8d
	callq	hipMemcpyToSymbol
.Ltmp49:
# %bb.75:                               # %_Z17hipMemcpyToSymbolIA2001_iE10hipError_tRKT_PKvmm13hipMemcpyKind.exit122
.Ltmp50:
	callq	hipDeviceSynchronize
.Ltmp51:
# %bb.76:
.Ltmp52:
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movabsq	$4294968296, %rdx               # imm = 0x1000003E8
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
.Ltmp53:
# %bb.77:
	testl	%eax, %eax
	jne	.LBB4_79
# %bb.78:
.Ltmp54:
	callq	_Z39__device_stub__vertex_disjoint_bfs_utilv
.Ltmp55:
.LBB4_79:
.Ltmp56:
	callq	hipDeviceSynchronize
.Ltmp57:
# %bb.80:
	movq	h_is_matched_edge(%rip), %rdi
.Ltmp59:
	movl	$d_is_matched_edge, %esi
	movl	$4004001, %edx                  # imm = 0x3D18A1
	xorl	%ecx, %ecx
	movl	$2, %r8d
	callq	hipMemcpyFromSymbol
.Ltmp60:
# %bb.81:                               # %_Z19hipMemcpyFromSymbolIA4004001_bE10hipError_tPvRKT_mm13hipMemcpyKind.exit
	movq	h_partner_vertex(%rip), %rdi
.Ltmp61:
	movl	$d_partner_vertex, %esi
	movl	$8004, %edx                     # imm = 0x1F44
	xorl	%ecx, %ecx
	movl	$2, %r8d
	callq	hipMemcpyFromSymbol
.Ltmp62:
# %bb.82:                               # %_Z19hipMemcpyFromSymbolIA2001_iE10hipError_tPvRKT_mm13hipMemcpyKind.exit
	leaq	192(%rsp), %rsi
	xorl	%edi, %edi
	callq	clock_gettime
.Ltmp63:
	callq	hipDeviceSynchronize
.Ltmp64:
# %bb.83:
.Ltmp66:
	callq	_Z14check_matchingv
.Ltmp67:
# %bb.84:
	movl	$.L.str.7, %edi
	movl	%eax, %esi
	xorl	%eax, %eax
	callq	printf
	movq	192(%rsp), %rax
	subq	208(%rsp), %rax
	imulq	$1000000000, %rax, %rax         # imm = 0x3B9ACA00
	addq	200(%rsp), %rax
	subq	216(%rsp), %rax
	cvtsi2sd	%rax, %xmm0
	divsd	.LCPI4_0(%rip), %xmm0
	movl	$.L.str.8, %edi
	movb	$1, %al
	callq	printf
.Ltmp69:
	callq	hipDeviceSynchronize
.Ltmp70:
# %bb.85:
	leaq	224(%rsp), %rdi
	callq	_ZNSt14basic_ifstreamIcSt11char_traitsIcEED1Ev
	xorl	%eax, %eax
	addq	$744, %rsp                      # imm = 0x2E8
	.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
.LBB4_86:
	.cfi_def_cfa_offset 800
.Ltmp71:
	jmp	.LBB4_110
.LBB4_87:
.Ltmp68:
	jmp	.LBB4_110
.LBB4_88:
.Ltmp22:
	movq	%rax, %rbx
	movq	128(%rsp), %rdi
	cmpq	%r15, %rdi
	je	.LBB4_91
# %bb.89:                               # %_ZNKSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE11_M_is_localEv.exit.i.i95
	callq	_ZdlPv
	jmp	.LBB4_91
.LBB4_90:
.Ltmp19:
	movq	%rax, %rbx
.LBB4_91:                               # %_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev.exit97
	leaq	16(%rsp), %rax
	movq	-16(%rax), %rdi
	cmpq	%rax, %rdi
	je	.LBB4_94
# %bb.92:                               # %_ZNKSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE11_M_is_localEv.exit.i.i98
	callq	_ZdlPv
	jmp	.LBB4_94
.LBB4_93:
.Ltmp16:
	movq	%rax, %rbx
.LBB4_94:                               # %_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev.exit100
	movq	64(%rsp), %rdi
	leaq	80(%rsp), %rax
	cmpq	%rax, %rdi
	je	.LBB4_96
# %bb.95:                               # %_ZNKSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE11_M_is_localEv.exit.i.i101
	callq	_ZdlPv
.LBB4_96:                               # %_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev.exit103
	movq	32(%rsp), %rdi
	leaq	48(%rsp), %rax
	cmpq	%rax, %rdi
	je	.LBB4_100
# %bb.97:                               # %_ZNKSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE11_M_is_localEv.exit.i.i104
	callq	_ZdlPv
	jmp	.LBB4_100
.LBB4_98:
.Ltmp11:
	jmp	.LBB4_106
.LBB4_99:
.Ltmp8:
	movq	%rax, %rbx
.LBB4_100:                              # %_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev.exit106
	movq	96(%rsp), %rdi
	leaq	112(%rsp), %rax
	cmpq	%rax, %rdi
	je	.LBB4_103
# %bb.101:                              # %_ZNKSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE11_M_is_localEv.exit.i.i107
	callq	_ZdlPv
	jmp	.LBB4_103
.LBB4_102:
.Ltmp5:
	movq	%rax, %rbx
.LBB4_103:                              # %_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev.exit109
	movq	160(%rsp), %rdi
	leaq	176(%rsp), %rax
	cmpq	%rax, %rdi
	je	.LBB4_111
# %bb.104:                              # %_ZNKSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE11_M_is_localEv.exit.i.i110
	callq	_ZdlPv
	jmp	.LBB4_111
.LBB4_105:
.Ltmp2:
.LBB4_106:
	movq	%rax, %rdi
	callq	__clang_call_terminate
.LBB4_107:
.Ltmp65:
	jmp	.LBB4_110
.LBB4_108:
.Ltmp58:
	jmp	.LBB4_110
.LBB4_109:
.Ltmp27:
.LBB4_110:
	movq	%rax, %rbx
.LBB4_111:
	leaq	224(%rsp), %rdi
	callq	_ZNSt14basic_ifstreamIcSt11char_traitsIcEED1Ev
	movq	%rbx, %rdi
	callq	_Unwind_Resume@PLT
.Lfunc_end4:
	.size	main, .Lfunc_end4-main
	.cfi_endproc
	.section	.gcc_except_table,"a",@progbits
	.p2align	2, 0x0
GCC_except_table4:
.Lexception0:
	.byte	255                             # @LPStart Encoding = omit
	.byte	3                               # @TType Encoding = udata4
	.uleb128 .Lttbase0-.Lttbaseref0
.Lttbaseref0:
	.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	1                               #   On action: 1
	.uleb128 .Ltmp3-.Lfunc_begin0           # >> Call Site 3 <<
	.uleb128 .Ltmp4-.Ltmp3                  #   Call between .Ltmp3 and .Ltmp4
	.uleb128 .Ltmp5-.Lfunc_begin0           #     jumps to .Ltmp5
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp4-.Lfunc_begin0           # >> Call Site 4 <<
	.uleb128 .Ltmp6-.Ltmp4                  #   Call between .Ltmp4 and .Ltmp6
	.byte	0                               #     has no landing pad
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp6-.Lfunc_begin0           # >> Call Site 5 <<
	.uleb128 .Ltmp7-.Ltmp6                  #   Call between .Ltmp6 and .Ltmp7
	.uleb128 .Ltmp8-.Lfunc_begin0           #     jumps to .Ltmp8
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp7-.Lfunc_begin0           # >> Call Site 6 <<
	.uleb128 .Ltmp9-.Ltmp7                  #   Call between .Ltmp7 and .Ltmp9
	.byte	0                               #     has no landing pad
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp9-.Lfunc_begin0           # >> Call Site 7 <<
	.uleb128 .Ltmp10-.Ltmp9                 #   Call between .Ltmp9 and .Ltmp10
	.uleb128 .Ltmp11-.Lfunc_begin0          #     jumps to .Ltmp11
	.byte	1                               #   On action: 1
	.uleb128 .Ltmp14-.Lfunc_begin0          # >> Call Site 8 <<
	.uleb128 .Ltmp13-.Ltmp14                #   Call between .Ltmp14 and .Ltmp13
	.uleb128 .Ltmp16-.Lfunc_begin0          #     jumps to .Ltmp16
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp13-.Lfunc_begin0          # >> Call Site 9 <<
	.uleb128 .Ltmp17-.Ltmp13                #   Call between .Ltmp13 and .Ltmp17
	.byte	0                               #     has no landing pad
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp17-.Lfunc_begin0          # >> Call Site 10 <<
	.uleb128 .Ltmp18-.Ltmp17                #   Call between .Ltmp17 and .Ltmp18
	.uleb128 .Ltmp19-.Lfunc_begin0          #     jumps to .Ltmp19
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp18-.Lfunc_begin0          # >> Call Site 11 <<
	.uleb128 .Ltmp20-.Ltmp18                #   Call between .Ltmp18 and .Ltmp20
	.byte	0                               #     has no landing pad
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp20-.Lfunc_begin0          # >> Call Site 12 <<
	.uleb128 .Ltmp21-.Ltmp20                #   Call between .Ltmp20 and .Ltmp21
	.uleb128 .Ltmp22-.Lfunc_begin0          #     jumps to .Ltmp22
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp23-.Lfunc_begin0          # >> Call Site 13 <<
	.uleb128 .Ltmp26-.Ltmp23                #   Call between .Ltmp23 and .Ltmp26
	.uleb128 .Ltmp27-.Lfunc_begin0          #     jumps to .Ltmp27
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp28-.Lfunc_begin0          # >> Call Site 14 <<
	.uleb128 .Ltmp57-.Ltmp28                #   Call between .Ltmp28 and .Ltmp57
	.uleb128 .Ltmp58-.Lfunc_begin0          #     jumps to .Ltmp58
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp59-.Lfunc_begin0          # >> Call Site 15 <<
	.uleb128 .Ltmp64-.Ltmp59                #   Call between .Ltmp59 and .Ltmp64
	.uleb128 .Ltmp65-.Lfunc_begin0          #     jumps to .Ltmp65
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp66-.Lfunc_begin0          # >> Call Site 16 <<
	.uleb128 .Ltmp67-.Ltmp66                #   Call between .Ltmp66 and .Ltmp67
	.uleb128 .Ltmp68-.Lfunc_begin0          #     jumps to .Ltmp68
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp69-.Lfunc_begin0          # >> Call Site 17 <<
	.uleb128 .Ltmp70-.Ltmp69                #   Call between .Ltmp69 and .Ltmp70
	.uleb128 .Ltmp71-.Lfunc_begin0          #     jumps to .Ltmp71
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp70-.Lfunc_begin0          # >> Call Site 18 <<
	.uleb128 .Lfunc_end4-.Ltmp70            #   Call between .Ltmp70 and .Lfunc_end4
	.byte	0                               #     has no landing pad
	.byte	0                               #   On action: cleanup
.Lcst_end0:
	.byte	1                               # >> Action Record 1 <<
                                        #   Catch TypeInfo 1
	.byte	0                               #   No further actions
	.p2align	2, 0x0
                                        # >> Catch TypeInfos <<
	.long	0                               # TypeInfo 1
.Lttbase0:
	.p2align	2, 0x0
                                        # -- End function
	.section	.text.__clang_call_terminate,"axG",@progbits,__clang_call_terminate,comdat
	.hidden	__clang_call_terminate          # -- Begin function __clang_call_terminate
	.weak	__clang_call_terminate
	.type	__clang_call_terminate,@function
__clang_call_terminate:                 # @__clang_call_terminate
	.cfi_startproc
# %bb.0:
	pushq	%rax
	.cfi_def_cfa_offset 16
	callq	__cxa_begin_catch
	callq	_ZSt9terminatev
.Lfunc_end5:
	.size	__clang_call_terminate, .Lfunc_end5-__clang_call_terminate
	.cfi_endproc
                                        # -- End function
	.section	.text._ZNSt8__detail18__to_chars_10_implIjEEvPcjT_,"axG",@progbits,_ZNSt8__detail18__to_chars_10_implIjEEvPcjT_,comdat
	.weak	_ZNSt8__detail18__to_chars_10_implIjEEvPcjT_ # -- Begin function _ZNSt8__detail18__to_chars_10_implIjEEvPcjT_
	.type	_ZNSt8__detail18__to_chars_10_implIjEEvPcjT_,@function
_ZNSt8__detail18__to_chars_10_implIjEEvPcjT_: # @_ZNSt8__detail18__to_chars_10_implIjEEvPcjT_
	.cfi_startproc
# %bb.0:
                                        # kill: def $esi killed $esi def $rsi
	cmpl	$100, %edx
	jb	.LBB6_1
# %bb.2:                                # %.lr.ph.preheader
	decl	%esi
.LBB6_3:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movl	%edx, %eax
	imulq	$1374389535, %rax, %rax         # imm = 0x51EB851F
	shrq	$37, %rax
	imull	$100, %eax, %ecx
	movl	%edx, %r8d
	subl	%ecx, %r8d
	movb	.L__const._ZNSt8__detail18__to_chars_10_implIjEEvPcjT_.__digits+1(%r8,%r8), %cl
	movl	%esi, %r9d
	movb	%cl, (%rdi,%r9)
	movb	.L__const._ZNSt8__detail18__to_chars_10_implIjEEvPcjT_.__digits(%r8,%r8), %cl
	leal	-1(%rsi), %r8d
	movb	%cl, (%rdi,%r8)
	addl	$-2, %esi
	cmpl	$9999, %edx                     # imm = 0x270F
	movl	%eax, %edx
	ja	.LBB6_3
	jmp	.LBB6_4
.LBB6_1:
	movl	%edx, %eax
.LBB6_4:                                # %._crit_edge
	cmpl	$10, %eax
	jb	.LBB6_6
# %bb.5:
	movl	%eax, %eax
	leaq	(%rax,%rax), %rcx
	movl	%ecx, %ecx
	movb	.L__const._ZNSt8__detail18__to_chars_10_implIjEEvPcjT_.__digits+1(%rcx), %cl
	movb	%cl, 1(%rdi)
	movb	.L__const._ZNSt8__detail18__to_chars_10_implIjEEvPcjT_.__digits(%rax,%rax), %al
	jmp	.LBB6_7
.LBB6_6:
	orb	$48, %al
.LBB6_7:
	movb	%al, (%rdi)
	retq
.Lfunc_end6:
	.size	_ZNSt8__detail18__to_chars_10_implIjEEvPcjT_, .Lfunc_end6-_ZNSt8__detail18__to_chars_10_implIjEEvPcjT_
	.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
	.cfi_offset %rbx, -16
	movq	__hip_gpubin_handle(%rip), %rbx
	testq	%rbx, %rbx
	jne	.LBB7_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB7_2:
	subq	$32, %rsp
	.cfi_adjust_cfa_offset 32
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z24vertex_disjoint_bfs_utilv, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	addq	$32, %rsp
	.cfi_adjust_cfa_offset -32
	movl	$d_flat_adj_list, %esi
	movl	$.L__unnamed_2, %edx
	movl	$.L__unnamed_2, %ecx
	movl	$8000000, %r9d                  # imm = 0x7A1200
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_adjust_cfa_offset 8
	pushq	$0
	.cfi_adjust_cfa_offset 8
	callq	__hipRegisterVar
	addq	$16, %rsp
	.cfi_adjust_cfa_offset -16
	movl	$d_degree, %esi
	movl	$.L__unnamed_3, %edx
	movl	$.L__unnamed_3, %ecx
	movl	$8004, %r9d                     # imm = 0x1F44
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_adjust_cfa_offset 8
	pushq	$0
	.cfi_adjust_cfa_offset 8
	callq	__hipRegisterVar
	addq	$16, %rsp
	.cfi_adjust_cfa_offset -16
	movl	$d_list_ptr, %esi
	movl	$.L__unnamed_4, %edx
	movl	$.L__unnamed_4, %ecx
	movl	$8008, %r9d                     # imm = 0x1F48
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_adjust_cfa_offset 8
	pushq	$0
	.cfi_adjust_cfa_offset 8
	callq	__hipRegisterVar
	addq	$16, %rsp
	.cfi_adjust_cfa_offset -16
	movl	$d_list_ptr_copy, %esi
	movl	$.L__unnamed_5, %edx
	movl	$.L__unnamed_5, %ecx
	movl	$8008, %r9d                     # imm = 0x1F48
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_adjust_cfa_offset 8
	pushq	$0
	.cfi_adjust_cfa_offset 8
	callq	__hipRegisterVar
	addq	$16, %rsp
	.cfi_adjust_cfa_offset -16
	movl	$d_is_matched_edge, %esi
	movl	$.L__unnamed_6, %edx
	movl	$.L__unnamed_6, %ecx
	movl	$4004001, %r9d                  # imm = 0x3D18A1
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_adjust_cfa_offset 8
	pushq	$0
	.cfi_adjust_cfa_offset 8
	callq	__hipRegisterVar
	addq	$16, %rsp
	.cfi_adjust_cfa_offset -16
	movl	$d_is_matched_vertex, %esi
	movl	$.L__unnamed_7, %edx
	movl	$.L__unnamed_7, %ecx
	movl	$2001, %r9d                     # imm = 0x7D1
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_adjust_cfa_offset 8
	pushq	$0
	.cfi_adjust_cfa_offset 8
	callq	__hipRegisterVar
	addq	$16, %rsp
	.cfi_adjust_cfa_offset -16
	movl	$d_partner_vertex, %esi
	movl	$.L__unnamed_8, %edx
	movl	$.L__unnamed_8, %ecx
	movl	$8004, %r9d                     # imm = 0x1F44
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_adjust_cfa_offset 8
	pushq	$0
	.cfi_adjust_cfa_offset 8
	callq	__hipRegisterVar
	addq	$16, %rsp
	.cfi_adjust_cfa_offset -16
	movl	$d_visited, %esi
	movl	$.L__unnamed_9, %edx
	movl	$.L__unnamed_9, %ecx
	movl	$8004, %r9d                     # imm = 0x1F44
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_adjust_cfa_offset 8
	pushq	$0
	.cfi_adjust_cfa_offset 8
	callq	__hipRegisterVar
	addq	$16, %rsp
	.cfi_adjust_cfa_offset -16
	movl	$d_bfs_parent, %esi
	movl	$.L__unnamed_10, %edx
	movl	$.L__unnamed_10, %ecx
	movl	$8004, %r9d                     # imm = 0x1F44
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_adjust_cfa_offset 8
	pushq	$0
	.cfi_adjust_cfa_offset 8
	callq	__hipRegisterVar
	addq	$16, %rsp
	.cfi_adjust_cfa_offset -16
	movl	$d_is_parent_change, %esi
	movl	$.L__unnamed_11, %edx
	movl	$.L__unnamed_11, %ecx
	movl	$2001, %r9d                     # imm = 0x7D1
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_adjust_cfa_offset 8
	pushq	$0
	.cfi_adjust_cfa_offset 8
	callq	__hipRegisterVar
	addq	$16, %rsp
	.cfi_adjust_cfa_offset -16
	movl	$d_frontier, %esi
	movl	$.L__unnamed_12, %edx
	movl	$.L__unnamed_12, %ecx
	movl	$8004, %r9d                     # imm = 0x1F44
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_adjust_cfa_offset 8
	pushq	$0
	.cfi_adjust_cfa_offset 8
	callq	__hipRegisterVar
	addq	$16, %rsp
	.cfi_adjust_cfa_offset -16
	movl	$d_next_frontier, %esi
	movl	$.L__unnamed_13, %edx
	movl	$.L__unnamed_13, %ecx
	movl	$8004, %r9d                     # imm = 0x1F44
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_adjust_cfa_offset 8
	pushq	$0
	.cfi_adjust_cfa_offset 8
	callq	__hipRegisterVar
	addq	$16, %rsp
	.cfi_adjust_cfa_offset -16
	movl	$num_aug_paths, %esi
	movl	$.L__unnamed_14, %edx
	movl	$.L__unnamed_14, %ecx
	movl	$4, %r9d
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_adjust_cfa_offset 8
	pushq	$0
	.cfi_adjust_cfa_offset 8
	callq	__hipRegisterVar
	addq	$16, %rsp
	.cfi_adjust_cfa_offset -16
	movl	$__hip_module_dtor, %edi
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end7:
	.size	__hip_module_ctor, .Lfunc_end7-__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	.LBB8_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
.LBB8_2:
	retq
.Lfunc_end8:
	.size	__hip_module_dtor, .Lfunc_end8-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	d_flat_adj_list,@object         # @d_flat_adj_list
	.local	d_flat_adj_list
	.comm	d_flat_adj_list,8000000,16
	.type	d_degree,@object                # @d_degree
	.local	d_degree
	.comm	d_degree,8004,16
	.type	d_list_ptr,@object              # @d_list_ptr
	.local	d_list_ptr
	.comm	d_list_ptr,8008,16
	.type	d_list_ptr_copy,@object         # @d_list_ptr_copy
	.local	d_list_ptr_copy
	.comm	d_list_ptr_copy,8008,16
	.type	d_is_matched_edge,@object       # @d_is_matched_edge
	.local	d_is_matched_edge
	.comm	d_is_matched_edge,4004001,16
	.type	d_is_matched_vertex,@object     # @d_is_matched_vertex
	.local	d_is_matched_vertex
	.comm	d_is_matched_vertex,2001,16
	.type	d_partner_vertex,@object        # @d_partner_vertex
	.local	d_partner_vertex
	.comm	d_partner_vertex,8004,16
	.type	d_visited,@object               # @d_visited
	.local	d_visited
	.comm	d_visited,8004,16
	.type	d_bfs_parent,@object            # @d_bfs_parent
	.local	d_bfs_parent
	.comm	d_bfs_parent,8004,16
	.type	d_is_parent_change,@object      # @d_is_parent_change
	.local	d_is_parent_change
	.comm	d_is_parent_change,2001,16
	.type	d_frontier,@object              # @d_frontier
	.local	d_frontier
	.comm	d_frontier,8004,16
	.type	d_next_frontier,@object         # @d_next_frontier
	.local	d_next_frontier
	.comm	d_next_frontier,8004,16
	.type	num_aug_paths,@object           # @num_aug_paths
	.local	num_aug_paths
	.comm	num_aug_paths,4,4
	.type	h_flat_adj_list,@object         # @h_flat_adj_list
	.bss
	.globl	h_flat_adj_list
	.p2align	3, 0x0
h_flat_adj_list:
	.quad	0
	.size	h_flat_adj_list, 8

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

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

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

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

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

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

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

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

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

	.type	fc,@object                      # @fc
	.data
	.globl	fc
	.p2align	2, 0x0
fc:
	.long	1000                            # 0x3e8
	.size	fc, 4

	.type	h_frontier,@object              # @h_frontier
	.bss
	.globl	h_frontier
	.p2align	3, 0x0
h_frontier:
	.quad	0
	.size	h_frontier, 8

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

	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"Matchings: "
	.size	.L.str, 12

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

	.type	_Z24vertex_disjoint_bfs_utilv,@object # @_Z24vertex_disjoint_bfs_utilv
	.section	.rodata,"a",@progbits
	.globl	_Z24vertex_disjoint_bfs_utilv
	.p2align	3, 0x0
_Z24vertex_disjoint_bfs_utilv:
	.quad	_Z39__device_stub__vertex_disjoint_bfs_utilv
	.size	_Z24vertex_disjoint_bfs_utilv, 8

	.type	.L.str.2,@object                # @.str.2
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str.2:
	.asciz	"Error! Not a matching!"
	.size	.L.str.2, 23

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	"FC_"
	.size	.L.str.3, 4

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

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

	.type	.L.str.6,@object                # @.str.6
.L.str.6:
	.asciz	"Done reading edges"
	.size	.L.str.6, 19

	.type	.L.str.7,@object                # @.str.7
.L.str.7:
	.asciz	"Number of matchings: %d \n"
	.size	.L.str.7, 26

	.type	.L.str.8,@object                # @.str.8
.L.str.8:
	.asciz	"Time elapsed %lf\n"
	.size	.L.str.8, 18

	.type	.L__const._ZNSt8__detail18__to_chars_10_implIjEEvPcjT_.__digits,@object # @__const._ZNSt8__detail18__to_chars_10_implIjEEvPcjT_.__digits
	.section	.rodata.str1.16,"aMS",@progbits,1
	.p2align	4, 0x0
.L__const._ZNSt8__detail18__to_chars_10_implIjEEvPcjT_.__digits:
	.asciz	"00010203040506070809101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899"
	.size	.L__const._ZNSt8__detail18__to_chars_10_implIjEEvPcjT_.__digits, 201

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

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"d_flat_adj_list"
	.size	.L__unnamed_2, 16

	.type	.L__unnamed_3,@object           # @2
.L__unnamed_3:
	.asciz	"d_degree"
	.size	.L__unnamed_3, 9

	.type	.L__unnamed_4,@object           # @3
.L__unnamed_4:
	.asciz	"d_list_ptr"
	.size	.L__unnamed_4, 11

	.type	.L__unnamed_5,@object           # @4
.L__unnamed_5:
	.asciz	"d_list_ptr_copy"
	.size	.L__unnamed_5, 16

	.type	.L__unnamed_6,@object           # @5
.L__unnamed_6:
	.asciz	"d_is_matched_edge"
	.size	.L__unnamed_6, 18

	.type	.L__unnamed_7,@object           # @6
.L__unnamed_7:
	.asciz	"d_is_matched_vertex"
	.size	.L__unnamed_7, 20

	.type	.L__unnamed_8,@object           # @7
.L__unnamed_8:
	.asciz	"d_partner_vertex"
	.size	.L__unnamed_8, 17

	.type	.L__unnamed_9,@object           # @8
.L__unnamed_9:
	.asciz	"d_visited"
	.size	.L__unnamed_9, 10

	.type	.L__unnamed_10,@object          # @9
.L__unnamed_10:
	.asciz	"d_bfs_parent"
	.size	.L__unnamed_10, 13

	.type	.L__unnamed_11,@object          # @10
.L__unnamed_11:
	.asciz	"d_is_parent_change"
	.size	.L__unnamed_11, 19

	.type	.L__unnamed_12,@object          # @11
.L__unnamed_12:
	.asciz	"d_frontier"
	.size	.L__unnamed_12, 11

	.type	.L__unnamed_13,@object          # @12
.L__unnamed_13:
	.asciz	"d_next_frontier"
	.size	.L__unnamed_13, 16

	.type	.L__unnamed_14,@object          # @13
.L__unnamed_14:
	.asciz	"num_aug_paths"
	.size	.L__unnamed_14, 14

	.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 _Z39__device_stub__vertex_disjoint_bfs_utilv
	.addrsig_sym __gxx_personality_v0
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Unwind_Resume
	.addrsig_sym d_flat_adj_list
	.addrsig_sym d_degree
	.addrsig_sym d_list_ptr
	.addrsig_sym d_list_ptr_copy
	.addrsig_sym d_is_matched_edge
	.addrsig_sym d_is_matched_vertex
	.addrsig_sym d_partner_vertex
	.addrsig_sym d_visited
	.addrsig_sym d_bfs_parent
	.addrsig_sym d_is_parent_change
	.addrsig_sym d_frontier
	.addrsig_sym d_next_frontier
	.addrsig_sym num_aug_paths
	.addrsig_sym _ZSt4cout
	.addrsig_sym _Z24vertex_disjoint_bfs_utilv
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z24vertex_disjoint_bfs_utilv ; -- Begin function _Z24vertex_disjoint_bfs_utilv
	.globl	_Z24vertex_disjoint_bfs_utilv
	.p2align	8
	.type	_Z24vertex_disjoint_bfs_utilv,@function
_Z24vertex_disjoint_bfs_utilv:          ; @_Z24vertex_disjoint_bfs_utilv
; %bb.0:
	v_lshl_or_b32 v2, s15, 10, v0
	s_mov_b32 s0, exec_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_add_nc_u32_e32 v0, 1, v2
	v_cmpx_lt_i32_e32 0x7cf, v2
	s_xor_b32 s0, exec_lo, s0
; %bb.1:                                ; %._Z25initialise_partner_vertexv.exit_crit_edge
	v_mov_b32_e32 v1, 0
; %bb.2:                                ; %Flow44
	s_and_not1_saveexec_b32 s0, s0
	s_cbranch_execz .LBB0_4
; %bb.3:                                ; %_Z13clear_visitedv.exit
	v_ashrrev_i32_e32 v3, 31, v2
	s_getpc_b64 s[2:3]
	s_add_u32 s2, s2, d_visited@rel32@lo+8
	s_addc_u32 s3, s3, d_visited@rel32@hi+16
	s_getpc_b64 s[4:5]
	s_add_u32 s4, s4, d_bfs_parent@rel32@lo+8
	s_addc_u32 s5, s5, d_bfs_parent@rel32@hi+16
	s_getpc_b64 s[6:7]
	s_add_u32 s6, s6, d_is_parent_change@rel32@lo+5
	s_addc_u32 s7, s7, d_is_parent_change@rel32@hi+13
	v_dual_mov_b32 v12, 0 :: v_dual_mov_b32 v13, 1
	v_lshlrev_b64 v[4:5], 2, v[2:3]
	v_ashrrev_i32_e32 v1, 31, v0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_u32 v6, vcc_lo, v4, s2
	v_add_co_ci_u32_e32 v7, vcc_lo, s3, v5, vcc_lo
	v_add_co_u32 v8, vcc_lo, v4, s4
	v_add_co_ci_u32_e32 v9, vcc_lo, s5, v5, vcc_lo
	v_add_co_u32 v10, vcc_lo, v2, s6
	s_getpc_b64 s[2:3]
	s_add_u32 s2, s2, d_partner_vertex@rel32@lo+8
	s_addc_u32 s3, s3, d_partner_vertex@rel32@hi+16
	v_add_co_ci_u32_e32 v11, vcc_lo, s7, v3, vcc_lo
	v_add_co_u32 v3, vcc_lo, v4, s2
	v_add_co_ci_u32_e32 v4, vcc_lo, s3, v5, vcc_lo
	v_mov_b32_e32 v5, -1
	global_store_b32 v[6:7], v12, off
	global_store_b32 v[8:9], v0, off
	global_store_b8 v[10:11], v13, off
	global_store_b32 v[3:4], v5, off
.LBB0_4:                                ; %_Z25initialise_partner_vertexv.exit
	s_or_b32 exec_lo, exec_lo, s0
	s_delay_alu instid0(SALU_CYCLE_1)
	s_mov_b32 s0, exec_lo
	v_cmpx_gt_i32_e32 0x3e8, v2
	s_cbranch_execz .LBB0_41
; %bb.5:
	v_lshlrev_b64 v[6:7], 2, v[0:1]
	s_getpc_b64 s[0:1]
	s_add_u32 s0, s0, d_visited@rel32@lo+4
	s_addc_u32 s1, s1, d_visited@rel32@hi+12
	s_mov_b32 s2, exec_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_add_co_u32 v4, vcc_lo, v6, s0
	v_add_co_ci_u32_e32 v5, vcc_lo, s1, v7, vcc_lo
	global_load_b32 v3, v[4:5], off
	s_waitcnt vmcnt(0)
	v_cmpx_eq_u32_e32 0, v3
	s_cbranch_execz .LBB0_27
; %bb.6:
	s_getpc_b64 s[0:1]
	s_add_u32 s0, s0, d_is_matched_vertex@rel32@lo+4
	s_addc_u32 s1, s1, d_is_matched_vertex@rel32@hi+12
	v_add_co_u32 v8, vcc_lo, v0, s0
	v_add_co_ci_u32_e32 v9, vcc_lo, s1, v1, vcc_lo
	global_load_u8 v3, v[8:9], off
	s_waitcnt vmcnt(0)
	v_cmp_eq_u16_e32 vcc_lo, 0, v3
	s_and_b32 exec_lo, exec_lo, vcc_lo
	s_cbranch_execz .LBB0_27
; %bb.7:
	s_getpc_b64 s[0:1]
	s_add_u32 s0, s0, d_frontier@rel32@lo+4
	s_addc_u32 s1, s1, d_frontier@rel32@hi+12
	v_add_co_u32 v9, vcc_lo, v6, s0
	v_add_co_ci_u32_e32 v10, vcc_lo, s1, v7, vcc_lo
	v_mov_b32_e32 v8, 1
	v_cmp_eq_u32_e32 vcc_lo, -1, v0
	s_mov_b32 s1, -1
	s_mov_b32 s3, exec_lo
	global_store_b32 v[9:10], v8, off
	v_cmpx_ne_u32_e32 -1, v0
	s_cbranch_execz .LBB0_23
; %bb.8:
	v_ashrrev_i32_e32 v3, 31, v2
	s_getpc_b64 s[4:5]
	s_add_u32 s4, s4, d_list_ptr@rel32@lo+4
	s_addc_u32 s5, s5, d_list_ptr@rel32@hi+12
	v_add_co_u32 v6, s0, v6, s4
	s_getpc_b64 s[6:7]
	s_add_u32 s6, s6, d_list_ptr@rel32@lo+12
	s_addc_u32 s7, s7, d_list_ptr@rel32@hi+20
	v_lshlrev_b64 v[2:3], 2, v[2:3]
	v_add_co_ci_u32_e64 v7, s0, s5, v7, s0
	s_mov_b32 s4, exec_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_co_u32 v9, s0, v2, s6
	v_add_co_ci_u32_e64 v10, s0, s7, v3, s0
	s_clause 0x1
	global_load_b32 v2, v[6:7], off
	global_load_b32 v13, v[9:10], off
	global_store_b32 v[4:5], v8, off
	s_waitcnt vmcnt(0)
	v_cmpx_lt_i32_e64 v2, v13
	s_cbranch_execz .LBB0_22
; %bb.9:                                ; %.lr.ph.i
	v_mad_u64_u32 v[4:5], null, 0x7d1, v0, 0
	v_ashrrev_i32_e32 v3, 31, v2
	s_getpc_b64 s[0:1]
	s_add_u32 s0, s0, d_flat_adj_list@rel32@lo+4
	s_addc_u32 s1, s1, d_flat_adj_list@rel32@hi+12
	s_mov_b32 s6, 0
                                        ; implicit-def: $sgpr5
                                        ; implicit-def: $sgpr7
	s_getpc_b64 s[12:13]
	s_add_u32 s12, s12, d_next_frontier@rel32@lo+4
	s_addc_u32 s13, s13, d_next_frontier@rel32@hi+12
	v_lshlrev_b64 v[7:8], 2, v[2:3]
	v_mad_u64_u32 v[9:10], null, 0x7d1, v1, v[5:6]
	v_mov_b32_e32 v3, 1
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_co_u32 v7, s0, v7, s0
	v_add_co_ci_u32_e64 v8, s0, s1, v8, s0
	s_delay_alu instid0(VALU_DEP_4)
	v_mov_b32_e32 v1, v9
.LBB0_10:                               ; =>This Inner Loop Header: Depth=1
	global_load_b32 v5, v[7:8], off
	s_getpc_b64 s[0:1]
	s_add_u32 s0, s0, d_visited@rel32@lo+4
	s_addc_u32 s1, s1, d_visited@rel32@hi+12
	s_mov_b32 s8, exec_lo
                                        ; implicit-def: $sgpr10
	s_waitcnt vmcnt(0)
	v_ashrrev_i32_e32 v6, 31, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[9:10], 2, v[5:6]
	v_add_co_u32 v11, s0, v9, s0
	s_delay_alu instid0(VALU_DEP_1)
	v_add_co_ci_u32_e64 v12, s0, s1, v10, s0
	global_atomic_swap_b32 v14, v[11:12], v3, off glc
	s_waitcnt vmcnt(0)
	v_cmp_ne_u32_e64 s0, 0, v14
	v_cmpx_eq_u32_e32 0, v14
	s_cbranch_execz .LBB0_16
; %bb.11:                               ;   in Loop: Header=BB0_10 Depth=1
	v_add_co_u32 v14, s1, v4, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_4) | instid1(VALU_DEP_1)
	v_add_co_ci_u32_e64 v15, s1, v1, v6, s1
	s_getpc_b64 s[10:11]
	s_add_u32 s10, s10, d_is_matched_edge@rel32@lo+4
	s_addc_u32 s11, s11, d_is_matched_edge@rel32@hi+12
	v_add_co_u32 v14, s1, v14, s10
	v_add_co_ci_u32_e64 v15, s1, s11, v15, s1
	s_getpc_b64 s[10:11]
	s_add_u32 s10, s10, d_bfs_parent@rel32@lo+4
	s_addc_u32 s11, s11, d_bfs_parent@rel32@hi+12
	s_mov_b32 s9, exec_lo
	global_store_b32 v[11:12], v3, off
	global_load_u8 v16, v[14:15], off
	v_add_co_u32 v14, s1, v9, s10
	s_delay_alu instid0(VALU_DEP_1)
	v_add_co_ci_u32_e64 v15, s1, s11, v10, s1
	s_mov_b32 s11, -1
                                        ; implicit-def: $sgpr10
	global_store_b32 v[14:15], v0, off
	s_waitcnt vmcnt(0)
	v_cmpx_eq_u16_e32 0, v16
	s_cbranch_execz .LBB0_15
; %bb.12:                               ;   in Loop: Header=BB0_10 Depth=1
	s_getpc_b64 s[10:11]
	s_add_u32 s10, s10, d_is_matched_vertex@rel32@lo+4
	s_addc_u32 s11, s11, d_is_matched_vertex@rel32@hi+12
	v_add_co_u32 v11, s1, v5, s10
	s_delay_alu instid0(VALU_DEP_1)
	v_add_co_ci_u32_e64 v12, s1, s11, v6, s1
	s_mov_b32 s11, 0
	s_mov_b32 s10, exec_lo
	global_load_u8 v11, v[11:12], off
	s_waitcnt vmcnt(0)
	v_cmpx_ne_u16_e32 0, v11
	s_cbranch_execz .LBB0_14
; %bb.13:                               ;   in Loop: Header=BB0_10 Depth=1
	v_add_co_u32 v9, s1, v9, s12
	s_delay_alu instid0(VALU_DEP_1)
	v_add_co_ci_u32_e64 v10, s1, s13, v10, s1
	s_mov_b32 s11, exec_lo
	global_store_b32 v[9:10], v3, off
.LBB0_14:                               ; %Flow35
                                        ;   in Loop: Header=BB0_10 Depth=1
	s_or_b32 exec_lo, exec_lo, s10
	s_mov_b32 s10, -1
	s_or_not1_b32 s11, s11, exec_lo
.LBB0_15:                               ; %Flow34
                                        ;   in Loop: Header=BB0_10 Depth=1
	s_or_b32 exec_lo, exec_lo, s9
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	s_and_not1_b32 s0, s0, exec_lo
	s_and_b32 s1, s11, exec_lo
	s_or_b32 s0, s0, s1
.LBB0_16:                               ; %Flow33
                                        ;   in Loop: Header=BB0_10 Depth=1
	s_or_b32 exec_lo, exec_lo, s8
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 s7, s7, exec_lo
	s_and_b32 s8, s10, exec_lo
	s_mov_b32 s1, -1
	s_or_b32 s7, s7, s8
	s_and_saveexec_b32 s8, s0
; %bb.17:                               ; %.critedge56.i
                                        ;   in Loop: Header=BB0_10 Depth=1
	v_add_nc_u32_e32 v2, 1, v2
	v_add_co_u32 v7, s1, v7, 4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_ci_u32_e64 v8, s1, 0, v8, s1
	v_cmp_ge_i32_e64 s0, v2, v13
	s_and_not1_b32 s7, s7, exec_lo
	s_delay_alu instid0(VALU_DEP_1)
	s_or_not1_b32 s1, s0, exec_lo
; %bb.18:                               ; %Flow36
                                        ;   in Loop: Header=BB0_10 Depth=1
	s_or_b32 exec_lo, exec_lo, s8
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b32 s0, exec_lo, s1
	s_or_b32 s6, s0, s6
	s_and_not1_b32 s0, s5, exec_lo
	s_and_b32 s1, s7, exec_lo
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 s5, s0, s1
	s_and_not1_b32 exec_lo, exec_lo, s6
	s_cbranch_execnz .LBB0_10
; %bb.19:                               ; %loop.exit.guard
	s_or_b32 exec_lo, exec_lo, s6
	s_mov_b32 s0, -1
	s_and_saveexec_b32 s1, s5
	s_delay_alu instid0(SALU_CYCLE_1)
	s_xor_b32 s1, exec_lo, s1
	s_cbranch_execz .LBB0_21
; %bb.20:                               ; %.critedge.i
	s_getpc_b64 s[6:7]
	s_add_u32 s6, s6, num_aug_paths@rel32@lo+4
	s_addc_u32 s7, s7, num_aug_paths@rel32@hi+12
	s_getpc_b64 s[8:9]
	s_add_u32 s8, s8, d_is_parent_change@rel32@lo+4
	s_addc_u32 s9, s9, d_is_parent_change@rel32@hi+12
	s_load_b32 s5, s[6:7], 0x0
	v_add_co_u32 v0, s0, v5, s8
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_4) | instid1(SALU_CYCLE_1)
	v_add_co_ci_u32_e64 v1, s0, s9, v6, s0
	v_dual_mov_b32 v2, 1 :: v_dual_mov_b32 v3, 0
	global_store_b8 v[0:1], v2, off
	s_waitcnt lgkmcnt(0)
	s_add_i32 s0, s5, 1
	v_mov_b32_e32 v4, s0
	s_xor_b32 s0, exec_lo, -1
	global_store_b32 v3, v4, s[6:7]
.LBB0_21:                               ; %Flow40
	s_or_b32 exec_lo, exec_lo, s1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_not1_b32 s1, s0, exec_lo
.LBB0_22:                               ; %Flow39
	s_or_b32 exec_lo, exec_lo, s4
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	s_and_not1_b32 s0, vcc_lo, exec_lo
	s_and_b32 s1, s1, exec_lo
	s_or_b32 vcc_lo, s0, s1
.LBB0_23:                               ; %Flow38
	s_or_b32 exec_lo, exec_lo, s3
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_b32 exec_lo, exec_lo, vcc_lo
	s_cbranch_execz .LBB0_26
; %bb.24:                               ; %.critedge52._crit_edge.i.preheader
	v_mov_b32_e32 v0, 0
	s_mov_b64 s[0:1], 4
.LBB0_25:                               ; %.critedge52._crit_edge.i
                                        ; =>This Inner Loop Header: Depth=1
	s_getpc_b64 s[4:5]
	s_add_u32 s4, s4, d_next_frontier@rel32@lo+4
	s_addc_u32 s5, s5, d_next_frontier@rel32@hi+12
	s_add_u32 s4, s0, s4
	s_addc_u32 s5, s1, s5
	global_load_b32 v1, v0, s[4:5]
	s_getpc_b64 s[4:5]
	s_add_u32 s4, s4, d_frontier@rel32@lo+4
	s_addc_u32 s5, s5, d_frontier@rel32@hi+12
	s_add_u32 s4, s0, s4
	s_addc_u32 s5, s1, s5
	s_add_u32 s0, s0, 4
	s_addc_u32 s1, s1, 0
	s_cmpk_lg_i32 s0, 0x1f44
	s_waitcnt vmcnt(0)
	global_store_b32 v0, v1, s[4:5]
	s_cbranch_scc1 .LBB0_25
.LBB0_26:                               ; %_Z19vertex_disjoint_bfsiii.exit
	s_waitcnt_vscnt null, 0x0
	s_barrier
	buffer_gl0_inv
.LBB0_27:
	s_or_b32 exec_lo, exec_lo, s2
	v_mov_b32_e32 v0, 0
	s_getpc_b64 s[0:1]
	s_add_u32 s0, s0, num_aug_paths@rel32@lo+4
	s_addc_u32 s1, s1, num_aug_paths@rel32@hi+12
	global_load_b32 v1, v0, s[0:1]
	s_mov_b32 s0, 1
	s_waitcnt vmcnt(0)
	v_cmp_gt_i32_e32 vcc_lo, 1, v1
	s_cbranch_vccnz .LBB0_41
; %bb.28:                               ; %.preheader.preheader
	v_dual_mov_b32 v1, -1 :: v_dual_mov_b32 v2, 1
	s_mov_b32 s1, 0
	s_getpc_b64 s[18:19]
	s_add_u32 s18, s18, d_partner_vertex@rel32@lo+4
	s_addc_u32 s19, s19, d_partner_vertex@rel32@hi+12
.LBB0_29:                               ; %.preheader
                                        ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB0_32 Depth 2
	v_mov_b32_e32 v3, s0
	s_getpc_b64 s[2:3]
	s_add_u32 s2, s2, d_is_parent_change@rel32@lo+4
	s_addc_u32 s3, s3, d_is_parent_change@rel32@hi+12
	global_load_u8 v3, v3, s[2:3]
	s_waitcnt vmcnt(0)
	v_cmp_eq_u16_e32 vcc_lo, 0, v3
	s_cbranch_vccnz .LBB0_40
; %bb.30:                               ; %.preheader.i
                                        ;   in Loop: Header=BB0_29 Depth=1
	s_getpc_b64 s[2:3]
	s_add_u32 s2, s2, d_bfs_parent@rel32@lo+4
	s_addc_u32 s3, s3, d_bfs_parent@rel32@hi+12
	s_lshl_b64 s[4:5], s[0:1], 2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_add_u32 s2, s4, s2
	s_addc_u32 s3, s5, s3
	global_load_b32 v3, v0, s[2:3]
	s_waitcnt vmcnt(0)
	v_cmp_eq_u32_e32 vcc_lo, s0, v3
	v_readfirstlane_b32 s4, v3
	s_cbranch_vccnz .LBB0_40
; %bb.31:                               ; %.lr.ph.i9.preheader
                                        ;   in Loop: Header=BB0_29 Depth=1
	s_mov_b32 s14, 1
	s_mov_b32 s2, s0
.LBB0_32:                               ; %.lr.ph.i9
                                        ;   Parent Loop BB0_29 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_mul_i32 s7, s2, 0x7d1
	s_and_b32 s10, s14, 1
	s_ashr_i32 s3, s2, 31
	s_mul_hi_i32 s6, s2, 0x7d1
	s_ashr_i32 s5, s4, 31
	s_add_u32 s7, s7, s4
	s_addc_u32 s11, s6, s5
	s_getpc_b64 s[8:9]
	s_add_u32 s8, s8, d_is_matched_edge@rel32@lo+4
	s_addc_u32 s9, s9, d_is_matched_edge@rel32@hi+12
	s_add_u32 s6, s7, s8
	s_addc_u32 s7, s11, s9
	s_mul_i32 s11, s4, 0x7d1
	s_mul_hi_i32 s12, s4, 0x7d1
	s_add_u32 s11, s11, s2
	s_addc_u32 s12, s12, s3
	s_add_u32 s8, s11, s8
	s_addc_u32 s9, s12, s9
	s_cmp_eq_u32 s10, 0
	s_mov_b32 s10, -1
	s_cbranch_scc1 .LBB0_34
; %bb.33:                               ;   in Loop: Header=BB0_32 Depth=2
	s_getpc_b64 s[10:11]
	s_add_u32 s10, s10, d_is_matched_vertex@rel32@lo+4
	s_addc_u32 s11, s11, d_is_matched_vertex@rel32@hi+12
	s_add_u32 s12, s2, s10
	s_addc_u32 s13, s3, s11
	s_add_u32 s16, s4, s10
	s_addc_u32 s17, s5, s11
	s_lshl_b64 s[20:21], s[2:3], 2
	v_dual_mov_b32 v3, s4 :: v_dual_mov_b32 v4, s2
	s_add_u32 s20, s20, s18
	s_addc_u32 s21, s21, s19
	s_lshl_b64 s[22:23], s[4:5], 2
	s_mov_b32 s10, 0
	s_clause 0x2
	global_store_b8 v0, v2, s[6:7]
	global_store_b8 v0, v2, s[12:13]
	global_store_b32 v0, v3, s[20:21]
	s_add_u32 s12, s22, s18
	s_addc_u32 s13, s23, s19
	s_clause 0x2
	global_store_b8 v0, v2, s[8:9]
	global_store_b8 v0, v2, s[16:17]
	global_store_b32 v0, v4, s[12:13]
.LBB0_34:                               ; %Flow27
                                        ;   in Loop: Header=BB0_32 Depth=2
	s_and_not1_b32 vcc_lo, exec_lo, s10
	s_cbranch_vccnz .LBB0_39
; %bb.35:                               ;   in Loop: Header=BB0_32 Depth=2
	s_getpc_b64 s[10:11]
	s_add_u32 s10, s10, d_partner_vertex@rel32@lo+4
	s_addc_u32 s11, s11, d_partner_vertex@rel32@hi+12
	s_lshl_b64 s[12:13], s[2:3], 2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_add_u32 s12, s12, s10
	s_addc_u32 s13, s13, s11
	global_load_b32 v3, v0, s[12:13]
	s_clause 0x1
	global_store_b8 v0, v0, s[6:7]
	global_store_b8 v0, v0, s[8:9]
	s_waitcnt vmcnt(0)
	v_cmp_ne_u32_e32 vcc_lo, s4, v3
	s_cbranch_vccnz .LBB0_37
; %bb.36:                               ;   in Loop: Header=BB0_32 Depth=2
	s_getpc_b64 s[6:7]
	s_add_u32 s6, s6, d_is_matched_vertex@rel32@lo+4
	s_addc_u32 s7, s7, d_is_matched_vertex@rel32@hi+12
	s_add_u32 s6, s2, s6
	s_addc_u32 s7, s3, s7
	s_clause 0x1
	global_store_b8 v0, v0, s[6:7]
	global_store_b32 v0, v1, s[12:13]
.LBB0_37:                               ;   in Loop: Header=BB0_32 Depth=2
	s_lshl_b64 s[6:7], s[4:5], 2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_add_u32 s6, s6, s10
	s_addc_u32 s7, s7, s11
	global_load_b32 v3, v0, s[6:7]
	s_waitcnt vmcnt(0)
	v_cmp_ne_u32_e32 vcc_lo, s2, v3
	s_cbranch_vccnz .LBB0_39
; %bb.38:                               ;   in Loop: Header=BB0_32 Depth=2
	s_getpc_b64 s[8:9]
	s_add_u32 s8, s8, d_is_matched_vertex@rel32@lo+4
	s_addc_u32 s9, s9, d_is_matched_vertex@rel32@hi+12
	s_add_u32 s4, s4, s8
	s_addc_u32 s5, s5, s9
	s_clause 0x1
	global_store_b8 v0, v0, s[4:5]
	global_store_b32 v0, v1, s[6:7]
.LBB0_39:                               ; %_Z13unmatch_edgesii.exit.i
                                        ;   in Loop: Header=BB0_32 Depth=2
	s_getpc_b64 s[4:5]
	s_add_u32 s4, s4, d_bfs_parent@rel32@lo+4
	s_addc_u32 s5, s5, d_bfs_parent@rel32@hi+12
	s_lshl_b64 s[2:3], s[2:3], 2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_add_u32 s2, s2, s4
	s_addc_u32 s3, s3, s5
	s_add_i32 s14, s14, 1
	global_load_b32 v3, v0, s[2:3]
	s_waitcnt vmcnt(0)
	v_ashrrev_i32_e32 v4, 31, v3
	v_readfirstlane_b32 s2, v3
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[4:5], 2, v[3:4]
	v_add_co_u32 v4, vcc_lo, v4, s4
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v5, vcc_lo, s5, v5, vcc_lo
	global_load_b32 v4, v[4:5], off
	s_waitcnt vmcnt(0)
	v_cmp_ne_u32_e32 vcc_lo, v4, v3
	v_readfirstlane_b32 s4, v4
	s_cbranch_vccnz .LBB0_32
.LBB0_40:                               ; %.loopexit.i
                                        ;   in Loop: Header=BB0_29 Depth=1
	s_add_i32 s0, s0, 1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmpk_lg_i32 s0, 0x7d1
	s_cbranch_scc1 .LBB0_29
.LBB0_41:                               ; %_Z16update_matchingsv.exit
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z24vertex_disjoint_bfs_utilv
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 0
		.amdhsa_user_sgpr_count 15
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 0
		.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 17
		.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_end0:
	.size	_Z24vertex_disjoint_bfs_utilv, .Lfunc_end0-_Z24vertex_disjoint_bfs_utilv
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 2268
; NumSgprs: 26
; NumVgprs: 17
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 3
; VGPRBlocks: 2
; NumSGPRsForWavesPerEU: 26
; NumVGPRsForWavesPerEU: 17
; 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
	.protected	d_flat_adj_list         ; @d_flat_adj_list
	.type	d_flat_adj_list,@object
	.section	.bss,"aw",@nobits
	.globl	d_flat_adj_list
	.p2align	4, 0x0
d_flat_adj_list:
	.zero	8000000
	.size	d_flat_adj_list, 8000000

	.protected	d_degree                ; @d_degree
	.type	d_degree,@object
	.globl	d_degree
	.p2align	4, 0x0
d_degree:
	.zero	8004
	.size	d_degree, 8004

	.protected	d_list_ptr              ; @d_list_ptr
	.type	d_list_ptr,@object
	.globl	d_list_ptr
	.p2align	4, 0x0
d_list_ptr:
	.zero	8008
	.size	d_list_ptr, 8008

	.protected	d_list_ptr_copy         ; @d_list_ptr_copy
	.type	d_list_ptr_copy,@object
	.globl	d_list_ptr_copy
	.p2align	4, 0x0
d_list_ptr_copy:
	.zero	8008
	.size	d_list_ptr_copy, 8008

	.protected	d_is_matched_edge       ; @d_is_matched_edge
	.type	d_is_matched_edge,@object
	.globl	d_is_matched_edge
	.p2align	4, 0x0
d_is_matched_edge:
	.zero	4004001
	.size	d_is_matched_edge, 4004001

	.protected	d_is_matched_vertex     ; @d_is_matched_vertex
	.type	d_is_matched_vertex,@object
	.globl	d_is_matched_vertex
	.p2align	4, 0x0
d_is_matched_vertex:
	.zero	2001
	.size	d_is_matched_vertex, 2001

	.protected	d_partner_vertex        ; @d_partner_vertex
	.type	d_partner_vertex,@object
	.globl	d_partner_vertex
	.p2align	4, 0x0
d_partner_vertex:
	.zero	8004
	.size	d_partner_vertex, 8004

	.protected	d_visited               ; @d_visited
	.type	d_visited,@object
	.globl	d_visited
	.p2align	4, 0x0
d_visited:
	.zero	8004
	.size	d_visited, 8004

	.protected	d_bfs_parent            ; @d_bfs_parent
	.type	d_bfs_parent,@object
	.globl	d_bfs_parent
	.p2align	4, 0x0
d_bfs_parent:
	.zero	8004
	.size	d_bfs_parent, 8004

	.protected	d_is_parent_change      ; @d_is_parent_change
	.type	d_is_parent_change,@object
	.globl	d_is_parent_change
	.p2align	4, 0x0
d_is_parent_change:
	.zero	2001
	.size	d_is_parent_change, 2001

	.protected	d_frontier              ; @d_frontier
	.type	d_frontier,@object
	.globl	d_frontier
	.p2align	4, 0x0
d_frontier:
	.zero	8004
	.size	d_frontier, 8004

	.protected	d_next_frontier         ; @d_next_frontier
	.type	d_next_frontier,@object
	.globl	d_next_frontier
	.p2align	4, 0x0
d_next_frontier:
	.zero	8004
	.size	d_next_frontier, 8004

	.protected	num_aug_paths           ; @num_aug_paths
	.type	num_aug_paths,@object
	.globl	num_aug_paths
	.p2align	2, 0x0
num_aug_paths:
	.long	0                               ; 0x0
	.size	num_aug_paths, 4

	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.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 d_flat_adj_list
	.addrsig_sym d_degree
	.addrsig_sym d_list_ptr
	.addrsig_sym d_is_matched_edge
	.addrsig_sym d_is_matched_vertex
	.addrsig_sym d_partner_vertex
	.addrsig_sym d_visited
	.addrsig_sym d_bfs_parent
	.addrsig_sym d_is_parent_change
	.addrsig_sym d_frontier
	.addrsig_sym d_next_frontier
	.addrsig_sym num_aug_paths
	.addrsig_sym __hip_cuid_
	.amdgpu_metadata
---
amdhsa.kernels:
  - .args:           []
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 4
    .kernarg_segment_size: 0
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z24vertex_disjoint_bfs_utilv
    .private_segment_fixed_size: 0
    .sgpr_count:     26
    .sgpr_spill_count: 0
    .symbol:         _Z24vertex_disjoint_bfs_utilv.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     17
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
