#include "Data.cuh"
#include "Data_kernel.cuh"

#include <fstream>
#include <iostream>
#include <limits>

Data::Data():
	rowSize(0), columnSize(0)//, flow(0)
{
}


Data::~Data()
{
    delete[] this->weightLeft;
    delete[] this->weightRight;
    delete[] this->weightUp;
    delete[] this->weightDown;
    delete[] this->weightS;
    delete[] this->weightT;
    delete[] this->height;
    delete[] this->capacity;
//    delete[] this->bfsTag;
    cudaFree(device_active);
    cudaFree(device_weightLeft);
    cudaFree(device_weightRight);
    cudaFree(device_weightUp);
    cudaFree(device_weightDown);
    cudaFree(device_weightS);
    cudaFree(device_weightT);
    cudaFree(device_height);
    cudaFree(device_capacity);
}


#define DEBUG
#ifdef DEBUG
void Data::Print()
{
	std::cout << "Weight:" << std::endl << std::endl;

    std::cout << "Left:" << std::endl;
	for (int i = 0; i < this->rowSize; i++) {
		for (int j = 0; j < this->columnSize; j++) {
			size_t idx = (i * this->columnSize + j);
            std::cout << this->weightLeft[idx] << " ";
		}
        std::cout << std::endl;
	}
    std::cout << std::endl;
    
    std::cout << "Right:" << std::endl;
	for (int i = 0; i < this->rowSize; i++) {
		for (int j = 0; j < this->columnSize; j++) {
			size_t idx = (i * this->columnSize + j);
            std::cout << this->weightRight[idx] << " ";
		}
        std::cout << std::endl;
	}
    std::cout << std::endl;
    
    std::cout << "Up:" << std::endl;
	for (int i = 0; i < this->rowSize; i++) {
		for (int j = 0; j < this->columnSize; j++) {
			size_t idx = (i * this->columnSize + j);
            std::cout << this->weightUp[idx] << " ";
		}
        std::cout << std::endl;
	}
    std::cout << std::endl;
    
    std::cout << "Down:" << std::endl;
	for (int i = 0; i < this->rowSize; i++) {
		for (int j = 0; j < this->columnSize; j++) {
			size_t idx = (i * this->columnSize + j);
            std::cout << this->weightDown[idx] << " ";
		}
        std::cout << std::endl;
	}
    std::cout << std::endl;

	std::cout << "Capacity:" << std::endl;
	for (int i = 0; i < this->rowSize; i++) {
		for (int j = 0; j < this->columnSize; j++) {
			size_t idx = (i * this->columnSize + j);
			std::cout << this->capacity[idx] << " ";
        }
        std::cout << std::endl;
	}
    std::cout << std::endl;

	std::cout << "Height:" << std::endl;
	for (int i = 0; i < this->rowSize; i++) {
		for (int j = 0; j < this->columnSize; j++) {
			size_t idx = (i * this->columnSize + j);
			std::cout << this->height[idx] << " ";
		}
        std::cout << std::endl;
	}
    std::cout << std::endl;
    std::cout << std::endl;
}
#endif


Data::Data(size_t rowSize, size_t columnSize) :
	rowSize(rowSize), columnSize(columnSize)//, flow(0)
{
	//weight = new int[4 * rowSize * columnSize](); // (): set to zero
    weightLeft = new int[rowSize * columnSize]();
    weightRight = new int[rowSize * columnSize]();
    weightUp = new int[rowSize * columnSize]();
    weightDown = new int[rowSize * columnSize]();
	weightS = new int[rowSize * columnSize]();
	weightT = new int[rowSize * columnSize]();
	height = new int[rowSize * columnSize]();
	capacity = new int[rowSize * columnSize]();

	//int *bfsTag;

    cudaMalloc( &device_active, sizeof(bool) );
    cudaMalloc( &device_weightLeft, sizeof(int) * rowSize * columnSize );
    cudaMalloc( &device_weightRight, sizeof(int) * rowSize * columnSize );
    cudaMalloc( &device_weightUp, sizeof(int) * rowSize * columnSize );
    cudaMalloc( &device_weightDown, sizeof(int) * rowSize * columnSize );
    cudaMalloc( &device_weightS, sizeof(int) * rowSize * columnSize );
    cudaMalloc( &device_weightT, sizeof(int) * rowSize * columnSize );
    cudaMalloc( &device_height, sizeof(int) * rowSize * columnSize );
    cudaMalloc( &device_capacity, sizeof(int) * rowSize * columnSize );

    cudaMemset( device_active, false, sizeof(bool) );
    cudaMemset( device_height, 0, sizeof(int) * rowSize * columnSize );
    cudaMemset( device_capacity, 0, sizeof(int) * rowSize * columnSize );


}


void Data::Read(const char * filename)
{
	using std::fstream;
	fstream fin;
	fin.open(filename, fstream::in);
	fin >> this->rowSize >> this->columnSize;

    this->weightLeft = new int[this->rowSize * this->columnSize]();
    this->weightRight = new int[this->rowSize * this->columnSize]();
    this->weightUp = new int[this->rowSize * this->columnSize]();
    this->weightDown = new int[this->rowSize * this->columnSize]();
	this->weightS = new int[this->rowSize * this->columnSize]();
	this->weightT = new int[this->rowSize * this->columnSize]();
	this->height = new int[this->rowSize * this->columnSize]();
	this->capacity = new int[this->rowSize * this->columnSize]();
    //this->bfsTag = new int[this->rowSize * this->columnSize]();

	for (int i = 0; i < this->rowSize * this->columnSize; i++) {
		// order: s, t, left, right, up, down
		fin >> this->weightS[i] >> this->weightT[i];
        fin >> this->weightLeft[i] >> this->weightRight[i] >> this->weightUp[i] >> this->weightDown[i];
	}
	fin.close();
    
    
    cudaMalloc( &device_active, sizeof(bool) );
    cudaMalloc( &device_weightLeft, sizeof(int) * rowSize * columnSize );
    cudaMalloc( &device_weightRight, sizeof(int) * rowSize * columnSize );
    cudaMalloc( &device_weightUp, sizeof(int) * rowSize * columnSize );
    cudaMalloc( &device_weightDown, sizeof(int) * rowSize * columnSize );
    cudaMalloc( &device_weightS, sizeof(int) * rowSize * columnSize );
    cudaMalloc( &device_weightT, sizeof(int) * rowSize * columnSize );
    cudaMalloc( &device_height, sizeof(int) * rowSize * columnSize );
    cudaMalloc( &device_capacity, sizeof(int) * rowSize * columnSize );

    cudaMemset( device_active, false, sizeof(bool) );
    cudaMemset( device_height, 0, sizeof(int) * rowSize * columnSize );
    cudaMemset( device_capacity, 0, sizeof(int) * rowSize * columnSize );
    
    cudaMemcpy( device_weightS, weightS, sizeof(int) * rowSize * columnSize, cudaMemcpyHostToDevice);
    cudaMemcpy( device_weightT, weightT, sizeof(int) * rowSize * columnSize, cudaMemcpyHostToDevice);
    cudaMemcpy( device_weightLeft, weightLeft, sizeof(int) * rowSize * columnSize, cudaMemcpyHostToDevice);
    cudaMemcpy( device_weightRight, weightRight, sizeof(int) * rowSize * columnSize, cudaMemcpyHostToDevice);
    cudaMemcpy( device_weightUp, weightUp, sizeof(int) * rowSize * columnSize, cudaMemcpyHostToDevice);
    cudaMemcpy( device_weightDown, weightDown, sizeof(int) * rowSize * columnSize, cudaMemcpyHostToDevice);    

    //cudaMemcpy(y, d_y, N*sizeof(float), cudaMemcpyDeviceToHost);
}







