#include <math.h>
#include <cuda.h>
#include <vector>
#include <fstream>
#include <string>
#include <sstream>
#include <iostream>

#define BLOCK_WIDTH 512
/*
extern __shared__ stands for shared memory on device, which has two "warps" of 32 threads.
Google CUDA shared memory and warps.
To replace extern __shared__ int __smem[]; which requires you to explicitly
know the data type is integer in advance. But input file could be int, float, or double.
Since we don't know the data type of shared meomry __smem[], we use
template<class T> where T stands for all possible data types. We also 
need to instantiate all possible data types later
In return (T *) __smem; it is data type conversion
Suggest to figure out difference between overload, override, redefine
*/
template<class T>
struct SharedMemory {
	__device__ inline operator T *() {
		extern __shared__ int __smem[];
		return (T *) __smem;
	}
	
	__device__ inline operator const T *() const {
		extern __shared__ int __smem[];
		return (T *) __smem;
	}
};

/////////////////////////////////////////////////////////////////////////////
//                   CUDA Kernel: Global memory
/////////////////////////////////////////////////////////////////////////////
template<class T, int blockSize>
__global__ void countGlobalMem(T *g_idata, int *g_odata, int N) {
	unsigned int i = blockSize * blockIdx.x + threadIdx.x;
	int gi = 0;
	if (i < N) {
		if (g_idata[i] == 1000) {
			atomicAdd(&g_odata[9], 1);
		} else {
			gi = (int) g_idata[i] / 100;
			atomicAdd(&g_odata[gi], 1);
		}
	}
}

/////////////////////////////////////////////////////////////////////////////
//                   CUDA Kernel: shared memory
/////////////////////////////////////////////////////////////////////////////
template<class T, int blockSize>
__global__ void countSharedMem(T *g_idata, int *g_odata, int N, int maxNum, int barrelSize) {
	/*
	Each block has a sdata
	*/
	extern __shared__ int sdata[];
	unsigned int tid = threadIdx.x;
	int numBarrel = maxNum/barrelSize;
	
	unsigned int i = blockSize * blockIdx.x + threadIdx.x;
	//gi is group/barrel index
	int gi = 0;
	if (i < N) {
		if (g_idata[i] == maxNum) {
			atomicAdd(&sdata[numBarrel-1], 1);
		} else {
			gi = (int) g_idata[i] / barrelSize;
			atomicAdd(&sdata[gi], 1);
		}
	}
	
	//wait until sdata[0~9] in all blocks are ready
	__syncthreads();
	/*
	every block has threadIdx.x from 0 to 511
	size of g_odata is numBarrel * blocks
	sum of all blocks is done in myCountTest(), note there
	is += when output to "q2b.txt"
	*/
	if (tid < numBarrel) {
		g_odata[blockIdx.x * numBarrel + tid] = sdata[tid];
	}	
}

//////////////////////////////////////////////////////////////////////////////
//                     CUDA Kernel: prefix sum (Naiive)
///////////////////////////////////////////////////////////////////////////////
int nextPowerOf2(int x)
{
    --x;
    x |= x >> 1;
    x |= x >> 2;
    x |= x >> 4;
    x |= x >> 8;
    x |= x >> 16;
    return ++x;
}

__global__ void scan(int *d_idata, int *d_odata, int N) {
	extern __shared__ int sdata[];
	
	//cunyi
	int i = blockIdx.x * blockDim.x + threadIdx.x;
	if (i < N) {
		sdata[threadIdx.x] = d_idata[i];
		//printf("\n sdata[%d]: %d", i, sdata[threadIdx.x]);
	}
	
	for (unsigned int stride = 1; stride <= threadIdx.x; stride *= 2) {
		__syncthreads();
		int in1 = sdata[threadIdx.x - stride];
		__syncthreads();
		sdata[threadIdx.x] += in1;
	}
	__syncthreads();
	if(i < N) {
		d_odata[threadIdx.x] = sdata[threadIdx.x];
		//printf("\n sdata[%d]: %d", i, d_odata[threadIdx.x]);
	}
	
}

///////////////////////////////////////////////////////////////////////////////
//                   Wrapper for countGlobalMem
///////////////////////////////////////////////////////////////////////////////
template<class T>
void countGMWrapper(int threads, int blocks, T *g_idata, int *g_odata ,int N) {
	/*
	1D block and 1D grid
	*/
	dim3 dimBlock(threads, 1, 1); 
	dim3 dimGrid(blocks, 1, 1);
	int smemSize = threads * sizeof (T);
	
	countGlobalMem<T, BLOCK_WIDTH><<<dimGrid, dimBlock, smemSize>>>(g_idata, g_odata, N);
}

///////////////////////////////////////////////////////////////////////////////
//                   Wrapper for countSharedMem
///////////////////////////////////////////////////////////////////////////////
template<class T>
void countSWrapper(int threads, int blocks, T *g_idata, int *g_odata ,int N, int maxNum, int barrelSize) {
	/*
	1D block and 1D grid
	*/
	dim3 dimBlock(threads, 1, 1); 
	dim3 dimGrid(blocks, 1, 1);
	int smemSize = threads * sizeof (T);
	
	countSharedMem<T, BLOCK_WIDTH><<<dimGrid, dimBlock, smemSize>>>(g_idata, g_odata, N, maxNum, barrelSize);
}

/////////////////////////////////////////////////////////////////////////////////
//                       Instantiate Template
/////////////////////////////////////////////////////////////////////////////////
template void
countGMWrapper<int>(int threads, int blocks, int *g_idata, int *g_odata, int N);

template void
countGMWrapper<float>(int threads, int blocks, float *g_idata, int *g_odata, int N);

template void
countGMWrapper<double>(int threads, int blocks, double *g_idata, int *g_odata, int N);

template void
countSWrapper<int>(int threads, int blocks, int *g_idata, int *g_odata ,int N, int maxNum, int barrelSize);

//////////////////////////////////////////////////////////////////////////////////
//						Test Function
//////////////////////////////////////////////////////////////////////////////////
void myCountTest(const char* filename) {
	int numBarrel = 10;
	//read test file and decide size of array
	std::vector<int> data;
    std::string line_;
    std::ifstream file_(filename);
    if(file_.is_open()) {
        while (getline(file_, line_)) {
            std::stringstream ss(line_);
            int i;
            while(ss>>i) {
                data.push_back(i);
                if (ss.peek() == ',' || ss.peek() == ' ') {
                    ss.ignore();
                }
            }
        }
        file_.close();
    } 
    int num_els = data.size();
	int numBlocks = num_els/BLOCK_WIDTH + 1;
	
	//Start to run Kernel_a
	int *d_in = NULL;
	int *d_out = NULL;
	cudaMalloc( (void **) &d_in, num_els * sizeof(int));
	cudaMalloc( (void **) &d_out, numBarrel * sizeof(int));
	int *in = (int *) malloc(num_els * sizeof(int));
	int *out = (int *) malloc(numBarrel * sizeof(int));	
	in = &data[0];
	std::vector<int> v(10);
	std::fill(v.begin(), v.end(), 0);
	std::copy(v.begin(), v.end(), out);
	cudaMemcpy(d_in, in, num_els * sizeof(int), cudaMemcpyHostToDevice);
	cudaMemcpy(d_out, out, numBarrel * sizeof(int), cudaMemcpyHostToDevice);
	countGMWrapper(BLOCK_WIDTH, numBlocks, d_in, d_out, num_els);
	cudaMemcpy(out, d_out, numBarrel * sizeof(int), cudaMemcpyDeviceToHost);
	std::ofstream fout1("q2a.txt", std::ios::app);
	for(int i = 0; i < numBarrel; i++) {
		if(fout1.is_open()) {
			fout1 << "\n Count[" <<i<<"]: " <<out[i];
		}
	}
	fout1.close();
	fout1.clear();
	cudaFree(d_out);
	//free(out);
	//d_in is not cleaned because we are going to run more cuda kernels using d_in
	free(in);
	//cudaFree(d_in);
	
	//Start to run Kernel_b, almost the same as kernel_a
	int *d_out_b = NULL;
	cudaMalloc( (void **) &d_out_b, numBarrel * numBlocks * sizeof(int));
	int *out_b = (int *) malloc(numBarrel * numBlocks * sizeof(int));
	//size of out_b is changed
	v.resize(numBarrel * numBlocks);
	std::fill(v.begin(), v.end(), 0);
	std::copy(v.begin(), v.end(), out_b);
	cudaMemcpy(d_out_b, out_b, numBarrel * numBlocks  * sizeof(int), cudaMemcpyHostToDevice);
	countSWrapper(BLOCK_WIDTH, numBlocks, d_in, d_out ,num_els, 1000, 100);
	cudaMemcpy(out_b, d_out_b, numBarrel * numBlocks * sizeof(int), cudaMemcpyDeviceToHost);
	std::ofstream fout2("q2b.txt", std::ios::app);
	int out_b_all;
	//int B[numBarrel];
	for(int i = 0; i < numBarrel; i++) {
		out_b_all = 0;
		for (int j = 0; j < numBlocks; j++)
			out_b_all += out_b[i + j * numBarrel];
		//B[i] = out_b_all;
		if(fout2.is_open()) {
			fout2 << "\n Count[" <<i<<"]: " <<out_b_all;
		}
	}
	fout2.close();
	fout2.clear();
	cudaFree(d_out_b);
	free(out_b);
	cudaFree(d_in);
	
	//start to run Kernel_c
	int n3 = nextPowerOf2(numBarrel);
	int *d_out_c = NULL;
	int *d_in_c = NULL;
	int *out_c = (int *) malloc(n3 * sizeof(int));
	v.resize(n3);
	std::fill(v.begin(), v.end(), 0);
	std::copy(v.begin(), v.end(), out_c);
	cudaMalloc( (void **) &d_in_c, n3 * sizeof(int));
	cudaMalloc( (void **) &d_out_c, n3 * sizeof(int));
	
	int *in_test = (int *) malloc(n3 * sizeof(int));
	std::vector<int> in_c;
	for (int i = 0; i < n3; i++) {
		if (i < numBarrel) {
			in_c.push_back(out[i]);
		} else {
			in_c.push_back(0);
		}
		//printf("\n c: %d", in_c[i]);
	}
	in_test = &in_c[0];
	cudaMemcpy(d_in_c, in_test, n3 * sizeof(int), cudaMemcpyHostToDevice);
	cudaMemcpy(d_out_c, out_c, n3 * sizeof(int), cudaMemcpyHostToDevice);
	scan<<<1, n3, n3*sizeof(int)>>>(d_in_c, d_out_c, n3);
	cudaMemcpy(out_c, d_out_c, n3 * sizeof(int), cudaMemcpyDeviceToHost);
	std::ofstream fout3("q2c.txt", std::ios::app);
	for(int i = 0; i < numBarrel; i++) {
		if(fout3.is_open()) {
			fout3 << "\n prescan[" <<i<<"]: " <<out_c[i];
		}
	}
	fout3.close();
	fout3.clear();
}

int main(int argc, char **argv) {
	myCountTest("inp.txt");
	return 0;
}

	.file	"tmpxft_00224826_00000000-6_prefix.cudafe1.cpp"
	.text
#APP
	.globl _ZSt21ios_base_library_initv

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB4210:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE4210:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.align 2
	.type	_ZNSt12_Vector_baseIiSaIiEE11_M_allocateEm.isra.0, @function
_ZNSt12_Vector_baseIiSaIiEE11_M_allocateEm.isra.0:
.LFB5112:
	.cfi_startproc
	testq	%rdi, %rdi
	je	.L3
	movabsq	$2305843009213693951, %rax
	cmpq	%rdi, %rax
	jnb	.L4
	pushq	%rax
	.cfi_def_cfa_offset 16
	movabsq	$4611686018427387903, %rax
	cmpq	%rdi, %rax
	jnb	.L5
	call	_ZSt28__throw_bad_array_new_lengthv@PLT
.L5:
	call	_ZSt17__throw_bad_allocv@PLT
.L4:
	.cfi_def_cfa_offset 8
	salq	$2, %rdi
	jmp	_Znwm@PLT
.L3:
	xorl	%eax, %eax
	ret
	.cfi_endproc
.LFE5112:
	.size	_ZNSt12_Vector_baseIiSaIiEE11_M_allocateEm.isra.0, .-_ZNSt12_Vector_baseIiSaIiEE11_M_allocateEm.isra.0
	.type	_ZSt4copyIN9__gnu_cxx17__normal_iteratorIPiSt6vectorIiSaIiEEEES2_ET0_T_S8_S7_.isra.0, @function
_ZSt4copyIN9__gnu_cxx17__normal_iteratorIPiSt6vectorIiSaIiEEEES2_ET0_T_S8_S7_.isra.0:
.LFB5116:
	.cfi_startproc
	movq	%rdi, %rax
	movq	%rdx, %rdi
	subq	%rax, %rsi
	movq	%rsi, %rdx
	cmpq	$4, %rsi
	jle	.L11
	movq	%rax, %rsi
	jmp	memmove@PLT
.L11:
	jne	.L10
	movl	(%rax), %eax
	movl	%eax, (%rdi)
.L10:
	ret
	.cfi_endproc
.LFE5116:
	.size	_ZSt4copyIN9__gnu_cxx17__normal_iteratorIPiSt6vectorIiSaIiEEEES2_ET0_T_S8_S7_.isra.0, .-_ZSt4copyIN9__gnu_cxx17__normal_iteratorIPiSt6vectorIiSaIiEEEES2_ET0_T_S8_S7_.isra.0
	.type	_ZL48__device_stub__Z14countGlobalMemIiLi512EEvPT_PiiPiS_i, @function