int Data::GetFlow()
{
    dim3 grid( (this->rowSize + 31) / 32, (this->columnSize + 31) / 32, 1);
    dim3 block(32, 1, 1);
    
	//int count = 0;// debug


    bool *active = new bool(true);
	//this->active = true;
	while (*active) {
    //for (int _ = 0; _ < 100; _ ++) {
        //count ++;
        //this->Print();
    
		//this->active = false;
        cudaMemset( device_active, false, sizeof(bool) );

        Data_PushLeftForLine<<< grid, block >>>(            
            this->device_active,
            this->device_weightLeft,
            this->device_weightRight,
            this->device_weightS,
            this->device_weightT,
            this->device_height,
            this->device_capacity,
            this->rowSize,
            this->columnSize
        );

        //cudaMemcpy(active, device_active, sizeof(bool), cudaMemcpyDeviceToHost);

        Data_PushUpForLine<<< grid, block >>>(            
            this->device_active,
            this->device_weightUp,
            this->device_weightDown,
            this->device_weightS,
            this->device_weightT,
            this->device_height,
            this->device_capacity,
            this->rowSize,
            this->columnSize
        );

        //cudaMemcpy(active, device_active, sizeof(bool), cudaMemcpyDeviceToHost);

        Data_PushRightForLine<<< grid, block >>>(            
            this->device_active,
            this->device_weightLeft,
            this->device_weightRight,
            this->device_weightS,
            this->device_weightT,
            this->device_height,
            this->device_capacity,
            this->rowSize,
            this->columnSize
        );

        //cudaMemcpy(active, device_active, sizeof(bool), cudaMemcpyDeviceToHost);

        Data_PushDownForLine<<< grid, block >>>(            
            this->device_active,
            this->device_weightUp,
            this->device_weightDown,
            this->device_weightS,
            this->device_weightT,
            this->device_height,
            this->device_capacity,
            this->rowSize,
            this->columnSize
        );
        cudaMemcpy(active, device_active, sizeof(bool), cudaMemcpyDeviceToHost);
	}
    
    //std::cout << count << std::endl;

//#ifdef DEBUG
    cudaMemcpy(weightS, device_weightS, sizeof(int) * rowSize * columnSize, cudaMemcpyDeviceToHost);
    cudaMemcpy(weightT, device_weightT, sizeof(int) * rowSize * columnSize, cudaMemcpyDeviceToHost);
    cudaMemcpy(weightLeft, device_weightLeft, sizeof(int) * rowSize * columnSize, cudaMemcpyDeviceToHost);
    cudaMemcpy(weightRight, device_weightRight, sizeof(int) * rowSize * columnSize, cudaMemcpyDeviceToHost);
    cudaMemcpy(weightUp, device_weightUp, sizeof(int) * rowSize * columnSize, cudaMemcpyDeviceToHost);
    cudaMemcpy(weightDown, device_weightDown, sizeof(int) * rowSize * columnSize, cudaMemcpyDeviceToHost);
    cudaMemcpy(height, device_height, sizeof(int) * rowSize * columnSize, cudaMemcpyDeviceToHost);
    cudaMemcpy(capacity, device_capacity, sizeof(int) * rowSize * columnSize, cudaMemcpyDeviceToHost);
//#endif

    
    // this->Print();
    // std::cout << "-----" << std::endl;
    
    
    return 0;
	//return this->flow;
}


void Data::BfsFromT()
{
    dim3 grid((this->rowSize * this->columnSize + 32767) / 32768, 32, 1);
    dim3 block(32, 32, 1);

    Data_BfsFromT<<< grid, block >>>(
        this->device_weightT,
        this->device_height,
        this->rowSize,
        this->columnSize
    );

    // debug
    //cudaMemcpy(height, device_height, sizeof(int) * rowSize * columnSize, cudaMemcpyDeviceToHost); this->Print();
    // debug

    int maxHeight = 1;
    for (bool *active = new bool(true); *active; maxHeight ++) {
        cudaMemset( device_active, false, sizeof(bool) );
        Data_BfsLevelK<<< grid, block >>>(
            this->device_active,
            this->device_weightUp,
            this->device_weightDown,
            this->device_weightLeft,
            this->device_weightRight,
            this->device_height,
            this->rowSize,
            this->columnSize,
            // parameter
            maxHeight
        );
        cudaMemcpy(active, device_active, sizeof(bool), cudaMemcpyDeviceToHost);
    }
    cudaMemcpy(height, device_height, sizeof(int) * rowSize * columnSize, cudaMemcpyDeviceToHost);
    
    // debug
    //this->Print();
    // debug
    
    // html debug
	int result[100][100] = { 0 };

	for (int i = 0; i < this->rowSize; i++) {
		for (int j = 0; j < this->columnSize; j++) {
			const size_t index = i * this->columnSize + j;
			int v = this->height[index];
			if (v != -1) {


				if (i != 0 && weightLeft[index] > 0 &&
					this->height[(i - 1) * this->columnSize + j] == -1) {
					result[i][j] = 1;
					result[i - 1][j] = 2;
				}

				if (i != this->rowSize && weightRight[index] > 0 &&
					this->height[(i + 1) * this->columnSize + j] == -1) {
					result[i][j] = 1;
					result[i + 1][j] = 2;
				}

				if (j != this->columnSize - 1 && weightUp[index] > 0 &&
					this->height[i * this->columnSize + j + 1] == -1) {
					result[i][j] = 1;
					result[i][j + 1] = 2;
				}

				if (j != 0 && weightDown[index] > 0 &&
					this->height[i * this->columnSize + j - 1] == -1) {
					result[i][j] = 1;
					result[i][j - 1] = 2;
				}


			}
		}
	}


	std::cout << this->rowSize << " " << this->columnSize << std::endl;
	for (int i = 0; i < this->rowSize; i++) {
		std::cout << "<tr>";
		for (int j = 0; j < this->columnSize; j++) {
			int v = this->height[i * this->columnSize + j];
			//if (v == 6) {
			//	std::cout << "(" << i << " " << j << ")" << std::endl;
			//}
			//std::cout << (v > 10000 ? -1 : v) << " ";

			//if (v == max_k) {
			//	std::cout << "<td class=\"yellow\">&nbsp;</td>";
			if (result[i][j] == 1) {
				std::cout << "<td class=\"yellow\">&nbsp;</td>";
			} else if (result[i][j] == 2) {
				std::cout << "<td class=\"purple\">&nbsp;</td>";
			} else if (this->weightS[i * this->columnSize + j] > 0) {
				std::cout << "<td class=\"red\">&nbsp;</td>";
			} else if (this->weightT[i * this->columnSize + j] > 0) {
				std::cout << "<td class=\"green\">&nbsp;</td>";

			} else {
				if (v == -1) {
					std::cout << "<td>&nbsp;</td>";
				} else {
					std::cout << "<td style=\"background-color: rgb(0, 0," << (5 * v) << ")\">&nbsp;</td>";
				}
			}
		}
		std::cout << "</tr>";
		//std::cout << std::endl;
	}
    // html debug
}


/*
void Data::BfsFromT()
{
	for (int i = 0; i < this->rowSize; i++) {
		for (int j = 0; j < this->columnSize; j++) {
			if (this->weightT[i * this->columnSize + j] > 0) {
				this->height[i * this->columnSize + j] = 1;
			} else {
				this->height[i * this->columnSize + j] = std::numeric_limits<int>::max();
			}
		}
	}

	//int max_k = -1;
	int maxHeight = 0;

	bool check = true;
	for (int k = 1; check; k++) {
		check = false;
		for (int i = 0; i < this->rowSize; i++) {
			for (int j = 0; j < this->columnSize; j++) {
				if (this->height[i * this->columnSize + j] == k) {
					//
					if (i != 0 && weightRight[(i - 1) * this->columnSize + j] > 0 &&
						this->height[(i - 1) * this->columnSize + j] > k) {
						this->height[(i - 1) * this->columnSize + j] = k + 1;
						check = true;
					}

					if (i != this->rowSize && weightLeft[(i + 1) * this->columnSize + j] > 0 &&
						this->height[(i + 1) * this->columnSize + j] > k) {
						this->height[(i + 1) * this->columnSize + j] = k + 1;
						check = true;
					}

					if (j != this->columnSize - 1 && weightDown[i * this->columnSize + j + 1] > 0 &&
						this->height[i * this->columnSize + j + 1] > k) {
						this->height[i * this->columnSize + j + 1] = k + 1;
						check = true;
					}

					if (j != 0 && weightUp[i * this->columnSize + j - 1] > 0 &&
						this->height[i * this->columnSize + j - 1] > k) {
						this->height[i * this->columnSize + j - 1] = k + 1;
						check = true;
					}
					//
				}
			}
		}
		//std::cout << "k = " << k << std::endl;
		maxHeight = k;
	}

	//for (int i = 0; i < this->rowSize * this->columnSize; i++) {
	//	std::cout << this->weight[i * 4] << " ";
	//	std::cout << this->weight[i * 4 + 1] << " ";
	//	std::cout << this->weight[i * 4 + 2] << " ";
	//	std::cout << this->weight[i * 4 + 3] << std::endl;
	//}


	// return if there is a path from t to s
	// for (int i = 0; i < this->rowSize; i++) {
		// for (int j = 0; j < this->columnSize; j++) {
			// const size_t index = i * this->columnSize + j;
			// if (this->height[index] <= maxHeight && this->weightS[index] > 0) {
				// return true;
			// }
		// }
	// }
	// return false;

#define HTML_DEBUG
#ifdef HTML_DEBUG

	int result[100][100] = { 0 };

	for (int i = 0; i < this->rowSize; i++) {
		for (int j = 0; j < this->columnSize; j++) {
			const size_t index = i * this->columnSize + j;
			int v = this->height[index];
			if (v <= maxHeight) {


				if (i != 0 && weightLeft[index] > 0 &&
					this->height[(i - 1) * this->columnSize + j] > maxHeight) {
					result[i][j] = 1;
					result[i - 1][j] = 2;
				}

				if (i != this->rowSize && weightRight[index] > 0 &&
					this->height[(i + 1) * this->columnSize + j] > maxHeight) {
					result[i][j] = 1;
					result[i + 1][j] = 2;
				}

				if (j != this->columnSize - 1 && weightUp[index] > 0 &&
					this->height[i * this->columnSize + j + 1] > maxHeight) {
					result[i][j] = 1;
					result[i][j + 1] = 2;
				}

				if (j != 0 && weightDown[index] > 0 &&
					this->height[i * this->columnSize + j - 1] > maxHeight) {
					result[i][j] = 1;
					result[i][j - 1] = 2;
				}


			}
		}
	}


	std::cout << this->rowSize << " " << this->columnSize << std::endl;
	for (int i = 0; i < this->rowSize; i++) {
		std::cout << "<tr>";
		for (int j = 0; j < this->columnSize; j++) {
			int v = this->height[i * this->columnSize + j];
			//if (v == 6) {
			//	std::cout << "(" << i << " " << j << ")" << std::endl;
			//}
			//std::cout << (v > 10000 ? -1 : v) << " ";

			//if (v == max_k) {
			//	std::cout << "<td class=\"yellow\">&nbsp;</td>";
			if (result[i][j] == 1) {
				std::cout << "<td class=\"yellow\">&nbsp;</td>";
			} else if (result[i][j] == 2) {
				std::cout << "<td class=\"purple\">&nbsp;</td>";
			} else if (this->weightS[i * this->columnSize + j] > 0) {
				std::cout << "<td class=\"red\">&nbsp;</td>";
			} else if (this->weightT[i * this->columnSize + j] > 0) {
				std::cout << "<td class=\"green\">&nbsp;</td>";

			} else {
				if (v > maxHeight) {
					std::cout << "<td>&nbsp;</td>";
				} else {
					std::cout << "<td style=\"background-color: rgb(0, 0," << (5 * v) << ")\">&nbsp;</td>";
				}
			}
		}
		std::cout << "</tr>";
		//std::cout << std::endl;
	}
#endif

	//std::cout << this->rowSize << " " << this->columnSize << std::endl;
	//for (int i = 0; i < this->rowSize; i++) {
	//	for (int j = 0; j < this->columnSize; j++) {
	//		int v = this->height[i * this->columnSize + j];
	//		if (v == 6) {
	//			std::cout << "(" << i << ", " << j << ")" << std::endl;
	//		}
	//	}
	//	std::cout << std::endl;
	//}
}
*/


	.file	"tmpxft_0022f9b5_00000000-6_Data.cudafe1.cpp"
	.text
#APP
	.globl _ZSt21ios_base_library_initv

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB3781:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE3781:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.type	_Z10cudaMallocIbE9cudaErrorPPT_m.isra.0, @function
_Z10cudaMallocIbE9cudaErrorPPT_m.isra.0:
.LFB4501:
	.cfi_startproc
	jmp	cudaMalloc@PLT
	.cfi_endproc
.LFE4501:
	.size	_Z10cudaMallocIbE9cudaErrorPPT_m.isra.0, .-_Z10cudaMallocIbE9cudaErrorPPT_m.isra.0
	.align 2
	.globl	_ZN4DataC2Ev
	.type	_ZN4DataC2Ev, @function
_ZN4DataC2Ev:
.LFB3767:
	.cfi_startproc
	endbr64
	xorl	%eax, %eax
	movq	%rax, (%rdi)
	movq	%rax, 8(%rdi)
	ret
	.cfi_endproc
.LFE3767:
	.size	_ZN4DataC2Ev, .-_ZN4DataC2Ev
	.globl	_ZN4DataC1Ev
	.set	_ZN4DataC1Ev,_ZN4DataC2Ev
	.align 2
	.globl	_ZN4DataD2Ev
	.type	_ZN4DataD2Ev, @function
_ZN4DataD2Ev:
.LFB3770:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA3770
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	movq	%rdi, %rbx
	movq	16(%rdi), %rdi
	testq	%rdi, %rdi
	je	.L5
	call	_ZdaPv@PLT
.L5:
	movq	24(%rbx), %rdi
	testq	%rdi, %rdi
	je	.L6
	call	_ZdaPv@PLT
.L6:
	movq	32(%rbx), %rdi
	testq	%rdi, %rdi
	je	.L7
	call	_ZdaPv@PLT
.L7:
	movq	40(%rbx), %rdi
	testq	%rdi, %rdi
	je	.L8
	call	_ZdaPv@PLT
.L8:
	movq	48(%rbx), %rdi
	testq	%rdi, %rdi
	je	.L9
	call	_ZdaPv@PLT
.L9:
	movq	56(%rbx), %rdi
	testq	%rdi, %rdi
	je	.L10
	call	_ZdaPv@PLT
.L10:
	movq	64(%rbx), %rdi
	testq	%rdi, %rdi
	je	.L11
	call	_ZdaPv@PLT
.L11:
	movq	72(%rbx), %rdi
	testq	%rdi, %rdi
	je	.L12
	call	_ZdaPv@PLT
.L12:
	movq	80(%rbx), %rdi
	call	cudaFree@PLT
	movq	88(%rbx), %rdi
	call	cudaFree@PLT
	movq	96(%rbx), %rdi
	call	cudaFree@PLT
	movq	104(%rbx), %rdi
	call	cudaFree@PLT
	movq	112(%rbx), %rdi
	call	cudaFree@PLT
	movq	120(%rbx), %rdi
	call	cudaFree@PLT
	movq	128(%rbx), %rdi
	call	cudaFree@PLT
	movq	136(%rbx), %rdi
	call	cudaFree@PLT
	movq	144(%rbx), %rdi
	call	cudaFree@PLT
	popq	%rbx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3770:
	.globl	__gxx_personality_v0
	.section	.gcc_except_table,"a",@progbits
.LLSDA3770:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE3770-.LLSDACSB3770
.LLSDACSB3770:
.LLSDACSE3770:
	.text
	.size	_ZN4DataD2Ev, .-_ZN4DataD2Ev
	.globl	_ZN4DataD1Ev
	.set	_ZN4DataD1Ev,_ZN4DataD2Ev
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"Weight:"
.LC1:
	.string	"Left:"
.LC2:
	.string	" "
.LC3:
	.string	"Right:"
.LC4:
	.string	"Up:"
.LC5:
	.string	"Down:"
.LC6:
	.string	"Capacity:"
.LC7:
	.string	"Height:"
	.text
	.align 2
	.globl	_ZN4Data5PrintEv
	.type	_ZN4Data5PrintEv, @function
_ZN4Data5PrintEv:
.LFB3772:
	.cfi_startproc
	endbr64
	pushq	%r14
	.cfi_def_cfa_offset 16
	.cfi_offset 14, -16
	leaq	.LC0(%rip), %rsi
	pushq	%r13
	.cfi_def_cfa_offset 24
	.cfi_offset 13, -24
	xorl	%r13d, %r13d
	pushq	%r12
	.cfi_def_cfa_offset 32
	.cfi_offset 12, -32
	leaq	.LC2(%rip), %r12
	pushq	%rbp
	.cfi_def_cfa_offset 40
	.cfi_offset 6, -40
	movq	%rdi, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset 3, -48
	leaq	_ZSt4cout(%rip), %rbx
	movq	%rbx, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
	leaq	.LC1(%rip), %rsi
	movq	%rbx, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