_ZL48__device_stub__Z14countGlobalMemIiLi512EEvPT_PiiPiS_i:
.LFB4234:
	.cfi_startproc
	subq	$136, %rsp
	.cfi_def_cfa_offset 144
	movq	%rdi, 24(%rsp)
	leaq	40(%rsp), %rcx
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movl	%edx, 12(%rsp)
	leaq	32(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 120(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 96(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 112(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L13
	pushq	40(%rsp)
	.cfi_def_cfa_offset 152
	leaq	_Z14countGlobalMemIiLi512EEvPT_Pii(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 160
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	112(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 152
	popq	%rdx
	.cfi_def_cfa_offset 144
.L13:
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L15
	call	__stack_chk_fail@PLT
.L15:
	addq	$136, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4234:
	.size	_ZL48__device_stub__Z14countGlobalMemIiLi512EEvPT_PiiPiS_i, .-_ZL48__device_stub__Z14countGlobalMemIiLi512EEvPT_PiiPiS_i
	.section	.text._Z14countGlobalMemIiLi512EEvPT_Pii,"axG",@progbits,_Z14countGlobalMemIiLi512EEvPT_Pii,comdat
	.weak	_Z14countGlobalMemIiLi512EEvPT_Pii
	.type	_Z14countGlobalMemIiLi512EEvPT_Pii, @function
_Z14countGlobalMemIiLi512EEvPT_Pii:
.LFB4627:
	.cfi_startproc
	endbr64
	jmp	_ZL48__device_stub__Z14countGlobalMemIiLi512EEvPT_PiiPiS_i
	.cfi_endproc
.LFE4627:
	.size	_Z14countGlobalMemIiLi512EEvPT_Pii, .-_Z14countGlobalMemIiLi512EEvPT_Pii
	.text
	.type	_ZL48__device_stub__Z14countGlobalMemIfLi512EEvPT_PiiPfPii, @function
_ZL48__device_stub__Z14countGlobalMemIfLi512EEvPT_PiiPfPii:
.LFB4236:
	.cfi_startproc
	subq	$136, %rsp
	.cfi_def_cfa_offset 144
	movq	%rdi, 24(%rsp)
	leaq	40(%rsp), %rcx
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movl	%edx, 12(%rsp)
	leaq	32(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 120(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 96(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 112(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L18
	pushq	40(%rsp)
	.cfi_def_cfa_offset 152
	leaq	_Z14countGlobalMemIfLi512EEvPT_Pii(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 160
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	112(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 152
	popq	%rdx
	.cfi_def_cfa_offset 144
.L18:
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L20
	call	__stack_chk_fail@PLT
.L20:
	addq	$136, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4236:
	.size	_ZL48__device_stub__Z14countGlobalMemIfLi512EEvPT_PiiPfPii, .-_ZL48__device_stub__Z14countGlobalMemIfLi512EEvPT_PiiPfPii
	.section	.text._Z14countGlobalMemIfLi512EEvPT_Pii,"axG",@progbits,_Z14countGlobalMemIfLi512EEvPT_Pii,comdat
	.weak	_Z14countGlobalMemIfLi512EEvPT_Pii
	.type	_Z14countGlobalMemIfLi512EEvPT_Pii, @function
_Z14countGlobalMemIfLi512EEvPT_Pii:
.LFB4628:
	.cfi_startproc
	endbr64
	jmp	_ZL48__device_stub__Z14countGlobalMemIfLi512EEvPT_PiiPfPii
	.cfi_endproc
.LFE4628:
	.size	_Z14countGlobalMemIfLi512EEvPT_Pii, .-_Z14countGlobalMemIfLi512EEvPT_Pii
	.text
	.type	_ZL48__device_stub__Z14countGlobalMemIdLi512EEvPT_PiiPdPii, @function
_ZL48__device_stub__Z14countGlobalMemIdLi512EEvPT_PiiPdPii:
.LFB4238:
	.cfi_startproc
	subq	$136, %rsp
	.cfi_def_cfa_offset 144
	movq	%rdi, 24(%rsp)
	leaq	40(%rsp), %rcx
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movl	%edx, 12(%rsp)
	leaq	32(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 120(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 96(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 112(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L23
	pushq	40(%rsp)
	.cfi_def_cfa_offset 152
	leaq	_Z14countGlobalMemIdLi512EEvPT_Pii(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 160
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	112(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 152
	popq	%rdx
	.cfi_def_cfa_offset 144
.L23:
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L25
	call	__stack_chk_fail@PLT
.L25:
	addq	$136, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4238:
	.size	_ZL48__device_stub__Z14countGlobalMemIdLi512EEvPT_PiiPdPii, .-_ZL48__device_stub__Z14countGlobalMemIdLi512EEvPT_PiiPdPii
	.section	.text._Z14countGlobalMemIdLi512EEvPT_Pii,"axG",@progbits,_Z14countGlobalMemIdLi512EEvPT_Pii,comdat
	.weak	_Z14countGlobalMemIdLi512EEvPT_Pii
	.type	_Z14countGlobalMemIdLi512EEvPT_Pii, @function
_Z14countGlobalMemIdLi512EEvPT_Pii:
.LFB4629:
	.cfi_startproc
	endbr64
	jmp	_ZL48__device_stub__Z14countGlobalMemIdLi512EEvPT_PiiPdPii
	.cfi_endproc
.LFE4629:
	.size	_Z14countGlobalMemIdLi512EEvPT_Pii, .-_Z14countGlobalMemIdLi512EEvPT_Pii
	.text
	.type	_ZL50__device_stub__Z14countSharedMemIiLi512EEvPT_PiiiiPiS_iii, @function
_ZL50__device_stub__Z14countSharedMemIiLi512EEvPT_PiiiiPiS_iii:
.LFB4240:
	.cfi_startproc
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movq	%rdi, 24(%rsp)
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movl	%edx, 12(%rsp)
	leaq	32(%rsp), %rdx
	movl	%ecx, 8(%rsp)
	leaq	40(%rsp), %rcx
	movl	%r8d, 4(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 96(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 128(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L28
	pushq	40(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z14countSharedMemIiLi512EEvPT_Piiii(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 176
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	112(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 168
	popq	%rdx
	.cfi_def_cfa_offset 160
.L28:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L30
	call	__stack_chk_fail@PLT
.L30:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4240:
	.size	_ZL50__device_stub__Z14countSharedMemIiLi512EEvPT_PiiiiPiS_iii, .-_ZL50__device_stub__Z14countSharedMemIiLi512EEvPT_PiiiiPiS_iii
	.section	.text._Z14countSharedMemIiLi512EEvPT_Piiii,"axG",@progbits,_Z14countSharedMemIiLi512EEvPT_Piiii,comdat
	.weak	_Z14countSharedMemIiLi512EEvPT_Piiii
	.type	_Z14countSharedMemIiLi512EEvPT_Piiii, @function
_Z14countSharedMemIiLi512EEvPT_Piiii:
.LFB4630:
	.cfi_startproc
	endbr64
	jmp	_ZL50__device_stub__Z14countSharedMemIiLi512EEvPT_PiiiiPiS_iii
	.cfi_endproc
.LFE4630:
	.size	_Z14countSharedMemIiLi512EEvPT_Piiii, .-_Z14countSharedMemIiLi512EEvPT_Piiii
	.text
	.globl	_Z12nextPowerOf2i
	.type	_Z12nextPowerOf2i, @function
_Z12nextPowerOf2i:
.LFB4192:
	.cfi_startproc
	endbr64
	decl	%edi
	movl	%edi, %eax
	sarl	%eax
	orl	%edi, %eax
	movl	%eax, %edi
	sarl	$2, %edi
	orl	%eax, %edi
	movl	%edi, %eax
	sarl	$4, %eax
	orl	%edi, %eax
	movl	%eax, %edx
	sarl	$8, %edx
	orl	%eax, %edx
	movl	%edx, %eax
	sarl	$16, %eax
	orl	%edx, %eax
	incl	%eax
	ret
	.cfi_endproc
.LFE4192:
	.size	_Z12nextPowerOf2i, .-_Z12nextPowerOf2i
	.section	.text._Z14countGMWrapperIiEviiPT_Pii,"axG",@progbits,_Z14countGMWrapperIiEviiPT_Pii,comdat
	.weak	_Z14countGMWrapperIiEviiPT_Pii
	.type	_Z14countGMWrapperIiEviiPT_Pii, @function
_Z14countGMWrapperIiEviiPT_Pii:
.LFB4559:
	.cfi_startproc
	endbr64
	movabsq	$4294967297, %rax
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	movl	%r8d, %r12d
	xorl	%r9d, %r9d
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	movq	%rcx, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	movq	%rdx, %rbx
	subq	$32, %rsp
	.cfi_def_cfa_offset 64
	movl	%edi, 8(%rsp)
	sall	$2, %edi
	movq	%rax, 12(%rsp)
	movslq	%edi, %r8
	movl	16(%rsp), %ecx
	movl	%esi, 20(%rsp)
	movq	8(%rsp), %rdx
	movq	%rax, 24(%rsp)
	movq	20(%rsp), %rdi
	movl	28(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L34
	addq	$32, %rsp
	.cfi_remember_state
	.cfi_def_cfa_offset 32
	movl	%r12d, %edx
	movq	%rbp, %rsi
	movq	%rbx, %rdi
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	jmp	_ZL48__device_stub__Z14countGlobalMemIiLi512EEvPT_PiiPiS_i
.L34:
	.cfi_restore_state
	addq	$32, %rsp
	.cfi_def_cfa_offset 32
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4559:
	.size	_Z14countGMWrapperIiEviiPT_Pii, .-_Z14countGMWrapperIiEviiPT_Pii
	.section	.text._Z14countGMWrapperIfEviiPT_Pii,"axG",@progbits,_Z14countGMWrapperIfEviiPT_Pii,comdat
	.weak	_Z14countGMWrapperIfEviiPT_Pii
	.type	_Z14countGMWrapperIfEviiPT_Pii, @function
_Z14countGMWrapperIfEviiPT_Pii:
.LFB4560:
	.cfi_startproc
	endbr64
	movabsq	$4294967297, %rax
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	movl	%r8d, %r12d
	xorl	%r9d, %r9d
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	movq	%rcx, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	movq	%rdx, %rbx
	subq	$32, %rsp
	.cfi_def_cfa_offset 64
	movl	%edi, 8(%rsp)
	sall	$2, %edi
	movq	%rax, 12(%rsp)
	movslq	%edi, %r8
	movl	16(%rsp), %ecx
	movl	%esi, 20(%rsp)
	movq	8(%rsp), %rdx
	movq	%rax, 24(%rsp)
	movq	20(%rsp), %rdi
	movl	28(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L37
	addq	$32, %rsp
	.cfi_remember_state
	.cfi_def_cfa_offset 32
	movl	%r12d, %edx
	movq	%rbp, %rsi
	movq	%rbx, %rdi
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	jmp	_ZL48__device_stub__Z14countGlobalMemIfLi512EEvPT_PiiPfPii
.L37:
	.cfi_restore_state
	addq	$32, %rsp
	.cfi_def_cfa_offset 32
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4560:
	.size	_Z14countGMWrapperIfEviiPT_Pii, .-_Z14countGMWrapperIfEviiPT_Pii
	.section	.text._Z14countGMWrapperIdEviiPT_Pii,"axG",@progbits,_Z14countGMWrapperIdEviiPT_Pii,comdat
	.weak	_Z14countGMWrapperIdEviiPT_Pii
	.type	_Z14countGMWrapperIdEviiPT_Pii, @function
_Z14countGMWrapperIdEviiPT_Pii:
.LFB4561:
	.cfi_startproc
	endbr64
	movabsq	$4294967297, %rax
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	movl	%r8d, %r12d
	xorl	%r9d, %r9d
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	movq	%rcx, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	movq	%rdx, %rbx
	subq	$32, %rsp
	.cfi_def_cfa_offset 64
	movl	%edi, 8(%rsp)
	sall	$3, %edi
	movq	%rax, 12(%rsp)
	movslq	%edi, %r8
	movl	16(%rsp), %ecx
	movl	%esi, 20(%rsp)
	movq	8(%rsp), %rdx
	movq	%rax, 24(%rsp)
	movq	20(%rsp), %rdi
	movl	28(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L40
	addq	$32, %rsp
	.cfi_remember_state
	.cfi_def_cfa_offset 32
	movl	%r12d, %edx
	movq	%rbp, %rsi
	movq	%rbx, %rdi
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	jmp	_ZL48__device_stub__Z14countGlobalMemIdLi512EEvPT_PiiPdPii
.L40:
	.cfi_restore_state
	addq	$32, %rsp
	.cfi_def_cfa_offset 32
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4561:
	.size	_Z14countGMWrapperIdEviiPT_Pii, .-_Z14countGMWrapperIdEviiPT_Pii
	.section	.text._Z13countSWrapperIiEviiPT_Piiii,"axG",@progbits,_Z13countSWrapperIiEviiPT_Piiii,comdat
	.weak	_Z13countSWrapperIiEviiPT_Piiii
	.type	_Z13countSWrapperIiEviiPT_Piiii, @function
_Z13countSWrapperIiEviiPT_Piiii:
.LFB4562:
	.cfi_startproc
	endbr64
	pushq	%r14
	.cfi_def_cfa_offset 16
	.cfi_offset 14, -16
	pushq	%r13
	.cfi_def_cfa_offset 24
	.cfi_offset 13, -24
	movl	%r9d, %r13d
	xorl	%r9d, %r9d
	pushq	%r12
	.cfi_def_cfa_offset 32
	.cfi_offset 12, -32
	movl	%r8d, %r12d
	leal	0(,%rdi,4), %r8d
	pushq	%rbp
	.cfi_def_cfa_offset 40
	.cfi_offset 6, -40
	movq	%rcx, %rbp
	movslq	%r8d, %r8
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset 3, -48
	movq	%rdx, %rbx
	movabsq	$4294967297, %rdx
	subq	$32, %rsp
	.cfi_def_cfa_offset 80
	movq	%rdx, 24(%rsp)
	movl	80(%rsp), %r14d
	movl	%edi, 8(%rsp)
	movq	%rdx, 12(%rsp)
	movl	16(%rsp), %ecx
	movl	%esi, 20(%rsp)
	movq	8(%rsp), %rdx
	movq	20(%rsp), %rdi
	movl	28(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L43
	addq	$32, %rsp
	.cfi_remember_state
	.cfi_def_cfa_offset 48
	movl	%r14d, %r8d
	movl	%r13d, %ecx
	movl	%r12d, %edx
	movq	%rbp, %rsi
	movq	%rbx, %rdi
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%rbp
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r13
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	jmp	_ZL50__device_stub__Z14countSharedMemIiLi512EEvPT_PiiiiPiS_iii
.L43:
	.cfi_restore_state
	addq	$32, %rsp
	.cfi_def_cfa_offset 48
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%rbp
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r13
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4562:
	.size	_Z13countSWrapperIiEviiPT_Piiii, .-_Z13countSWrapperIiEviiPT_Piiii
	.text
	.globl	_Z26__device_stub__Z4scanPiS_iPiS_i
	.type	_Z26__device_stub__Z4scanPiS_iPiS_i, @function
_Z26__device_stub__Z4scanPiS_iPiS_i:
.LFB4232:
	.cfi_startproc
	endbr64
	subq	$136, %rsp
	.cfi_def_cfa_offset 144
	movq	%rdi, 24(%rsp)
	leaq	40(%rsp), %rcx
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movl	%edx, 12(%rsp)
	leaq	32(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 120(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 96(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 112(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L46
	pushq	40(%rsp)
	.cfi_def_cfa_offset 152
	leaq	_Z4scanPiS_i(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 160
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	112(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 152
	popq	%rdx
	.cfi_def_cfa_offset 144
.L46:
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L48
	call	__stack_chk_fail@PLT
.L48:
	addq	$136, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4232:
	.size	_Z26__device_stub__Z4scanPiS_iPiS_i, .-_Z26__device_stub__Z4scanPiS_iPiS_i
	.globl	_Z4scanPiS_i
	.type	_Z4scanPiS_i, @function
_Z4scanPiS_i:
.LFB4233:
	.cfi_startproc
	endbr64
	jmp	_Z26__device_stub__Z4scanPiS_iPiS_i
	.cfi_endproc
.LFE4233:
	.size	_Z4scanPiS_i, .-_Z4scanPiS_i
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"_Z14countSharedMemIiLi512EEvPT_Piiii"
.LC1:
	.string	"_Z14countGlobalMemIdLi512EEvPT_Pii"
.LC2:
	.string	"_Z14countGlobalMemIfLi512EEvPT_Pii"
.LC3:
	.string	"_Z14countGlobalMemIiLi512EEvPT_Pii"
.LC4:
	.string	"_Z4scanPiS_i"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB4243:
	.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	.LC0(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z14countSharedMemIiLi512EEvPT_Piiii(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC1(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z14countGlobalMemIdLi512EEvPT_Pii(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC2(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z14countGlobalMemIfLi512EEvPT_Pii(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC3(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z14countGlobalMemIiLi512EEvPT_Pii(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC4(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z4scanPiS_i(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE4243:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.text._ZNSt6vectorIiSaIiEED2Ev,"axG",@progbits,_ZNSt6vectorIiSaIiEED5Ev,comdat
	.align 2
	.weak	_ZNSt6vectorIiSaIiEED2Ev
	.type	_ZNSt6vectorIiSaIiEED2Ev, @function
_ZNSt6vectorIiSaIiEED2Ev:
.LFB4573:
	.cfi_startproc
	endbr64
	movq	(%rdi), %rax
	testq	%rax, %rax
	je	.L53
	movq	16(%rdi), %rsi
	movq	%rax, %rdi
	subq	%rax, %rsi
	jmp	_ZdlPvm@PLT
.L53:
	ret
	.cfi_endproc
.LFE4573:
	.size	_ZNSt6vectorIiSaIiEED2Ev, .-_ZNSt6vectorIiSaIiEED2Ev
	.weak	_ZNSt6vectorIiSaIiEED1Ev
	.set	_ZNSt6vectorIiSaIiEED1Ev,_ZNSt6vectorIiSaIiEED2Ev
	.section	.text._ZNKSt6vectorIiSaIiEE12_M_check_lenEmPKc,"axG",@progbits,_ZNKSt6vectorIiSaIiEE12_M_check_lenEmPKc,comdat
	.align 2
	.weak	_ZNKSt6vectorIiSaIiEE12_M_check_lenEmPKc
	.type	_ZNKSt6vectorIiSaIiEE12_M_check_lenEmPKc, @function
_ZNKSt6vectorIiSaIiEE12_M_check_lenEmPKc:
.LFB4949:
	.cfi_startproc
	endbr64
	movq	%rdi, %rax
	movq	%rdx, %rdi
	movq	8(%rax), %rcx
	subq	(%rax), %rcx
	movabsq	$2305843009213693951, %rax
	sarq	$2, %rcx
	movq	%rax, %rdx
	subq	%rcx, %rdx
	cmpq	%rsi, %rdx
	jnb	.L56
	pushq	%rax
	.cfi_def_cfa_offset 16
	call	_ZSt20__throw_length_errorPKc@PLT
.L56:
	.cfi_def_cfa_offset 8
	cmpq	%rcx, %rsi
	cmovb	%rcx, %rsi
	addq	%rsi, %rcx
	jc	.L55
	movabsq	$2305843009213693951, %rax
	cmpq	%rax, %rcx
	cmovbe	%rcx, %rax
.L55:
	ret
	.cfi_endproc
.LFE4949:
	.size	_ZNKSt6vectorIiSaIiEE12_M_check_lenEmPKc, .-_ZNKSt6vectorIiSaIiEE12_M_check_lenEmPKc
	.section	.rodata._ZNSt6vectorIiSaIiEE9push_backERKi.str1.1,"aMS",@progbits,1
.LC5:
	.string	"vector::_M_realloc_insert"
	.section	.text._ZNSt6vectorIiSaIiEE9push_backERKi,"axG",@progbits,_ZNSt6vectorIiSaIiEE9push_backERKi,comdat
	.align 2
	.weak	_ZNSt6vectorIiSaIiEE9push_backERKi
	.type	_ZNSt6vectorIiSaIiEE9push_backERKi, @function
_ZNSt6vectorIiSaIiEE9push_backERKi:
.LFB4599:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	movq	%rdi, %r14
	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	$24, %rsp
	.cfi_def_cfa_offset 80
	movq	8(%rdi), %rbx
	cmpq	16(%rdi), %rbx
	je	.L65
	movl	(%rsi), %eax
	addq	$4, %rbx
	movl	%eax, -4(%rbx)
	movq	%rbx, 8(%rdi)
	jmp	.L64
.L65:
	movq	%rsi, 8(%rsp)
	leaq	.LC5(%rip), %rdx
	movl	$1, %esi
	movq	%rbx, %r15
	call	_ZNKSt6vectorIiSaIiEE12_M_check_lenEmPKc
	movq	(%r14), %r12
	movq	8(%r14), %r13
	movq	%rax, %rdi
	movq	%rax, (%rsp)
	call	_ZNSt12_Vector_baseIiSaIiEE11_M_allocateEm.isra.0
	movq	8(%rsp), %r8
	subq	%r12, %r15
	movq	%rax, %rbp
	movl	(%r8), %eax
	movl	%eax, 0(%rbp,%r15)
	testq	%r15, %r15
	jle	.L67
	movq	%r15, %rdx
	movq	%r12, %rsi
	movq	%rbp, %rdi
	call	memmove@PLT
.L67:
	subq	%rbx, %r13
	leaq	4(%rbp,%r15), %rdi
	testq	%r13, %r13
	jle	.L68
	movq	%r13, %rdx
	movq	%rbx, %rsi
	call	memmove@PLT
	movq	%rax, %rdi
.L68:
	addq	%rdi, %r13
	testq	%r12, %r12
	je	.L69
	movq	16(%r14), %rsi
	movq	%r12, %rdi
	subq	%r12, %rsi
	call	_ZdlPvm@PLT
.L69:
	movq	(%rsp), %rax
	movq	%rbp, (%r14)
	movq	%r13, 8(%r14)
	leaq	0(%rbp,%rax,4), %rax
	movq	%rax, 16(%r14)
.L64:
	addq	$24, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4599:
	.size	_ZNSt6vectorIiSaIiEE9push_backERKi, .-_ZNSt6vectorIiSaIiEE9push_backERKi
	.section	.text._ZNSt27__uninitialized_default_n_1ILb1EE18__uninit_default_nIPimEET_S3_T0_,"axG",@progbits,_ZNSt27__uninitialized_default_n_1ILb1EE18__uninit_default_nIPimEET_S3_T0_,comdat
	.weak	_ZNSt27__uninitialized_default_n_1ILb1EE18__uninit_default_nIPimEET_S3_T0_
	.type	_ZNSt27__uninitialized_default_n_1ILb1EE18__uninit_default_nIPimEET_S3_T0_, @function
_ZNSt27__uninitialized_default_n_1ILb1EE18__uninit_default_nIPimEET_S3_T0_:
.LFB5075:
	.cfi_startproc
	endbr64
	movq	%rdi, %rax
	testq	%rsi, %rsi
	je	.L74
	xorl	%r8d, %r8d
	leaq	4(%rdi), %rax
	movl	%r8d, (%rdi)
	decq	%rsi
	je	.L74
	leaq	(%rax,%rsi,4), %rdx
	movq	%rax, %rcx
.L76:
	cmpq	%rcx, %rdx
	je	.L83
	xorl	%esi, %esi
	addq	$4, %rcx
	movl	%esi, -4(%rcx)
	jmp	.L76
.L83:
	subq	%rdi, %rdx
	leaq	-4(%rax,%rdx), %rax
.L74:
	ret
	.cfi_endproc
.LFE5075:
	.size	_ZNSt27__uninitialized_default_n_1ILb1EE18__uninit_default_nIPimEET_S3_T0_, .-_ZNSt27__uninitialized_default_n_1ILb1EE18__uninit_default_nIPimEET_S3_T0_
	.section	.rodata._ZNSt6vectorIiSaIiEE17_M_default_appendEm.str1.1,"aMS",@progbits,1
.LC6:
	.string	"vector::_M_default_append"
	.section	.text._ZNSt6vectorIiSaIiEE17_M_default_appendEm,"axG",@progbits,_ZNSt6vectorIiSaIiEE17_M_default_appendEm,comdat
	.align 2
	.weak	_ZNSt6vectorIiSaIiEE17_M_default_appendEm
	.type	_ZNSt6vectorIiSaIiEE17_M_default_appendEm, @function
_ZNSt6vectorIiSaIiEE17_M_default_appendEm:
.LFB4845:
	.cfi_startproc
	endbr64
	testq	%rsi, %rsi
	je	.L97
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	movq	%rsi, %r14
	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
	movq	%rdi, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	pushq	%rdx
	.cfi_def_cfa_offset 64
	movq	8(%rdi), %rdi
	movq	16(%rbp), %rax
	movq	0(%rbp), %r15
	subq	%rdi, %rax
	sarq	$2, %rax
	cmpq	%rsi, %rax
	jb	.L86
	call	_ZNSt27__uninitialized_default_n_1ILb1EE18__uninit_default_nIPimEET_S3_T0_
	movq	%rax, 8(%rbp)
	jmp	.L84
.L86:
	movq	%rdi, %rbx
	leaq	.LC6(%rip), %rdx
	movq	%rbp, %rdi
	call	_ZNKSt6vectorIiSaIiEE12_M_check_lenEmPKc
	subq	%r15, %rbx
	movq	%rax, %rdi
	movq	%rax, %r13
	call	_ZNSt12_Vector_baseIiSaIiEE11_M_allocateEm.isra.0
	movq	%r14, %rsi
	leaq	(%rax,%rbx), %rdi
	movq	%rax, %r12
	call	_ZNSt27__uninitialized_default_n_1ILb1EE18__uninit_default_nIPimEET_S3_T0_
	testq	%rbx, %rbx
	je	.L87
	movq	%rbx, %rdx
	movq	%r15, %rsi
	movq	%r12, %rdi
	call	memmove@PLT
	jmp	.L88
.L87:
	testq	%r15, %r15
	je	.L89
.L88:
	movq	16(%rbp), %rsi
	movq	%r15, %rdi
	subq	%r15, %rsi
	call	_ZdlPvm@PLT
.L89:
	sarq	$2, %rbx
	movq	%r12, 0(%rbp)
	addq	%r14, %rbx
	leaq	(%r12,%rbx,4), %rax
	movq	%rax, 8(%rbp)
	leaq	(%r12,%r13,4), %rax
	movq	%rax, 16(%rbp)
.L84:
	popq	%rax
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
.L97:
	.cfi_restore 3
	.cfi_restore 6
	.cfi_restore 12
	.cfi_restore 13
	.cfi_restore 14
	.cfi_restore 15
	ret
	.cfi_endproc
.LFE4845:
	.size	_ZNSt6vectorIiSaIiEE17_M_default_appendEm, .-_ZNSt6vectorIiSaIiEE17_M_default_appendEm
	.section	.text._ZNSt6vectorIiSaIiEE6resizeEm,"axG",@progbits,_ZNSt6vectorIiSaIiEE6resizeEm,comdat
	.align 2
	.weak	_ZNSt6vectorIiSaIiEE6resizeEm
	.type	_ZNSt6vectorIiSaIiEE6resizeEm, @function
_ZNSt6vectorIiSaIiEE6resizeEm:
.LFB4624:
	.cfi_startproc
	endbr64
	movq	8(%rdi), %rdx
	movq	(%rdi), %rcx
	movq	%rdx, %rax
	subq	%rcx, %rax
	sarq	$2, %rax
	cmpq	%rsi, %rax
	jnb	.L101
	subq	%rax, %rsi
	jmp	_ZNSt6vectorIiSaIiEE17_M_default_appendEm
.L101:
	cmpq	%rax, %rsi
	jnb	.L100
	leaq	(%rcx,%rsi,4), %rax
	cmpq	%rdx, %rax
	je	.L100
	movq	%rax, 8(%rdi)
.L100:
	ret
	.cfi_endproc
.LFE4624:
	.size	_ZNSt6vectorIiSaIiEE6resizeEm, .-_ZNSt6vectorIiSaIiEE6resizeEm
	.section	.rodata.str1.1
.LC7:
	.string	"q2a.txt"
.LC8:
	.string	"\n Count["
.LC9:
	.string	"]: "
.LC10:
	.string	"q2b.txt"
.LC11:
	.string	"q2c.txt"
.LC12:
	.string	"\n prescan["
	.text
	.globl	_Z11myCountTestPKc
	.type	_Z11myCountTestPKc, @function
_Z11myCountTestPKc:
.LFB4195:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA4195
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	xorl	%r10d, %r10d
	movq	%rdi, %rsi
	xorl	%r11d, %r11d
	movl	$8, %edx
	movq	%rsp, %rbp
	.cfi_def_cfa_register 6
	pushq	%r15
	pushq	%r14
	pushq	%r13
	pushq	%r12
	pushq	%rbx
	subq	$2312, %rsp
	.cfi_offset 15, -24
	.cfi_offset 14, -32
	.cfi_offset 13, -40
	.cfi_offset 12, -48
	.cfi_offset 3, -56
	movq	%fs:40, %rax
	movq	%rax, -56(%rbp)
	xorl	%eax, %eax
	leaq	-2128(%rbp), %rax
	movq	%r10, -2216(%rbp)
	movq	%rax, -2144(%rbp)
	leaq	-576(%rbp), %rax
	movq	%rax, %rdi
	movq	%r10, -2208(%rbp)
	movq	%r10, -2200(%rbp)
	movq	%r11, -2136(%rbp)
	movb	$0, -2128(%rbp)
	movq	%rax, -2312(%rbp)
.LEHB0:
	call	_ZNSt14basic_ifstreamIcSt11char_traitsIcEEC1EPKcSt13_Ios_Openmode@PLT
.LEHE0:
	leaq	-456(%rbp), %rdi
	leaq	-2144(%rbp), %rbx
	call	_ZNKSt12__basic_fileIcE7is_openEv@PLT
	testb	%al, %al
	je	.L105
.L104:
	movq	-2312(%rbp), %rdi
	movq	%rbx, %rsi
.LEHB1:
	call	_ZSt7getlineIcSt11char_traitsIcESaIcEERSt13basic_istreamIT_T0_ES7_RNSt7__cxx1112basic_stringIS4_S5_T1_EE@PLT
	movq	(%rax), %rdx
	movq	-24(%rdx), %rdx
	testb	$5, 32(%rax,%rdx)
	jne	.L175
	leaq	-1088(%rbp), %rax
	movl	$24, %edx
	movq	%rbx, %rsi
	movq	%rax, %rdi
	movq	%rax, %r14
	call	_ZNSt7__cxx1118basic_stringstreamIcSt11char_traitsIcESaIcEEC1ERKNS_12basic_stringIcS2_S3_EESt13_Ios_Openmode@PLT
.LEHE1:
	leaq	-2168(%rbp), %r12
.L106:
	movq	%r12, %rsi
	movq	%r14, %rdi
.LEHB2:
	call	_ZNSirsERi@PLT
	movq	(%rax), %rdx
	movq	-24(%rdx), %rdx
	testb	$5, 32(%rax,%rdx)
	jne	.L176
	leaq	-2216(%rbp), %rdi
	movq	%r12, %rsi
	call	_ZNSt6vectorIiSaIiEE9push_backERKi
	movq	%r14, %rdi
	call	_ZNSi4peekEv@PLT
	cmpl	$44, %eax
	jne	.L107
.L109:
	movq	%r14, %rdi
	call	_ZNSi6ignoreEv@PLT
	jmp	.L106
.L107:
	movq	%r14, %rdi
	call	_ZNSi4peekEv@PLT
.LEHE2:
	cmpl	$32, %eax
	jne	.L106
	jmp	.L109
.L176:
	movq	%r14, %rdi
	call	_ZNSt7__cxx1118basic_stringstreamIcSt11char_traitsIcESaIcEED1Ev@PLT
	jmp	.L104
.L175:
	movq	-2312(%rbp), %rdi
.LEHB3:
	call	_ZNSt14basic_ifstreamIcSt11char_traitsIcEE5closeEv@PLT
.L105:
	movq	-2216(%rbp), %r14
	movq	-2208(%rbp), %rbx
	movl	$512, %ecx
	xorl	%r9d, %r9d
	leaq	-2280(%rbp), %rdi
	movq	%r9, -2280(%rbp)
	subq	%r14, %rbx
	movq	%r9, -2272(%rbp)
	sarq	$2, %rbx
	movl	%ebx, %eax
	cltd
	idivl	%ecx
	leal	1(%rax), %r13d
	movslq	%ebx, %rax
	salq	$2, %rax
	movq	%rax, %rsi
	movq	%rax, -2296(%rbp)
	call	cudaMalloc@PLT
	leaq	-2272(%rbp), %rdi
	movl	$40, %esi
	call	cudaMalloc@PLT
	movl	$40, %edi
	call	malloc@PLT
	xorl	%r8d, %r8d
	movl	$40, %edi
	movq	%r8, -2184(%rbp)
	movq	%rax, %r12
	movq	%r8, -2176(%rbp)
	call	_Znwm@PLT
.LEHE3:
	movq	%rax, %r15
	movq	%rax, -2192(%rbp)
	movl	$10, %esi
	leaq	40(%rax), %rax
	movq	%r15, %rdi
	movq	%rax, -2176(%rbp)
	call	_ZNSt27__uninitialized_default_n_1ILb1EE18__uninit_default_nIPimEET_S3_T0_
	movq	%rax, -2184(%rbp)
	movq	%rax, %rsi
	movq	%r15, %rax
.L112:
	cmpq	%rax, %rsi
	je	.L177
	xorl	%edi, %edi
	addq	$4, %rax
	movl	%edi, -4(%rax)
	jmp	.L112
.L177:
	movq	%r12, %rdx
	movq	%r15, %rdi
	call	_ZSt4copyIN9__gnu_cxx17__normal_iteratorIPiSt6vectorIiSaIiEEEES2_ET0_T_S8_S7_.isra.0
	movq	-2296(%rbp), %rdx
	movl	$1, %ecx
	movq	%r14, %rsi
	movq	-2280(%rbp), %rdi
.LEHB4:
	call	cudaMemcpy@PLT
	movq	-2272(%rbp), %rdi
	movl	$1, %ecx
	movl	$40, %edx
	movq	%r12, %rsi
	call	cudaMemcpy@PLT
	movq	-2272(%rbp), %rcx
	movq	-2280(%rbp), %rdx
	movl	%ebx, %r8d
	movl	%r13d, %esi
	movl	$512, %edi
	call	_Z14countGMWrapperIiEviiPT_Pii
	movq	-2272(%rbp), %rsi
	movl	$2, %ecx
	movl	$40, %edx
	movq	%r12, %rdi
	call	cudaMemcpy@PLT
	leaq	-2112(%rbp), %rax
	movl	$1, %edx
	leaq	.LC7(%rip), %rsi
	movq	%rax, %rdi
	movq	%rax, -2320(%rbp)
	call	_ZNSt14basic_ofstreamIcSt11char_traitsIcEEC1EPKcSt13_Ios_Openmode@PLT
.LEHE4:
	xorl	%esi, %esi
	leaq	.LC8(%rip), %r15
	movq	%rsi, -2296(%rbp)
.L115:
	leaq	-2000(%rbp), %rdi
	call	_ZNKSt12__basic_fileIcE7is_openEv@PLT
	testb	%al, %al
	je	.L114
	movq	-2320(%rbp), %rdi
	movq	%r15, %rsi
.LEHB5:
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	-2296(%rbp), %esi
	movq	%rax, %rdi
	call	_ZNSolsEi@PLT
	movq	%rax, %rdi
	leaq	.LC9(%rip), %rsi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%rax, %rdi
	movq	-2296(%rbp), %rax
	movl	(%r12,%rax,4), %esi
	call	_ZNSolsEi@PLT
.L114:
	incq	-2296(%rbp)
	cmpq	$10, -2296(%rbp)
	jne	.L115
	movq	-2320(%rbp), %rdi
	call	_ZNSt14basic_ofstreamIcSt11char_traitsIcEE5closeEv@PLT
	leaq	-1864(%rbp), %rdi
	xorl	%esi, %esi
	call	_ZNSt9basic_iosIcSt11char_traitsIcEE5clearESt12_Ios_Iostate@PLT
	movq	-2272(%rbp), %rdi
	call	cudaFree@PLT
	imull	$10, %r13d, %r15d
	movq	%r14, %rdi
	call	free@PLT
	xorl	%ecx, %ecx
	leaq	-2264(%rbp), %rdi
	movq	%rcx, -2264(%rbp)
	movslq	%r15d, %r15
	leaq	0(,%r15,4), %r14
	movq	%r14, %rsi
	call	cudaMalloc@PLT
	movq	%r14, %rdi
	call	malloc@PLT
	movq	%r15, %rsi
	movq	%rax, -2296(%rbp)
	leaq	-2192(%rbp), %rax
	movq	%rax, %rdi
	movq	%rax, -2328(%rbp)
	call	_ZNSt6vectorIiSaIiEE6resizeEm
	movq	-2192(%rbp), %rdi
	movq	-2184(%rbp), %rsi
	movq	%rdi, %rax
.L116:
	cmpq	%rax, %rsi
	je	.L178
	xorl	%edx, %edx
	addq	$4, %rax
	movl	%edx, -4(%rax)
	jmp	.L116
.L178:
	movq	-2296(%rbp), %rdx
	call	_ZSt4copyIN9__gnu_cxx17__normal_iteratorIPiSt6vectorIiSaIiEEEES2_ET0_T_S8_S7_.isra.0
	movq	-2296(%rbp), %rsi
	movl	$1, %ecx
	movq	%r14, %rdx
	movq	-2264(%rbp), %rdi
	call	cudaMemcpy@PLT
	pushq	%r15
	movq	-2272(%rbp), %rcx
	movl	%ebx, %r8d
	movl	%r13d, %esi
	pushq	$100
	movl	$1000, %r9d
	movl	$512, %edi
	movq	-2280(%rbp), %rdx
	.cfi_escape 0x2e,0x10
	call	_Z13countSWrapperIiEviiPT_Piiii
	movq	-2264(%rbp), %rsi
	popq	%r10
	movl	$2, %ecx
	movq	%r14, %rdx
	movq	-2296(%rbp), %rdi
	popq	%r11
	.cfi_escape 0x2e,0
	call	cudaMemcpy@PLT
	leaq	-1600(%rbp), %rax
	movl	$1, %edx
	leaq	.LC10(%rip), %rsi
	movq	%rax, %rdi
	movq	%rax, -2304(%rbp)
	call	_ZNSt14basic_ofstreamIcSt11char_traitsIcEEC1EPKcSt13_Ios_Openmode@PLT
.LEHE5:
	xorl	%ebx, %ebx
	leaq	.LC8(%rip), %r15
.L118:
	movq	-2296(%rbp), %rax
	xorl	%r14d, %r14d
	leaq	(%rax,%rbx,4), %rdx
	xorl	%eax, %eax
.L122:
	cmpl	%r13d, %eax
	jge	.L179
	addl	(%rdx), %r14d
	incl	%eax
	addq	$40, %rdx
	jmp	.L122
.L179:
	leaq	-1488(%rbp), %rdi
	call	_ZNKSt12__basic_fileIcE7is_openEv@PLT
	testb	%al, %al
	je	.L120
	movq	-2304(%rbp), %rdi
	movq	%r15, %rsi
.LEHB6:
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%rax, %rdi
	movl	%ebx, %esi
	call	_ZNSolsEi@PLT
	movq	%rax, %rdi
	leaq	.LC9(%rip), %rsi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%rax, %rdi
	movl	%r14d, %esi
	call	_ZNSolsEi@PLT
.L120:
	incq	%rbx
	cmpq	$10, %rbx
	jne	.L118
	movq	-2304(%rbp), %rdi
	call	_ZNSt14basic_ofstreamIcSt11char_traitsIcEE5closeEv@PLT
	leaq	-1352(%rbp), %rdi
	xorl	%esi, %esi
	call	_ZNSt9basic_iosIcSt11char_traitsIcEE5clearESt12_Ios_Iostate@PLT
	movq	-2264(%rbp), %rdi
	call	cudaFree@PLT
	movq	-2296(%rbp), %rdi
	call	free@PLT
	movq	-2280(%rbp), %rdi
	call	cudaFree@PLT
	xorl	%r9d, %r9d
	movl	$64, %edi
	movq	%r9, -2256(%rbp)
	movq	%r9, -2248(%rbp)
	call	malloc@PLT
	movq	-2328(%rbp), %rdi
	movl	$16, %esi
	movq	%rax, -2296(%rbp)
	call	_ZNSt6vectorIiSaIiEE6resizeEm
	movq	-2192(%rbp), %rdi
	movq	-2184(%rbp), %rsi
	movq	%rdi, %rax
.L123:
	cmpq	%rax, %rsi
	je	.L180
	xorl	%r8d, %r8d
	addq	$4, %rax
	movl	%r8d, -4(%rax)
	jmp	.L123
.L180:
	movq	-2296(%rbp), %rdx
	call	_ZSt4copyIN9__gnu_cxx17__normal_iteratorIPiSt6vectorIiSaIiEEEES2_ET0_T_S8_S7_.isra.0
	leaq	-2248(%rbp), %rdi
	movl	$64, %esi
	call	cudaMalloc@PLT
	leaq	-2256(%rbp), %rdi
	movl	$64, %esi
	call	cudaMalloc@PLT
.LEHE6:
	xorl	%ecx, %ecx
	leaq	-2168(%rbp), %rax
	movq	%r12, -2328(%rbp)
	xorl	%r13d, %r13d
	movq	%rcx, -2168(%rbp)
	movq	%rcx, -2160(%rbp)
	movq	%rcx, -2152(%rbp)
	movq	%rax, -2344(%rbp)
.L130:
	cmpl	$9, %r13d
	jg	.L125
	movq	-2328(%rbp), %rsi
	movq	-2344(%rbp), %rdi
.LEHB7:
	call	_ZNSt6vectorIiSaIiEE9push_backERKi
	jmp	.L126
.L125:
	movq	-2160(%rbp), %rbx
	movq	-2152(%rbp), %r15
	cmpq	%r15, %rbx
	je	.L127
	xorl	%edx, %edx
	addq	$4, %rbx
	movl	%edx, -4(%rbx)
	movq	%rbx, -2160(%rbp)
	jmp	.L126
.L127:
	leaq	-2168(%rbp), %rdi
	leaq	.LC5(%rip), %rdx
	movl	$1, %esi
	call	_ZNKSt6vectorIiSaIiEE12_M_check_lenEmPKc
	movq	-2168(%rbp), %r12
	movq	%rax, %rdi
	movq	%rax, -2336(%rbp)
	subq	%r12, %rbx
	call	_ZNSt12_Vector_baseIiSaIiEE11_M_allocateEm.isra.0
	movq	%rax, %r14
	xorl	%eax, %eax
	movl	%eax, (%r14,%rbx)
	testq	%rbx, %rbx
	jle	.L128
	movq	%rbx, %rdx
	movq	%r12, %rsi
	movq	%r14, %rdi
	call	memmove@PLT
.L128:
	leaq	4(%r14,%rbx), %rbx
	testq	%r12, %r12
	je	.L129
	movq	%r15, %rsi
	movq	%r12, %rdi
	subq	%r12, %rsi
	call	_ZdlPvm@PLT
.L129:
	movq	-2336(%rbp), %rax
	movq	%r14, -2168(%rbp)
	movq	%rbx, -2160(%rbp)
	leaq	(%r14,%rax,4), %rax
	movq	%rax, -2152(%rbp)
.L126:
	incl	%r13d
	addq	$4, -2328(%rbp)
	cmpl	$16, %r13d
	jne	.L130
	movq	-2168(%rbp), %rsi
	movl	$1, %ecx
	movl	$64, %edx
	movq	-2248(%rbp), %rdi
	call	cudaMemcpy@PLT
	movq	-2296(%rbp), %rsi
	movl	$1, %ecx
	movl	$64, %edx
	movq	-2256(%rbp), %rdi
	call	cudaMemcpy@PLT
	movl	$268435457, %edx
	xorl	%r9d, %r9d
	movl	$64, %r8d
	movl	$1, %ecx
	salq	$4, %rdx
	movl	$1, %esi
	movabsq	$4294967297, %rdi
	call	__cudaPushCallConfiguration@PLT
	movq	-2256(%rbp), %rsi
	testl	%eax, %eax
	jne	.L131
	movq	-2248(%rbp), %rdi
	movl	$16, %edx
	call	_Z26__device_stub__Z4scanPiS_iPiS_i
.L131:
	movq	-2256(%rbp), %rsi
	movl	$2, %ecx
	movl	$64, %edx
	movq	-2296(%rbp), %rdi
	call	cudaMemcpy@PLT
	leaq	-1088(%rbp), %rbx
	movl	$1, %edx
	leaq	.LC11(%rip), %rsi
	movq	%rbx, %rdi
	call	_ZNSt14basic_ofstreamIcSt11char_traitsIcEEC1EPKcSt13_Ios_Openmode@PLT
.LEHE7:
	xorl	%r12d, %r12d
	leaq	.LC12(%rip), %r13
.L133:
	leaq	-976(%rbp), %rdi
	call	_ZNKSt12__basic_fileIcE7is_openEv@PLT
	testb	%al, %al
	je	.L132
	movq	%r13, %rsi
	movq	%rbx, %rdi
.LEHB8:
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%rax, %rdi
	movl	%r12d, %esi
	call	_ZNSolsEi@PLT
	movq	%rax, %rdi
	leaq	.LC9(%rip), %rsi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%rax, %rdi
	movq	-2296(%rbp), %rax
	movl	(%rax,%r12,4), %esi
	call	_ZNSolsEi@PLT
.L132:
	incq	%r12
	cmpq	$10, %r12
	jne	.L133
	movq	%rbx, %rdi
	call	_ZNSt14basic_ofstreamIcSt11char_traitsIcEE5closeEv@PLT
	leaq	-840(%rbp), %rdi
	xorl	%esi, %esi
	call	_ZNSt9basic_iosIcSt11char_traitsIcEE5clearESt12_Ios_Iostate@PLT
.LEHE8:
	movq	%rbx, %rdi
	call	_ZNSt14basic_ofstreamIcSt11char_traitsIcEED1Ev@PLT
	leaq	-2168(%rbp), %rdi
	call	_ZNSt6vectorIiSaIiEED1Ev
	movq	-2304(%rbp), %rdi
	call	_ZNSt14basic_ofstreamIcSt11char_traitsIcEED1Ev@PLT
	leaq	-2112(%rbp), %rdi
	call	_ZNSt14basic_ofstreamIcSt11char_traitsIcEED1Ev@PLT
	leaq	-2192(%rbp), %rdi
	call	_ZNSt6vectorIiSaIiEED1Ev
	leaq	-576(%rbp), %rdi
	call	_ZNSt14basic_ifstreamIcSt11char_traitsIcEED1Ev@PLT
	leaq	-2144(%rbp), %rdi
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE10_M_disposeEv@PLT
	leaq	-2216(%rbp), %rdi
	call	_ZNSt6vectorIiSaIiEED1Ev
	movq	-56(%rbp), %rax
	subq	%fs:40, %rax
	je	.L143
	jmp	.L173
.L146:
	endbr64
	movq	%rax, %rbx
.L134:
	movq	%r14, %rdi
	call	_ZNSt7__cxx1118basic_stringstreamIcSt11char_traitsIcESaIcEED1Ev@PLT
	jmp	.L135
.L152:
	endbr64
	movq	%rax, %r15
.L136:
	movq	%rbx, %rdi
	movq	%r15, %rbx
	call	_ZNSt14basic_ofstreamIcSt11char_traitsIcEED1Ev@PLT
	jmp	.L137
.L151:
	endbr64
	movq	%rax, %rbx
.L137:
	leaq	-2168(%rbp), %rdi
	call	_ZNSt6vectorIiSaIiEED1Ev
	jmp	.L138
.L150:
	endbr64
	movq	%rax, %rbx
.L138:
	movq	-2304(%rbp), %rdi
	call	_ZNSt14basic_ofstreamIcSt11char_traitsIcEED1Ev@PLT
	jmp	.L139
.L149:
	endbr64
	movq	%rax, %rbx
.L139:
	movq	-2320(%rbp), %rdi
	call	_ZNSt14basic_ofstreamIcSt11char_traitsIcEED1Ev@PLT
	jmp	.L140
.L148:
	endbr64
	movq	%rax, %rbx
.L140:
	leaq	-2192(%rbp), %rdi
	call	_ZNSt6vectorIiSaIiEED1Ev
	jmp	.L135
.L147:
	endbr64
	movq	%rax, %rbx
.L135:
	movq	-2312(%rbp), %rdi
	call	_ZNSt14basic_ifstreamIcSt11char_traitsIcEED1Ev@PLT
	jmp	.L141
.L145:
	endbr64
	movq	%rax, %rbx
.L141:
	leaq	-2144(%rbp), %rdi
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE10_M_disposeEv@PLT
	leaq	-2216(%rbp), %rdi
	call	_ZNSt6vectorIiSaIiEED1Ev
	movq	-56(%rbp), %rax
	subq	%fs:40, %rax
	jne	.L173
	movq	%rbx, %rdi
.LEHB9:
	call	_Unwind_Resume@PLT
.LEHE9:
.L173:
	call	__stack_chk_fail@PLT
.L143:
	leaq	-40(%rbp), %rsp
	popq	%rbx
	popq	%r12
	popq	%r13
	popq	%r14
	popq	%r15
	popq	%rbp
	.cfi_def_cfa 7, 8
	ret
	.cfi_endproc
.LFE4195:
	.globl	__gxx_personality_v0
	.section	.gcc_except_table,"a",@progbits
.LLSDA4195:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE4195-.LLSDACSB4195
.LLSDACSB4195:
	.uleb128 .LEHB0-.LFB4195
	.uleb128 .LEHE0-.LEHB0
	.uleb128 .L145-.LFB4195
	.uleb128 0
	.uleb128 .LEHB1-.LFB4195
	.uleb128 .LEHE1-.LEHB1
	.uleb128 .L147-.LFB4195
	.uleb128 0
	.uleb128 .LEHB2-.LFB4195
	.uleb128 .LEHE2-.LEHB2
	.uleb128 .L146-.LFB4195
	.uleb128 0
	.uleb128 .LEHB3-.LFB4195
	.uleb128 .LEHE3-.LEHB3
	.uleb128 .L147-.LFB4195
	.uleb128 0
	.uleb128 .LEHB4-.LFB4195
	.uleb128 .LEHE4-.LEHB4
	.uleb128 .L148-.LFB4195
	.uleb128 0
	.uleb128 .LEHB5-.LFB4195
	.uleb128 .LEHE5-.LEHB5
	.uleb128 .L149-.LFB4195
	.uleb128 0
	.uleb128 .LEHB6-.LFB4195
	.uleb128 .LEHE6-.LEHB6
	.uleb128 .L150-.LFB4195
	.uleb128 0
	.uleb128 .LEHB7-.LFB4195
	.uleb128 .LEHE7-.LEHB7
	.uleb128 .L151-.LFB4195
	.uleb128 0
	.uleb128 .LEHB8-.LFB4195
	.uleb128 .LEHE8-.LEHB8
	.uleb128 .L152-.LFB4195
	.uleb128 0
	.uleb128 .LEHB9-.LFB4195
	.uleb128 .LEHE9-.LEHB9
	.uleb128 0
	.uleb128 0
.LLSDACSE4195:
	.text
	.size	_Z11myCountTestPKc, .-_Z11myCountTestPKc
	.section	.rodata.str1.1
.LC13:
	.string	"inp.txt"
	.section	.text.startup
	.globl	main
	.type	main, @function
main:
.LFB4207:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	leaq	.LC13(%rip), %rdi
	call	_Z11myCountTestPKc
	xorl	%eax, %eax
	popq	%rdx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4207:
	.size	main, .-main
	.section	.nvFatBinSegment,"aw"
	.align 8
	.type	_ZL15__fatDeviceText, @object
	.size	_ZL15__fatDeviceText, 24
_ZL15__fatDeviceText:
	.long	1180844977
	.long	1
	.quad	fatbinData
	.quad	0
	.local	_ZL20__cudaFatCubinHandle
	.comm	_ZL20__cudaFatCubinHandle,8,8
	.hidden	DW.ref.__gxx_personality_v0
	.weak	DW.ref.__gxx_personality_v0
	.section	.data.rel.local.DW.ref.__gxx_personality_v0,"awG",@progbits,DW.ref.__gxx_personality_v0,comdat
	.align 8
	.type	DW.ref.__gxx_personality_v0, @object
	.size	DW.ref.__gxx_personality_v0, 8
DW.ref.__gxx_personality_v0:
	.quad	__gxx_personality_v0
	.ident	"GCC: (Ubuntu 13.3.0-6ubuntu2~24.04) 13.3.0"
	.section	.note.GNU-stack,"",@progbits
	.section	.note.gnu.property,"a"
	.align 8
	.long	1f - 0f
	.long	4f - 1f
	.long	5
0:
	.string	"GNU"
1:
	.align 8
	.long	0xc0000002
	.long	3f - 2f
2:
	.long	0x3
3:
	.align 8
4:

#include <math.h>
#include <hip/hip_runtime.h>
#include <vector>
#include <fstream>
#include <string>
#include <sstream>
#include <iostream>

#define BLOCK_WIDTH 512
/*
extern __shared__ stands for shared memory on device, which has two "warps" of 32 threads.
Google CUDA shared memory and warps.
To replace extern __shared__ int __smem[]; which requires you to explicitly
know the data type is integer in advance. But input file could be int, float, or double.
Since we don't know the data type of shared meomry __smem[], we use
template<class T> where T stands for all possible data types. We also 
need to instantiate all possible data types later
In return (T *) __smem; it is data type conversion
Suggest to figure out difference between overload, override, redefine
*/
template<class T>
struct SharedMemory {
	__device__ inline operator T *() {
		extern __shared__ int __smem[];
		return (T *) __smem;
	}
	
	__device__ inline operator const T *() const {
		extern __shared__ int __smem[];
		return (T *) __smem;
	}
};

/////////////////////////////////////////////////////////////////////////////
//                   CUDA Kernel: Global memory
/////////////////////////////////////////////////////////////////////////////
template<class T, int blockSize>
__global__ void countGlobalMem(T *g_idata, int *g_odata, int N) {
	unsigned int i = blockSize * blockIdx.x + threadIdx.x;
	int gi = 0;
	if (i < N) {
		if (g_idata[i] == 1000) {
			atomicAdd(&g_odata[9], 1);
		} else {
			gi = (int) g_idata[i] / 100;
			atomicAdd(&g_odata[gi], 1);
		}
	}
}

/////////////////////////////////////////////////////////////////////////////
//                   CUDA Kernel: shared memory
/////////////////////////////////////////////////////////////////////////////
template<class T, int blockSize>
__global__ void countSharedMem(T *g_idata, int *g_odata, int N, int maxNum, int barrelSize) {
	/*
	Each block has a sdata
	*/
	extern __shared__ int sdata[];
	unsigned int tid = threadIdx.x;
	int numBarrel = maxNum/barrelSize;
	
	unsigned int i = blockSize * blockIdx.x + threadIdx.x;
	//gi is group/barrel index
	int gi = 0;
	if (i < N) {
		if (g_idata[i] == maxNum) {
			atomicAdd(&sdata[numBarrel-1], 1);
		} else {
			gi = (int) g_idata[i] / barrelSize;
			atomicAdd(&sdata[gi], 1);
		}
	}
	
	//wait until sdata[0~9] in all blocks are ready
	__syncthreads();
	/*
	every block has threadIdx.x from 0 to 511
	size of g_odata is numBarrel * blocks
	sum of all blocks is done in myCountTest(), note there
	is += when output to "q2b.txt"
	*/
	if (tid < numBarrel) {
		g_odata[blockIdx.x * numBarrel + tid] = sdata[tid];
	}	
}

//////////////////////////////////////////////////////////////////////////////
//                     CUDA Kernel: prefix sum (Naiive)
///////////////////////////////////////////////////////////////////////////////
int nextPowerOf2(int x)
{
    --x;
    x |= x >> 1;
    x |= x >> 2;
    x |= x >> 4;
    x |= x >> 8;
    x |= x >> 16;
    return ++x;
}

__global__ void scan(int *d_idata, int *d_odata, int N) {
	extern __shared__ int sdata[];
	
	//cunyi
	int i = blockIdx.x * blockDim.x + threadIdx.x;
	if (i < N) {
		sdata[threadIdx.x] = d_idata[i];
		//printf("\n sdata[%d]: %d", i, sdata[threadIdx.x]);
	}
	
	for (unsigned int stride = 1; stride <= threadIdx.x; stride *= 2) {
		__syncthreads();
		int in1 = sdata[threadIdx.x - stride];
		__syncthreads();
		sdata[threadIdx.x] += in1;
	}
	__syncthreads();
	if(i < N) {
		d_odata[threadIdx.x] = sdata[threadIdx.x];
		//printf("\n sdata[%d]: %d", i, d_odata[threadIdx.x]);
	}
	
}

///////////////////////////////////////////////////////////////////////////////
//                   Wrapper for countGlobalMem
///////////////////////////////////////////////////////////////////////////////
template<class T>
void countGMWrapper(int threads, int blocks, T *g_idata, int *g_odata ,int N) {
	/*
	1D block and 1D grid
	*/
	dim3 dimBlock(threads, 1, 1); 
	dim3 dimGrid(blocks, 1, 1);
	int smemSize = threads * sizeof (T);
	
	countGlobalMem<T, BLOCK_WIDTH><<<dimGrid, dimBlock, smemSize>>>(g_idata, g_odata, N);
}

///////////////////////////////////////////////////////////////////////////////
//                   Wrapper for countSharedMem
///////////////////////////////////////////////////////////////////////////////
template<class T>
void countSWrapper(int threads, int blocks, T *g_idata, int *g_odata ,int N, int maxNum, int barrelSize) {
	/*
	1D block and 1D grid
	*/
	dim3 dimBlock(threads, 1, 1); 
	dim3 dimGrid(blocks, 1, 1);
	int smemSize = threads * sizeof (T);
	
	countSharedMem<T, BLOCK_WIDTH><<<dimGrid, dimBlock, smemSize>>>(g_idata, g_odata, N, maxNum, barrelSize);
}

/////////////////////////////////////////////////////////////////////////////////
//                       Instantiate Template
/////////////////////////////////////////////////////////////////////////////////
template void
countGMWrapper<int>(int threads, int blocks, int *g_idata, int *g_odata, int N);

template void
countGMWrapper<float>(int threads, int blocks, float *g_idata, int *g_odata, int N);

template void
countGMWrapper<double>(int threads, int blocks, double *g_idata, int *g_odata, int N);

template void
countSWrapper<int>(int threads, int blocks, int *g_idata, int *g_odata ,int N, int maxNum, int barrelSize);

//////////////////////////////////////////////////////////////////////////////////
//						Test Function
//////////////////////////////////////////////////////////////////////////////////
void myCountTest(const char* filename) {
	int numBarrel = 10;
	//read test file and decide size of array
	std::vector<int> data;
    std::string line_;
    std::ifstream file_(filename);
    if(file_.is_open()) {
        while (getline(file_, line_)) {
            std::stringstream ss(line_);
            int i;
            while(ss>>i) {
                data.push_back(i);
                if (ss.peek() == ',' || ss.peek() == ' ') {
                    ss.ignore();
                }
            }
        }
        file_.close();
    } 
    int num_els = data.size();
	int numBlocks = num_els/BLOCK_WIDTH + 1;
	
	//Start to run Kernel_a
	int *d_in = NULL;
	int *d_out = NULL;
	hipMalloc( (void **) &d_in, num_els * sizeof(int));
	hipMalloc( (void **) &d_out, numBarrel * sizeof(int));
	int *in = (int *) malloc(num_els * sizeof(int));
	int *out = (int *) malloc(numBarrel * sizeof(int));	
	in = &data[0];
	std::vector<int> v(10);
	std::fill(v.begin(), v.end(), 0);
	std::copy(v.begin(), v.end(), out);
	hipMemcpy(d_in, in, num_els * sizeof(int), hipMemcpyHostToDevice);
	hipMemcpy(d_out, out, numBarrel * sizeof(int), hipMemcpyHostToDevice);
	countGMWrapper(BLOCK_WIDTH, numBlocks, d_in, d_out, num_els);
	hipMemcpy(out, d_out, numBarrel * sizeof(int), hipMemcpyDeviceToHost);
	std::ofstream fout1("q2a.txt", std::ios::app);
	for(int i = 0; i < numBarrel; i++) {
		if(fout1.is_open()) {
			fout1 << "\n Count[" <<i<<"]: " <<out[i];
		}
	}
	fout1.close();
	fout1.clear();
	hipFree(d_out);
	//free(out);
	//d_in is not cleaned because we are going to run more cuda kernels using d_in
	free(in);
	//cudaFree(d_in);
	
	//Start to run Kernel_b, almost the same as kernel_a
	int *d_out_b = NULL;
	hipMalloc( (void **) &d_out_b, numBarrel * numBlocks * sizeof(int));
	int *out_b = (int *) malloc(numBarrel * numBlocks * sizeof(int));
	//size of out_b is changed
	v.resize(numBarrel * numBlocks);
	std::fill(v.begin(), v.end(), 0);
	std::copy(v.begin(), v.end(), out_b);
	hipMemcpy(d_out_b, out_b, numBarrel * numBlocks  * sizeof(int), hipMemcpyHostToDevice);
	countSWrapper(BLOCK_WIDTH, numBlocks, d_in, d_out ,num_els, 1000, 100);
	hipMemcpy(out_b, d_out_b, numBarrel * numBlocks * sizeof(int), hipMemcpyDeviceToHost);
	std::ofstream fout2("q2b.txt", std::ios::app);
	int out_b_all;
	//int B[numBarrel];
	for(int i = 0; i < numBarrel; i++) {
		out_b_all = 0;
		for (int j = 0; j < numBlocks; j++)
			out_b_all += out_b[i + j * numBarrel];
		//B[i] = out_b_all;
		if(fout2.is_open()) {
			fout2 << "\n Count[" <<i<<"]: " <<out_b_all;
		}
	}
	fout2.close();
	fout2.clear();
	hipFree(d_out_b);
	free(out_b);
	hipFree(d_in);
	
	//start to run Kernel_c
	int n3 = nextPowerOf2(numBarrel);
	int *d_out_c = NULL;
	int *d_in_c = NULL;
	int *out_c = (int *) malloc(n3 * sizeof(int));
	v.resize(n3);
	std::fill(v.begin(), v.end(), 0);
	std::copy(v.begin(), v.end(), out_c);
	hipMalloc( (void **) &d_in_c, n3 * sizeof(int));
	hipMalloc( (void **) &d_out_c, n3 * sizeof(int));
	
	int *in_test = (int *) malloc(n3 * sizeof(int));
	std::vector<int> in_c;
	for (int i = 0; i < n3; i++) {
		if (i < numBarrel) {
			in_c.push_back(out[i]);
		} else {
			in_c.push_back(0);
		}
		//printf("\n c: %d", in_c[i]);
	}
	in_test = &in_c[0];
	hipMemcpy(d_in_c, in_test, n3 * sizeof(int), hipMemcpyHostToDevice);
	hipMemcpy(d_out_c, out_c, n3 * sizeof(int), hipMemcpyHostToDevice);
	scan<<<1, n3, n3*sizeof(int)>>>(d_in_c, d_out_c, n3);
	hipMemcpy(out_c, d_out_c, n3 * sizeof(int), hipMemcpyDeviceToHost);
	std::ofstream fout3("q2c.txt", std::ios::app);
	for(int i = 0; i < numBarrel; i++) {
		if(fout3.is_open()) {
			fout3 << "\n prescan[" <<i<<"]: " <<out_c[i];
		}
	}
	fout3.close();
	fout3.clear();
}

int main(int argc, char **argv) {
	myCountTest("inp.txt");
	return 0;
}

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

                                        # End of file scope inline assembly
	.globl	_Z12nextPowerOf2i               # -- Begin function _Z12nextPowerOf2i
	.type	_Z12nextPowerOf2i,@function
_Z12nextPowerOf2i:                      # @_Z12nextPowerOf2i
	.cfi_startproc
# %bb.0:
                                        # kill: def $edi killed $edi def $rdi
	leal	-1(%rdi), %eax
	movl	%eax, %ecx
	sarl	%ecx
	orl	%eax, %ecx
	movl	%ecx, %eax
	sarl	$2, %eax
	orl	%ecx, %eax
	movl	%eax, %ecx
	sarl	$4, %ecx
	orl	%eax, %ecx
	movl	%ecx, %edx
	sarl	$8, %edx
	orl	%ecx, %edx
	movl	%edx, %eax
	sarl	$16, %eax
	orl	%edx, %eax
	incl	%eax
	retq
.Lfunc_end0:
	.size	_Z12nextPowerOf2i, .Lfunc_end0-_Z12nextPowerOf2i
	.cfi_endproc
                                        # -- End function
	.globl	_Z19__device_stub__scanPiS_i    # -- Begin function _Z19__device_stub__scanPiS_i
	.type	_Z19__device_stub__scanPiS_i,@function
_Z19__device_stub__scanPiS_i:           # @_Z19__device_stub__scanPiS_i
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$112, %rsp
	.cfi_def_cfa_offset 160
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	40(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	32(%rsp), %rcx
	movq	%rsi, (%rcx)
	leaq	12(%rsp), %rsi
	movl	%edx, (%rsi)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%rbx)
	movq	%rsi, 16(%rbx)
	leaq	64(%rsp), %r14
	leaq	48(%rsp), %r15
	leaq	24(%rsp), %r12
	leaq	16(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z4scanPiS_i, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$128, %rsp
	.cfi_adjust_cfa_offset -128
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end1:
	.size	_Z19__device_stub__scanPiS_i, .Lfunc_end1-_Z19__device_stub__scanPiS_i
	.cfi_endproc
                                        # -- End function
	.section	.text._Z14countGMWrapperIiEviiPT_Pii,"axG",@progbits,_Z14countGMWrapperIiEviiPT_Pii,comdat
	.weak	_Z14countGMWrapperIiEviiPT_Pii  # -- Begin function _Z14countGMWrapperIiEviiPT_Pii
	.type	_Z14countGMWrapperIiEviiPT_Pii,@function
_Z14countGMWrapperIiEviiPT_Pii:         # @_Z14countGMWrapperIiEviiPT_Pii
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset %rbx, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movl	%r8d, %ebx
	movq	%rcx, %r14
	movq	%rdx, %r15
	movl	%edi, %edx
	btsq	$32, %rdx
	movl	%esi, %eax
	btsq	$32, %rax
	shll	$2, %edi
	movslq	%edi, %r8
	movq	%rax, %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	je	.LBB2_2
# %bb.1:
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.LBB2_2:
	.cfi_def_cfa_offset 32
	movq	%r15, %rdi
	movq	%r14, %rsi
	movl	%ebx, %edx
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	jmp	_Z29__device_stub__countGlobalMemIiLi512EEvPT_Pii # TAILCALL
.Lfunc_end2:
	.size	_Z14countGMWrapperIiEviiPT_Pii, .Lfunc_end2-_Z14countGMWrapperIiEviiPT_Pii
	.cfi_endproc
                                        # -- End function
	.section	.text._Z29__device_stub__countGlobalMemIiLi512EEvPT_Pii,"axG",@progbits,_Z29__device_stub__countGlobalMemIiLi512EEvPT_Pii,comdat
	.weak	_Z29__device_stub__countGlobalMemIiLi512EEvPT_Pii # -- Begin function _Z29__device_stub__countGlobalMemIiLi512EEvPT_Pii
	.type	_Z29__device_stub__countGlobalMemIiLi512EEvPT_Pii,@function
_Z29__device_stub__countGlobalMemIiLi512EEvPT_Pii: # @_Z29__device_stub__countGlobalMemIiLi512EEvPT_Pii
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$112, %rsp
	.cfi_def_cfa_offset 160
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	40(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	32(%rsp), %rcx
	movq	%rsi, (%rcx)
	leaq	12(%rsp), %rsi
	movl	%edx, (%rsi)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%rbx)
	movq	%rsi, 16(%rbx)
	leaq	64(%rsp), %r14
	leaq	48(%rsp), %r15
	leaq	24(%rsp), %r12
	leaq	16(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z14countGlobalMemIiLi512EEvPT_Pii, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$128, %rsp
	.cfi_adjust_cfa_offset -128
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end3:
	.size	_Z29__device_stub__countGlobalMemIiLi512EEvPT_Pii, .Lfunc_end3-_Z29__device_stub__countGlobalMemIiLi512EEvPT_Pii
	.cfi_endproc
                                        # -- End function
	.section	.text._Z14countGMWrapperIfEviiPT_Pii,"axG",@progbits,_Z14countGMWrapperIfEviiPT_Pii,comdat
	.weak	_Z14countGMWrapperIfEviiPT_Pii  # -- Begin function _Z14countGMWrapperIfEviiPT_Pii
	.type	_Z14countGMWrapperIfEviiPT_Pii,@function
_Z14countGMWrapperIfEviiPT_Pii:         # @_Z14countGMWrapperIfEviiPT_Pii
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset %rbx, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movl	%r8d, %ebx
	movq	%rcx, %r14
	movq	%rdx, %r15
	movl	%edi, %edx
	btsq	$32, %rdx
	movl	%esi, %eax
	btsq	$32, %rax
	shll	$2, %edi
	movslq	%edi, %r8
	movq	%rax, %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	je	.LBB4_2
# %bb.1:
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.LBB4_2:
	.cfi_def_cfa_offset 32
	movq	%r15, %rdi
	movq	%r14, %rsi
	movl	%ebx, %edx
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	jmp	_Z29__device_stub__countGlobalMemIfLi512EEvPT_Pii # TAILCALL
.Lfunc_end4:
	.size	_Z14countGMWrapperIfEviiPT_Pii, .Lfunc_end4-_Z14countGMWrapperIfEviiPT_Pii
	.cfi_endproc
                                        # -- End function
	.section	.text._Z29__device_stub__countGlobalMemIfLi512EEvPT_Pii,"axG",@progbits,_Z29__device_stub__countGlobalMemIfLi512EEvPT_Pii,comdat
	.weak	_Z29__device_stub__countGlobalMemIfLi512EEvPT_Pii # -- Begin function _Z29__device_stub__countGlobalMemIfLi512EEvPT_Pii
	.type	_Z29__device_stub__countGlobalMemIfLi512EEvPT_Pii,@function
_Z29__device_stub__countGlobalMemIfLi512EEvPT_Pii: # @_Z29__device_stub__countGlobalMemIfLi512EEvPT_Pii
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$112, %rsp
	.cfi_def_cfa_offset 160
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	40(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	32(%rsp), %rcx
	movq	%rsi, (%rcx)
	leaq	12(%rsp), %rsi
	movl	%edx, (%rsi)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%rbx)
	movq	%rsi, 16(%rbx)
	leaq	64(%rsp), %r14
	leaq	48(%rsp), %r15
	leaq	24(%rsp), %r12
	leaq	16(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z14countGlobalMemIfLi512EEvPT_Pii, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$128, %rsp
	.cfi_adjust_cfa_offset -128
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end5:
	.size	_Z29__device_stub__countGlobalMemIfLi512EEvPT_Pii, .Lfunc_end5-_Z29__device_stub__countGlobalMemIfLi512EEvPT_Pii
	.cfi_endproc
                                        # -- End function
	.section	.text._Z14countGMWrapperIdEviiPT_Pii,"axG",@progbits,_Z14countGMWrapperIdEviiPT_Pii,comdat
	.weak	_Z14countGMWrapperIdEviiPT_Pii  # -- Begin function _Z14countGMWrapperIdEviiPT_Pii
	.type	_Z14countGMWrapperIdEviiPT_Pii,@function
_Z14countGMWrapperIdEviiPT_Pii:         # @_Z14countGMWrapperIdEviiPT_Pii
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset %rbx, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movl	%r8d, %ebx
	movq	%rcx, %r14
	movq	%rdx, %r15
	movl	%edi, %edx
	btsq	$32, %rdx
	movl	%esi, %eax
	btsq	$32, %rax
	shll	$3, %edi
	movslq	%edi, %r8
	movq	%rax, %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	je	.LBB6_2
# %bb.1:
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.LBB6_2:
	.cfi_def_cfa_offset 32
	movq	%r15, %rdi
	movq	%r14, %rsi
	movl	%ebx, %edx
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	jmp	_Z29__device_stub__countGlobalMemIdLi512EEvPT_Pii # TAILCALL
.Lfunc_end6:
	.size	_Z14countGMWrapperIdEviiPT_Pii, .Lfunc_end6-_Z14countGMWrapperIdEviiPT_Pii
	.cfi_endproc
                                        # -- End function
	.section	.text._Z29__device_stub__countGlobalMemIdLi512EEvPT_Pii,"axG",@progbits,_Z29__device_stub__countGlobalMemIdLi512EEvPT_Pii,comdat
	.weak	_Z29__device_stub__countGlobalMemIdLi512EEvPT_Pii # -- Begin function _Z29__device_stub__countGlobalMemIdLi512EEvPT_Pii
	.type	_Z29__device_stub__countGlobalMemIdLi512EEvPT_Pii,@function
_Z29__device_stub__countGlobalMemIdLi512EEvPT_Pii: # @_Z29__device_stub__countGlobalMemIdLi512EEvPT_Pii
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$112, %rsp
	.cfi_def_cfa_offset 160
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	40(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	32(%rsp), %rcx
	movq	%rsi, (%rcx)
	leaq	12(%rsp), %rsi
	movl	%edx, (%rsi)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%rbx)
	movq	%rsi, 16(%rbx)
	leaq	64(%rsp), %r14
	leaq	48(%rsp), %r15
	leaq	24(%rsp), %r12
	leaq	16(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z14countGlobalMemIdLi512EEvPT_Pii, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$128, %rsp
	.cfi_adjust_cfa_offset -128
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end7:
	.size	_Z29__device_stub__countGlobalMemIdLi512EEvPT_Pii, .Lfunc_end7-_Z29__device_stub__countGlobalMemIdLi512EEvPT_Pii
	.cfi_endproc
                                        # -- End function
	.section	.text._Z13countSWrapperIiEviiPT_Piiii,"axG",@progbits,_Z13countSWrapperIiEviiPT_Piiii,comdat
	.weak	_Z13countSWrapperIiEviiPT_Piiii # -- Begin function _Z13countSWrapperIiEviiPT_Piiii
	.type	_Z13countSWrapperIiEviiPT_Piiii,@function
_Z13countSWrapperIiEviiPT_Piiii:        # @_Z13countSWrapperIiEviiPT_Piiii
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movl	%r9d, %ebx
	movl	%r8d, %ebp
	movq	%rcx, %r14
	movq	%rdx, %r12
	movl	48(%rsp), %r15d
	movl	%edi, %edx
	btsq	$32, %rdx
	movl	%esi, %eax
	btsq	$32, %rax
	shll	$2, %edi
	movslq	%edi, %r8
	movq	%rax, %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	je	.LBB8_2
# %bb.1:
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.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
.LBB8_2:
	.cfi_def_cfa_offset 48
	movq	%r12, %rdi
	movq	%r14, %rsi
	movl	%ebp, %edx
	movl	%ebx, %ecx
	movl	%r15d, %r8d
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.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
	jmp	_Z29__device_stub__countSharedMemIiLi512EEvPT_Piiii # TAILCALL
.Lfunc_end8:
	.size	_Z13countSWrapperIiEviiPT_Piiii, .Lfunc_end8-_Z13countSWrapperIiEviiPT_Piiii
	.cfi_endproc
                                        # -- End function
	.section	.text._Z29__device_stub__countSharedMemIiLi512EEvPT_Piiii,"axG",@progbits,_Z29__device_stub__countSharedMemIiLi512EEvPT_Piiii,comdat
	.weak	_Z29__device_stub__countSharedMemIiLi512EEvPT_Piiii # -- Begin function _Z29__device_stub__countSharedMemIiLi512EEvPT_Piiii
	.type	_Z29__device_stub__countSharedMemIiLi512EEvPT_Piiii,@function
_Z29__device_stub__countSharedMemIiLi512EEvPT_Piiii: # @_Z29__device_stub__countSharedMemIiLi512EEvPT_Piiii
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$128, %rsp
	.cfi_def_cfa_offset 176
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	40(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	32(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	12(%rsp), %rsi
	movl	%edx, (%rsi)
	leaq	8(%rsp), %rdx
	movl	%ecx, (%rdx)
	leaq	4(%rsp), %rcx
	movl	%r8d, (%rcx)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%rbx)
	leaq	64(%rsp), %r14
	leaq	48(%rsp), %r15
	leaq	24(%rsp), %r12
	leaq	16(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z14countSharedMemIiLi512EEvPT_Piiii, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$144, %rsp
	.cfi_adjust_cfa_offset -144
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end9:
	.size	_Z29__device_stub__countSharedMemIiLi512EEvPT_Piiii, .Lfunc_end9-_Z29__device_stub__countSharedMemIiLi512EEvPT_Piiii
	.cfi_endproc
                                        # -- End function
	.text
	.globl	_Z11myCountTestPKc              # -- Begin function _Z11myCountTestPKc
	.type	_Z11myCountTestPKc,@function
_Z11myCountTestPKc:                     # @_Z11myCountTestPKc
.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	$2232, %rsp                     # imm = 0x8B8
	.cfi_def_cfa_offset 2288
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movq	%rdi, %rsi
	xorps	%xmm0, %xmm0
	movaps	%xmm0, 64(%rsp)
	xorl	%eax, %eax
	movq	%rax, 80(%rsp)
	leaq	160(%rsp), %rcx
	movq	%rcx, -16(%rcx)
	movq	%rax, -8(%rcx)
	movb	$0, (%rcx)
.Ltmp0:
	leaq	1712(%rsp), %rbx
	movq	%rbx, %rdi
	movl	$8, %edx
	callq	_ZNSt14basic_ifstreamIcSt11char_traitsIcEEC1EPKcSt13_Ios_Openmode
.Ltmp1:
# %bb.1:
	leaq	1832(%rsp), %rdi
	callq	_ZNKSt12__basic_fileIcE7is_openEv
	testb	%al, %al
	je	.LBB10_18
# %bb.2:                                # %.preheader196
	leaq	304(%rsp), %r14
	leaq	144(%rsp), %r15
	leaq	176(%rsp), %r12
	leaq	1200(%rsp), %r13
	leaq	64(%rsp), %rbp
.LBB10_3:                               # =>This Loop Header: Depth=1
                                        #     Child Loop BB10_7 Depth 2
	movq	1712(%rsp), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
.Ltmp3:
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
.Ltmp4:
# %bb.4:                                # %.noexc
                                        #   in Loop: Header=BB10_3 Depth=1
.Ltmp5:
	movsbl	%al, %edx
	movq	%rbx, %rdi
	movq	%r15, %rsi
	callq	_ZSt7getlineIcSt11char_traitsIcESaIcEERSt13basic_istreamIT_T0_ES7_RNSt7__cxx1112basic_stringIS4_S5_T1_EES4_
.Ltmp6:
# %bb.5:                                #   in Loop: Header=BB10_3 Depth=1
	movq	(%rax), %rcx
	movq	-24(%rcx), %rcx
	testb	$5, 32(%rax,%rcx)
	jne	.LBB10_19
# %bb.6:                                #   in Loop: Header=BB10_3 Depth=1
.Ltmp121:
	movq	%r12, %rdi
	movq	%r15, %rsi
	movl	$24, %edx
	callq	_ZNSt7__cxx1118basic_stringstreamIcSt11char_traitsIcESaIcEEC1ERKNS_12basic_stringIcS2_S3_EESt13_Ios_Openmode
.Ltmp122:
.LBB10_7:                               #   Parent Loop BB10_3 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
.Ltmp124:
	movq	%r12, %rdi
	movq	%r13, %rsi
	callq	_ZNSirsERi
.Ltmp125:
# %bb.8:                                #   in Loop: Header=BB10_7 Depth=2
	movq	(%rax), %rcx
	movq	-24(%rcx), %rcx
	testb	$5, 32(%rax,%rcx)
	jne	.LBB10_17
# %bb.9:                                #   in Loop: Header=BB10_7 Depth=2
	movq	72(%rsp), %rsi
	cmpq	80(%rsp), %rsi
	je	.LBB10_11
# %bb.10:                               #   in Loop: Header=BB10_7 Depth=2
	movl	1200(%rsp), %eax
	movl	%eax, (%rsi)
	addq	$4, %rsi
	movq	%rsi, 72(%rsp)
	jmp	.LBB10_12
.LBB10_11:                              #   in Loop: Header=BB10_7 Depth=2
.Ltmp126:
	movq	%rbp, %rdi
	movq	%r13, %rdx
	callq	_ZNSt6vectorIiSaIiEE17_M_realloc_insertIJRKiEEEvN9__gnu_cxx17__normal_iteratorIPiS1_EEDpOT_
.Ltmp127:
.LBB10_12:                              # %_ZNSt6vectorIiSaIiEE9push_backERKi.exit
                                        #   in Loop: Header=BB10_7 Depth=2
.Ltmp128:
	movq	%r12, %rdi
	callq	_ZNSi4peekEv
.Ltmp129:
# %bb.13:                               #   in Loop: Header=BB10_7 Depth=2
	cmpl	$44, %eax
	je	.LBB10_16
# %bb.14:                               #   in Loop: Header=BB10_7 Depth=2
.Ltmp130:
	movq	%r12, %rdi
	callq	_ZNSi4peekEv
.Ltmp131:
# %bb.15:                               #   in Loop: Header=BB10_7 Depth=2
	cmpl	$32, %eax
	jne	.LBB10_7
.LBB10_16:                              #   in Loop: Header=BB10_7 Depth=2
.Ltmp132:
	movq	%r12, %rdi
	callq	_ZNSi6ignoreEv
.Ltmp133:
	jmp	.LBB10_7
.LBB10_17:                              #   in Loop: Header=BB10_3 Depth=1
	movl	$_ZTTNSt7__cxx1118basic_stringstreamIcSt11char_traitsIcESaIcEEE, %esi
	movq	%r12, %rdi
	callq	_ZNSt7__cxx1118basic_stringstreamIcSt11char_traitsIcESaIcEED2Ev
	movq	%r14, %rdi
	callq	_ZNSt8ios_baseD2Ev
	jmp	.LBB10_3
.LBB10_18:
	xorl	%eax, %eax
	xorl	%ebx, %ebx
	jmp	.LBB10_21
.LBB10_19:
.Ltmp8:
	leaq	1712(%rsp), %rdi
	callq	_ZNSt14basic_ifstreamIcSt11char_traitsIcEE5closeEv
.Ltmp9:
# %bb.20:                               # %._crit_edge221
	movq	64(%rsp), %rax
	movq	72(%rsp), %rbx
.LBB10_21:
	subq	%rax, %rbx
	movq	%rbx, %rax
	shrq	$2, %rax
	leal	511(%rax), %r15d
	testl	%eax, %eax
	movq	%rax, 40(%rsp)                  # 8-byte Spill
	cmovnsl	%eax, %r15d
	xorl	%eax, %eax
	leaq	24(%rsp), %rdi
	movq	%rax, (%rdi)
	movq	%rax, 8(%rsp)
	shlq	$30, %rbx
	sarq	$30, %rbx
.Ltmp11:
	movq	%rbx, %rsi
	callq	hipMalloc
.Ltmp12:
# %bb.22:
.Ltmp13:
	leaq	8(%rsp), %rdi
	movl	$40, %esi
	callq	hipMalloc
.Ltmp14:
# %bb.23:
	movl	$40, %edi
	callq	malloc
	movq	%rax, %r14
	movq	64(%rsp), %r12
.Ltmp16:
	movl	$40, %edi
	callq	_Znwm
.Ltmp17:
# %bb.24:
	movq	%rax, 88(%rsp)
	movq	%rax, %rcx
	addq	$40, %rcx
	movq	%rcx, 104(%rsp)
	movq	%rcx, 96(%rsp)
	xorl	%ecx, %ecx
	movq	%rcx, 32(%rax)
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rax)
	movups	%xmm0, (%rax)
	movq	%rcx, 32(%r14)
	movups	%xmm0, 16(%r14)
	movups	%xmm0, (%r14)
	movq	24(%rsp), %rdi
.Ltmp19:
	movq	%r12, %rsi
	movq	%rbx, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
.Ltmp20:
# %bb.25:
	movq	8(%rsp), %rdi
.Ltmp21:
	movl	$40, %edx
	movq	%r14, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
.Ltmp22:
# %bb.26:
	sarl	$9, %r15d
	incl	%r15d
	movq	24(%rsp), %rdx
	movq	8(%rsp), %rcx
.Ltmp23:
	movl	$512, %edi                      # imm = 0x200
	movl	%r15d, %esi
	movq	40(%rsp), %r8                   # 8-byte Reload
                                        # kill: def $r8d killed $r8d killed $r8
	callq	_Z14countGMWrapperIiEviiPT_Pii
.Ltmp24:
# %bb.27:
	movq	8(%rsp), %rsi
.Ltmp25:
	movl	$40, %edx
	movq	%r14, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
.Ltmp26:
# %bb.28:
.Ltmp28:
	movq	%r12, 32(%rsp)                  # 8-byte Spill
	leaq	176(%rsp), %rdi
	movl	$.L.str, %esi
	movl	$1, %edx
	callq	_ZNSt14basic_ofstreamIcSt11char_traitsIcEEC1EPKcSt13_Ios_Openmode
.Ltmp29:
# %bb.29:                               # %.preheader195
	leaq	288(%rsp), %rbx
	xorl	%ebp, %ebp
	leaq	176(%rsp), %r13
.LBB10_30:                              # =>This Inner Loop Header: Depth=1
	movq	%rbx, %rdi
	callq	_ZNKSt12__basic_fileIcE7is_openEv
	testb	%al, %al
	je	.LBB10_35
# %bb.31:                               #   in Loop: Header=BB10_30 Depth=1
.Ltmp31:
	movl	$.L.str.1, %esi
	movl	$8, %edx
	movq	%r13, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
.Ltmp32:
# %bb.32:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc.exit
                                        #   in Loop: Header=BB10_30 Depth=1
.Ltmp33:
	movq	%r13, %rdi
	movl	%ebp, %esi
	callq	_ZNSolsEi
.Ltmp34:
# %bb.33:                               #   in Loop: Header=BB10_30 Depth=1
.Ltmp35:
	movq	%rax, %r12
	movl	$.L.str.2, %esi
	movl	$3, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
.Ltmp36:
# %bb.34:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc.exit148
                                        #   in Loop: Header=BB10_30 Depth=1
	movl	(%r14,%rbp,4), %esi
.Ltmp37:
	movq	%r12, %rdi
	callq	_ZNSolsEi
.Ltmp38:
.LBB10_35:                              #   in Loop: Header=BB10_30 Depth=1
	incq	%rbp
	cmpq	$10, %rbp
	jne	.LBB10_30
# %bb.36:
.Ltmp40:
	leaq	176(%rsp), %rbx
	movq	%rbx, %rdi
	callq	_ZNSt14basic_ofstreamIcSt11char_traitsIcEE5closeEv
.Ltmp41:
# %bb.37:
	movq	176(%rsp), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
.Ltmp42:
	xorl	%esi, %esi
	callq	_ZNSt9basic_iosIcSt11char_traitsIcEE5clearESt12_Ios_Iostate
.Ltmp43:
# %bb.38:
	movq	8(%rsp), %rdi
.Ltmp44:
	callq	hipFree
.Ltmp45:
# %bb.39:
	movq	32(%rsp), %rdi                  # 8-byte Reload
	callq	free
	leaq	56(%rsp), %rdi
	movq	$0, (%rdi)
	leal	(%r15,%r15), %eax
	leal	(%rax,%rax,4), %eax
	movslq	%eax, %rbx
	leaq	(,%rbx,4), %rbp
.Ltmp47:
	movq	%rbp, %rsi
	callq	hipMalloc
.Ltmp48:
# %bb.40:
	movq	%rbp, %rdi
	callq	malloc
	movq	%rax, %r12
.Ltmp50:
	leaq	88(%rsp), %rdi
	movq	%rbx, %rsi
	callq	_ZNSt6vectorIiSaIiEE6resizeEm
.Ltmp51:
# %bb.41:
	movq	88(%rsp), %rbx
	movq	96(%rsp), %r13
	cmpq	%r13, %rbx
	je	.LBB10_43
# %bb.42:                               # %.lr.ph.i.i.i.i150.preheader
	movq	%r13, %rdx
	subq	%rbx, %rdx
	addq	$-4, %rdx
	andq	$-4, %rdx
	addq	$4, %rdx
	movq	%rbx, %rdi
	xorl	%esi, %esi
	callq	memset@PLT
.LBB10_43:                              # %_ZSt4fillIN9__gnu_cxx17__normal_iteratorIPiSt6vectorIiSaIiEEEEiEvT_S7_RKT0_.exit153
	subq	%rbx, %r13
	cmpq	$5, %r13
	jl	.LBB10_105
# %bb.44:
	movq	%r12, %rdi
	movq	%rbx, %rsi
	movq	%r13, %rdx
	callq	memmove@PLT
.LBB10_45:                              # %_ZSt4copyIN9__gnu_cxx17__normal_iteratorIPiSt6vectorIiSaIiEEEES2_ET0_T_S8_S7_.exit154
	movq	56(%rsp), %rdi
.Ltmp52:
	movq	%r12, %rsi
	movq	%rbp, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
.Ltmp53:
# %bb.46:
	movq	24(%rsp), %rdx
	movq	8(%rsp), %rcx
.Ltmp54:
	movl	$100, (%rsp)
	movl	$512, %edi                      # imm = 0x200
	movl	%r15d, %esi
	movq	40(%rsp), %r8                   # 8-byte Reload
                                        # kill: def $r8d killed $r8d killed $r8
	movl	$1000, %r9d                     # imm = 0x3E8
	callq	_Z13countSWrapperIiEviiPT_Piiii
.Ltmp55:
# %bb.47:
	movq	56(%rsp), %rsi
.Ltmp56:
	movq	%r12, %rdi
	movq	%rbp, %rdx
	movl	$2, %ecx
	callq	hipMemcpy
.Ltmp57:
# %bb.48:
.Ltmp59:
	movq	%r12, 136(%rsp)                 # 8-byte Spill
	leaq	1200(%rsp), %rdi
	movl	$.L.str.3, %esi
	movl	$1, %edx
	callq	_ZNSt14basic_ofstreamIcSt11char_traitsIcEEC1EPKcSt13_Ios_Openmode
.Ltmp60:
# %bb.49:                               # %.preheader194
	movl	%r15d, %eax
	movq	%rax, 32(%rsp)                  # 8-byte Spill
	xorl	%r12d, %r12d
	leaq	1200(%rsp), %rbx
	movq	136(%rsp), %rbp                 # 8-byte Reload
.LBB10_50:                              # %.preheader193
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB10_52 Depth 2
	movl	$0, %r13d
	cmpl	$-511, 40(%rsp)                 # 4-byte Folded Reload
                                        # imm = 0xFE01
	jl	.LBB10_53
# %bb.51:                               # %.lr.ph.preheader
                                        #   in Loop: Header=BB10_50 Depth=1
	xorl	%r13d, %r13d
	movq	32(%rsp), %rax                  # 8-byte Reload
	movq	%rbp, %rcx
.LBB10_52:                              # %.lr.ph
                                        #   Parent Loop BB10_50 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	addl	(%rcx), %r13d
	addq	$40, %rcx
	decq	%rax
	jne	.LBB10_52
.LBB10_53:                              # %._crit_edge
                                        #   in Loop: Header=BB10_50 Depth=1
	leaq	1312(%rsp), %rdi
	callq	_ZNKSt12__basic_fileIcE7is_openEv
	testb	%al, %al
	je	.LBB10_58
# %bb.54:                               #   in Loop: Header=BB10_50 Depth=1
.Ltmp62:
	movl	$.L.str.1, %esi
	movl	$8, %edx
	movq	%rbx, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
.Ltmp63:
# %bb.55:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc.exit156
                                        #   in Loop: Header=BB10_50 Depth=1
.Ltmp64:
	movq	%rbx, %rdi
	movl	%r12d, %esi
	callq	_ZNSolsEi
.Ltmp65:
# %bb.56:                               #   in Loop: Header=BB10_50 Depth=1
.Ltmp66:
	movq	%rax, %r15
	movl	$.L.str.2, %esi
	movl	$3, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
.Ltmp67:
# %bb.57:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc.exit158
                                        #   in Loop: Header=BB10_50 Depth=1
.Ltmp68:
	movq	%r15, %rdi
	movl	%r13d, %esi
	callq	_ZNSolsEi
.Ltmp69:
.LBB10_58:                              #   in Loop: Header=BB10_50 Depth=1
	incq	%r12
	addq	$4, %rbp
	cmpq	$10, %r12
	jne	.LBB10_50
# %bb.59:
.Ltmp71:
	leaq	1200(%rsp), %r15
	movq	%r15, %rdi
	callq	_ZNSt14basic_ofstreamIcSt11char_traitsIcEE5closeEv
.Ltmp72:
# %bb.60:
	movq	1200(%rsp), %rax
	movq	-24(%rax), %rdi
	addq	%r15, %rdi
.Ltmp73:
	xorl	%esi, %esi
	callq	_ZNSt9basic_iosIcSt11char_traitsIcEE5clearESt12_Ios_Iostate
.Ltmp74:
# %bb.61:
	movq	56(%rsp), %rdi
.Ltmp75:
	callq	hipFree
.Ltmp76:
# %bb.62:
	movq	136(%rsp), %rdi                 # 8-byte Reload
	callq	free
	movq	24(%rsp), %rdi
.Ltmp77:
	callq	hipFree
.Ltmp78:
# %bb.63:
	xorl	%eax, %eax
	movq	%rax, 16(%rsp)
	movq	%rax, 48(%rsp)
	movl	$64, %edi
	callq	malloc
	movq	%rax, %r12
.Ltmp80:
	leaq	88(%rsp), %rdi
	movl	$16, %esi
	callq	_ZNSt6vectorIiSaIiEE6resizeEm
.Ltmp81:
# %bb.64:
	movq	88(%rsp), %rbx
	movq	96(%rsp), %r13
	cmpq	%r13, %rbx
	je	.LBB10_66
# %bb.65:                               # %.lr.ph.i.i.i.i160.preheader
	movq	%r13, %rdx
	subq	%rbx, %rdx
	addq	$-4, %rdx
	andq	$-4, %rdx
	addq	$4, %rdx
	movq	%rbx, %rdi
	xorl	%esi, %esi
	callq	memset@PLT
.LBB10_66:                              # %_ZSt4fillIN9__gnu_cxx17__normal_iteratorIPiSt6vectorIiSaIiEEEEiEvT_S7_RKT0_.exit163
	subq	%rbx, %r13
	cmpq	$5, %r13
	jl	.LBB10_107
# %bb.67:
	movq	%r12, %rdi
	movq	%rbx, %rsi
	movq	%r13, %rdx
	callq	memmove@PLT
.LBB10_68:                              # %_ZSt4copyIN9__gnu_cxx17__normal_iteratorIPiSt6vectorIiSaIiEEEES2_ET0_T_S8_S7_.exit164
.Ltmp82:
	leaq	48(%rsp), %rdi
	movl	$64, %esi
	callq	hipMalloc
.Ltmp83:
# %bb.69:
.Ltmp84:
	leaq	16(%rsp), %rdi
	movl	$64, %esi
	callq	hipMalloc
.Ltmp85:
# %bb.70:
	xorps	%xmm0, %xmm0
	leaq	112(%rsp), %rbx
	movaps	%xmm0, (%rbx)
	movq	$0, 16(%rbx)
	xorl	%ebp, %ebp
	leaq	688(%rsp), %r13
.LBB10_71:                              # =>This Inner Loop Header: Depth=1
	cmpq	$9, %rbp
	ja	.LBB10_74
# %bb.72:                               #   in Loop: Header=BB10_71 Depth=1
	movq	120(%rsp), %rsi
	cmpq	128(%rsp), %rsi
	je	.LBB10_77
# %bb.73:                               #   in Loop: Header=BB10_71 Depth=1
	movl	(%r14), %eax
	movl	%eax, (%rsi)
	jmp	.LBB10_76
.LBB10_74:                              #   in Loop: Header=BB10_71 Depth=1
	movl	$0, 688(%rsp)
	movq	120(%rsp), %rsi
	cmpq	128(%rsp), %rsi
	je	.LBB10_78
# %bb.75:                               #   in Loop: Header=BB10_71 Depth=1
	movl	$0, (%rsi)
.LBB10_76:                              # %_ZNSt6vectorIiSaIiEE9push_backERKi.exit167
                                        #   in Loop: Header=BB10_71 Depth=1
	addq	$4, %rsi
	movq	%rsi, 120(%rsp)
	jmp	.LBB10_79
.LBB10_77:                              #   in Loop: Header=BB10_71 Depth=1
.Ltmp90:
	movq	%rbx, %rdi
	movq	%r14, %rdx
	callq	_ZNSt6vectorIiSaIiEE17_M_realloc_insertIJRKiEEEvN9__gnu_cxx17__normal_iteratorIPiS1_EEDpOT_
.Ltmp91:
	jmp	.LBB10_79
.LBB10_78:                              #   in Loop: Header=BB10_71 Depth=1
.Ltmp87:
	movq	%rbx, %rdi
	movq	%r13, %rdx
	callq	_ZNSt6vectorIiSaIiEE17_M_realloc_insertIJiEEEvN9__gnu_cxx17__normal_iteratorIPiS1_EEDpOT_
.Ltmp88:
.LBB10_79:                              # %_ZNSt6vectorIiSaIiEE9push_backERKi.exit167
                                        #   in Loop: Header=BB10_71 Depth=1
	incq	%rbp
	addq	$4, %r14
	cmpq	$16, %rbp
	jne	.LBB10_71
# %bb.80:
	movq	112(%rsp), %rsi
	movq	48(%rsp), %rdi
.Ltmp93:
	movl	$64, %edx
	movl	$1, %ecx
	callq	hipMemcpy
.Ltmp94:
# %bb.81:
	movq	16(%rsp), %rdi
.Ltmp95:
	movl	$64, %edx
	movq	%r12, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
.Ltmp96:
# %bb.82:
.Ltmp97:
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movabsq	$4294967312, %rdx               # imm = 0x100000010
	movl	$64, %r8d
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
.Ltmp98:
# %bb.83:
	testl	%eax, %eax
	jne	.LBB10_85
# %bb.84:
	movq	48(%rsp), %rdi
	movq	16(%rsp), %rsi
.Ltmp99:
	movl	$16, %edx
	callq	_Z19__device_stub__scanPiS_i
.Ltmp100:
.LBB10_85:
	movq	16(%rsp), %rsi
.Ltmp101:
	movl	$64, %edx
	movq	%r12, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
.Ltmp102:
# %bb.86:
.Ltmp104:
	leaq	688(%rsp), %rdi
	movl	$.L.str.4, %esi
	movl	$1, %edx
	callq	_ZNSt14basic_ofstreamIcSt11char_traitsIcEEC1EPKcSt13_Ios_Openmode
.Ltmp105:
# %bb.87:                               # %.preheader
	leaq	800(%rsp), %rbx
	xorl	%r14d, %r14d
	leaq	688(%rsp), %r13
.LBB10_88:                              # =>This Inner Loop Header: Depth=1
	movq	%rbx, %rdi
	callq	_ZNKSt12__basic_fileIcE7is_openEv
	testb	%al, %al
	je	.LBB10_93
# %bb.89:                               #   in Loop: Header=BB10_88 Depth=1
.Ltmp107:
	movl	$.L.str.5, %esi
	movl	$10, %edx
	movq	%r13, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
.Ltmp108:
# %bb.90:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc.exit170
                                        #   in Loop: Header=BB10_88 Depth=1
.Ltmp109:
	movq	%r13, %rdi
	movl	%r14d, %esi
	callq	_ZNSolsEi
.Ltmp110:
# %bb.91:                               #   in Loop: Header=BB10_88 Depth=1
.Ltmp111:
	movq	%rax, %rbp
	movl	$.L.str.2, %esi
	movl	$3, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
.Ltmp112:
# %bb.92:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc.exit172
                                        #   in Loop: Header=BB10_88 Depth=1
	movl	(%r12,%r14,4), %esi
.Ltmp113:
	movq	%rbp, %rdi
	callq	_ZNSolsEi
.Ltmp114:
.LBB10_93:                              #   in Loop: Header=BB10_88 Depth=1
	incq	%r14
	cmpq	$10, %r14
	jne	.LBB10_88
# %bb.94:
.Ltmp116:
	leaq	688(%rsp), %rbx
	movq	%rbx, %rdi
	callq	_ZNSt14basic_ofstreamIcSt11char_traitsIcEE5closeEv
.Ltmp117:
# %bb.95:
	movq	688(%rsp), %rax
	addq	-24(%rax), %rbx
.Ltmp118:
	movq	%rbx, %rdi
	xorl	%esi, %esi
	callq	_ZNSt9basic_iosIcSt11char_traitsIcEE5clearESt12_Ios_Iostate
.Ltmp119:
# %bb.96:
	movq	_ZTTSt14basic_ofstreamIcSt11char_traitsIcEE(%rip), %rbx
	leaq	696(%rsp), %rdi
	movq	%rbx, -8(%rdi)
	movq	_ZTTSt14basic_ofstreamIcSt11char_traitsIcEE+24(%rip), %r14
	movq	-24(%rbx), %rax
	movq	%r14, 688(%rsp,%rax)
	callq	_ZNSt13basic_filebufIcSt11char_traitsIcEED2Ev
	leaq	936(%rsp), %rdi
	callq	_ZNSt8ios_baseD2Ev
	movq	112(%rsp), %rdi
	testq	%rdi, %rdi
	je	.LBB10_98
# %bb.97:
	callq	_ZdlPv
.LBB10_98:                              # %_ZNSt6vectorIiSaIiEED2Ev.exit
	leaq	1208(%rsp), %rdi
	movq	%rbx, -8(%rdi)
	movq	-24(%rbx), %rax
	movq	%r14, 1200(%rsp,%rax)
	callq	_ZNSt13basic_filebufIcSt11char_traitsIcEED2Ev
	movl	$248, %r12d
	addq	%r12, %r15
	movq	%r15, %rdi
	callq	_ZNSt8ios_baseD2Ev
	leaq	184(%rsp), %rdi
	movq	%rbx, -8(%rdi)
	movq	-24(%rbx), %rax
	movq	%r14, 176(%rsp,%rax)
	callq	_ZNSt13basic_filebufIcSt11char_traitsIcEED2Ev
	leaq	176(%rsp), %rdi
	addq	%r12, %rdi
	callq	_ZNSt8ios_baseD2Ev
	movq	88(%rsp), %rdi
	testq	%rdi, %rdi
	je	.LBB10_100
# %bb.99:
	callq	_ZdlPv
.LBB10_100:                             # %_ZNSt6vectorIiSaIiEED2Ev.exit174
	leaq	1712(%rsp), %rdi
	callq	_ZNSt14basic_ifstreamIcSt11char_traitsIcEED1Ev
	movq	144(%rsp), %rdi
	leaq	160(%rsp), %rax
	cmpq	%rax, %rdi
	je	.LBB10_102
# %bb.101:                              # %_ZNKSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE11_M_is_localEv.exit.i.i
	callq	_ZdlPv
.LBB10_102:                             # %_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev.exit
	movq	64(%rsp), %rdi
	testq	%rdi, %rdi
	je	.LBB10_104
# %bb.103:
	callq	_ZdlPv
.LBB10_104:                             # %_ZNSt6vectorIiSaIiEED2Ev.exit176
	addq	$2232, %rsp                     # imm = 0x8B8
	.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
.LBB10_105:
	.cfi_def_cfa_offset 2288
	cmpq	$4, %r13
	jne	.LBB10_45
# %bb.106:
	movl	(%rbx), %eax
	movl	%eax, (%r12)
	jmp	.LBB10_45
.LBB10_107:
	cmpq	$4, %r13
	jne	.LBB10_68
# %bb.108:
	movl	(%rbx), %eax
	movl	%eax, (%r12)
	jmp	.LBB10_68
.LBB10_109:                             # %.loopexit.split-lp
.Ltmp10:
	jmp	.LBB10_130
.LBB10_110:
.Ltmp106:
	jmp	.LBB10_126
.LBB10_111:
.Ltmp61:
	jmp	.LBB10_139
.LBB10_112:
.Ltmp49:
	jmp	.LBB10_139
.LBB10_113:
.Ltmp30:
	jmp	.LBB10_123
.LBB10_114:
.Ltmp18:
	jmp	.LBB10_130
.LBB10_115:
.Ltmp2:
	movq	%rax, %rbx
	jmp	.LBB10_145
.LBB10_116:
.Ltmp120:
	jmp	.LBB10_132
.LBB10_117:
.Ltmp15:
	jmp	.LBB10_130
.LBB10_118:
.Ltmp86:
	jmp	.LBB10_136
.LBB10_119:
.Ltmp46:
	jmp	.LBB10_139
.LBB10_120:
.Ltmp79:
	jmp	.LBB10_136
.LBB10_121:
.Ltmp58:
	jmp	.LBB10_139
.LBB10_122:
.Ltmp27:
.LBB10_123:
	movq	%rax, %rbx
	jmp	.LBB10_141
.LBB10_124:
.Ltmp103:
	jmp	.LBB10_126
.LBB10_125:
.Ltmp92:
	jmp	.LBB10_126
.LBB10_127:
.Ltmp89:
.LBB10_126:
	movq	%rax, %rbx
	jmp	.LBB10_133
.LBB10_128:
.Ltmp123:
	jmp	.LBB10_130
.LBB10_129:                             # %.loopexit
.Ltmp7:
.LBB10_130:
	movq	%rax, %rbx
	jmp	.LBB10_144
.LBB10_131:
.Ltmp115:
.LBB10_132:
	movq	%rax, %rbx
	leaq	688(%rsp), %rdi
	callq	_ZNSt14basic_ofstreamIcSt11char_traitsIcEED1Ev
.LBB10_133:
	movq	112(%rsp), %rdi
	testq	%rdi, %rdi
	je	.LBB10_137
# %bb.134:
	callq	_ZdlPv
	jmp	.LBB10_137
.LBB10_135:
.Ltmp70:
.LBB10_136:
	movq	%rax, %rbx
.LBB10_137:
	leaq	1200(%rsp), %rdi
	callq	_ZNSt14basic_ofstreamIcSt11char_traitsIcEED1Ev
	jmp	.LBB10_140
.LBB10_138:
.Ltmp39:
.LBB10_139:
	movq	%rax, %rbx
.LBB10_140:
	leaq	176(%rsp), %rdi
	callq	_ZNSt14basic_ofstreamIcSt11char_traitsIcEED1Ev
.LBB10_141:
	movq	88(%rsp), %rdi
	testq	%rdi, %rdi
	je	.LBB10_144
# %bb.142:
	callq	_ZdlPv
	jmp	.LBB10_144
.LBB10_143:
.Ltmp134:
	movq	%rax, %rbx
	leaq	176(%rsp), %rdi
	movl	$_ZTTNSt7__cxx1118basic_stringstreamIcSt11char_traitsIcESaIcEEE, %esi
	callq	_ZNSt7__cxx1118basic_stringstreamIcSt11char_traitsIcESaIcEED2Ev
	movq	%r14, %rdi
	callq	_ZNSt8ios_baseD2Ev
.LBB10_144:
	leaq	1712(%rsp), %rdi
	callq	_ZNSt14basic_ifstreamIcSt11char_traitsIcEED1Ev
.LBB10_145:
	movq	144(%rsp), %rdi
	leaq	160(%rsp), %rax
	cmpq	%rax, %rdi
	je	.LBB10_147
# %bb.146:                              # %_ZNKSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE11_M_is_localEv.exit.i.i181
	callq	_ZdlPv
.LBB10_147:                             # %_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev.exit183
	movq	64(%rsp), %rdi
	testq	%rdi, %rdi
	je	.LBB10_149
# %bb.148:
	callq	_ZdlPv
.LBB10_149:                             # %_ZNSt6vectorIiSaIiEED2Ev.exit185
	movq	%rbx, %rdi
	callq	_Unwind_Resume@PLT
.Lfunc_end10:
	.size	_Z11myCountTestPKc, .Lfunc_end10-_Z11myCountTestPKc
	.cfi_endproc
	.section	.gcc_except_table,"a",@progbits
	.p2align	2, 0x0
GCC_except_table10:
.Lexception0:
	.byte	255                             # @LPStart Encoding = omit
	.byte	255                             # @TType Encoding = omit
	.byte	1                               # Call site Encoding = uleb128
	.uleb128 .Lcst_end0-.Lcst_begin0
.Lcst_begin0:
	.uleb128 .Ltmp0-.Lfunc_begin0           # >> Call Site 1 <<
	.uleb128 .Ltmp1-.Ltmp0                  #   Call between .Ltmp0 and .Ltmp1
	.uleb128 .Ltmp2-.Lfunc_begin0           #     jumps to .Ltmp2
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp3-.Lfunc_begin0           # >> Call Site 2 <<
	.uleb128 .Ltmp6-.Ltmp3                  #   Call between .Ltmp3 and .Ltmp6
	.uleb128 .Ltmp7-.Lfunc_begin0           #     jumps to .Ltmp7
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp121-.Lfunc_begin0         # >> Call Site 3 <<
	.uleb128 .Ltmp122-.Ltmp121              #   Call between .Ltmp121 and .Ltmp122
	.uleb128 .Ltmp123-.Lfunc_begin0         #     jumps to .Ltmp123
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp124-.Lfunc_begin0         # >> Call Site 4 <<
	.uleb128 .Ltmp133-.Ltmp124              #   Call between .Ltmp124 and .Ltmp133
	.uleb128 .Ltmp134-.Lfunc_begin0         #     jumps to .Ltmp134
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp8-.Lfunc_begin0           # >> Call Site 5 <<
	.uleb128 .Ltmp9-.Ltmp8                  #   Call between .Ltmp8 and .Ltmp9
	.uleb128 .Ltmp10-.Lfunc_begin0          #     jumps to .Ltmp10
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp11-.Lfunc_begin0          # >> Call Site 6 <<
	.uleb128 .Ltmp14-.Ltmp11                #   Call between .Ltmp11 and .Ltmp14
	.uleb128 .Ltmp15-.Lfunc_begin0          #     jumps to .Ltmp15
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp16-.Lfunc_begin0          # >> Call Site 7 <<
	.uleb128 .Ltmp17-.Ltmp16                #   Call between .Ltmp16 and .Ltmp17
	.uleb128 .Ltmp18-.Lfunc_begin0          #     jumps to .Ltmp18
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp19-.Lfunc_begin0          # >> Call Site 8 <<
	.uleb128 .Ltmp26-.Ltmp19                #   Call between .Ltmp19 and .Ltmp26
	.uleb128 .Ltmp27-.Lfunc_begin0          #     jumps to .Ltmp27
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp28-.Lfunc_begin0          # >> Call Site 9 <<
	.uleb128 .Ltmp29-.Ltmp28                #   Call between .Ltmp28 and .Ltmp29
	.uleb128 .Ltmp30-.Lfunc_begin0          #     jumps to .Ltmp30
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp31-.Lfunc_begin0          # >> Call Site 10 <<
	.uleb128 .Ltmp38-.Ltmp31                #   Call between .Ltmp31 and .Ltmp38
	.uleb128 .Ltmp39-.Lfunc_begin0          #     jumps to .Ltmp39
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp40-.Lfunc_begin0          # >> Call Site 11 <<
	.uleb128 .Ltmp45-.Ltmp40                #   Call between .Ltmp40 and .Ltmp45
	.uleb128 .Ltmp46-.Lfunc_begin0          #     jumps to .Ltmp46
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp47-.Lfunc_begin0          # >> Call Site 12 <<
	.uleb128 .Ltmp48-.Ltmp47                #   Call between .Ltmp47 and .Ltmp48
	.uleb128 .Ltmp49-.Lfunc_begin0          #     jumps to .Ltmp49
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp50-.Lfunc_begin0          # >> Call Site 13 <<
	.uleb128 .Ltmp51-.Ltmp50                #   Call between .Ltmp50 and .Ltmp51
	.uleb128 .Ltmp58-.Lfunc_begin0          #     jumps to .Ltmp58
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp51-.Lfunc_begin0          # >> Call Site 14 <<
	.uleb128 .Ltmp52-.Ltmp51                #   Call between .Ltmp51 and .Ltmp52
	.byte	0                               #     has no landing pad
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp52-.Lfunc_begin0          # >> Call Site 15 <<
	.uleb128 .Ltmp57-.Ltmp52                #   Call between .Ltmp52 and .Ltmp57
	.uleb128 .Ltmp58-.Lfunc_begin0          #     jumps to .Ltmp58
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp59-.Lfunc_begin0          # >> Call Site 16 <<
	.uleb128 .Ltmp60-.Ltmp59                #   Call between .Ltmp59 and .Ltmp60
	.uleb128 .Ltmp61-.Lfunc_begin0          #     jumps to .Ltmp61
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp62-.Lfunc_begin0          # >> Call Site 17 <<
	.uleb128 .Ltmp69-.Ltmp62                #   Call between .Ltmp62 and .Ltmp69
	.uleb128 .Ltmp70-.Lfunc_begin0          #     jumps to .Ltmp70
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp71-.Lfunc_begin0          # >> Call Site 18 <<
	.uleb128 .Ltmp78-.Ltmp71                #   Call between .Ltmp71 and .Ltmp78
	.uleb128 .Ltmp79-.Lfunc_begin0          #     jumps to .Ltmp79
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp80-.Lfunc_begin0          # >> Call Site 19 <<
	.uleb128 .Ltmp81-.Ltmp80                #   Call between .Ltmp80 and .Ltmp81
	.uleb128 .Ltmp86-.Lfunc_begin0          #     jumps to .Ltmp86
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp81-.Lfunc_begin0          # >> Call Site 20 <<
	.uleb128 .Ltmp82-.Ltmp81                #   Call between .Ltmp81 and .Ltmp82
	.byte	0                               #     has no landing pad
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp82-.Lfunc_begin0          # >> Call Site 21 <<
	.uleb128 .Ltmp85-.Ltmp82                #   Call between .Ltmp82 and .Ltmp85
	.uleb128 .Ltmp86-.Lfunc_begin0          #     jumps to .Ltmp86
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp90-.Lfunc_begin0          # >> Call Site 22 <<
	.uleb128 .Ltmp91-.Ltmp90                #   Call between .Ltmp90 and .Ltmp91
	.uleb128 .Ltmp92-.Lfunc_begin0          #     jumps to .Ltmp92
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp87-.Lfunc_begin0          # >> Call Site 23 <<
	.uleb128 .Ltmp88-.Ltmp87                #   Call between .Ltmp87 and .Ltmp88
	.uleb128 .Ltmp89-.Lfunc_begin0          #     jumps to .Ltmp89
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp93-.Lfunc_begin0          # >> Call Site 24 <<
	.uleb128 .Ltmp102-.Ltmp93               #   Call between .Ltmp93 and .Ltmp102
	.uleb128 .Ltmp103-.Lfunc_begin0         #     jumps to .Ltmp103
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp104-.Lfunc_begin0         # >> Call Site 25 <<
	.uleb128 .Ltmp105-.Ltmp104              #   Call between .Ltmp104 and .Ltmp105
	.uleb128 .Ltmp106-.Lfunc_begin0         #     jumps to .Ltmp106
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp107-.Lfunc_begin0         # >> Call Site 26 <<
	.uleb128 .Ltmp114-.Ltmp107              #   Call between .Ltmp107 and .Ltmp114
	.uleb128 .Ltmp115-.Lfunc_begin0         #     jumps to .Ltmp115
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp116-.Lfunc_begin0         # >> Call Site 27 <<
	.uleb128 .Ltmp119-.Ltmp116              #   Call between .Ltmp116 and .Ltmp119
	.uleb128 .Ltmp120-.Lfunc_begin0         #     jumps to .Ltmp120
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp119-.Lfunc_begin0         # >> Call Site 28 <<
	.uleb128 .Lfunc_end10-.Ltmp119          #   Call between .Ltmp119 and .Lfunc_end10
	.byte	0                               #     has no landing pad
	.byte	0                               #   On action: cleanup
.Lcst_end0:
	.p2align	2, 0x0
                                        # -- End function
	.section	.text._ZNSt6vectorIiSaIiEE6resizeEm,"axG",@progbits,_ZNSt6vectorIiSaIiEE6resizeEm,comdat
	.weak	_ZNSt6vectorIiSaIiEE6resizeEm   # -- Begin function _ZNSt6vectorIiSaIiEE6resizeEm
	.p2align	1, 0x90
	.type	_ZNSt6vectorIiSaIiEE6resizeEm,@function
_ZNSt6vectorIiSaIiEE6resizeEm:          # @_ZNSt6vectorIiSaIiEE6resizeEm
	.cfi_startproc
# %bb.0:
	movq	(%rdi), %rdx
	movq	8(%rdi), %rcx
	movq	%rcx, %r8
	subq	%rdx, %r8
	sarq	$2, %r8
	movq	%rsi, %rax
	subq	%r8, %rax
	jbe	.LBB11_1
# %bb.5:
	movq	%rax, %rsi
	jmp	_ZNSt6vectorIiSaIiEE17_M_default_appendEm # TAILCALL
.LBB11_1:
	jae	.LBB11_4
# %bb.2:
	leaq	(%rdx,%rsi,4), %rax
	cmpq	%rax, %rcx
	je	.LBB11_4
# %bb.3:
	movq	%rax, 8(%rdi)
.LBB11_4:                               # %_ZNSt6vectorIiSaIiEE15_M_erase_at_endEPi.exit
	retq
.Lfunc_end11:
	.size	_ZNSt6vectorIiSaIiEE6resizeEm, .Lfunc_end11-_ZNSt6vectorIiSaIiEE6resizeEm
	.cfi_endproc
                                        # -- End function
	.text
	.globl	main                            # -- Begin function main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%rax
	.cfi_def_cfa_offset 16
	movl	$.L.str.6, %edi
	callq	_Z11myCountTestPKc
	xorl	%eax, %eax
	popq	%rcx
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end12:
	.size	main, .Lfunc_end12-main
	.cfi_endproc
                                        # -- End function
	.section	.text._ZNSt6vectorIiSaIiEE17_M_realloc_insertIJRKiEEEvN9__gnu_cxx17__normal_iteratorIPiS1_EEDpOT_,"axG",@progbits,_ZNSt6vectorIiSaIiEE17_M_realloc_insertIJRKiEEEvN9__gnu_cxx17__normal_iteratorIPiS1_EEDpOT_,comdat
	.weak	_ZNSt6vectorIiSaIiEE17_M_realloc_insertIJRKiEEEvN9__gnu_cxx17__normal_iteratorIPiS1_EEDpOT_ # -- Begin function _ZNSt6vectorIiSaIiEE17_M_realloc_insertIJRKiEEEvN9__gnu_cxx17__normal_iteratorIPiS1_EEDpOT_
	.p2align	1, 0x90
	.type	_ZNSt6vectorIiSaIiEE17_M_realloc_insertIJRKiEEEvN9__gnu_cxx17__normal_iteratorIPiS1_EEDpOT_,@function
_ZNSt6vectorIiSaIiEE17_M_realloc_insertIJRKiEEEvN9__gnu_cxx17__normal_iteratorIPiS1_EEDpOT_: # @_ZNSt6vectorIiSaIiEE17_M_realloc_insertIJRKiEEEvN9__gnu_cxx17__normal_iteratorIPiS1_EEDpOT_
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	subq	$24, %rsp
	.cfi_def_cfa_offset 80
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movq	%rdx, (%rsp)                    # 8-byte Spill
	movq	%rsi, %r12
	movq	%rdi, %rbx
	movl	$1, %esi
	movl	$.L.str.8, %edx
	callq	_ZNKSt6vectorIiSaIiEE12_M_check_lenEmPKc
	movq	(%rbx), %r13
	movq	8(%rbx), %r15
	movq	%r12, 8(%rsp)                   # 8-byte Spill
	subq	%r13, %r12
	movq	%r12, %r14
	sarq	$2, %r14
	testq	%rax, %rax
	movq	%rax, 16(%rsp)                  # 8-byte Spill
	je	.LBB13_1
# %bb.2:
	movq	%rbx, %rdi
	movq	%rax, %rsi
	xorl	%edx, %edx
	callq	_ZNSt15__new_allocatorIiE8allocateEmPKv
	movq	%rax, %rbp
	jmp	.LBB13_3
.LBB13_1:
	xorl	%ebp, %ebp
.LBB13_3:                               # %_ZNSt12_Vector_baseIiSaIiEE11_M_allocateEm.exit
	movq	(%rsp), %rax                    # 8-byte Reload
	movl	(%rax), %eax
	movl	%eax, (%rbp,%r14,4)
	testq	%r12, %r12
	jle	.LBB13_5
# %bb.4:
	movq	%rbp, %rdi
	movq	%r13, %rsi
	movq	%r12, %rdx
	callq	memmove@PLT
.LBB13_5:                               # %_ZNSt6vectorIiSaIiEE11_S_relocateEPiS2_S2_RS0_.exit
	addq	%rbp, %r12
	addq	$4, %r12
	movq	8(%rsp), %rsi                   # 8-byte Reload
	subq	%rsi, %r15
	testq	%r15, %r15
	jle	.LBB13_7
# %bb.6:
	movq	%r12, %rdi
	movq	%r15, %rdx
	callq	memmove@PLT
.LBB13_7:                               # %_ZNSt6vectorIiSaIiEE11_S_relocateEPiS2_S2_RS0_.exit16
	testq	%r13, %r13
	je	.LBB13_9
# %bb.8:
	movq	%r13, %rdi
	callq	_ZdlPv
.LBB13_9:                               # %_ZNSt12_Vector_baseIiSaIiEE13_M_deallocateEPim.exit
	addq	%r15, %r12
	movq	%rbp, (%rbx)
	movq	%r12, 8(%rbx)
	movq	16(%rsp), %rax                  # 8-byte Reload
	leaq	(,%rax,4), %rax
	addq	%rbp, %rax
	movq	%rax, 16(%rbx)
	addq	$24, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end13:
	.size	_ZNSt6vectorIiSaIiEE17_M_realloc_insertIJRKiEEEvN9__gnu_cxx17__normal_iteratorIPiS1_EEDpOT_, .Lfunc_end13-_ZNSt6vectorIiSaIiEE17_M_realloc_insertIJRKiEEEvN9__gnu_cxx17__normal_iteratorIPiS1_EEDpOT_
	.cfi_endproc
                                        # -- End function
	.section	.text._ZNKSt6vectorIiSaIiEE12_M_check_lenEmPKc,"axG",@progbits,_ZNKSt6vectorIiSaIiEE12_M_check_lenEmPKc,comdat
	.weak	_ZNKSt6vectorIiSaIiEE12_M_check_lenEmPKc # -- Begin function _ZNKSt6vectorIiSaIiEE12_M_check_lenEmPKc
	.p2align	1, 0x90
	.type	_ZNKSt6vectorIiSaIiEE12_M_check_lenEmPKc,@function
_ZNKSt6vectorIiSaIiEE12_M_check_lenEmPKc: # @_ZNKSt6vectorIiSaIiEE12_M_check_lenEmPKc
	.cfi_startproc
# %bb.0:
	movabsq	$2305843009213693951, %rcx      # imm = 0x1FFFFFFFFFFFFFFF
	movq	8(%rdi), %r8
	subq	(%rdi), %r8
	sarq	$2, %r8
	movq	%rcx, %rax
	subq	%r8, %rax
	cmpq	%rsi, %rax
	jb	.LBB14_2
# %bb.1:
	cmpq	%rsi, %r8
	cmovaq	%r8, %rsi
	leaq	(%rsi,%r8), %rax
	cmpq	%rcx, %rax
	cmovaeq	%rcx, %rax
	addq	%r8, %rsi
	cmovbq	%rcx, %rax
	retq
.LBB14_2:
	pushq	%rax
	.cfi_def_cfa_offset 16
	movq	%rdx, %rdi
	callq	_ZSt20__throw_length_errorPKc
.Lfunc_end14:
	.size	_ZNKSt6vectorIiSaIiEE12_M_check_lenEmPKc, .Lfunc_end14-_ZNKSt6vectorIiSaIiEE12_M_check_lenEmPKc
	.cfi_endproc
                                        # -- End function
	.section	.text._ZNSt15__new_allocatorIiE8allocateEmPKv,"axG",@progbits,_ZNSt15__new_allocatorIiE8allocateEmPKv,comdat
	.weak	_ZNSt15__new_allocatorIiE8allocateEmPKv # -- Begin function _ZNSt15__new_allocatorIiE8allocateEmPKv
	.p2align	1, 0x90
	.type	_ZNSt15__new_allocatorIiE8allocateEmPKv,@function
_ZNSt15__new_allocatorIiE8allocateEmPKv: # @_ZNSt15__new_allocatorIiE8allocateEmPKv
	.cfi_startproc
# %bb.0:
	pushq	%rax
	.cfi_def_cfa_offset 16
	movq	%rsi, %rax
	shrq	$61, %rax
	jne	.LBB15_1
# %bb.3:
	shlq	$2, %rsi
	movq	%rsi, %rdi
	popq	%rax
	.cfi_def_cfa_offset 8
	jmp	_Znwm                           # TAILCALL
.LBB15_1:
	.cfi_def_cfa_offset 16
	shrq	$62, %rsi
	je	.LBB15_2
# %bb.4:
	callq	_ZSt28__throw_bad_array_new_lengthv
.LBB15_2:
	callq	_ZSt17__throw_bad_allocv
.Lfunc_end15:
	.size	_ZNSt15__new_allocatorIiE8allocateEmPKv, .Lfunc_end15-_ZNSt15__new_allocatorIiE8allocateEmPKv
	.cfi_endproc
                                        # -- End function
	.section	.text._ZNSt6vectorIiSaIiEE17_M_default_appendEm,"axG",@progbits,_ZNSt6vectorIiSaIiEE17_M_default_appendEm,comdat
	.weak	_ZNSt6vectorIiSaIiEE17_M_default_appendEm # -- Begin function _ZNSt6vectorIiSaIiEE17_M_default_appendEm
	.p2align	1, 0x90
	.type	_ZNSt6vectorIiSaIiEE17_M_default_appendEm,@function
_ZNSt6vectorIiSaIiEE17_M_default_appendEm: # @_ZNSt6vectorIiSaIiEE17_M_default_appendEm
	.cfi_startproc
# %bb.0:
	testq	%rsi, %rsi
	je	.LBB16_16
# %bb.1:
	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
	pushq	%rax
	.cfi_def_cfa_offset 64
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movq	%rsi, %r14
	movq	%rdi, %rbx
	movq	8(%rdi), %r15
	movq	16(%rdi), %rax
	subq	%r15, %rax
	sarq	$2, %rax
	cmpq	%rsi, %rax
	jae	.LBB16_2
# %bb.5:
	movq	(%rbx), %r12
	subq	%r12, %r15
	movq	%r15, %r13
	sarq	$2, %r13
	movl	$.L.str.10, %edx
	movq	%rbx, %rdi
	movq	%r14, %rsi
	callq	_ZNKSt6vectorIiSaIiEE12_M_check_lenEmPKc
	testq	%rax, %rax
	movq	%rax, (%rsp)                    # 8-byte Spill
	je	.LBB16_6
# %bb.7:
	movq	%rbx, %rdi
	movq	%rax, %rsi
	xorl	%edx, %edx
	callq	_ZNSt15__new_allocatorIiE8allocateEmPKv
	movq	%rax, %rbp
	jmp	.LBB16_8
.LBB16_2:
	movl	$0, (%r15)
	leaq	4(%r15), %rdi
	cmpq	$1, %r14
	je	.LBB16_4
# %bb.3:                                # %_ZSt6fill_nIPimiET_S1_T0_RKT1_.exit.loopexit.i.i.i
	leaq	-4(,%r14,4), %rdx
	xorl	%esi, %esi
	callq	memset@PLT
	leaq	(%r15,%r14,4), %rdi
.LBB16_4:                               # %_ZSt27__uninitialized_default_n_aIPimiET_S1_T0_RSaIT1_E.exit
	movq	%rdi, 8(%rbx)
	jmp	.LBB16_15
.LBB16_6:
	xorl	%ebp, %ebp
.LBB16_8:
	leaq	(,%r13,4), %r13
	addq	%rbp, %r13
	movl	$0, (%r13)
	cmpq	$1, %r14
	je	.LBB16_10
# %bb.9:                                # %_ZSt6fill_nIPimiET_S1_T0_RKT1_.exit.loopexit.i.i.i30
	leaq	4(%r13), %rdi
	leaq	-4(,%r14,4), %rdx
	xorl	%esi, %esi
	callq	memset@PLT
.LBB16_10:                              # %_ZSt27__uninitialized_default_n_aIPimiET_S1_T0_RSaIT1_E.exit32
	testq	%r15, %r15
	jle	.LBB16_12
# %bb.11:
	movq	%rbp, %rdi
	movq	%r12, %rsi
	movq	%r15, %rdx
	callq	memmove@PLT
.LBB16_12:                              # %_ZNSt6vectorIiSaIiEE11_S_relocateEPiS2_S2_RS0_.exit
	testq	%r12, %r12
	je	.LBB16_14
# %bb.13:
	movq	%r12, %rdi
	callq	_ZdlPv
.LBB16_14:                              # %_ZNSt12_Vector_baseIiSaIiEE13_M_deallocateEPim.exit35
	movq	%rbp, (%rbx)
	leaq	(,%r14,4), %rax
	addq	%r13, %rax
	movq	%rax, 8(%rbx)
	movq	(%rsp), %rax                    # 8-byte Reload
	leaq	(,%rax,4), %rax
	addq	%rbp, %rax
	movq	%rax, 16(%rbx)
.LBB16_15:
	addq	$8, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	.cfi_restore %rbx
	.cfi_restore %r12
	.cfi_restore %r13
	.cfi_restore %r14
	.cfi_restore %r15
	.cfi_restore %rbp
.LBB16_16:
	retq
.Lfunc_end16:
	.size	_ZNSt6vectorIiSaIiEE17_M_default_appendEm, .Lfunc_end16-_ZNSt6vectorIiSaIiEE17_M_default_appendEm
	.cfi_endproc
                                        # -- End function
	.section	.text._ZNSt6vectorIiSaIiEE17_M_realloc_insertIJiEEEvN9__gnu_cxx17__normal_iteratorIPiS1_EEDpOT_,"axG",@progbits,_ZNSt6vectorIiSaIiEE17_M_realloc_insertIJiEEEvN9__gnu_cxx17__normal_iteratorIPiS1_EEDpOT_,comdat
	.weak	_ZNSt6vectorIiSaIiEE17_M_realloc_insertIJiEEEvN9__gnu_cxx17__normal_iteratorIPiS1_EEDpOT_ # -- Begin function _ZNSt6vectorIiSaIiEE17_M_realloc_insertIJiEEEvN9__gnu_cxx17__normal_iteratorIPiS1_EEDpOT_
	.p2align	1, 0x90
	.type	_ZNSt6vectorIiSaIiEE17_M_realloc_insertIJiEEEvN9__gnu_cxx17__normal_iteratorIPiS1_EEDpOT_,@function
_ZNSt6vectorIiSaIiEE17_M_realloc_insertIJiEEEvN9__gnu_cxx17__normal_iteratorIPiS1_EEDpOT_: # @_ZNSt6vectorIiSaIiEE17_M_realloc_insertIJiEEEvN9__gnu_cxx17__normal_iteratorIPiS1_EEDpOT_
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	subq	$24, %rsp
	.cfi_def_cfa_offset 80
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movq	%rdx, (%rsp)                    # 8-byte Spill
	movq	%rsi, %r12
	movq	%rdi, %rbx
	movl	$1, %esi
	movl	$.L.str.8, %edx
	callq	_ZNKSt6vectorIiSaIiEE12_M_check_lenEmPKc
	movq	(%rbx), %r13
	movq	8(%rbx), %r15
	movq	%r12, 8(%rsp)                   # 8-byte Spill
	subq	%r13, %r12
	movq	%r12, %r14
	sarq	$2, %r14
	testq	%rax, %rax
	movq	%rax, 16(%rsp)                  # 8-byte Spill
	je	.LBB17_1
# %bb.2:
	movq	%rbx, %rdi
	movq	%rax, %rsi
	xorl	%edx, %edx
	callq	_ZNSt15__new_allocatorIiE8allocateEmPKv
	movq	%rax, %rbp
	jmp	.LBB17_3
.LBB17_1:
	xorl	%ebp, %ebp
.LBB17_3:                               # %_ZNSt12_Vector_baseIiSaIiEE11_M_allocateEm.exit
	movq	(%rsp), %rax                    # 8-byte Reload
	movl	(%rax), %eax
	movl	%eax, (%rbp,%r14,4)
	testq	%r12, %r12
	jle	.LBB17_5
# %bb.4:
	movq	%rbp, %rdi
	movq	%r13, %rsi
	movq	%r12, %rdx
	callq	memmove@PLT
.LBB17_5:                               # %_ZNSt6vectorIiSaIiEE11_S_relocateEPiS2_S2_RS0_.exit
	addq	%rbp, %r12
	addq	$4, %r12
	movq	8(%rsp), %rsi                   # 8-byte Reload
	subq	%rsi, %r15
	testq	%r15, %r15
	jle	.LBB17_7
# %bb.6:
	movq	%r12, %rdi
	movq	%r15, %rdx
	callq	memmove@PLT
.LBB17_7:                               # %_ZNSt6vectorIiSaIiEE11_S_relocateEPiS2_S2_RS0_.exit16
	testq	%r13, %r13
	je	.LBB17_9
# %bb.8:
	movq	%r13, %rdi
	callq	_ZdlPv
.LBB17_9:                               # %_ZNSt12_Vector_baseIiSaIiEE13_M_deallocateEPim.exit
	addq	%r15, %r12
	movq	%rbp, (%rbx)
	movq	%r12, 8(%rbx)
	movq	16(%rsp), %rax                  # 8-byte Reload
	leaq	(,%rax,4), %rax
	addq	%rbp, %rax
	movq	%rax, 16(%rbx)
	addq	$24, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end17:
	.size	_ZNSt6vectorIiSaIiEE17_M_realloc_insertIJiEEEvN9__gnu_cxx17__normal_iteratorIPiS1_EEDpOT_, .Lfunc_end17-_ZNSt6vectorIiSaIiEE17_M_realloc_insertIJiEEEvN9__gnu_cxx17__normal_iteratorIPiS1_EEDpOT_
	.cfi_endproc
                                        # -- End function
	.text
	.p2align	4, 0x90                         # -- Begin function __hip_module_ctor
	.type	__hip_module_ctor,@function
__hip_module_ctor:                      # @__hip_module_ctor
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	subq	$32, %rsp
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -16
	movq	__hip_gpubin_handle(%rip), %rbx
	testq	%rbx, %rbx
	jne	.LBB18_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB18_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z4scanPiS_i, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z14countGlobalMemIiLi512EEvPT_Pii, %esi
	movl	$.L__unnamed_2, %edx
	movl	$.L__unnamed_2, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z14countGlobalMemIfLi512EEvPT_Pii, %esi
	movl	$.L__unnamed_3, %edx
	movl	$.L__unnamed_3, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z14countGlobalMemIdLi512EEvPT_Pii, %esi
	movl	$.L__unnamed_4, %edx
	movl	$.L__unnamed_4, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z14countSharedMemIiLi512EEvPT_Piiii, %esi
	movl	$.L__unnamed_5, %edx
	movl	$.L__unnamed_5, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end18:
	.size	__hip_module_ctor, .Lfunc_end18-__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	.LBB19_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
.LBB19_2:
	retq
.Lfunc_end19:
	.size	__hip_module_dtor, .Lfunc_end19-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_Z4scanPiS_i,@object            # @_Z4scanPiS_i
	.section	.rodata,"a",@progbits
	.globl	_Z4scanPiS_i
	.p2align	3, 0x0
_Z4scanPiS_i:
	.quad	_Z19__device_stub__scanPiS_i
	.size	_Z4scanPiS_i, 8

	.type	_Z14countGlobalMemIiLi512EEvPT_Pii,@object # @_Z14countGlobalMemIiLi512EEvPT_Pii
	.section	.rodata._Z14countGlobalMemIiLi512EEvPT_Pii,"aG",@progbits,_Z14countGlobalMemIiLi512EEvPT_Pii,comdat
	.weak	_Z14countGlobalMemIiLi512EEvPT_Pii
	.p2align	3, 0x0
_Z14countGlobalMemIiLi512EEvPT_Pii:
	.quad	_Z29__device_stub__countGlobalMemIiLi512EEvPT_Pii
	.size	_Z14countGlobalMemIiLi512EEvPT_Pii, 8

	.type	_Z14countGlobalMemIfLi512EEvPT_Pii,@object # @_Z14countGlobalMemIfLi512EEvPT_Pii
	.section	.rodata._Z14countGlobalMemIfLi512EEvPT_Pii,"aG",@progbits,_Z14countGlobalMemIfLi512EEvPT_Pii,comdat
	.weak	_Z14countGlobalMemIfLi512EEvPT_Pii
	.p2align	3, 0x0
_Z14countGlobalMemIfLi512EEvPT_Pii:
	.quad	_Z29__device_stub__countGlobalMemIfLi512EEvPT_Pii
	.size	_Z14countGlobalMemIfLi512EEvPT_Pii, 8

	.type	_Z14countGlobalMemIdLi512EEvPT_Pii,@object # @_Z14countGlobalMemIdLi512EEvPT_Pii
	.section	.rodata._Z14countGlobalMemIdLi512EEvPT_Pii,"aG",@progbits,_Z14countGlobalMemIdLi512EEvPT_Pii,comdat
	.weak	_Z14countGlobalMemIdLi512EEvPT_Pii
	.p2align	3, 0x0
_Z14countGlobalMemIdLi512EEvPT_Pii:
	.quad	_Z29__device_stub__countGlobalMemIdLi512EEvPT_Pii
	.size	_Z14countGlobalMemIdLi512EEvPT_Pii, 8

	.type	_Z14countSharedMemIiLi512EEvPT_Piiii,@object # @_Z14countSharedMemIiLi512EEvPT_Piiii
	.section	.rodata._Z14countSharedMemIiLi512EEvPT_Piiii,"aG",@progbits,_Z14countSharedMemIiLi512EEvPT_Piiii,comdat
	.weak	_Z14countSharedMemIiLi512EEvPT_Piiii
	.p2align	3, 0x0
_Z14countSharedMemIiLi512EEvPT_Piiii:
	.quad	_Z29__device_stub__countSharedMemIiLi512EEvPT_Piiii
	.size	_Z14countSharedMemIiLi512EEvPT_Piiii, 8

	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"q2a.txt"
	.size	.L.str, 8

	.type	.L.str.1,@object                # @.str.1
.L.str.1:
	.asciz	"\n Count["
	.size	.L.str.1, 9

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

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	"q2b.txt"
	.size	.L.str.3, 8

	.type	.L.str.4,@object                # @.str.4
.L.str.4:
	.asciz	"q2c.txt"
	.size	.L.str.4, 8

	.type	.L.str.5,@object                # @.str.5
.L.str.5:
	.asciz	"\n prescan["
	.size	.L.str.5, 11

	.type	.L.str.6,@object                # @.str.6
.L.str.6:
	.asciz	"inp.txt"
	.size	.L.str.6, 8

	.type	.L.str.8,@object                # @.str.8
.L.str.8:
	.asciz	"vector::_M_realloc_insert"
	.size	.L.str.8, 26

	.type	.L.str.10,@object               # @.str.10
.L.str.10:
	.asciz	"vector::_M_default_append"
	.size	.L.str.10, 26

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z4scanPiS_i"
	.size	.L__unnamed_1, 13

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"_Z14countGlobalMemIiLi512EEvPT_Pii"
	.size	.L__unnamed_2, 35

	.type	.L__unnamed_3,@object           # @2
.L__unnamed_3:
	.asciz	"_Z14countGlobalMemIfLi512EEvPT_Pii"
	.size	.L__unnamed_3, 35

	.type	.L__unnamed_4,@object           # @3
.L__unnamed_4:
	.asciz	"_Z14countGlobalMemIdLi512EEvPT_Pii"
	.size	.L__unnamed_4, 35

	.type	.L__unnamed_5,@object           # @4
.L__unnamed_5:
	.asciz	"_Z14countSharedMemIiLi512EEvPT_Piiii"
	.size	.L__unnamed_5, 37

	.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 _Z19__device_stub__scanPiS_i
	.addrsig_sym _Z29__device_stub__countGlobalMemIiLi512EEvPT_Pii
	.addrsig_sym _Z29__device_stub__countGlobalMemIfLi512EEvPT_Pii
	.addrsig_sym _Z29__device_stub__countGlobalMemIdLi512EEvPT_Pii
	.addrsig_sym _Z29__device_stub__countSharedMemIiLi512EEvPT_Piiii
	.addrsig_sym __gxx_personality_v0
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Unwind_Resume
	.addrsig_sym _Z4scanPiS_i
	.addrsig_sym _Z14countGlobalMemIiLi512EEvPT_Pii
	.addrsig_sym _Z14countGlobalMemIfLi512EEvPT_Pii
	.addrsig_sym _Z14countGlobalMemIdLi512EEvPT_Pii
	.addrsig_sym _Z14countSharedMemIiLi512EEvPT_Piiii
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z4scanPiS_i            ; -- Begin function _Z4scanPiS_i
	.globl	_Z4scanPiS_i
	.p2align	8
	.type	_Z4scanPiS_i,@function
_Z4scanPiS_i:                           ; @_Z4scanPiS_i
; %bb.0:
	s_clause 0x2
	s_load_b32 s4, s[0:1], 0x24
	s_load_b32 s5, s[0:1], 0x10
	s_load_b128 s[0:3], s[0:1], 0x0
	v_lshl_add_u32 v3, v0, 2, 0
	s_waitcnt lgkmcnt(0)
	s_and_b32 s4, s4, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[1:2], null, s15, s4, v[0:1]
	v_cmp_gt_i32_e32 vcc_lo, s5, v1
	s_and_saveexec_b32 s4, vcc_lo
	s_cbranch_execz .LBB0_2
; %bb.1:
	v_ashrrev_i32_e32 v2, 31, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[1:2], 2, v[1:2]
	v_add_co_u32 v1, s0, s0, v1
	s_delay_alu instid0(VALU_DEP_1)
	v_add_co_ci_u32_e64 v2, s0, s1, v2, s0
	global_load_b32 v1, v[1:2], off
	s_waitcnt vmcnt(0)
	ds_store_b32 v3, v1
.LBB0_2:
	s_or_b32 exec_lo, exec_lo, s4
	v_cmp_ne_u32_e64 s0, 0, v0
	s_mov_b32 s4, 0
	s_delay_alu instid0(VALU_DEP_1)
	s_and_saveexec_b32 s1, s0
	s_cbranch_execz .LBB0_5
; %bb.3:                                ; %.lr.ph
	s_mov_b32 s5, 1
.LBB0_4:                                ; =>This Inner Loop Header: Depth=1
	s_delay_alu instid0(SALU_CYCLE_1)
	v_subrev_nc_u32_e32 v1, s5, v0
	s_waitcnt lgkmcnt(0)
	s_barrier
	buffer_gl0_inv
	s_lshl_b32 s5, s5, 1
	v_lshl_add_u32 v1, v1, 2, 0
	v_cmp_gt_u32_e64 s0, s5, v0
	ds_load_b32 v1, v1
	s_waitcnt lgkmcnt(0)
	s_barrier
	buffer_gl0_inv
	ds_load_b32 v2, v3
	s_or_b32 s4, s0, s4
	s_waitcnt lgkmcnt(0)
	v_add_nc_u32_e32 v1, v2, v1
	ds_store_b32 v3, v1
	s_and_not1_b32 exec_lo, exec_lo, s4
	s_cbranch_execnz .LBB0_4
.LBB0_5:                                ; %Flow22
	s_or_b32 exec_lo, exec_lo, s1
	s_waitcnt lgkmcnt(0)
	s_barrier
	buffer_gl0_inv
	s_and_saveexec_b32 s0, vcc_lo
	s_cbranch_execz .LBB0_7
; %bb.6:
	v_lshlrev_b32_e32 v0, 2, v0
	s_delay_alu instid0(VALU_DEP_1)
	v_add_nc_u32_e32 v1, 0, v0
	ds_load_b32 v1, v1
	s_waitcnt lgkmcnt(0)
	global_store_b32 v0, v1, s[2:3]
.LBB0_7:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z4scanPiS_i
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 280
		.amdhsa_user_sgpr_count 15
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 4
		.amdhsa_next_free_sgpr 16
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end0:
	.size	_Z4scanPiS_i, .Lfunc_end0-_Z4scanPiS_i
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 336
; NumSgprs: 18
; NumVgprs: 4
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 4
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.section	.text._Z14countGlobalMemIiLi512EEvPT_Pii,"axG",@progbits,_Z14countGlobalMemIiLi512EEvPT_Pii,comdat
	.protected	_Z14countGlobalMemIiLi512EEvPT_Pii ; -- Begin function _Z14countGlobalMemIiLi512EEvPT_Pii
	.globl	_Z14countGlobalMemIiLi512EEvPT_Pii
	.p2align	8
	.type	_Z14countGlobalMemIiLi512EEvPT_Pii,@function
_Z14countGlobalMemIiLi512EEvPT_Pii:     ; @_Z14countGlobalMemIiLi512EEvPT_Pii
; %bb.0:
	s_load_b32 s2, s[0:1], 0x10
	v_lshl_add_u32 v0, s15, 9, v0
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_1)
	v_cmp_gt_u32_e32 vcc_lo, s2, v0
	s_and_saveexec_b32 s2, vcc_lo
	s_cbranch_execz .LBB1_4
; %bb.1:
	s_load_b128 s[0:3], s[0:1], 0x0
	v_mov_b32_e32 v1, 0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[0:1], 2, v[0:1]
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v0, vcc_lo, s0, v0
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v1, vcc_lo, s1, v1, vcc_lo
	s_mov_b32 s0, exec_lo
	global_load_b32 v2, v[0:1], off
	v_mov_b32_e32 v0, 9
	v_mov_b32_e32 v1, 0
	s_waitcnt vmcnt(0)
	v_cmpx_ne_u32_e32 0x3e8, v2
; %bb.2:
	v_mul_hi_i32 v0, 0x51eb851f, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshrrev_b32_e32 v1, 31, v0
	v_ashrrev_i32_e32 v0, 5, v0
	v_add_nc_u32_e32 v0, v0, v1
	s_delay_alu instid0(VALU_DEP_1)
	v_ashrrev_i32_e32 v1, 31, v0
; %bb.3:                                ; %.sink.split
	s_or_b32 exec_lo, exec_lo, s0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_lshlrev_b64 v[0:1], 2, v[0:1]
	v_mov_b32_e32 v2, 1
	v_add_co_u32 v0, vcc_lo, s2, v0
	s_delay_alu instid0(VALU_DEP_3)
	v_add_co_ci_u32_e32 v1, vcc_lo, s3, v1, vcc_lo
	global_atomic_add_u32 v[0:1], v2, off
.LBB1_4:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z14countGlobalMemIiLi512EEvPT_Pii
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 20
		.amdhsa_user_sgpr_count 15
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 3
		.amdhsa_next_free_sgpr 16
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.section	.text._Z14countGlobalMemIiLi512EEvPT_Pii,"axG",@progbits,_Z14countGlobalMemIiLi512EEvPT_Pii,comdat
.Lfunc_end1:
	.size	_Z14countGlobalMemIiLi512EEvPT_Pii, .Lfunc_end1-_Z14countGlobalMemIiLi512EEvPT_Pii
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 204
; NumSgprs: 18
; NumVgprs: 3
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 3
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.section	.text._Z14countGlobalMemIfLi512EEvPT_Pii,"axG",@progbits,_Z14countGlobalMemIfLi512EEvPT_Pii,comdat
	.protected	_Z14countGlobalMemIfLi512EEvPT_Pii ; -- Begin function _Z14countGlobalMemIfLi512EEvPT_Pii
	.globl	_Z14countGlobalMemIfLi512EEvPT_Pii
	.p2align	8
	.type	_Z14countGlobalMemIfLi512EEvPT_Pii,@function
_Z14countGlobalMemIfLi512EEvPT_Pii:     ; @_Z14countGlobalMemIfLi512EEvPT_Pii
; %bb.0:
	s_load_b32 s2, s[0:1], 0x10
	v_lshl_add_u32 v0, s15, 9, v0
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_1)
	v_cmp_gt_u32_e32 vcc_lo, s2, v0
	s_and_saveexec_b32 s2, vcc_lo
	s_cbranch_execz .LBB2_4
; %bb.1:
	s_load_b128 s[0:3], s[0:1], 0x0
	v_mov_b32_e32 v1, 0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[0:1], 2, v[0:1]
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v0, vcc_lo, s0, v0
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v1, vcc_lo, s1, v1, vcc_lo
	s_mov_b32 s0, exec_lo
	global_load_b32 v2, v[0:1], off
	v_mov_b32_e32 v0, 9
	v_mov_b32_e32 v1, 0
	s_waitcnt vmcnt(0)
	v_cmpx_neq_f32_e32 0x447a0000, v2
; %bb.2:
	v_cvt_i32_f32_e32 v0, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_hi_i32 v0, 0x51eb851f, v0
	v_lshrrev_b32_e32 v1, 31, v0
	v_ashrrev_i32_e32 v0, 5, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v0, v0, v1
	v_ashrrev_i32_e32 v1, 31, v0
; %bb.3:                                ; %.sink.split
	s_or_b32 exec_lo, exec_lo, s0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_lshlrev_b64 v[0:1], 2, v[0:1]
	v_mov_b32_e32 v2, 1
	v_add_co_u32 v0, vcc_lo, s2, v0
	s_delay_alu instid0(VALU_DEP_3)
	v_add_co_ci_u32_e32 v1, vcc_lo, s3, v1, vcc_lo
	global_atomic_add_u32 v[0:1], v2, off
.LBB2_4:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z14countGlobalMemIfLi512EEvPT_Pii
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 20
		.amdhsa_user_sgpr_count 15
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 3
		.amdhsa_next_free_sgpr 16
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.section	.text._Z14countGlobalMemIfLi512EEvPT_Pii,"axG",@progbits,_Z14countGlobalMemIfLi512EEvPT_Pii,comdat
.Lfunc_end2:
	.size	_Z14countGlobalMemIfLi512EEvPT_Pii, .Lfunc_end2-_Z14countGlobalMemIfLi512EEvPT_Pii
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 208
; NumSgprs: 18
; NumVgprs: 3
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 3
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.section	.text._Z14countGlobalMemIdLi512EEvPT_Pii,"axG",@progbits,_Z14countGlobalMemIdLi512EEvPT_Pii,comdat
	.protected	_Z14countGlobalMemIdLi512EEvPT_Pii ; -- Begin function _Z14countGlobalMemIdLi512EEvPT_Pii
	.globl	_Z14countGlobalMemIdLi512EEvPT_Pii
	.p2align	8
	.type	_Z14countGlobalMemIdLi512EEvPT_Pii,@function
_Z14countGlobalMemIdLi512EEvPT_Pii:     ; @_Z14countGlobalMemIdLi512EEvPT_Pii
; %bb.0:
	s_load_b32 s2, s[0:1], 0x10
	v_lshl_add_u32 v0, s15, 9, v0
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_1)
	v_cmp_gt_u32_e32 vcc_lo, s2, v0
	s_and_saveexec_b32 s2, vcc_lo
	s_cbranch_execz .LBB3_4
; %bb.1:
	s_load_b128 s[0:3], s[0:1], 0x0
	v_mov_b32_e32 v1, 0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[0:1], 3, v[0:1]
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v0, vcc_lo, s0, v0
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v1, vcc_lo, s1, v1, vcc_lo
	s_mov_b32 s0, exec_lo
	global_load_b64 v[2:3], v[0:1], off
	v_mov_b32_e32 v0, 9
	v_mov_b32_e32 v1, 0
	s_waitcnt vmcnt(0)
	v_cmpx_neq_f64_e32 0x408f4000, v[2:3]
; %bb.2:
	v_cvt_i32_f64_e32 v0, v[2:3]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_hi_i32 v0, 0x51eb851f, v0
	v_lshrrev_b32_e32 v1, 31, v0
	v_ashrrev_i32_e32 v0, 5, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v0, v0, v1
	v_ashrrev_i32_e32 v1, 31, v0
; %bb.3:                                ; %.sink.split
	s_or_b32 exec_lo, exec_lo, s0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_lshlrev_b64 v[0:1], 2, v[0:1]
	v_mov_b32_e32 v2, 1
	v_add_co_u32 v0, vcc_lo, s2, v0
	s_delay_alu instid0(VALU_DEP_3)
	v_add_co_ci_u32_e32 v1, vcc_lo, s3, v1, vcc_lo
	global_atomic_add_u32 v[0:1], v2, off
.LBB3_4:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z14countGlobalMemIdLi512EEvPT_Pii
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 20
		.amdhsa_user_sgpr_count 15
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 4
		.amdhsa_next_free_sgpr 16
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.section	.text._Z14countGlobalMemIdLi512EEvPT_Pii,"axG",@progbits,_Z14countGlobalMemIdLi512EEvPT_Pii,comdat
.Lfunc_end3:
	.size	_Z14countGlobalMemIdLi512EEvPT_Pii, .Lfunc_end3-_Z14countGlobalMemIdLi512EEvPT_Pii
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 208
; NumSgprs: 18
; NumVgprs: 4
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 4
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.section	.text._Z14countSharedMemIiLi512EEvPT_Piiii,"axG",@progbits,_Z14countSharedMemIiLi512EEvPT_Piiii,comdat
	.protected	_Z14countSharedMemIiLi512EEvPT_Piiii ; -- Begin function _Z14countSharedMemIiLi512EEvPT_Piiii
	.globl	_Z14countSharedMemIiLi512EEvPT_Piiii
	.p2align	8
	.type	_Z14countSharedMemIiLi512EEvPT_Piiii,@function
_Z14countSharedMemIiLi512EEvPT_Piiii:   ; @_Z14countSharedMemIiLi512EEvPT_Piiii
; %bb.0:
	s_load_b256 s[0:7], s[0:1], 0x0
	s_waitcnt lgkmcnt(0)
	s_ashr_i32 s7, s6, 31
	s_ashr_i32 s10, s5, 31
	s_add_i32 s6, s6, s7
	s_add_i32 s11, s5, s10
	s_xor_b32 s6, s6, s7
	s_xor_b32 s11, s11, s10
	v_cvt_f32_u32_e32 v1, s6
	s_sub_i32 s8, 0, s6
	s_xor_b32 s10, s10, s7
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_iflag_f32_e32 v1, v1
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v1, 0x4f7ffffe, v1
	v_cvt_u32_f32_e32 v1, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_readfirstlane_b32 s9, v1
	v_lshl_add_u32 v1, s15, 9, v0
	s_mul_i32 s8, s8, s9
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	v_cmp_gt_u32_e32 vcc_lo, s4, v1
	s_mul_hi_u32 s8, s9, s8
	s_add_i32 s9, s9, s8
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_mul_hi_u32 s8, s11, s9
	s_mul_i32 s12, s8, s6
	s_delay_alu instid0(SALU_CYCLE_1)
	s_sub_i32 s11, s11, s12
	s_add_i32 s12, s8, 1
	s_sub_i32 s13, s11, s6
	s_cmp_ge_u32 s11, s6
	s_cselect_b32 s8, s12, s8
	s_cselect_b32 s11, s13, s11
	s_add_i32 s12, s8, 1
	s_cmp_ge_u32 s11, s6
	s_cselect_b32 s8, s12, s8
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_xor_b32 s4, s8, s10
	s_sub_i32 s4, s4, s10
	s_and_saveexec_b32 s8, vcc_lo
	s_cbranch_execz .LBB4_6
; %bb.1:
	v_mov_b32_e32 v2, 0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[1:2], 2, v[1:2]
	v_add_co_u32 v1, vcc_lo, s0, v1
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v2, vcc_lo, s1, v2, vcc_lo
	s_mov_b32 s0, exec_lo
	global_load_b32 v2, v[1:2], off
                                        ; implicit-def: $vgpr1
	s_waitcnt vmcnt(0)
	v_cmpx_ne_u32_e64 s5, v2
	s_xor_b32 s0, exec_lo, s0
	s_cbranch_execz .LBB4_3
; %bb.2:
	v_ashrrev_i32_e32 v1, 31, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v2, v2, v1
	v_xor_b32_e32 v2, v2, v1
	v_xor_b32_e32 v1, s7, v1
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_hi_u32 v3, v2, s9
	v_mul_lo_u32 v4, v3, s6
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_sub_nc_u32_e32 v2, v2, v4
	v_add_nc_u32_e32 v4, 1, v3
	v_subrev_nc_u32_e32 v5, s6, v2
	v_cmp_le_u32_e32 vcc_lo, s6, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_cndmask_b32 v3, v3, v4 :: v_dual_cndmask_b32 v2, v2, v5
	v_add_nc_u32_e32 v4, 1, v3
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cmp_le_u32_e32 vcc_lo, s6, v2
	v_cndmask_b32_e32 v2, v3, v4, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v2, v2, v1
	v_sub_nc_u32_e32 v1, v2, v1
.LBB4_3:                                ; %Flow
	s_and_not1_saveexec_b32 s0, s0
; %bb.4:
	s_add_i32 s1, s4, -1
	s_delay_alu instid0(SALU_CYCLE_1)
	v_mov_b32_e32 v1, s1
; %bb.5:                                ; %.sink.split
	s_or_b32 exec_lo, exec_lo, s0
	s_delay_alu instid0(VALU_DEP_1)
	v_lshl_add_u32 v1, v1, 2, 0
	v_mov_b32_e32 v2, 1
	ds_add_u32 v1, v2
.LBB4_6:                                ; %Flow37
	s_or_b32 exec_lo, exec_lo, s8
	s_waitcnt lgkmcnt(0)
	s_barrier
	buffer_gl0_inv
	s_mov_b32 s0, exec_lo
	v_cmpx_gt_u32_e64 s4, v0
	s_cbranch_execz .LBB4_8
; %bb.7:
	v_lshl_add_u32 v3, v0, 2, 0
	v_mad_u64_u32 v[1:2], null, s4, s15, v[0:1]
	v_mov_b32_e32 v2, 0
	ds_load_b32 v3, v3
	v_lshlrev_b64 v[0:1], 2, v[1:2]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v0, vcc_lo, s2, v0
	v_add_co_ci_u32_e32 v1, vcc_lo, s3, v1, vcc_lo
	s_waitcnt lgkmcnt(0)
	global_store_b32 v[0:1], v3, off
.LBB4_8:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z14countSharedMemIiLi512EEvPT_Piiii
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 28
		.amdhsa_user_sgpr_count 15
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 6
		.amdhsa_next_free_sgpr 16
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.section	.text._Z14countSharedMemIiLi512EEvPT_Piiii,"axG",@progbits,_Z14countSharedMemIiLi512EEvPT_Piiii,comdat
.Lfunc_end4:
	.size	_Z14countSharedMemIiLi512EEvPT_Piiii, .Lfunc_end4-_Z14countSharedMemIiLi512EEvPT_Piiii
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 500
; NumSgprs: 18
; NumVgprs: 6
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 6
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym sdata
	.addrsig_sym __hip_cuid_
	.amdgpu_metadata
---
amdhsa.kernels:
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .offset:         16
        .size:           4
        .value_kind:     by_value
      - .offset:         24
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         28
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         32
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         36
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         38
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         40
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         42
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         44
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         46
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         64
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         88
        .size:           2
        .value_kind:     hidden_grid_dims
      - .offset:         144
        .size:           4
        .value_kind:     hidden_dynamic_lds_size
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 280
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z4scanPiS_i
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z4scanPiS_i.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     4
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .offset:         16
        .size:           4
        .value_kind:     by_value
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 20
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z14countGlobalMemIiLi512EEvPT_Pii
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z14countGlobalMemIiLi512EEvPT_Pii.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     3
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .offset:         16
        .size:           4
        .value_kind:     by_value
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 20
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z14countGlobalMemIfLi512EEvPT_Pii
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z14countGlobalMemIfLi512EEvPT_Pii.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     3
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .offset:         16
        .size:           4
        .value_kind:     by_value
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 20
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z14countGlobalMemIdLi512EEvPT_Pii
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z14countGlobalMemIdLi512EEvPT_Pii.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     4
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .offset:         16
        .size:           4
        .value_kind:     by_value
      - .offset:         20
        .size:           4
        .value_kind:     by_value
      - .offset:         24
        .size:           4
        .value_kind:     by_value
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 28
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z14countSharedMemIiLi512EEvPT_Piiii
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z14countSharedMemIiLi512EEvPT_Piiii.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     6
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