.L40:
	cmpq	0(%rbp), %r13
	jnb	.L65
	xorl	%r14d, %r14d
.L42:
	movq	8(%rbp), %rax
	cmpq	%rax, %r14
	jnb	.L66
	imulq	%r13, %rax
	movq	16(%rbp), %rdx
	movq	%rbx, %rdi
	addq	%r14, %rax
	incq	%r14
	movl	(%rdx,%rax,4), %esi
	call	_ZNSolsEi@PLT
	movq	%r12, %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	jmp	.L42
.L66:
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rbx, %rdi
	incq	%r13
	call	_ZNSolsEPFRSoS_E@PLT
	jmp	.L40
.L65:
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rbx, %rdi
	xorl	%r13d, %r13d
	leaq	.LC2(%rip), %r12
	call	_ZNSolsEPFRSoS_E@PLT
	leaq	.LC3(%rip), %rsi
	movq	%rbx, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
.L43:
	cmpq	0(%rbp), %r13
	jnb	.L67
	xorl	%r14d, %r14d
.L45:
	movq	8(%rbp), %rax
	cmpq	%rax, %r14
	jnb	.L68
	imulq	%r13, %rax
	movq	24(%rbp), %rdx
	movq	%rbx, %rdi
	addq	%r14, %rax
	incq	%r14
	movl	(%rdx,%rax,4), %esi
	call	_ZNSolsEi@PLT
	movq	%r12, %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	jmp	.L45
.L68:
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rbx, %rdi
	incq	%r13
	call	_ZNSolsEPFRSoS_E@PLT
	jmp	.L43
.L67:
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rbx, %rdi
	xorl	%r13d, %r13d
	leaq	.LC2(%rip), %r12
	call	_ZNSolsEPFRSoS_E@PLT
	leaq	.LC4(%rip), %rsi
	movq	%rbx, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
.L46:
	cmpq	0(%rbp), %r13
	jnb	.L69
	xorl	%r14d, %r14d
.L48:
	movq	8(%rbp), %rax
	cmpq	%rax, %r14
	jnb	.L70
	imulq	%r13, %rax
	movq	32(%rbp), %rdx
	movq	%rbx, %rdi
	addq	%r14, %rax
	incq	%r14
	movl	(%rdx,%rax,4), %esi
	call	_ZNSolsEi@PLT
	movq	%r12, %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	jmp	.L48
.L70:
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rbx, %rdi
	incq	%r13
	call	_ZNSolsEPFRSoS_E@PLT
	jmp	.L46
.L69:
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rbx, %rdi
	xorl	%r13d, %r13d
	leaq	.LC2(%rip), %r12
	call	_ZNSolsEPFRSoS_E@PLT
	leaq	.LC5(%rip), %rsi
	movq	%rbx, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
.L49:
	cmpq	0(%rbp), %r13
	jnb	.L71
	xorl	%r14d, %r14d
.L51:
	movq	8(%rbp), %rax
	cmpq	%rax, %r14
	jnb	.L72
	imulq	%r13, %rax
	movq	40(%rbp), %rdx
	movq	%rbx, %rdi
	addq	%r14, %rax
	incq	%r14
	movl	(%rdx,%rax,4), %esi
	call	_ZNSolsEi@PLT
	movq	%r12, %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	jmp	.L51
.L72:
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rbx, %rdi
	incq	%r13
	call	_ZNSolsEPFRSoS_E@PLT
	jmp	.L49
.L71:
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rbx, %rdi
	xorl	%r13d, %r13d
	leaq	.LC2(%rip), %r12
	call	_ZNSolsEPFRSoS_E@PLT
	leaq	.LC6(%rip), %rsi
	movq	%rbx, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
.L52:
	cmpq	0(%rbp), %r13
	jnb	.L73
	xorl	%r14d, %r14d
.L54:
	movq	8(%rbp), %rax
	cmpq	%rax, %r14
	jnb	.L74
	imulq	%r13, %rax
	movq	72(%rbp), %rdx
	movq	%rbx, %rdi
	addq	%r14, %rax
	incq	%r14
	movl	(%rdx,%rax,4), %esi
	call	_ZNSolsEi@PLT
	movq	%r12, %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	jmp	.L54
.L74:
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rbx, %rdi
	incq	%r13
	call	_ZNSolsEPFRSoS_E@PLT
	jmp	.L52
.L73:
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rbx, %rdi
	xorl	%r13d, %r13d
	leaq	.LC2(%rip), %r12
	call	_ZNSolsEPFRSoS_E@PLT
	leaq	.LC7(%rip), %rsi
	movq	%rbx, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
.L55:
	cmpq	0(%rbp), %r13
	jnb	.L75
	xorl	%r14d, %r14d
.L57:
	movq	8(%rbp), %rax
	cmpq	%rax, %r14
	jnb	.L76
	imulq	%r13, %rax
	movq	64(%rbp), %rdx
	movq	%rbx, %rdi
	addq	%r14, %rax
	incq	%r14
	movl	(%rdx,%rax,4), %esi
	call	_ZNSolsEi@PLT
	movq	%r12, %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	jmp	.L57
.L76:
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rbx, %rdi
	incq	%r13
	call	_ZNSolsEPFRSoS_E@PLT
	jmp	.L55
.L75:
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rbx, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %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	_ZNSolsEPFRSoS_E@PLT
	.cfi_endproc
.LFE3772:
	.size	_ZN4Data5PrintEv, .-_ZN4Data5PrintEv
	.align 2
	.globl	_ZN4DataC2Emm
	.type	_ZN4DataC2Emm, @function
_ZN4DataC2Emm:
.LFB3774:
	.cfi_startproc
	endbr64
	movabsq	$2305843009213693950, %rax
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	movq	%rsi, %r13
	imulq	%rdx, %r13
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	movq	%rdi, %rbx
	pushq	%rcx
	.cfi_def_cfa_offset 48
	movq	%rsi, (%rdi)
	movq	%rdx, 8(%rdi)
	cmpq	%r13, %rax
	jb	.L78
	leaq	0(,%r13,4), %rbp
	xorl	%r12d, %r12d
	movq	%rbp, %rdi
	call	_Znam@PLT
	jmp	.L79
.L78:
	call	__cxa_throw_bad_array_new_length@PLT
.L79:
	cmpq	%r12, %r13
	je	.L96
	xorl	%edx, %edx
	movl	%edx, (%rax,%r12,4)
	incq	%r12
	jmp	.L79
.L96:
	movq	%rax, 16(%rbx)
	movq	%rbp, %rdi
	call	_Znam@PLT
	xorl	%edx, %edx
.L81:
	cmpq	%r12, %rdx
	je	.L97
	xorl	%r11d, %r11d
	movl	%r11d, (%rax,%rdx,4)
	incq	%rdx
	jmp	.L81
.L97:
	movq	%rax, 24(%rbx)
	movq	%rbp, %rdi
	call	_Znam@PLT
	xorl	%edx, %edx
.L83:
	cmpq	%r12, %rdx
	je	.L98
	xorl	%r10d, %r10d
	movl	%r10d, (%rax,%rdx,4)
	incq	%rdx
	jmp	.L83
.L98:
	movq	%rax, 32(%rbx)
	movq	%rbp, %rdi
	call	_Znam@PLT
	xorl	%edx, %edx
.L85:
	cmpq	%r12, %rdx
	je	.L99
	xorl	%r9d, %r9d
	movl	%r9d, (%rax,%rdx,4)
	incq	%rdx
	jmp	.L85
.L99:
	movq	%rax, 40(%rbx)
	movq	%rbp, %rdi
	call	_Znam@PLT
	xorl	%edx, %edx
.L87:
	cmpq	%r12, %rdx
	je	.L100
	xorl	%r8d, %r8d
	movl	%r8d, (%rax,%rdx,4)
	incq	%rdx
	jmp	.L87
.L100:
	movq	%rax, 48(%rbx)
	movq	%rbp, %rdi
	call	_Znam@PLT
	xorl	%edx, %edx
.L89:
	cmpq	%r12, %rdx
	je	.L101
	xorl	%edi, %edi
	movl	%edi, (%rax,%rdx,4)
	incq	%rdx
	jmp	.L89
.L101:
	movq	%rax, 56(%rbx)
	movq	%rbp, %rdi
	call	_Znam@PLT
	xorl	%edx, %edx
.L91:
	cmpq	%r12, %rdx
	je	.L102
	xorl	%esi, %esi
	movl	%esi, (%rax,%rdx,4)
	incq	%rdx
	jmp	.L91
.L102:
	movq	%rax, 64(%rbx)
	movq	%rbp, %rdi
	call	_Znam@PLT
	xorl	%edx, %edx
.L93:
	cmpq	%r12, %rdx
	je	.L103
	xorl	%ecx, %ecx
	movl	%ecx, (%rax,%rdx,4)
	incq	%rdx
	jmp	.L93
.L103:
	movq	%rax, 72(%rbx)
	leaq	80(%rbx), %rdi
	movl	$1, %esi
	call	_Z10cudaMallocIbE9cudaErrorPPT_m.isra.0
	leaq	88(%rbx), %rdi
	movq	%rbp, %rsi
	call	_Z10cudaMallocIbE9cudaErrorPPT_m.isra.0
	leaq	96(%rbx), %rdi
	movq	%rbp, %rsi
	call	_Z10cudaMallocIbE9cudaErrorPPT_m.isra.0
	leaq	104(%rbx), %rdi
	movq	%rbp, %rsi
	call	_Z10cudaMallocIbE9cudaErrorPPT_m.isra.0
	leaq	112(%rbx), %rdi
	movq	%rbp, %rsi
	call	_Z10cudaMallocIbE9cudaErrorPPT_m.isra.0
	leaq	120(%rbx), %rdi
	movq	%rbp, %rsi
	call	_Z10cudaMallocIbE9cudaErrorPPT_m.isra.0
	leaq	128(%rbx), %rdi
	movq	%rbp, %rsi
	call	_Z10cudaMallocIbE9cudaErrorPPT_m.isra.0
	leaq	136(%rbx), %rdi
	movq	%rbp, %rsi
	call	_Z10cudaMallocIbE9cudaErrorPPT_m.isra.0
	leaq	144(%rbx), %rdi
	movq	%rbp, %rsi
	call	_Z10cudaMallocIbE9cudaErrorPPT_m.isra.0
	movq	80(%rbx), %rdi
	movl	$1, %edx
	xorl	%esi, %esi
	call	cudaMemset@PLT
	movq	136(%rbx), %rdi
	movq	%rbp, %rdx
	xorl	%esi, %esi
	call	cudaMemset@PLT
	movq	144(%rbx), %rdi
	movq	%rbp, %rdx
	popq	%rax
	.cfi_def_cfa_offset 40
	xorl	%esi, %esi
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%rbp
	.cfi_def_cfa_offset 24
	popq	%r12
	.cfi_def_cfa_offset 16
	popq	%r13
	.cfi_def_cfa_offset 8
	jmp	cudaMemset@PLT
	.cfi_endproc
.LFE3774:
	.size	_ZN4DataC2Emm, .-_ZN4DataC2Emm
	.globl	_ZN4DataC1Emm
	.set	_ZN4DataC1Emm,_ZN4DataC2Emm
	.align 2
	.globl	_ZN4Data7GetFlowEv
	.type	_ZN4Data7GetFlowEv, @function
_ZN4Data7GetFlowEv:
.LFB3777:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	movq	%rdi, %rbx
	subq	$40, %rsp
	.cfi_def_cfa_offset 64
	movq	8(%rdi), %rax
	movq	(%rdi), %rcx
	movl	$1, %edi
	movl	$1, 16(%rsp)
	addq	$31, %rax
	leaq	31(%rcx), %rdx
	movl	$1, 28(%rsp)
	shrq	$5, %rax
	shrq	$5, %rdx
	movl	%eax, 12(%rsp)
	movl	$134217729, %eax
	salq	$5, %rax
	movl	%edx, 8(%rsp)
	movq	%rax, 20(%rsp)
	call	_Znwm@PLT
	movb	$1, (%rax)
	movq	%rax, %rbp
.L105:
	cmpb	$0, 0(%rbp)
	je	.L112
	movq	80(%rbx), %rdi
	xorl	%esi, %esi
	movl	$1, %edx
	call	cudaMemset@PLT
	movl	28(%rsp), %ecx
	movl	16(%rsp), %esi
	xorl	%r9d, %r9d
	movq	20(%rsp), %rdx
	movq	8(%rsp), %rdi
	xorl	%r8d, %r8d
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L106
	movq	120(%rbx), %rcx
	movq	96(%rbx), %rdx
	movq	88(%rbx), %rsi
	movq	80(%rbx), %rdi
	pushq	%r10
	.cfi_def_cfa_offset 72
	movq	8(%rbx), %rax
	pushq	%rax
	.cfi_def_cfa_offset 80
	movq	(%rbx), %rax
	pushq	%rax
	.cfi_def_cfa_offset 88
	pushq	144(%rbx)
	.cfi_def_cfa_offset 96
	movq	136(%rbx), %r9
	movq	128(%rbx), %r8
	call	_Z20Data_PushLeftForLinePbPiS0_S0_S0_S0_S0_ii@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 64
.L106:
	movl	28(%rsp), %ecx
	movq	20(%rsp), %rdx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movq	8(%rsp), %rdi
	movl	16(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L107
	movq	120(%rbx), %rcx
	movq	112(%rbx), %rdx
	movq	104(%rbx), %rsi
	movq	80(%rbx), %rdi
	pushq	%r9
	.cfi_def_cfa_offset 72
	movq	8(%rbx), %rax
	pushq	%rax
	.cfi_def_cfa_offset 80
	movq	(%rbx), %rax
	pushq	%rax
	.cfi_def_cfa_offset 88
	pushq	144(%rbx)
	.cfi_def_cfa_offset 96
	movq	136(%rbx), %r9
	movq	128(%rbx), %r8
	call	_Z18Data_PushUpForLinePbPiS0_S0_S0_S0_S0_ii@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 64
.L107:
	movl	28(%rsp), %ecx
	movq	20(%rsp), %rdx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movq	8(%rsp), %rdi
	movl	16(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L108
	movq	120(%rbx), %rcx
	movq	96(%rbx), %rdx
	movq	88(%rbx), %rsi
	movq	80(%rbx), %rdi
	pushq	%r8
	.cfi_def_cfa_offset 72
	movq	8(%rbx), %rax
	pushq	%rax
	.cfi_def_cfa_offset 80
	movq	(%rbx), %rax
	pushq	%rax
	.cfi_def_cfa_offset 88
	pushq	144(%rbx)
	.cfi_def_cfa_offset 96
	movq	136(%rbx), %r9
	movq	128(%rbx), %r8
	call	_Z21Data_PushRightForLinePbPiS0_S0_S0_S0_S0_ii@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 64
.L108:
	movl	28(%rsp), %ecx
	movq	20(%rsp), %rdx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movq	8(%rsp), %rdi
	movl	16(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L109
	movq	120(%rbx), %rcx
	movq	112(%rbx), %rdx
	movq	104(%rbx), %rsi
	movq	80(%rbx), %rdi
	pushq	%rax
	.cfi_def_cfa_offset 72
	movq	8(%rbx), %rax
	pushq	%rax
	.cfi_def_cfa_offset 80
	movq	(%rbx), %rax
	pushq	%rax
	.cfi_def_cfa_offset 88
	pushq	144(%rbx)
	.cfi_def_cfa_offset 96
	movq	136(%rbx), %r9
	movq	128(%rbx), %r8
	call	_Z20Data_PushDownForLinePbPiS0_S0_S0_S0_S0_ii@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 64
.L109:
	movq	80(%rbx), %rsi
	movl	$2, %ecx
	movl	$1, %edx
	movq	%rbp, %rdi
	call	cudaMemcpy@PLT
	jmp	.L105
.L112:
	movq	(%rbx), %rdx
	imulq	8(%rbx), %rdx
	movl	$2, %ecx
	movq	120(%rbx), %rsi
	movq	48(%rbx), %rdi
	salq	$2, %rdx
	call	cudaMemcpy@PLT
	movq	(%rbx), %rdx
	imulq	8(%rbx), %rdx
	movl	$2, %ecx
	movq	128(%rbx), %rsi
	movq	56(%rbx), %rdi
	salq	$2, %rdx
	call	cudaMemcpy@PLT
	movq	(%rbx), %rdx
	imulq	8(%rbx), %rdx
	movl	$2, %ecx
	movq	88(%rbx), %rsi
	movq	16(%rbx), %rdi
	salq	$2, %rdx
	call	cudaMemcpy@PLT
	movq	(%rbx), %rdx
	imulq	8(%rbx), %rdx
	movl	$2, %ecx
	movq	96(%rbx), %rsi
	movq	24(%rbx), %rdi
	salq	$2, %rdx
	call	cudaMemcpy@PLT
	movq	(%rbx), %rdx
	imulq	8(%rbx), %rdx
	movl	$2, %ecx
	movq	104(%rbx), %rsi
	movq	32(%rbx), %rdi
	salq	$2, %rdx
	call	cudaMemcpy@PLT
	movq	(%rbx), %rdx
	imulq	8(%rbx), %rdx
	movl	$2, %ecx
	movq	112(%rbx), %rsi
	movq	40(%rbx), %rdi
	salq	$2, %rdx
	call	cudaMemcpy@PLT
	movq	(%rbx), %rdx
	imulq	8(%rbx), %rdx
	movl	$2, %ecx
	movq	136(%rbx), %rsi
	movq	64(%rbx), %rdi
	salq	$2, %rdx
	call	cudaMemcpy@PLT
	movq	(%rbx), %rdx
	imulq	8(%rbx), %rdx
	movl	$2, %ecx
	movq	144(%rbx), %rsi
	movq	72(%rbx), %rdi
	salq	$2, %rdx
	call	cudaMemcpy@PLT
	addq	$40, %rsp
	.cfi_def_cfa_offset 24
	xorl	%eax, %eax
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3777:
	.size	_ZN4Data7GetFlowEv, .-_ZN4Data7GetFlowEv
	.section	.rodata.str1.1
.LC8:
	.string	"<tr>"
.LC9:
	.string	"<td class=\"yellow\">&nbsp;</td>"
.LC10:
	.string	"<td class=\"purple\">&nbsp;</td>"
.LC11:
	.string	"<td class=\"red\">&nbsp;</td>"
.LC12:
	.string	"<td class=\"green\">&nbsp;</td>"
.LC13:
	.string	"<td>&nbsp;</td>"
.LC14:
	.string	"<td style=\"background-color: rgb(0, 0,"
.LC15:
	.string	")\">&nbsp;</td>"
.LC16:
	.string	"</tr>"
	.text
	.align 2
	.globl	_ZN4Data8BfsFromTEv
	.type	_ZN4Data8BfsFromTEv, @function
_ZN4Data8BfsFromTEv:
.LFB3778:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	leaq	-36864(%rsp), %r11
	.cfi_def_cfa 11, 36920
.LPSRL0:
	subq	$4096, %rsp
	orq	$0, (%rsp)
	cmpq	%r11, %rsp
	jne	.LPSRL0
	.cfi_def_cfa_register 7
	subq	$3192, %rsp
	.cfi_def_cfa_offset 40112
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$1, %ecx
	movabsq	$137438953504, %rdx
	movq	%fs:40, %rax
	movq	%rax, 40040(%rsp)
	movq	(%rdi), %rax
	movq	%rdi, %r15
	movq	%rdx, 28(%rsp)
	movl	$1, 36(%rsp)
	imulq	8(%rdi), %rax
	addq	$32767, %rax
	shrq	$15, %rax
	movl	%eax, 16(%rsp)
	movl	$134217729, %eax
	salq	$5, %rax
	movq	%rax, 20(%rsp)
	movq	16(%rsp), %rdi
	movl	24(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L114
	movq	136(%r15), %rsi
	movq	128(%r15), %rdi
	movl	8(%r15), %ecx
	movl	(%r15), %edx
	call	_Z13Data_BfsFromTPiS_ii@PLT
.L114:
	movl	$1, %edi
	movl	$1, %ebp
	call	_Znwm@PLT
	movb	$1, (%rax)
	movq	%rax, %rbx
.L115:
	cmpb	$0, (%rbx)
	je	.L148
	movq	80(%r15), %rdi
	xorl	%esi, %esi
	movl	$1, %edx
	call	cudaMemset@PLT
	movl	36(%rsp), %ecx
	movl	24(%rsp), %esi
	xorl	%r9d, %r9d
	movq	28(%rsp), %rdx
	movq	16(%rsp), %rdi
	xorl	%r8d, %r8d
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L116
	movq	88(%r15), %rcx
	movq	112(%r15), %rdx
	movq	104(%r15), %rsi
	movq	80(%r15), %rdi
	pushq	%rax
	.cfi_def_cfa_offset 40120
	movq	8(%r15), %rax
	movq	136(%r15), %r9
	pushq	%rbp
	.cfi_def_cfa_offset 40128
	movq	96(%r15), %r8
	pushq	%rax
	.cfi_def_cfa_offset 40136
	movq	(%r15), %rax
	pushq	%rax
	.cfi_def_cfa_offset 40144
	call	_Z14Data_BfsLevelKPbPiS0_S0_S0_S0_iii@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 40112
.L116:
	movq	80(%r15), %rsi
	movl	$2, %ecx
	movl	$1, %edx
	incl	%ebp
	movq	%rbx, %rdi
	call	cudaMemcpy@PLT
	jmp	.L115
.L148:
	movq	(%r15), %rdx
	imulq	8(%r15), %rdx
	movl	$2, %ecx
	xorl	%ebx, %ebx
	movq	136(%r15), %rsi
	movq	64(%r15), %rdi
	movq	$-400, %r13
	salq	$2, %rdx
	call	cudaMemcpy@PLT
	xorl	%eax, %eax
	leaq	40(%rsp), %rdi
	movq	(%r15), %r12
	movl	$10000, %ecx
	xorl	%esi, %esi
	rep stosl
	leaq	40(%rsp), %rax
	movq	%rax, 8(%rsp)
.L118:
	cmpq	%rsi, %r12
	je	.L126
	movq	8(%r15), %rbp
	leal	1(%rsi), %edi
	leaq	-1(%rsi), %r9
	movq	8(%rsp), %rax
	movslq	%edi, %rdi
	xorl	%r8d, %r8d
	leaq	0(,%rbp,4), %rcx
	leaq	-1(%rbp), %r10
	subq	%rbx, %rax
	movq	%rcx, %rdx
	imulq	%rcx, %r9
	imulq	%rsi, %rdx
	imulq	%rdi, %rcx
	imulq	$400, %rdi, %rdi
.L127:
	cmpq	%r8, %rbp
	je	.L149
	movq	64(%r15), %r11
	cmpl	$-1, (%r11,%rdx)
	je	.L120
	testq	%rsi, %rsi
	je	.L121
	movq	16(%r15), %r14
	cmpl	$0, (%r14,%rdx)
	jle	.L121
	cmpl	$-1, (%r11,%r9)
	jne	.L121
	movl	$1, (%rax)
	movl	$2, (%rax,%r13)
.L121:
	movq	24(%r15), %r14
	cmpl	$0, (%r14,%rdx)
	jle	.L122
	cmpl	$-1, (%r11,%rcx)
	jne	.L122
	leaq	(%rax,%rbx), %r14
	movl	$1, (%rax)
	movl	$2, (%r14,%rdi)
.L122:
	cmpq	%r8, %r10
	je	.L123
	movq	32(%r15), %r14
	cmpl	$0, (%r14,%rdx)
	jle	.L123
	cmpl	$-1, 4(%r11,%rdx)
	jne	.L123
	movl	$1, (%rax)
	movl	$2, 4(%rax)
.L123:
	testq	%r8, %r8
	je	.L120
	movq	40(%r15), %r14
	cmpl	$0, (%r14,%rdx)
	jle	.L120
	cmpl	$-1, -4(%r11,%rdx)
	jne	.L120
	movl	$1, (%rax)
	movl	$2, -4(%rax)
.L120:
	incq	%r8
	addq	$4, %rax
	addq	$4, %rdx
	addq	$4, %r9
	addq	$4, %rcx
	jmp	.L127
.L149:
	incq	%rsi
	subq	$400, %rbx
	jmp	.L118
.L126:
	movq	%r12, %rsi
	leaq	_ZSt4cout(%rip), %rdi
	xorl	%ebp, %ebp
	call	_ZNSo9_M_insertImEERSoT_@PLT
	leaq	.LC2(%rip), %rsi
	leaq	.LC8(%rip), %r14
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	8(%r15), %rsi
	movq	%rax, %rdi
	call	_ZNSo9_M_insertImEERSoT_@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
.L128:
	cmpq	(%r15), %rbp
	jnb	.L150
	imulq	$400, %rbp, %r13
	movq	%r14, %rsi
	leaq	_ZSt4cout(%rip), %rdi
	xorl	%r12d, %r12d
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	8(%rsp), %rax
	addq	%rax, %r13
.L129:
	movq	8(%r15), %rax
	cmpq	%rax, %r12
	jnb	.L151
	imulq	%rbp, %rax
	movq	64(%r15), %rdx
	movl	0(%r13,%r12,4), %ecx
	leaq	.LC9(%rip), %rsi
	addq	%r12, %rax
	movl	(%rdx,%rax,4), %ebx
	cmpl	$1, %ecx
	je	.L147
	leaq	.LC10(%rip), %rsi
	cmpl	$2, %ecx
	je	.L147
	movq	48(%r15), %rcx
	leaq	.LC11(%rip), %rsi
	cmpl	$0, (%rcx,%rax,4)
	jg	.L147
	movq	56(%r15), %rcx
	leaq	.LC12(%rip), %rsi
	cmpl	$0, (%rcx,%rax,4)
	jg	.L147
	cmpl	$-1, %ebx
	jne	.L135
	leaq	.LC13(%rip), %rsi
.L147:
	leaq	_ZSt4cout(%rip), %rdi
	jmp	.L146
.L135:
	leaq	.LC14(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	leal	(%rbx,%rbx,4), %esi
	movq	%rax, %rdi
	call	_ZNSolsEi@PLT
	leaq	.LC15(%rip), %rsi
	movq	%rax, %rdi
.L146:
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	incq	%r12
	jmp	.L129
.L151:
	leaq	.LC16(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	incq	%rbp
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	jmp	.L128
.L150:
	movq	40040(%rsp), %rax
	subq	%fs:40, %rax
	je	.L138
	call	__stack_chk_fail@PLT
.L138:
	addq	$40056, %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
.LFE3778:
	.size	_ZN4Data8BfsFromTEv, .-_ZN4Data8BfsFromTEv
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB3804:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	movq	%rax, %rdi
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rdx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE3804:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.text
	.align 2
	.globl	_ZN4Data4ReadEPKc
	.type	_ZN4Data4ReadEPKc, @function
_ZN4Data4ReadEPKc:
.LFB3776:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA3776
	endbr64
	pushq	%r14
	.cfi_def_cfa_offset 16
	.cfi_offset 14, -16
	pushq	%r13
	.cfi_def_cfa_offset 24
	.cfi_offset 13, -24
	pushq	%r12
	.cfi_def_cfa_offset 32
	.cfi_offset 12, -32
	pushq	%rbp
	.cfi_def_cfa_offset 40
	.cfi_offset 6, -40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset 3, -48
	movq	%rdi, %rbx
	subq	$560, %rsp
	.cfi_def_cfa_offset 608
	movq	%rsi, 8(%rsp)
	leaq	24(%rsp), %r12
	movq	%r12, %rdi
	movq	%fs:40, %rax
	movq	%rax, 552(%rsp)
	xorl	%eax, %eax
.LEHB0:
	call	_ZNSt13basic_fstreamIcSt11char_traitsIcEEC1Ev@PLT
.LEHE0:
	movq	8(%rsp), %rsi
	movl	$8, %edx
	movq	%r12, %rdi
.LEHB1:
	call	_ZNSt13basic_fstreamIcSt11char_traitsIcEE4openEPKcSt13_Ios_Openmode@PLT
	movq	%rbx, %rsi
	movq	%r12, %rdi
	call	_ZNSi10_M_extractImEERSiRT_@PLT
	movq	%rax, %rdi
	leaq	8(%rbx), %rsi
	call	_ZNSi10_M_extractImEERSiRT_@PLT
	movabsq	$2305843009213693950, %rax
	movq	(%rbx), %rbp
	imulq	8(%rbx), %rbp
	cmpq	%rbp, %rax
	jb	.L155
	leaq	0(,%rbp,4), %rdi
	call	_Znam@PLT
	xorl	%edx, %edx
	jmp	.L156
.L155:
	movq	552(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L201
	call	__cxa_throw_bad_array_new_length@PLT
.L156:
	cmpq	%rdx, %rbp
	je	.L204
	xorl	%r13d, %r13d
	movl	%r13d, (%rax,%rdx,4)
	incq	%rdx
	jmp	.L156
.L204:
	movq	(%rbx), %rbp
	imulq	8(%rbx), %rbp
	movq	%rax, 16(%rbx)
	movabsq	$2305843009213693950, %rax
	cmpq	%rbp, %rax
	jb	.L159
	leaq	0(,%rbp,4), %rdi
	call	_Znam@PLT
	xorl	%edx, %edx
	jmp	.L160
.L159:
	movq	552(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L201
	call	__cxa_throw_bad_array_new_length@PLT
.L160:
	cmpq	%rdx, %rbp
	je	.L205
	xorl	%r11d, %r11d
	movl	%r11d, (%rax,%rdx,4)
	incq	%rdx
	jmp	.L160
.L205:
	movq	(%rbx), %rbp
	imulq	8(%rbx), %rbp
	movq	%rax, 24(%rbx)
	movabsq	$2305843009213693950, %rax
	cmpq	%rbp, %rax
	jb	.L163
	leaq	0(,%rbp,4), %rdi
	call	_Znam@PLT
	xorl	%edx, %edx
	jmp	.L164
.L163:
	movq	552(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L201
	call	__cxa_throw_bad_array_new_length@PLT
.L164:
	cmpq	%rdx, %rbp
	je	.L206
	xorl	%r10d, %r10d
	movl	%r10d, (%rax,%rdx,4)
	incq	%rdx
	jmp	.L164
.L206:
	movq	(%rbx), %rbp
	imulq	8(%rbx), %rbp
	movq	%rax, 32(%rbx)
	movabsq	$2305843009213693950, %rax
	cmpq	%rbp, %rax
	jb	.L167
	leaq	0(,%rbp,4), %rdi
	call	_Znam@PLT
	xorl	%edx, %edx
	jmp	.L168
.L167:
	movq	552(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L201
	call	__cxa_throw_bad_array_new_length@PLT
.L168:
	cmpq	%rdx, %rbp
	je	.L207
	xorl	%r9d, %r9d
	movl	%r9d, (%rax,%rdx,4)
	incq	%rdx
	jmp	.L168
.L207:
	movq	(%rbx), %rbp
	imulq	8(%rbx), %rbp
	movq	%rax, 40(%rbx)
	movabsq	$2305843009213693950, %rax
	cmpq	%rbp, %rax
	jb	.L171
	leaq	0(,%rbp,4), %rdi
	call	_Znam@PLT
	xorl	%edx, %edx
	jmp	.L172
.L171:
	movq	552(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L201
	call	__cxa_throw_bad_array_new_length@PLT
.L172:
	cmpq	%rdx, %rbp
	je	.L208
	xorl	%r8d, %r8d
	movl	%r8d, (%rax,%rdx,4)
	incq	%rdx
	jmp	.L172
.L208:
	movq	(%rbx), %rbp
	imulq	8(%rbx), %rbp
	movq	%rax, 48(%rbx)
	movabsq	$2305843009213693950, %rax
	cmpq	%rbp, %rax
	jb	.L175
	leaq	0(,%rbp,4), %rdi
	call	_Znam@PLT
	xorl	%edx, %edx
	jmp	.L176
.L175:
	movq	552(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L201
	call	__cxa_throw_bad_array_new_length@PLT
.L176:
	cmpq	%rdx, %rbp
	je	.L209
	xorl	%edi, %edi
	movl	%edi, (%rax,%rdx,4)
	incq	%rdx
	jmp	.L176
.L209:
	movq	(%rbx), %rbp
	imulq	8(%rbx), %rbp
	movq	%rax, 56(%rbx)
	movabsq	$2305843009213693950, %rax
	cmpq	%rbp, %rax
	jb	.L179
	leaq	0(,%rbp,4), %rdi
	call	_Znam@PLT
	xorl	%edx, %edx
	jmp	.L180
.L179:
	movq	552(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L201
	call	__cxa_throw_bad_array_new_length@PLT
.L180:
	cmpq	%rdx, %rbp
	je	.L210
	xorl	%esi, %esi
	movl	%esi, (%rax,%rdx,4)
	incq	%rdx
	jmp	.L180
.L210:
	movq	(%rbx), %rbp
	imulq	8(%rbx), %rbp
	movq	%rax, 64(%rbx)
	movabsq	$2305843009213693950, %rax
	cmpq	%rbp, %rax
	jb	.L183
	leaq	0(,%rbp,4), %rdi
	call	_Znam@PLT
	xorl	%edx, %edx
	jmp	.L184
.L183:
	movq	552(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L201
	call	__cxa_throw_bad_array_new_length@PLT
.L184:
	cmpq	%rdx, %rbp
	je	.L211
	xorl	%ecx, %ecx
	movl	%ecx, (%rax,%rdx,4)
	incq	%rdx
	jmp	.L184
.L211:
	movq	%rax, 72(%rbx)
	xorl	%r14d, %r14d
	leaq	24(%rsp), %r13
.L187:
	movq	(%rbx), %rax
	imulq	8(%rbx), %rax
	cmpq	%rax, %r14
	jnb	.L212
	movq	48(%rbx), %rsi
	leaq	0(,%r14,4), %rbp
	movq	%r13, %rdi
	addq	%rbp, %rsi
	call	_ZNSirsERi@PLT
	movq	56(%rbx), %rsi
	movq	%rax, %rdi
	addq	%rbp, %rsi
	call	_ZNSirsERi@PLT
	movq	16(%rbx), %rsi
	movq	%r13, %rdi
	addq	%rbp, %rsi
	call	_ZNSirsERi@PLT
	movq	24(%rbx), %rsi
	movq	%rax, %rdi
	addq	%rbp, %rsi
	call	_ZNSirsERi@PLT
	movq	32(%rbx), %rsi
	movq	%rax, %rdi
	addq	%rbp, %rsi
	call	_ZNSirsERi@PLT
	movq	40(%rbx), %rsi
	movq	%rax, %rdi
	addq	%rbp, %rsi
	call	_ZNSirsERi@PLT
	incq	%r14
	jmp	.L187
.L212:
	movq	%r13, %rdi
	call	_ZNSt13basic_fstreamIcSt11char_traitsIcEE5closeEv@PLT
	leaq	80(%rbx), %rdi
	movl	$1, %esi
	call	_Z10cudaMallocIbE9cudaErrorPPT_m.isra.0
	movq	(%rbx), %rsi
	imulq	8(%rbx), %rsi
	leaq	88(%rbx), %rdi
	salq	$2, %rsi
	call	_Z10cudaMallocIbE9cudaErrorPPT_m.isra.0
	movq	(%rbx), %rsi
	imulq	8(%rbx), %rsi
	leaq	96(%rbx), %rdi
	salq	$2, %rsi
	call	_Z10cudaMallocIbE9cudaErrorPPT_m.isra.0
	movq	(%rbx), %rsi
	imulq	8(%rbx), %rsi
	leaq	104(%rbx), %rdi
	salq	$2, %rsi
	call	_Z10cudaMallocIbE9cudaErrorPPT_m.isra.0
	movq	(%rbx), %rsi
	imulq	8(%rbx), %rsi
	leaq	112(%rbx), %rdi
	salq	$2, %rsi
	call	_Z10cudaMallocIbE9cudaErrorPPT_m.isra.0
	movq	(%rbx), %rsi
	imulq	8(%rbx), %rsi
	leaq	120(%rbx), %rdi
	salq	$2, %rsi
	call	_Z10cudaMallocIbE9cudaErrorPPT_m.isra.0
	movq	(%rbx), %rsi
	imulq	8(%rbx), %rsi
	leaq	128(%rbx), %rdi
	salq	$2, %rsi
	call	_Z10cudaMallocIbE9cudaErrorPPT_m.isra.0
	movq	(%rbx), %rsi
	imulq	8(%rbx), %rsi
	leaq	136(%rbx), %rdi
	salq	$2, %rsi
	call	_Z10cudaMallocIbE9cudaErrorPPT_m.isra.0
	movq	(%rbx), %rsi
	imulq	8(%rbx), %rsi
	leaq	144(%rbx), %rdi
	salq	$2, %rsi
	call	_Z10cudaMallocIbE9cudaErrorPPT_m.isra.0
	movq	80(%rbx), %rdi
	movl	$1, %edx
	xorl	%esi, %esi
	call	cudaMemset@PLT
	movq	(%rbx), %rdx
	imulq	8(%rbx), %rdx
	xorl	%esi, %esi
	movq	136(%rbx), %rdi
	salq	$2, %rdx
	call	cudaMemset@PLT
	movq	(%rbx), %rdx
	imulq	8(%rbx), %rdx
	xorl	%esi, %esi
	movq	144(%rbx), %rdi
	salq	$2, %rdx
	call	cudaMemset@PLT
	movq	(%rbx), %rdx
	imulq	8(%rbx), %rdx
	movl	$1, %ecx
	movq	48(%rbx), %rsi
	movq	120(%rbx), %rdi
	salq	$2, %rdx
	call	cudaMemcpy@PLT
	movq	(%rbx), %rdx
	imulq	8(%rbx), %rdx
	movl	$1, %ecx
	movq	56(%rbx), %rsi
	movq	128(%rbx), %rdi
	salq	$2, %rdx
	call	cudaMemcpy@PLT
	movq	(%rbx), %rdx
	imulq	8(%rbx), %rdx
	movl	$1, %ecx
	movq	16(%rbx), %rsi
	movq	88(%rbx), %rdi
	salq	$2, %rdx
	call	cudaMemcpy@PLT
	movq	(%rbx), %rdx
	imulq	8(%rbx), %rdx
	movl	$1, %ecx
	movq	24(%rbx), %rsi
	movq	96(%rbx), %rdi
	salq	$2, %rdx
	call	cudaMemcpy@PLT
	movq	(%rbx), %rdx
	imulq	8(%rbx), %rdx
	movl	$1, %ecx
	movq	32(%rbx), %rsi
	movq	104(%rbx), %rdi
	salq	$2, %rdx
	call	cudaMemcpy@PLT
	movq	(%rbx), %rdx
	imulq	8(%rbx), %rdx
	movl	$1, %ecx
	movq	40(%rbx), %rsi
	movq	112(%rbx), %rdi
	salq	$2, %rdx
	call	cudaMemcpy@PLT
.LEHE1:
	movq	%r13, %rdi
	call	_ZNSt13basic_fstreamIcSt11char_traitsIcEED1Ev@PLT
	movq	552(%rsp), %rax
	subq	%fs:40, %rax
	je	.L191
	jmp	.L201
.L200:
	endbr64
	movq	%rax, %rbx
.L189:
	movq	%r12, %rdi
	call	_ZNSt13basic_fstreamIcSt11char_traitsIcEED1Ev@PLT
	movq	552(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L201
	movq	%rbx, %rdi
.LEHB2:
	call	_Unwind_Resume@PLT
.LEHE2:
.L201:
	call	__stack_chk_fail@PLT
.L191:
	addq	$560, %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
.LFE3776:
	.section	.gcc_except_table
.LLSDA3776:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE3776-.LLSDACSB3776
.LLSDACSB3776:
	.uleb128 .LEHB0-.LFB3776
	.uleb128 .LEHE0-.LEHB0
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB1-.LFB3776
	.uleb128 .LEHE1-.LEHB1
	.uleb128 .L200-.LFB3776
	.uleb128 0
	.uleb128 .LEHB2-.LFB3776
	.uleb128 .LEHE2-.LEHB2
	.uleb128 0
	.uleb128 0
.LLSDACSE3776:
	.text
	.size	_ZN4Data4ReadEPKc, .-_ZN4Data4ReadEPKc
	.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 <hip/hip_runtime.h>
class Data
{
private:
	//bool active;
	size_t rowSize, columnSize;
	int *weightLeft, *weightRight, *weightUp, *weightDown;
	int *weightS, *weightT;
	int *height, *capacity;

	//int flow;
	//int *bfsTag;

    // variables on device (GPU)
    bool *device_active;
	int *device_weightLeft, *device_weightRight, *device_weightUp, *device_weightDown;
	int *device_weightS, *device_weightT;
	int *device_height, *device_capacity;

public:
	Data();
	Data(size_t rowSize, size_t columnSize);
    ~Data();
	void Read(const char * filename);
	int GetFlow();
	void BfsFromT();
    void Print();
};

	.text
	.file	"Data.hip"
	.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 __hip_cuid_

	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.amdgpu_metadata
---
amdhsa.kernels:  []
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
