﻿char *title = "Floyd's algorithm";
char *description = "Алгоритм Флойда - поиск всех кратчайших путей в графе";
/*
Алгоритм Флойда является одним из методов поиска кратчайших путей в графе. 
В отличии от алгоритма Дейкстры, который позволяет при доведении до конца построить 
ориентированное дерево кратчайших путей от некоторой вершины, метод Флойда позволяет 
найти длины всех кратчайших путей в графе. Конечно эта задача может быть решена 
и многократным применением алгоритма Дейкстры (каждый раз последовательно выбираем 
вершину от первой до N-ной, пока не получим кратчайшие пути от всех вершин графа), 
однако реализация подобной процедуры потребовала бы значительных вычислительных затрат.

Прежде чем представлять алгоритмы, необходимо ввести некоторые обозначения. 
Перенумеруем вершины исходного графа целыми числами от 1 до N. Обозначим через di,jm длину кратчайшего пути 
из вершинм i в вершину j, который в качестве промежуточных может содержать только первые m вершин графа. 
(Напомним, что промежуточной вершиной пути является любая принадлежащая ему вершина, не совпадающая 
с его начальной или конечной вершинами.) Если между вершинами i и j не существует ни одного пути указанного типа, 
то условно будем считать, что di,jm=∞. Из данного определения величин di,jm следует, что величина di,j0, 
представляет длину кратчайшего пути из вершины i в вершину j, не имеющего промежуточных вершин, 
т. е. длину кратчайшей дуги, соединяющей i с j (если такие дуги присутствуют в графе). 
для любой вершины i положим di,im= 0. Отметим далее, что величина di,jmпредставляет длину кратчайшего пути 
между вершинами i и j.

Обозначим через Dm матрицу размера NxN, элемент (i, j) которой совпадает с di,jm. 
Если в исходном графе нам известна длина каждой дуги, то мы можем сформировать матрицу D0. 
Наша цель состоит в определении матрицы DN, представляющей кратчайшие пути между всеми вершинами рассматриваемого графа.

В алгоритме Флойда в качестве исходной выступает матрица D0. 
Вначале из этой матрицы вычисляется матрица D1. 
Затем по матрице D1 вычисляется матрицав D2 и т. д. 
Процесс повторяется до тех пор, пока по матрице DN-1 не будет вычислена матрица DN.

Рассмотрим основную идею, лежащую в основе алгоритма Флойда. 
Суть алгоритма Флойда заключается в проверке того, не окажется ли путь из вершины i в вершину j короче, 
если он будет проходить через некоторую промежуточную вершину m. Предположим, что нам известны:

кратчайший путь из вершины i в вершину m, в котором в качестве промежуточных допускается использование только первых (m - 1) вершин;
кратчайший путь из вершины m в вершину j, в котором в качестве промежуточных допускается использование только первых (m - 1) вершин;
кратчайший путь из вершины i в вершину j, в котором в качестве промежуточных допускается использование только первых (m - 1) вершин.

Поскольку по предположению исходный граф не может содержать контуров отрицательной длины, 
один из двух путей — путь, совпадающий с представленным в пункте 3, или путь, являющийся объединением 
путей из пунктов 1 и 2 — должен быть кратчайшим путем из вершины i в вершину j, 
в котором в качестве промежуточных допускается использование только первых m вершин. Таким образом,

di,jm=min{ di,mm-1+ dm,jm-1; di,jm-1}

Из соотношения видно, что для вычисления элементов матрицы Dm необходимо располагать лишь элементами матрицы Dm-1. 
Более того, соответствующие вычисления могут быть проведены без обращения к исходному графу. 
Теперь мм в состоянии дать формальное описание алгоритма Флойда для нахождения на графе кратчайших путей 
между всеми парами вершин. 

Алгоритм

Перенумеровать вершины графа от 1 до N целыми числами, определить матрицу D0, каждый элемент di,j  
которой есть длина кратчайшей дуги между вершинами i и j. Если такой дуги нет, положить значение элемента 
равным ∞. Кроме того, положить значения диагонального элемента di,iравным 0.
Для целого m, последовательно принимающего значения 1...N определить по элементам матрицы Dm-1 элементы Dm
Алгоритм заканчивается получением матрицы всех кратчайших путей DN, N – число вершин графа.
 
Напомним, для определения по известным элементам матрицы Dm-1 элементов матрицы  Dm 
в алгоритме Флойда применяется рекурсивное соотношение:

di,jm=min{ di,mm-1+ dm,jm-1; di,jm-1}

di,jm – элемент матрицы Dm, di,jm-1 – элементы матрицы Dm-1 найденой на предыдущем шаге алгоритма.
*/
#include <iostream>
#include <cstdio>
#include <cstdlib>
#include <stdio.h>
#include <string.h>
#include <time.h>
#include <ctype.h>
#include <limits.h>
#include <cuda.h>
#include <cuda_runtime.h>
#include <device_launch_parameters.h>

#define assert( bool ) 
int strempty(const char *p)
{
	if (!p)
		return (1);
	for (; *p; p++)
		if (!isspace(*p))
			return (0);
	return (1);
}
char *mystrtok(char **m, char *s, char c)
{
	char *p = s ? s : *m;
	if (!*p)
		return 0;
	*m = strchr(p, c);
	if (*m)
		*(*m)++ = 0;
	else
		*m = p + strlen(p);
	return p;
}
#ifndef max
#define max( a, b ) ( ((a) > (b)) ? (a) : (b) )
#endif

#ifndef min
#define min( a, b ) ( ((a) < (b)) ? (a) : (b) )
#endif

__global__ void global_init(long *prev, long *next, int *intermedian, int m, int n){
	for (int id = blockDim.x*blockIdx.x + threadIdx.x; id < n*n; id += blockDim.x*gridDim.x) {
		intermedian[id] = -1;
	}
}

__global__ void global_floyd(long *prev, long *next, int *intermedian, int m, int n){
	for (int id = blockDim.x*blockIdx.x + threadIdx.x; id < n*n; id += blockDim.x*gridDim.x) {
		int i = id/n;
		int j = id%n;
		if(i != j && prev[i*n+m] != LONG_MAX && prev[m*n+j] != LONG_MAX && (prev[i*n+m] + prev[m*n+j]) < prev[i*n+j]){
			next[id] = prev[i*n+m] + prev[m*n+j];
			intermedian[id] = m;
		}
		else {
			next[id] = prev[id];
		}
	}
}

__host__ void host_floyd(int gridSize, int blockSize, long *matrix, int *intermedian, int n)
{
	cudaError_t err;
	long *device_prev;
	long *device_next;
	int *device_intermedian;

	err = cudaMalloc((void**)&device_prev, n*n * sizeof(long));
	err = cudaMalloc((void**)&device_next, n*n * sizeof(long));
	err = cudaMalloc((void**)&device_intermedian, n*n * sizeof(int));

	err = cudaMemcpy(device_prev, matrix, n*n*sizeof(long), cudaMemcpyHostToDevice);

	int blocks = (gridSize > 0)? gridSize : min(max(1, (int)pow((double)n*n, 0.333333333333333)), 15);
	int threads = (blockSize > 0)? blockSize : min(max(1, (int)pow((double)n*n, 0.333333333333333)), 15);
	global_init <<< blocks, threads >>>(device_prev, device_next, device_intermedian, -1, n);

	for(int m = 0; m < n ; m++){
		global_floyd <<< blocks, threads >>>(device_prev, device_next, device_intermedian, m, n);
		err = cudaMemcpy(matrix, device_next, n*n*sizeof(long), cudaMemcpyDeviceToHost);
		for (int i = 0; i < n; i++){
			for (int j = 0; j < n; j++){
				printf("%ld%s", matrix[i*n + j], ((j == n - 1) ? "\n" : ";"));
			}
		}
		printf("\n");
		long * t = device_prev; device_prev = device_next; device_next = t;
	}

	err = cudaMemcpy(matrix, device_prev, n*n*sizeof(long), cudaMemcpyDeviceToHost);
	err = cudaMemcpy(intermedian, device_intermedian, n*n*sizeof(int), cudaMemcpyDeviceToHost);

	err = cudaFree(device_prev);
	err = cudaFree(device_next);
	err = cudaFree(device_intermedian);

	err = err;
}

int main(int argc, char* argv[])
{
	int gridSize = 0;
	int blockSize = 0;

	printf("Title :\t%s\n", title); fflush(stdout);

	if (argc < 4) {
		printf("Usage :\t%s [-g <gridSize>] [-b <blockSize>] <inputfilename> <outputfilename> <intermedianfilename>\n", argv[0]); fflush(stdout);
		printf("\tinputfilename - source matrix of path prices or empty\n"); fflush(stdout);
		printf("\toutputfilename - output floyd's matrix of path prices\n"); fflush(stdout);
		printf("\tintermedianfilename - output matrix of intermedian points or empty\n"); fflush(stdout);
		exit(-1);
	}

	int argId = 1;
	for(; argId < argc && argv[argId][0]=='-' ; argId++){
		switch(argv[argId][1]){
		case 'g':
			gridSize = atoi(argv[++argId]);
			break;
		case 'b':
			blockSize = atoi(argv[++argId]);
			break;
		}
	}

	char *inputFileName = argv[argId++];
	char *outputFileName = argv[argId++];
	char *intermedianFileName = argv[argId++];

	printf("Input File Name :\t%s\n", inputFileName); fflush(stdout);
	printf("Output File Name :\t%s\n", outputFileName); fflush(stdout);
	printf("Intermedian File Name :\t%s\n", intermedianFileName); fflush(stdout);

	char buffer[4096];
	char *tok;
	char *p;
	int n;         /* Ранг текущего массива */
	long *matrix;  /* Массив цен */
	int *intermedian;  /* Массивов промежуточных точек */
	int i, j;

	FILE *fs = fopen(inputFileName, "r");
	if (fs == NULL) {
		fprintf(stderr, "File open error (%s)\n", inputFileName); fflush(stderr);
		exit(-1);
	}

	n = 0;

	/* Заполняем массив числами из файла */
	/* Операция выполняетя только на хост процессе */
	/* Операция выполняетя в два прохода по файлу */
	/* На первом проходе определяется ранг матрицы */
	/* На втором проходе считываются данные */
	for (i = 0; (tok = fgets(buffer, sizeof(buffer), fs)) != NULL; i++)
	{
		j = 0;
		for (tok = mystrtok(&p, tok, ';'); tok != NULL; tok = mystrtok(&p, NULL, ';'))
		{
			j++;
		}
		n = max(n, j);
	}
	n = max(n, i);

	matrix = (long *)malloc(n*n*sizeof(long));
	intermedian = (int *)malloc(n*n*sizeof(int));

	fseek(fs, 0, SEEK_SET);

	for (i = 0; (tok = fgets(buffer, sizeof(buffer), fs)) != NULL; i++)
	{
		j = 0;
		for (tok = mystrtok(&p, tok, ';'); tok != NULL; tok = mystrtok(&p, NULL, ';'))
		{
			/* Пустые элементы - это запрещённые пути */
			matrix[n*i + j++] = strempty(tok) ? LONG_MAX : atol(tok);
		}
		for (; j < n; j++) matrix[n*i + j] = LONG_MAX;
	}
	for (j = 0; j < (n - i)*n; j++) matrix[n*i + j] = LONG_MAX;
	for (i = 0; i < n; i++) matrix[n*i + i] = LONG_MAX; /* Запрещаем петли */

	fclose(fs);

	printf("Matrix rank :\t%d\n", n);
	for (i = 0; i < n; i++){
		for (j = 0; j < n; j++){
			printf("%ld%s", matrix[i*n + j], ((j == n - 1) ? "\n" : "\t"));
		}
	}
	fflush(stdout);

	// Find/set the device.
	int device_size = 0;
	cudaGetDeviceCount(&device_size);
	for (i = 0; i < device_size; ++i)
	{
		cudaDeviceProp cudaDeviceProp;
		cudaGetDeviceProperties(&cudaDeviceProp, i);
		printf("Running on GPU %d (%s)\n", i, cudaDeviceProp.name); 

		printf("Device has ECC support enabled %d\n",cudaDeviceProp.ECCEnabled);
		printf("Number of asynchronous engines %d\n",cudaDeviceProp.asyncEngineCount);
		printf("Device can map host memory with cudaHostAlloc/cudaHostGetDevicePointer %d\n",cudaDeviceProp.canMapHostMemory);
		printf("Clock frequency in kilohertz %d\n",cudaDeviceProp.clockRate);
		printf("Compute mode (See cudaComputeMode) %d\n",cudaDeviceProp.computeMode);
		printf("Device can possibly execute multiple kernels concurrently %d\n",cudaDeviceProp.concurrentKernels);
		printf("Device can concurrently copy memory and execute a kernel. Deprecated. Use instead asyncEngineCount. %d\n",cudaDeviceProp.deviceOverlap);
		printf("Device is integrated as opposed to discrete %d\n",cudaDeviceProp.integrated);
		printf("Specified whether there is a run time limit on kernels %d\n",cudaDeviceProp.kernelExecTimeoutEnabled);
		printf("Size of L2 cache in bytes %d\n",cudaDeviceProp.l2CacheSize);
		printf("Major compute capability %d\n",cudaDeviceProp.major);
		printf("Maximum size of each dimension of a grid %d\n",cudaDeviceProp.maxGridSize[0]);
		printf("Maximum size of each dimension of a grid %d\n",cudaDeviceProp.maxGridSize[1]);
		printf("Maximum size of each dimension of a grid %d\n",cudaDeviceProp.maxGridSize[2]);
		printf("Maximum 1D surface size %d\n",cudaDeviceProp.maxSurface1D);
		printf("Maximum 1D layered surface dimensions %d\n",cudaDeviceProp.maxSurface1DLayered[0]);
		printf("Maximum 1D layered surface dimensions %d\n",cudaDeviceProp.maxSurface1DLayered[1]);
		printf("Maximum 2D surface dimensions %d\n",cudaDeviceProp.maxSurface2D[0]);
		printf("Maximum 2D surface dimensions %d\n",cudaDeviceProp.maxSurface2D[1]);
		printf("Maximum 2D layered surface dimensions %d\n",cudaDeviceProp.maxSurface2DLayered[0]);
		printf("Maximum 2D layered surface dimensions %d\n",cudaDeviceProp.maxSurface2DLayered[1]);
		printf("Maximum 2D layered surface dimensions %d\n",cudaDeviceProp.maxSurface2DLayered[2]);
		printf("Maximum 3D surface dimensions %d\n",cudaDeviceProp.maxSurface3D[0]);
		printf("Maximum 3D surface dimensions %d\n",cudaDeviceProp.maxSurface3D[1]);
		printf("Maximum 3D surface dimensions %d\n",cudaDeviceProp.maxSurface3D[2]);
		printf("Maximum Cubemap surface dimensions %d\n",cudaDeviceProp.maxSurfaceCubemap);
		printf("Maximum Cubemap layered surface dimensions %d\n",cudaDeviceProp.maxSurfaceCubemapLayered[0]);
		printf("Maximum Cubemap layered surface dimensions %d\n",cudaDeviceProp.maxSurfaceCubemapLayered[1]);
		printf("Maximum 1D texture size %d\n",cudaDeviceProp.maxTexture1D);
		printf("Maximum 1D layered texture dimensions %d\n",cudaDeviceProp.maxTexture1DLayered[0]);
		printf("Maximum 1D layered texture dimensions %d\n",cudaDeviceProp.maxTexture1DLayered[1]);
		printf("Maximum size for 1D textures bound to linear memory %d\n",cudaDeviceProp.maxTexture1DLinear);
		printf("Maximum 1D mipmapped texture size %d\n",cudaDeviceProp.maxTexture1DMipmap);
		printf("Maximum 2D texture dimensions %d\n",cudaDeviceProp.maxTexture2D[0]);
		printf("Maximum 2D texture dimensions %d\n",cudaDeviceProp.maxTexture2D[1]);
		printf("Maximum 2D texture dimensions if texture gather operations have to be performed %d\n",cudaDeviceProp.maxTexture2DGather[0]);
		printf("Maximum 2D texture dimensions if texture gather operations have to be performed %d\n",cudaDeviceProp.maxTexture2DGather[1]);
		printf("Maximum 2D layered texture dimensions %d\n",cudaDeviceProp.maxTexture2DLayered[0]);
		printf("Maximum 2D layered texture dimensions %d\n",cudaDeviceProp.maxTexture2DLayered[1]);
		printf("Maximum 2D layered texture dimensions %d\n",cudaDeviceProp.maxTexture2DLayered[2]);
		printf("Maximum dimensions (width, height, pitch) for 2D textures bound to pitched memory %d\n",cudaDeviceProp.maxTexture2DLinear[0]);
		printf("Maximum dimensions (width, height, pitch) for 2D textures bound to pitched memory %d\n",cudaDeviceProp.maxTexture2DLinear[1]);
		printf("Maximum dimensions (width, height, pitch) for 2D textures bound to pitched memory %d\n",cudaDeviceProp.maxTexture2DLinear[2]);
		printf("Maximum 2D mipmapped texture dimensions %d\n",cudaDeviceProp.maxTexture2DMipmap[0]);
		printf("Maximum 2D mipmapped texture dimensions %d\n",cudaDeviceProp.maxTexture2DMipmap[1]);
		printf("Maximum 3D texture dimensions %d\n",cudaDeviceProp.maxTexture3D[0]);
		printf("Maximum 3D texture dimensions %d\n",cudaDeviceProp.maxTexture3D[1]);
		printf("Maximum 3D texture dimensions %d\n",cudaDeviceProp.maxTexture3D[2]);
		printf("Maximum alternate 3D texture dimensions %d\n",cudaDeviceProp.maxTexture3DAlt[0]);
		printf("Maximum alternate 3D texture dimensions %d\n",cudaDeviceProp.maxTexture3DAlt[1]);
		printf("Maximum alternate 3D texture dimensions %d\n",cudaDeviceProp.maxTexture3DAlt[2]);
		printf("Maximum Cubemap texture dimensions %d\n",cudaDeviceProp.maxTextureCubemap);
		printf("Maximum Cubemap layered texture dimensions %d\n",cudaDeviceProp.maxTextureCubemapLayered[0]);
		printf("Maximum Cubemap layered texture dimensions %d\n",cudaDeviceProp.maxTextureCubemapLayered[1]);
		printf("Maximum size of each dimension of a block %d\n",cudaDeviceProp.maxThreadsDim[0]);
		printf("Maximum size of each dimension of a block %d\n",cudaDeviceProp.maxThreadsDim[1]);
		printf("Maximum size of each dimension of a block %d\n",cudaDeviceProp.maxThreadsDim[2]);
		printf("Maximum number of threads per block %d\n",cudaDeviceProp.maxThreadsPerBlock);
		printf("Maximum resident threads per multiprocessor %d\n",cudaDeviceProp.maxThreadsPerMultiProcessor);
		printf("Maximum pitch in bytes allowed by memory copies %d\n",cudaDeviceProp.memPitch);
		printf("Global memory bus width in bits %d\n",cudaDeviceProp.memoryBusWidth);
		printf("Peak memory clock frequency in kilohertz %d\n",cudaDeviceProp.memoryClockRate);
		printf("Minor compute capability %d\n",cudaDeviceProp.minor);
		printf("Number of multiprocessors on device %d\n",cudaDeviceProp.multiProcessorCount);
		printf("PCI bus ID of the device %d\n",cudaDeviceProp.pciBusID);
		printf("PCI device ID of the device %d\n",cudaDeviceProp.pciDeviceID);
		printf("PCI domain ID of the device %d\n",cudaDeviceProp.pciDomainID);
		printf("32-bit registers available per block %d\n",cudaDeviceProp.regsPerBlock);
		printf("Shared memory available per block in bytes %d\n",cudaDeviceProp.sharedMemPerBlock);
		printf("Device supports stream priorities %d\n",cudaDeviceProp.streamPrioritiesSupported);
		printf("Alignment requirements for surfaces %d\n",cudaDeviceProp.surfaceAlignment);
		printf("1 if device is a Tesla device using TCC driver, 0 otherwise %d\n",cudaDeviceProp.tccDriver);
		printf("Alignment requirement for textures %d\n",cudaDeviceProp.textureAlignment);
		printf("Pitch alignment requirement for texture references bound to pitched memory %d\n",cudaDeviceProp.texturePitchAlignment);
		printf("Constant memory available on device in bytes %d\n",cudaDeviceProp.totalConstMem);
		printf("Global memory available on device in bytes %d\n",cudaDeviceProp.totalGlobalMem);
		printf("Device shares a unified address space with the host %d\n",cudaDeviceProp.unifiedAddressing);
		printf("Warp size in threads %d\n",cudaDeviceProp.warpSize);

		fflush(stdout);
	}

	host_floyd(gridSize, blockSize, matrix, intermedian, n);

	cudaDeviceReset();

	/* Bыводим результаты */
	fs = fopen(outputFileName, "w");
	if (fs == NULL) {
		fprintf(stderr, "File open error (%s)\n", outputFileName); fflush(stderr);
		exit(-1);
	}
	for (i = 0; i < n; i++){
		for (j = 0; j < n; j++){
			if (matrix[i*n + j] != LONG_MAX)
				fprintf(fs, "%ld%s", matrix[i*n + j], ((j == n - 1) ? "\n" : ";"));
			else
				fprintf(fs, "%s", ((j == n - 1) ? "\n" : ";"));
		}
	}
	fclose(fs);

	/* Bыводим результаты */
	fs = fopen(intermedianFileName, "w");
	if (fs == NULL) {
		fprintf(stderr, "File open error (%s)\n", intermedianFileName); fflush(stderr);
		exit(-1);
	}
	for (i = 0; i < n; i++){
		for (j = 0; j < n; j++){
			if (intermedian[i*n + j] >= 0)
				fprintf(fs, "%d%s", intermedian[i*n + j], ((j == n - 1) ? "\n" : ";"));
			else
				fprintf(fs, "%s", ((j == n - 1) ? "\n" : ";"));
		}
	}
	fclose(fs);

	free(matrix);
	free(intermedian);

	exit(0);
}
	.file	"tmpxft_002d93c8_00000000-6_floyd.cudafe1.cpp"
	.text
#APP
	.globl _ZSt21ios_base_library_initv

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB3641:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE3641:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z8stremptyPKc
	.type	_Z8stremptyPKc, @function
_Z8stremptyPKc:
.LFB3635:
	.cfi_startproc
	endbr64
	testq	%rdi, %rdi
	jne	.L10
	movl	$1, %eax
	ret
.L10:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	movq	%rdi, %rbx
.L4:
	movsbl	(%rbx), %edi
	testb	%dil, %dil
	je	.L15
	call	isspace@PLT
	testl	%eax, %eax
	je	.L2
	incq	%rbx
	jmp	.L4
.L15:
	movl	$1, %eax
.L2:
	popq	%rbx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3635:
	.size	_Z8stremptyPKc, .-_Z8stremptyPKc
	.globl	_Z8mystrtokPPcS_c
	.type	_Z8mystrtokPPcS_c, @function
_Z8mystrtokPPcS_c:
.LFB3636:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	movq	%rdi, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	movq	%rsi, %rbx
	pushq	%rcx
	.cfi_def_cfa_offset 32
	testq	%rsi, %rsi
	jne	.L17
	movq	(%rdi), %rbx
.L17:
	cmpb	$0, (%rbx)
	je	.L21
	movsbl	%dl, %esi
	movq	%rbx, %rdi
	call	strchr@PLT
	movq	%rax, 0(%rbp)
	testq	%rax, %rax
	je	.L19
	leaq	1(%rax), %rdx
	movq	%rdx, 0(%rbp)
	movb	$0, (%rax)
	jmp	.L18
.L19:
	movq	%rbx, %rdi
	call	strlen@PLT
	addq	%rbx, %rax
	movq	%rax, 0(%rbp)
	jmp	.L18
.L21:
	xorl	%ebx, %ebx
.L18:
	movq	%rbx, %rax
	popq	%rdx
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3636:
	.size	_Z8mystrtokPPcS_c, .-_Z8mystrtokPPcS_c
	.globl	_Z37__device_stub__Z11global_initPlS_PiiiPlS_Piii
	.type	_Z37__device_stub__Z11global_initPlS_PiiiPlS_Piii, @function
_Z37__device_stub__Z11global_initPlS_PiiiPlS_Piii:
.LFB3663:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movq	%rdi, 24(%rsp)
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	32(%rsp), %rdx
	movl	%ecx, 4(%rsp)
	leaq	40(%rsp), %rcx
	movl	%r8d, (%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	8(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 120(%rsp)
	movq	%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	.L24
	pushq	40(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z11global_initPlS_Piii(%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
.L24:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L26
	call	__stack_chk_fail@PLT
.L26:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3663:
	.size	_Z37__device_stub__Z11global_initPlS_PiiiPlS_Piii, .-_Z37__device_stub__Z11global_initPlS_PiiiPlS_Piii
	.globl	_Z11global_initPlS_Piii
	.type	_Z11global_initPlS_Piii, @function
_Z11global_initPlS_Piii:
.LFB3664:
	.cfi_startproc
	endbr64
	jmp	_Z37__device_stub__Z11global_initPlS_PiiiPlS_Piii
	.cfi_endproc
.LFE3664:
	.size	_Z11global_initPlS_Piii, .-_Z11global_initPlS_Piii
	.globl	_Z38__device_stub__Z12global_floydPlS_PiiiPlS_Piii
	.type	_Z38__device_stub__Z12global_floydPlS_PiiiPlS_Piii, @function
_Z38__device_stub__Z12global_floydPlS_PiiiPlS_Piii:
.LFB3665:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movq	%rdi, 24(%rsp)
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	32(%rsp), %rdx
	movl	%ecx, 4(%rsp)
	leaq	40(%rsp), %rcx
	movl	%r8d, (%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	8(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 120(%rsp)
	movq	%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	.L29
	pushq	40(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z12global_floydPlS_Piii(%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
.L29:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L31
	call	__stack_chk_fail@PLT
.L31:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3665:
	.size	_Z38__device_stub__Z12global_floydPlS_PiiiPlS_Piii, .-_Z38__device_stub__Z12global_floydPlS_PiiiPlS_Piii
	.globl	_Z12global_floydPlS_Piii
	.type	_Z12global_floydPlS_Piii, @function
_Z12global_floydPlS_Piii:
.LFB3666:
	.cfi_startproc
	endbr64
	jmp	_Z38__device_stub__Z12global_floydPlS_PiiiPlS_Piii
	.cfi_endproc
.LFE3666:
	.size	_Z12global_floydPlS_Piii, .-_Z12global_floydPlS_Piii
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"\n"
.LC1:
	.string	";"
.LC3:
	.string	"%ld%s"
	.text
	.globl	_Z10host_floydiiPlPii
	.type	_Z10host_floydiiPlPii, @function
_Z10host_floydiiPlPii:
.LFB3637:
	.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
	movl	%r8d, %r13d
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	imull	%r8d, %r13d
	movl	%esi, %r12d
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movl	%edi, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movl	%r8d, %ebx
	movslq	%r13d, %r13
	subq	$136, %rsp
	.cfi_def_cfa_offset 192
	movq	%rdx, 8(%rsp)
	leaq	72(%rsp), %rdi
	movq	%rcx, 32(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 120(%rsp)
	xorl	%eax, %eax
	leaq	0(,%r13,8), %rax
	movq	%rax, %rsi
	movq	%rax, (%rsp)
	call	cudaMalloc@PLT
	movq	(%rsp), %rsi
	leaq	80(%rsp), %rdi
	call	cudaMalloc@PLT
	leaq	0(,%r13,4), %rax
	leaq	88(%rsp), %rdi
	movq	%rax, %rsi
	movq	%rax, 16(%rsp)
	call	cudaMalloc@PLT
	movq	(%rsp), %rdx
	movq	8(%rsp), %rsi
	movl	$1, %ecx
	movq	72(%rsp), %rdi
	call	cudaMemcpy@PLT
	testl	%ebp, %ebp
	jg	.L35
	cvtsi2sdl	%ebx, %xmm0
	movsd	.LC2(%rip), %xmm1
	mulsd	%xmm0, %xmm0
	movq	%xmm0, %rbp
	call	pow@PLT
	cvttsd2sil	%xmm0, %eax
	testl	%eax, %eax
	jg	.L36
.L38:
	movsd	.LC2(%rip), %xmm1
	movq	%rbp, %xmm0
	call	pow@PLT
	cvttsd2sil	%xmm0, %eax
	testl	%eax, %eax
	jg	.L37
	movl	$1, %ebp
	jmp	.L35
.L36:
	movsd	.LC2(%rip), %xmm1
	movq	%rbp, %xmm0
	call	pow@PLT
	cvttsd2sil	%xmm0, %eax
	cmpl	$14, %eax
	jle	.L38
	movl	$15, %ebp
	jmp	.L35
.L37:
	movsd	.LC2(%rip), %xmm1
	movq	%rbp, %xmm0
	call	pow@PLT
	cvttsd2sil	%xmm0, %ebp
.L35:
	testl	%r12d, %r12d
	jg	.L39
	cvtsi2sdl	%ebx, %xmm0
	movsd	.LC2(%rip), %xmm1
	mulsd	%xmm0, %xmm0
	movq	%xmm0, %r12
	call	pow@PLT
	cvttsd2sil	%xmm0, %eax
	testl	%eax, %eax
	jg	.L40
.L42:
	movsd	.LC2(%rip), %xmm1
	movq	%r12, %xmm0
	call	pow@PLT
	cvttsd2sil	%xmm0, %eax
	testl	%eax, %eax
	jg	.L41
	movl	$1, %r12d
	jmp	.L39
.L40:
	movsd	.LC2(%rip), %xmm1
	movq	%r12, %xmm0
	call	pow@PLT
	cvttsd2sil	%xmm0, %eax
	cmpl	$14, %eax
	jle	.L42
	movl	$15, %r12d
	jmp	.L39
.L41:
	movsd	.LC2(%rip), %xmm1
	movq	%r12, %xmm0
	call	pow@PLT
	cvttsd2sil	%xmm0, %r12d
.L39:
	movabsq	$4294967297, %rax
	movl	%r12d, 108(%rsp)
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movq	%rax, 112(%rsp)
	movl	116(%rsp), %ecx
	movq	%rax, 100(%rsp)
	movq	108(%rsp), %rdx
	movl	%ebp, 96(%rsp)
	movl	104(%rsp), %esi
	movq	96(%rsp), %rdi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L43
	movq	88(%rsp), %rdx
	movq	80(%rsp), %rsi
	movl	%ebx, %r8d
	orl	$-1, %ecx
	movq	72(%rsp), %rdi
	call	_Z37__device_stub__Z11global_initPlS_PiiiPlS_Piii
.L43:
	movabsq	$4294967297, %r15
	movslq	%ebx, %r14
	xorl	%r13d, %r13d
	leaq	0(,%r14,8), %rax
	movq	%rax, 24(%rsp)
.L44:
	cmpl	%ebx, %r13d
	jge	.L55
	movq	%r15, 112(%rsp)
	movl	116(%rsp), %ecx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movq	%r15, 100(%rsp)
	movl	104(%rsp), %esi
	movl	%r12d, 108(%rsp)
	movq	108(%rsp), %rdx
	movl	%ebp, 96(%rsp)
	movq	96(%rsp), %rdi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L45
	movq	88(%rsp), %rdx
	movq	80(%rsp), %rsi
	movl	%ebx, %r8d
	movl	%r13d, %ecx
	movq	72(%rsp), %rdi
	call	_Z38__device_stub__Z12global_floydPlS_PiiiPlS_Piii
.L45:
	movq	(%rsp), %rdx
	movq	80(%rsp), %rsi
	movl	$2, %ecx
	movq	8(%rsp), %rdi
	call	cudaMemcpy@PLT
	movq	8(%rsp), %r10
	xorl	%r9d, %r9d
.L46:
	xorl	%r8d, %r8d
.L48:
	leal	-1(%rbx), %eax
	movq	(%r10,%r8,8), %rdx
	leaq	.LC0(%rip), %rcx
	movl	$2, %edi
	cmpl	%r8d, %eax
	leaq	.LC1(%rip), %rax
	movq	%r9, 56(%rsp)
	leaq	.LC3(%rip), %rsi
	cmovne	%rax, %rcx
	xorl	%eax, %eax
	movq	%r10, 48(%rsp)
	movq	%r8, 40(%rsp)
	call	__printf_chk@PLT
	movq	40(%rsp), %r8
	movq	48(%rsp), %r10
	movq	56(%rsp), %r9
	incq	%r8
	cmpq	%r14, %r8
	jne	.L48
	movq	24(%rsp), %rax
	incq	%r9
	addq	%rax, %r10
	cmpq	%r14, %r9
	jne	.L46
	leaq	.LC0(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	incl	%r13d
	call	__printf_chk@PLT
	movq	72(%rsp), %rax
	movq	80(%rsp), %rdx
	movq	%rax, 80(%rsp)
	movq	%rdx, 72(%rsp)
	jmp	.L44
.L55:
	movq	(%rsp), %rdx
	movq	72(%rsp), %rsi
	movl	$2, %ecx
	movq	8(%rsp), %rdi
	call	cudaMemcpy@PLT
	movq	16(%rsp), %rdx
	movq	88(%rsp), %rsi
	movl	$2, %ecx
	movq	32(%rsp), %rdi
	call	cudaMemcpy@PLT
	movq	72(%rsp), %rdi
	call	cudaFree@PLT
	movq	80(%rsp), %rdi
	call	cudaFree@PLT
	movq	88(%rsp), %rdi
	call	cudaFree@PLT
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L51
	call	__stack_chk_fail@PLT
.L51:
	addq	$136, %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
.LFE3637:
	.size	_Z10host_floydiiPlPii, .-_Z10host_floydiiPlPii
	.section	.rodata.str1.1
.LC4:
	.string	"\t"
.LC5:
	.string	"Title :\t%s\n"
.LC6:
	.string	"Usage :\t%s [-g <gridSize>] [-b <blockSize>] <inputfilename> <outputfilename> <intermedianfilename>\n"
.LC7:
	.string	"\tinputfilename - source matrix of path prices or empty\n"
.LC8:
	.string	"\toutputfilename - output floyd's matrix of path prices\n"
.LC9:
	.string	"\tintermedianfilename - output matrix of intermedian points or empty\n"
.LC10:
	.string	"Input File Name :\t%s\n"
.LC11:
	.string	"Output File Name :\t%s\n"
.LC12:
	.string	"Intermedian File Name :\t%s\n"
.LC13:
	.string	"r"
.LC14:
	.string	"File open error (%s)\n"
.LC15:
	.string	"Matrix rank :\t%d\n"
.LC16:
	.string	"Running on GPU %d (%s)\n"
.LC17:
	.string	"Device has ECC support enabled %d\n"
.LC18:
	.string	"Number of asynchronous engines %d\n"
.LC19:
	.string	"Device can map host memory with cudaHostAlloc/cudaHostGetDevicePointer %d\n"
.LC20:
	.string	"Clock frequency in kilohertz %d\n"
.LC21:
	.string	"Compute mode (See cudaComputeMode) %d\n"
.LC22:
	.string	"Device can possibly execute multiple kernels concurrently %d\n"
.LC23:
	.string	"Device can concurrently copy memory and execute a kernel. Deprecated. Use instead asyncEngineCount. %d\n"
.LC24:
	.string	"Device is integrated as opposed to discrete %d\n"
.LC25:
	.string	"Specified whether there is a run time limit on kernels %d\n"
.LC26:
	.string	"Size of L2 cache in bytes %d\n"
.LC27:
	.string	"Major compute capability %d\n"
.LC28:
	.string	"Maximum size of each dimension of a grid %d\n"
.LC29:
	.string	"Maximum 1D surface size %d\n"
.LC30:
	.string	"Maximum 1D layered surface dimensions %d\n"
.LC31:
	.string	"Maximum 2D surface dimensions %d\n"
.LC32:
	.string	"Maximum 2D layered surface dimensions %d\n"
.LC33:
	.string	"Maximum 3D surface dimensions %d\n"
.LC34:
	.string	"Maximum Cubemap surface dimensions %d\n"
.LC35:
	.string	"Maximum Cubemap layered surface dimensions %d\n"
.LC36:
	.string	"Maximum 1D texture size %d\n"
.LC37:
	.string	"Maximum 1D layered texture dimensions %d\n"
.LC38:
	.string	"Maximum size for 1D textures bound to linear memory %d\n"
.LC39:
	.string	"Maximum 1D mipmapped texture size %d\n"
.LC40:
	.string	"Maximum 2D texture dimensions %d\n"
.LC41:
	.string	"Maximum 2D texture dimensions if texture gather operations have to be performed %d\n"
.LC42:
	.string	"Maximum 2D layered texture dimensions %d\n"
.LC43:
	.string	"Maximum dimensions (width, height, pitch) for 2D textures bound to pitched memory %d\n"
.LC44:
	.string	"Maximum 2D mipmapped texture dimensions %d\n"
.LC45:
	.string	"Maximum 3D texture dimensions %d\n"
.LC46:
	.string	"Maximum alternate 3D texture dimensions %d\n"
.LC47:
	.string	"Maximum Cubemap texture dimensions %d\n"
.LC48:
	.string	"Maximum Cubemap layered texture dimensions %d\n"
.LC49:
	.string	"Maximum size of each dimension of a block %d\n"
.LC50:
	.string	"Maximum number of threads per block %d\n"
.LC51:
	.string	"Maximum resident threads per multiprocessor %d\n"
.LC52:
	.string	"Maximum pitch in bytes allowed by memory copies %d\n"
.LC53:
	.string	"Global memory bus width in bits %d\n"
.LC54:
	.string	"Peak memory clock frequency in kilohertz %d\n"
.LC55:
	.string	"Minor compute capability %d\n"
.LC56:
	.string	"Number of multiprocessors on device %d\n"
.LC57:
	.string	"PCI bus ID of the device %d\n"
.LC58:
	.string	"PCI device ID of the device %d\n"
.LC59:
	.string	"PCI domain ID of the device %d\n"
.LC60:
	.string	"32-bit registers available per block %d\n"
.LC61:
	.string	"Shared memory available per block in bytes %d\n"
.LC62:
	.string	"Device supports stream priorities %d\n"
.LC63:
	.string	"Alignment requirements for surfaces %d\n"
.LC64:
	.string	"1 if device is a Tesla device using TCC driver, 0 otherwise %d\n"
.LC65:
	.string	"Alignment requirement for textures %d\n"
.LC66:
	.string	"Pitch alignment requirement for texture references bound to pitched memory %d\n"
.LC67:
	.string	"Constant memory available on device in bytes %d\n"
.LC68:
	.string	"Global memory available on device in bytes %d\n"
.LC69:
	.string	"Device shares a unified address space with the host %d\n"
.LC70:
	.string	"Warp size in threads %d\n"
.LC71:
	.string	"w"
.LC72:
	.string	"%s"
.LC73:
	.string	"%d%s"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB3638:
	.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
	subq	$4096, %rsp
	.cfi_def_cfa_offset 4152
	orq	$0, (%rsp)
	subq	$1128, %rsp
	.cfi_def_cfa_offset 5280
	movq	title(%rip), %rdx
	movq	%fs:40, %rax
	movq	%rax, 5208(%rsp)
	xorl	%eax, %eax
	movl	%edi, %ebp
	movq	%rsi, %rbx
	movl	$2, %edi
	leaq	.LC5(%rip), %rsi
	call	__printf_chk@PLT
	movq	stdout(%rip), %rdi
	call	fflush@PLT
	cmpl	$3, %ebp
	jg	.L96
	movq	(%rbx), %rdx
	leaq	.LC6(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	stdout(%rip), %rdi
	call	fflush@PLT
	leaq	.LC7(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	stdout(%rip), %rdi
	call	fflush@PLT
	leaq	.LC8(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	stdout(%rip), %rdi
	call	fflush@PLT
	movl	$2, %edi
	leaq	.LC9(%rip), %rsi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	stdout(%rip), %rdi
	jmp	.L115
.L96:
	xorl	%edx, %edx
	movl	$1, %eax
	movl	%edx, 44(%rsp)
	movl	%edx, 40(%rsp)
.L57:
	movslq	%eax, %rcx
	leal	1(%rax), %r12d
	movq	(%rbx,%rcx,8), %rdx
	movq	8(%rbx,%rcx,8), %rdi
	cmpb	$45, (%rdx)
	jne	.L60
	movb	1(%rdx), %dl
	cmpb	$98, %dl
	je	.L58
	cmpb	$103, %dl
	jne	.L59
	call	atoi@PLT
	movl	%eax, 40(%rsp)
	jmp	.L114
.L58:
	call	atoi@PLT
	movl	%eax, 44(%rsp)
.L114:
	movl	%r12d, %eax
.L59:
	incl	%eax
	cmpl	%eax, %ebp
	jg	.L57
.L60:
	cltq
	leaq	.LC10(%rip), %rsi
	movl	$2, %edi
	xorl	%r12d, %r12d
	leaq	0(,%rax,8), %rdx
	movq	(%rbx,%rax,8), %r13
	movq	8(%rbx,%rdx), %rax
	movq	%rax, 16(%rsp)
	movq	16(%rbx,%rdx), %rax
	movq	%r13, %rdx
	xorl	%ebx, %ebx
	movq	%rax, 24(%rsp)
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	stdout(%rip), %rdi
	call	fflush@PLT
	movq	16(%rsp), %rdx
	leaq	.LC11(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movq	stdout(%rip), %rdi
	call	fflush@PLT
	movq	24(%rsp), %rdx
	leaq	.LC12(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movq	stdout(%rip), %rdi
	call	fflush@PLT
	leaq	.LC13(%rip), %rsi
	movq	%r13, %rdi
	call	fopen@PLT
	movq	%rax, 8(%rsp)
	leaq	1112(%rsp), %rax
	cmpq	$0, 8(%rsp)
	movq	%rax, 48(%rsp)
	jne	.L62
	movq	%r13, %rcx
.L117:
	movq	stderr(%rip), %rdi
	leaq	.LC14(%rip), %rdx
	movl	$2, %esi
	xorl	%eax, %eax
.L118:
	call	__fprintf_chk@PLT
	movq	stderr(%rip), %rdi
.L115:
	call	fflush@PLT
	orl	$-1, %edi
.L116:
	call	exit@PLT
.L120:
	cmpl	%r14d, %ebx
	cmovl	%r14d, %ebx
	incl	%r12d
.L62:
	movq	8(%rsp), %rcx
	movq	48(%rsp), %rdi
	movl	$4096, %esi
	movl	$4096, %edx
	call	__fgets_chk@PLT
	movq	%rax, %rsi
	testq	%rax, %rax
	je	.L119
	leaq	72(%rsp), %r13
	movl	$59, %edx
	xorl	%r14d, %r14d
	movq	%r13, %rdi
	call	_Z8mystrtokPPcS_c
.L63:
	testq	%rax, %rax
	je	.L120
	movl	$59, %edx
	xorl	%esi, %esi
	movq	%r13, %rdi
	incl	%r14d
	call	_Z8mystrtokPPcS_c
	jmp	.L63
.L119:
	movabsq	$9223372036854775807, %rbp
	cmpl	%r12d, %ebx
	movl	%r12d, %r14d
	cmovge	%ebx, %r14d
	xorl	%r13d, %r13d
	xorl	%r12d, %r12d
	movl	%r14d, %ebx
	imull	%r14d, %ebx
	movslq	%ebx, %rbx
	leaq	0(,%rbx,8), %rdi
	call	malloc@PLT
	leaq	0(,%rbx,4), %rdi
	movq	%rax, (%rsp)
	call	malloc@PLT
	movq	8(%rsp), %rdi
	xorl	%edx, %edx
	xorl	%esi, %esi
	movq	%rax, 32(%rsp)
	call	fseek@PLT
.L66:
	movq	8(%rsp), %rcx
	movq	48(%rsp), %rdi
	movl	$4096, %esi
	movslq	%r13d, %rbx
	movl	$4096, %edx
	call	__fgets_chk@PLT
	testq	%rax, %rax
	movq	%rax, %rsi
	je	.L121
	leaq	72(%rsp), %rdi
	movl	$59, %edx
	call	_Z8mystrtokPPcS_c
	movq	%rax, %rdi
	movq	(%rsp), %rax
	leaq	(%rax,%rbx,8), %r15
	xorl	%ebx, %ebx
.L67:
	testq	%rdi, %rdi
	je	.L122
	movq	%rdi, 56(%rsp)
	call	_Z8stremptyPKc
	movl	%eax, %edx
	movq	%rbp, %rax
	testl	%edx, %edx
	jne	.L68
	movq	56(%rsp), %rdi
	call	atol@PLT
.L68:
	movq	%rax, (%r15,%rbx,8)
	leaq	72(%rsp), %rdi
	movl	$59, %edx
	xorl	%esi, %esi
	incq	%rbx
	call	_Z8mystrtokPPcS_c
	movq	%rax, %rdi
	jmp	.L67
.L122:
	movslq	%ebx, %rbx
.L70:
	cmpl	%ebx, %r14d
	jle	.L123
	movq	%rbp, (%r15,%rbx,8)
	incq	%rbx
	jmp	.L70
.L123:
	incl	%r12d
	addl	%r14d, %r13d
	jmp	.L66
.L121:
	movl	%r14d, %edx
	movq	(%rsp), %rax
	movabsq	$9223372036854775807, %rsi
	subl	%r12d, %edx
	imull	%r14d, %edx
	leaq	(%rax,%rbx,8), %rcx
	xorl	%eax, %eax
.L73:
	cmpl	%eax, %edx
	jle	.L124
	movq	%rsi, (%rcx,%rax,8)
	incq	%rax
	jmp	.L73
.L124:
	movq	(%rsp), %r12
	movslq	%r14d, %r15
	xorl	%eax, %eax
	movabsq	$9223372036854775807, %rcx
	leaq	8(,%r15,8), %rbx
	movq	%r12, %rdx
.L75:
	cmpl	%r14d, %eax
	je	.L125
	movq	%rcx, (%rdx)
	incl	%eax
	addq	%rbx, %rdx
	jmp	.L75
.L125:
	movq	8(%rsp), %rdi
	xorl	%r13d, %r13d
	call	fclose@PLT
	movl	%r14d, %edx
	movl	$2, %edi
	xorl	%eax, %eax
	leaq	.LC15(%rip), %rsi
	call	__printf_chk@PLT
	leaq	-8(%rbx), %rax
	movq	(%rsp), %rbx
	movq	%rax, 48(%rsp)
.L77:
	cmpl	%r14d, %r13d
	je	.L126
	xorl	%ebp, %ebp
.L79:
	leal	-1(%r14), %eax
	movq	(%rbx,%rbp,8), %rdx
	leaq	.LC0(%rip), %rcx
	movl	$2, %edi
	cmpl	%ebp, %eax
	leaq	.LC4(%rip), %rax
	leaq	.LC3(%rip), %rsi
	cmovne	%rax, %rcx
	xorl	%eax, %eax
	incq	%rbp
	call	__printf_chk@PLT
	cmpl	%ebp, %r14d
	jg	.L79
	movq	48(%rsp), %rax
	incl	%r13d
	addq	%rax, %rbx
	jmp	.L77
.L126:
	movq	stdout(%rip), %rdi
	xorl	%ebx, %ebx
	leaq	.LC16(%rip), %r13
	call	fflush@PLT
	xorl	%eax, %eax
	leaq	68(%rsp), %rdi
	movl	%eax, 68(%rsp)
	call	cudaGetDeviceCount@PLT
.L80:
	cmpl	%ebx, 68(%rsp)
	jle	.L127
	leaq	80(%rsp), %rbp
	movl	%ebx, %esi
	movq	%rbp, %rdi
	call	cudaGetDeviceProperties_v2@PLT
	movq	%rbp, %rcx
	movl	%ebx, %edx
	movq	%r13, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	leaq	.LC28(%rip), %rbp
	call	__printf_chk@PLT
	movl	660(%rsp), %edx
	leaq	.LC17(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movl	680(%rsp), %edx
	leaq	.LC18(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movl	480(%rsp), %edx
	leaq	.LC19(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movl	428(%rsp), %edx
	leaq	.LC20(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movl	484(%rsp), %edx
	leaq	.LC21(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movl	656(%rsp), %edx
	leaq	.LC22(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movl	464(%rsp), %edx
	leaq	.LC23(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movl	476(%rsp), %edx
	leaq	.LC24(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movl	472(%rsp), %edx
	leaq	.LC25(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movl	696(%rsp), %edx
	leaq	.LC26(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movl	440(%rsp), %edx
	leaq	.LC27(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movq	%rbp, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movl	416(%rsp), %edx
	call	__printf_chk@PLT
	movq	%rbp, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movl	420(%rsp), %edx
	call	__printf_chk@PLT
	movq	%rbp, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movl	424(%rsp), %edx
	leaq	.LC30(%rip), %rbp
	call	__printf_chk@PLT
	movl	592(%rsp), %edx
	leaq	.LC29(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movq	%rbp, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movl	616(%rsp), %edx
	call	__printf_chk@PLT
	movq	%rbp, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movl	620(%rsp), %edx
	leaq	.LC31(%rip), %rbp
	call	__printf_chk@PLT
	movq	%rbp, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movl	596(%rsp), %edx
	call	__printf_chk@PLT
	movq	%rbp, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movl	600(%rsp), %edx
	leaq	.LC32(%rip), %rbp
	call	__printf_chk@PLT
	movq	%rbp, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movl	624(%rsp), %edx
	call	__printf_chk@PLT
	movq	%rbp, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movl	628(%rsp), %edx
	call	__printf_chk@PLT
	movq	%rbp, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movl	632(%rsp), %edx
	leaq	.LC33(%rip), %rbp
	call	__printf_chk@PLT
	movq	%rbp, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movl	604(%rsp), %edx
	call	__printf_chk@PLT
	movq	%rbp, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movl	608(%rsp), %edx
	call	__printf_chk@PLT
	movq	%rbp, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movl	612(%rsp), %edx
	leaq	.LC35(%rip), %rbp
	call	__printf_chk@PLT
	movl	636(%rsp), %edx
	leaq	.LC34(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movq	%rbp, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movl	640(%rsp), %edx
	call	__printf_chk@PLT
	movq	%rbp, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movl	644(%rsp), %edx
	leaq	.LC37(%rip), %rbp
	call	__printf_chk@PLT
	movl	488(%rsp), %edx
	leaq	.LC36(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movq	%rbp, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movl	564(%rsp), %edx
	call	__printf_chk@PLT
	movq	%rbp, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movl	568(%rsp), %edx
	leaq	.LC40(%rip), %rbp
	call	__printf_chk@PLT
	movl	496(%rsp), %edx
	leaq	.LC38(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movl	492(%rsp), %edx
	leaq	.LC39(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movq	%rbp, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movl	500(%rsp), %edx
	call	__printf_chk@PLT
	movq	%rbp, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movl	504(%rsp), %edx
	leaq	.LC41(%rip), %rbp
	call	__printf_chk@PLT
	movq	%rbp, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movl	528(%rsp), %edx
	call	__printf_chk@PLT
	movq	%rbp, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movl	532(%rsp), %edx
	leaq	.LC42(%rip), %rbp
	call	__printf_chk@PLT
	movq	%rbp, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movl	572(%rsp), %edx
	call	__printf_chk@PLT
	movq	%rbp, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movl	576(%rsp), %edx
	call	__printf_chk@PLT
	movq	%rbp, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movl	580(%rsp), %edx
	leaq	.LC43(%rip), %rbp
	call	__printf_chk@PLT
	movq	%rbp, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movl	516(%rsp), %edx
	call	__printf_chk@PLT
	movq	%rbp, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movl	520(%rsp), %edx
	call	__printf_chk@PLT
	movq	%rbp, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movl	524(%rsp), %edx
	leaq	.LC44(%rip), %rbp
	call	__printf_chk@PLT
	movq	%rbp, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movl	508(%rsp), %edx
	call	__printf_chk@PLT
	movq	%rbp, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movl	512(%rsp), %edx
	leaq	.LC45(%rip), %rbp
	call	__printf_chk@PLT
	movq	%rbp, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movl	536(%rsp), %edx
	call	__printf_chk@PLT
	movq	%rbp, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movl	540(%rsp), %edx
	call	__printf_chk@PLT
	movq	%rbp, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movl	544(%rsp), %edx
	leaq	.LC46(%rip), %rbp
	call	__printf_chk@PLT
	movq	%rbp, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movl	548(%rsp), %edx
	call	__printf_chk@PLT
	movq	%rbp, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movl	552(%rsp), %edx
	call	__printf_chk@PLT
	movq	%rbp, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movl	556(%rsp), %edx
	leaq	.LC48(%rip), %rbp
	call	__printf_chk@PLT
	movl	560(%rsp), %edx
	leaq	.LC47(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movq	%rbp, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movl	584(%rsp), %edx
	call	__printf_chk@PLT
	movq	%rbp, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movl	588(%rsp), %edx
	leaq	.LC49(%rip), %rbp
	call	__printf_chk@PLT
	movq	%rbp, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movl	404(%rsp), %edx
	call	__printf_chk@PLT
	movq	%rbp, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movl	408(%rsp), %edx
	call	__printf_chk@PLT
	movq	%rbp, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movl	412(%rsp), %edx
	call	__printf_chk@PLT
	movl	400(%rsp), %edx
	leaq	.LC50(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movl	704(%rsp), %edx
	leaq	.LC51(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movq	392(%rsp), %rdx
	leaq	.LC52(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movl	692(%rsp), %edx
	leaq	.LC53(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movl	688(%rsp), %edx
	leaq	.LC54(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movl	444(%rsp), %edx
	leaq	.LC55(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movl	468(%rsp), %edx
	leaq	.LC56(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movl	664(%rsp), %edx
	leaq	.LC57(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movl	668(%rsp), %edx
	leaq	.LC58(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movl	672(%rsp), %edx
	leaq	.LC59(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movl	384(%rsp), %edx
	leaq	.LC60(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movq	376(%rsp), %rdx
	leaq	.LC61(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movl	708(%rsp), %edx
	leaq	.LC62(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movq	648(%rsp), %rdx
	leaq	.LC63(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movl	676(%rsp), %edx
	leaq	.LC64(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movq	448(%rsp), %rdx
	leaq	.LC65(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movq	456(%rsp), %rdx
	leaq	.LC66(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movq	432(%rsp), %rdx
	leaq	.LC67(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movq	368(%rsp), %rdx
	leaq	.LC68(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movl	684(%rsp), %edx
	leaq	.LC69(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movl	388(%rsp), %edx
	leaq	.LC70(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movq	stdout(%rip), %rdi
	call	fflush@PLT
	incl	%ebx
	jmp	.L80
.L127:
	movq	32(%rsp), %rcx
	movq	(%rsp), %rdx
	movl	%r14d, %r8d
	xorl	%ebp, %ebp
	movl	44(%rsp), %esi
	movl	40(%rsp), %edi
	movabsq	$9223372036854775807, %r13
	call	_Z10host_floydiiPlPii
	call	cudaDeviceReset@PLT
	movq	16(%rsp), %rdi
	leaq	.LC71(%rip), %rsi
	call	fopen@PLT
	movq	%rax, 8(%rsp)
	testq	%rax, %rax
	jne	.L82
	movq	16(%rsp), %rcx
	movq	stderr(%rip), %rdi
	movl	$2, %esi
	leaq	.LC14(%rip), %rdx
	jmp	.L118
.L87:
	movq	(%r12,%rbx,8), %rcx
	leal	-1(%r14), %eax
	cmpq	%r13, %rcx
	je	.L83
	cmpl	%ebx, %eax
	leaq	.LC1(%rip), %r8
	leaq	.LC0(%rip), %rax
	movq	8(%rsp), %rdi
	cmove	%rax, %r8
	leaq	.LC3(%rip), %rdx
	movl	$2, %esi
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
.L85:
	incq	%rbx
	cmpl	%ebx, %r14d
	jg	.L87
	movq	48(%rsp), %rax
	incl	%ebp
	addq	%rax, %r12
.L82:
	cmpl	%r14d, %ebp
	je	.L128
	xorl	%ebx, %ebx
	jmp	.L87
.L83:
	cmpl	%ebx, %eax
	leaq	.LC1(%rip), %rcx
	leaq	.LC0(%rip), %rax
	movq	8(%rsp), %rdi
	cmove	%rax, %rcx
	leaq	.LC72(%rip), %rdx
	movl	$2, %esi
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
	jmp	.L85
.L128:
	movq	8(%rsp), %rdi
	call	fclose@PLT
	movq	24(%rsp), %rdi
	leaq	.LC71(%rip), %rsi
	call	fopen@PLT
	movq	%rax, %rbx
	testq	%rax, %rax
	je	.L88
	leaq	0(,%r15,4), %rax
	movq	32(%rsp), %r13
	xorl	%r12d, %r12d
	leaq	.LC1(%rip), %rbp
	movq	%rax, 8(%rsp)
	jmp	.L89
.L88:
	movq	24(%rsp), %rcx
	jmp	.L117
.L94:
	movl	0(%r13,%r15,4), %ecx
	leal	-1(%r14), %eax
	testl	%ecx, %ecx
	js	.L90
	cmpl	%r15d, %eax
	leaq	.LC0(%rip), %r8
	movl	$2, %esi
	movq	%rbx, %rdi
	cmovne	%rbp, %r8
	leaq	.LC73(%rip), %rdx
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
.L92:
	incq	%r15
	cmpl	%r15d, %r14d
	jg	.L94
	movq	8(%rsp), %rax
	incl	%r12d
	addq	%rax, %r13
.L89:
	cmpl	%r14d, %r12d
	je	.L129
	xorl	%r15d, %r15d
	jmp	.L94
.L90:
	cmpl	%r15d, %eax
	leaq	.LC0(%rip), %rcx
	movl	$2, %esi
	movq	%rbx, %rdi
	cmovne	%rbp, %rcx
	leaq	.LC72(%rip), %rdx
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
	jmp	.L92
.L129:
	movq	%rbx, %rdi
	call	fclose@PLT
	movq	(%rsp), %rdi
	call	free@PLT
	movq	32(%rsp), %rdi
	call	free@PLT
	xorl	%edi, %edi
	jmp	.L116
	.cfi_endproc
.LFE3638:
	.size	main, .-main
	.section	.rodata.str1.1
.LC74:
	.string	"_Z12global_floydPlS_Piii"
.LC75:
	.string	"_Z11global_initPlS_Piii"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB3668:
	.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	.LC74(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z12global_floydPlS_Piii(%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	.LC75(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z11global_initPlS_Piii(%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
.LFE3668:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.nvFatBinSegment,"aw"
	.align 8
	.type	_ZL15__fatDeviceText, @object
	.size	_ZL15__fatDeviceText, 24
_ZL15__fatDeviceText:
	.long	1180844977
	.long	1
	.quad	fatbinData
	.quad	0
	.local	_ZL20__cudaFatCubinHandle
	.comm	_ZL20__cudaFatCubinHandle,8,8
	.globl	description
	.section	.rodata.str1.1
.LC76:
	.string	"\320\220\320\273\320\263\320\276\321\200\320\270\321\202\320\274 \320\244\320\273\320\276\320\271\320\264\320\260 - \320\277\320\276\320\270\321\201\320\272 \320\262\321\201\320\265\321\205 \320\272\321\200\320\260\321\202\321\207\320\260\320\271\321\210\320\270\321\205 \320\277\321\203\321\202\320\265\320\271 \320\262 \320\263\321\200\320\260\321\204\320\265"
	.section	.data.rel.local,"aw"
	.align 8
	.type	description, @object
	.size	description, 8
description:
	.quad	.LC76
	.globl	title
	.section	.rodata.str1.1
.LC77:
	.string	"Floyd's algorithm"
	.section	.data.rel.local
	.align 8
	.type	title, @object
	.size	title, 8
title:
	.quad	.LC77
	.section	.rodata.cst8,"aM",@progbits,8
	.align 8
.LC2:
	.long	1431655759
	.long	1070945621
	.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:

﻿char *title = "Floyd's algorithm";
char *description = "Алгоритм Флойда - поиск всех кратчайших путей в графе";
/*
Алгоритм Флойда является одним из методов поиска кратчайших путей в графе. 
В отличии от алгоритма Дейкстры, который позволяет при доведении до конца построить 
ориентированное дерево кратчайших путей от некоторой вершины, метод Флойда позволяет 
найти длины всех кратчайших путей в графе. Конечно эта задача может быть решена 
и многократным применением алгоритма Дейкстры (каждый раз последовательно выбираем 
вершину от первой до N-ной, пока не получим кратчайшие пути от всех вершин графа), 
однако реализация подобной процедуры потребовала бы значительных вычислительных затрат.

Прежде чем представлять алгоритмы, необходимо ввести некоторые обозначения. 
Перенумеруем вершины исходного графа целыми числами от 1 до N. Обозначим через di,jm длину кратчайшего пути 
из вершинм i в вершину j, который в качестве промежуточных может содержать только первые m вершин графа. 
(Напомним, что промежуточной вершиной пути является любая принадлежащая ему вершина, не совпадающая 
с его начальной или конечной вершинами.) Если между вершинами i и j не существует ни одного пути указанного типа, 
то условно будем считать, что di,jm=∞. Из данного определения величин di,jm следует, что величина di,j0, 
представляет длину кратчайшего пути из вершины i в вершину j, не имеющего промежуточных вершин, 
т. е. длину кратчайшей дуги, соединяющей i с j (если такие дуги присутствуют в графе). 
для любой вершины i положим di,im= 0. Отметим далее, что величина di,jmпредставляет длину кратчайшего пути 
между вершинами i и j.

Обозначим через Dm матрицу размера NxN, элемент (i, j) которой совпадает с di,jm. 
Если в исходном графе нам известна длина каждой дуги, то мы можем сформировать матрицу D0. 
Наша цель состоит в определении матрицы DN, представляющей кратчайшие пути между всеми вершинами рассматриваемого графа.

В алгоритме Флойда в качестве исходной выступает матрица D0. 
Вначале из этой матрицы вычисляется матрица D1. 
Затем по матрице D1 вычисляется матрицав D2 и т. д. 
Процесс повторяется до тех пор, пока по матрице DN-1 не будет вычислена матрица DN.

Рассмотрим основную идею, лежащую в основе алгоритма Флойда. 
Суть алгоритма Флойда заключается в проверке того, не окажется ли путь из вершины i в вершину j короче, 
если он будет проходить через некоторую промежуточную вершину m. Предположим, что нам известны:

кратчайший путь из вершины i в вершину m, в котором в качестве промежуточных допускается использование только первых (m - 1) вершин;
кратчайший путь из вершины m в вершину j, в котором в качестве промежуточных допускается использование только первых (m - 1) вершин;
кратчайший путь из вершины i в вершину j, в котором в качестве промежуточных допускается использование только первых (m - 1) вершин.

Поскольку по предположению исходный граф не может содержать контуров отрицательной длины, 
один из двух путей — путь, совпадающий с представленным в пункте 3, или путь, являющийся объединением 
путей из пунктов 1 и 2 — должен быть кратчайшим путем из вершины i в вершину j, 
в котором в качестве промежуточных допускается использование только первых m вершин. Таким образом,

di,jm=min{ di,mm-1+ dm,jm-1; di,jm-1}

Из соотношения видно, что для вычисления элементов матрицы Dm необходимо располагать лишь элементами матрицы Dm-1. 
Более того, соответствующие вычисления могут быть проведены без обращения к исходному графу. 
Теперь мм в состоянии дать формальное описание алгоритма Флойда для нахождения на графе кратчайших путей 
между всеми парами вершин. 

Алгоритм

Перенумеровать вершины графа от 1 до N целыми числами, определить матрицу D0, каждый элемент di,j  
которой есть длина кратчайшей дуги между вершинами i и j. Если такой дуги нет, положить значение элемента 
равным ∞. Кроме того, положить значения диагонального элемента di,iравным 0.
Для целого m, последовательно принимающего значения 1...N определить по элементам матрицы Dm-1 элементы Dm
Алгоритм заканчивается получением матрицы всех кратчайших путей DN, N – число вершин графа.
 
Напомним, для определения по известным элементам матрицы Dm-1 элементов матрицы  Dm 
в алгоритме Флойда применяется рекурсивное соотношение:

di,jm=min{ di,mm-1+ dm,jm-1; di,jm-1}

di,jm – элемент матрицы Dm, di,jm-1 – элементы матрицы Dm-1 найденой на предыдущем шаге алгоритма.
*/
#include <iostream>
#include <cstdio>
#include <cstdlib>
#include <stdio.h>
#include <string.h>
#include <time.h>
#include <ctype.h>
#include <limits.h>
#include <hip/hip_runtime.h>



#define assert( bool ) 
int strempty(const char *p)
{
	if (!p)
		return (1);
	for (; *p; p++)
		if (!isspace(*p))
			return (0);
	return (1);
}
char *mystrtok(char **m, char *s, char c)
{
	char *p = s ? s : *m;
	if (!*p)
		return 0;
	*m = strchr(p, c);
	if (*m)
		*(*m)++ = 0;
	else
		*m = p + strlen(p);
	return p;
}
#ifndef max
#define max( a, b ) ( ((a) > (b)) ? (a) : (b) )
#endif

#ifndef min
#define min( a, b ) ( ((a) < (b)) ? (a) : (b) )
#endif

__global__ void global_init(long *prev, long *next, int *intermedian, int m, int n){
	for (int id = blockDim.x*blockIdx.x + threadIdx.x; id < n*n; id += blockDim.x*gridDim.x) {
		intermedian[id] = -1;
	}
}

__global__ void global_floyd(long *prev, long *next, int *intermedian, int m, int n){
	for (int id = blockDim.x*blockIdx.x + threadIdx.x; id < n*n; id += blockDim.x*gridDim.x) {
		int i = id/n;
		int j = id%n;
		if(i != j && prev[i*n+m] != LONG_MAX && prev[m*n+j] != LONG_MAX && (prev[i*n+m] + prev[m*n+j]) < prev[i*n+j]){
			next[id] = prev[i*n+m] + prev[m*n+j];
			intermedian[id] = m;
		}
		else {
			next[id] = prev[id];
		}
	}
}

__host__ void host_floyd(int gridSize, int blockSize, long *matrix, int *intermedian, int n)
{
	hipError_t err;
	long *device_prev;
	long *device_next;
	int *device_intermedian;

	err = hipMalloc((void**)&device_prev, n*n * sizeof(long));
	err = hipMalloc((void**)&device_next, n*n * sizeof(long));
	err = hipMalloc((void**)&device_intermedian, n*n * sizeof(int));

	err = hipMemcpy(device_prev, matrix, n*n*sizeof(long), hipMemcpyHostToDevice);

	int blocks = (gridSize > 0)? gridSize : min(max(1, (int)pow((double)n*n, 0.333333333333333)), 15);
	int threads = (blockSize > 0)? blockSize : min(max(1, (int)pow((double)n*n, 0.333333333333333)), 15);
	global_init <<< blocks, threads >>>(device_prev, device_next, device_intermedian, -1, n);

	for(int m = 0; m < n ; m++){
		global_floyd <<< blocks, threads >>>(device_prev, device_next, device_intermedian, m, n);
		err = hipMemcpy(matrix, device_next, n*n*sizeof(long), hipMemcpyDeviceToHost);
		for (int i = 0; i < n; i++){
			for (int j = 0; j < n; j++){
				printf("%ld%s", matrix[i*n + j], ((j == n - 1) ? "\n" : ";"));
			}
		}
		printf("\n");
		long * t = device_prev; device_prev = device_next; device_next = t;
	}

	err = hipMemcpy(matrix, device_prev, n*n*sizeof(long), hipMemcpyDeviceToHost);
	err = hipMemcpy(intermedian, device_intermedian, n*n*sizeof(int), hipMemcpyDeviceToHost);

	err = hipFree(device_prev);
	err = hipFree(device_next);
	err = hipFree(device_intermedian);

	err = err;
}

int main(int argc, char* argv[])
{
	int gridSize = 0;
	int blockSize = 0;

	printf("Title :\t%s\n", title); fflush(stdout);

	if (argc < 4) {
		printf("Usage :\t%s [-g <gridSize>] [-b <blockSize>] <inputfilename> <outputfilename> <intermedianfilename>\n", argv[0]); fflush(stdout);
		printf("\tinputfilename - source matrix of path prices or empty\n"); fflush(stdout);
		printf("\toutputfilename - output floyd's matrix of path prices\n"); fflush(stdout);
		printf("\tintermedianfilename - output matrix of intermedian points or empty\n"); fflush(stdout);
		exit(-1);
	}

	int argId = 1;
	for(; argId < argc && argv[argId][0]=='-' ; argId++){
		switch(argv[argId][1]){
		case 'g':
			gridSize = atoi(argv[++argId]);
			break;
		case 'b':
			blockSize = atoi(argv[++argId]);
			break;
		}
	}

	char *inputFileName = argv[argId++];
	char *outputFileName = argv[argId++];
	char *intermedianFileName = argv[argId++];

	printf("Input File Name :\t%s\n", inputFileName); fflush(stdout);
	printf("Output File Name :\t%s\n", outputFileName); fflush(stdout);
	printf("Intermedian File Name :\t%s\n", intermedianFileName); fflush(stdout);

	char buffer[4096];
	char *tok;
	char *p;
	int n;         /* Ранг текущего массива */
	long *matrix;  /* Массив цен */
	int *intermedian;  /* Массивов промежуточных точек */
	int i, j;

	FILE *fs = fopen(inputFileName, "r");
	if (fs == NULL) {
		fprintf(stderr, "File open error (%s)\n", inputFileName); fflush(stderr);
		exit(-1);
	}

	n = 0;

	/* Заполняем массив числами из файла */
	/* Операция выполняетя только на хост процессе */
	/* Операция выполняетя в два прохода по файлу */
	/* На первом проходе определяется ранг матрицы */
	/* На втором проходе считываются данные */
	for (i = 0; (tok = fgets(buffer, sizeof(buffer), fs)) != NULL; i++)
	{
		j = 0;
		for (tok = mystrtok(&p, tok, ';'); tok != NULL; tok = mystrtok(&p, NULL, ';'))
		{
			j++;
		}
		n = max(n, j);
	}
	n = max(n, i);

	matrix = (long *)malloc(n*n*sizeof(long));
	intermedian = (int *)malloc(n*n*sizeof(int));

	fseek(fs, 0, SEEK_SET);

	for (i = 0; (tok = fgets(buffer, sizeof(buffer), fs)) != NULL; i++)
	{
		j = 0;
		for (tok = mystrtok(&p, tok, ';'); tok != NULL; tok = mystrtok(&p, NULL, ';'))
		{
			/* Пустые элементы - это запрещённые пути */
			matrix[n*i + j++] = strempty(tok) ? LONG_MAX : atol(tok);
		}
		for (; j < n; j++) matrix[n*i + j] = LONG_MAX;
	}
	for (j = 0; j < (n - i)*n; j++) matrix[n*i + j] = LONG_MAX;
	for (i = 0; i < n; i++) matrix[n*i + i] = LONG_MAX; /* Запрещаем петли */

	fclose(fs);

	printf("Matrix rank :\t%d\n", n);
	for (i = 0; i < n; i++){
		for (j = 0; j < n; j++){
			printf("%ld%s", matrix[i*n + j], ((j == n - 1) ? "\n" : "\t"));
		}
	}
	fflush(stdout);

	// Find/set the device.
	int device_size = 0;
	hipGetDeviceCount(&device_size);
	for (i = 0; i < device_size; ++i)
	{
		hipDeviceProp_t hipDeviceProp_t;
		hipGetDeviceProperties(&hipDeviceProp_t, i);
		printf("Running on GPU %d (%s)\n", i, hipDeviceProp_t.name); 

		printf("Device has ECC support enabled %d\n",hipDeviceProp_t.ECCEnabled);
		printf("Number of asynchronous engines %d\n",hipDeviceProp_t.asyncEngineCount);
		printf("Device can map host memory with cudaHostAlloc/cudaHostGetDevicePointer %d\n",hipDeviceProp_t.canMapHostMemory);
		printf("Clock frequency in kilohertz %d\n",hipDeviceProp_t.clockRate);
		printf("Compute mode (See cudaComputeMode) %d\n",hipDeviceProp_t.computeMode);
		printf("Device can possibly execute multiple kernels concurrently %d\n",hipDeviceProp_t.concurrentKernels);
		printf("Device can concurrently copy memory and execute a kernel. Deprecated. Use instead asyncEngineCount. %d\n",hipDeviceProp_t.deviceOverlap);
		printf("Device is integrated as opposed to discrete %d\n",hipDeviceProp_t.integrated);
		printf("Specified whether there is a run time limit on kernels %d\n",hipDeviceProp_t.kernelExecTimeoutEnabled);
		printf("Size of L2 cache in bytes %d\n",hipDeviceProp_t.l2CacheSize);
		printf("Major compute capability %d\n",hipDeviceProp_t.major);
		printf("Maximum size of each dimension of a grid %d\n",hipDeviceProp_t.maxGridSize[0]);
		printf("Maximum size of each dimension of a grid %d\n",hipDeviceProp_t.maxGridSize[1]);
		printf("Maximum size of each dimension of a grid %d\n",hipDeviceProp_t.maxGridSize[2]);
		printf("Maximum 1D surface size %d\n",hipDeviceProp_t.maxSurface1D);
		printf("Maximum 1D layered surface dimensions %d\n",hipDeviceProp_t.maxSurface1DLayered[0]);
		printf("Maximum 1D layered surface dimensions %d\n",hipDeviceProp_t.maxSurface1DLayered[1]);
		printf("Maximum 2D surface dimensions %d\n",hipDeviceProp_t.maxSurface2D[0]);
		printf("Maximum 2D surface dimensions %d\n",hipDeviceProp_t.maxSurface2D[1]);
		printf("Maximum 2D layered surface dimensions %d\n",hipDeviceProp_t.maxSurface2DLayered[0]);
		printf("Maximum 2D layered surface dimensions %d\n",hipDeviceProp_t.maxSurface2DLayered[1]);
		printf("Maximum 2D layered surface dimensions %d\n",hipDeviceProp_t.maxSurface2DLayered[2]);
		printf("Maximum 3D surface dimensions %d\n",hipDeviceProp_t.maxSurface3D[0]);
		printf("Maximum 3D surface dimensions %d\n",hipDeviceProp_t.maxSurface3D[1]);
		printf("Maximum 3D surface dimensions %d\n",hipDeviceProp_t.maxSurface3D[2]);
		printf("Maximum Cubemap surface dimensions %d\n",hipDeviceProp_t.maxSurfaceCubemap);
		printf("Maximum Cubemap layered surface dimensions %d\n",hipDeviceProp_t.maxSurfaceCubemapLayered[0]);
		printf("Maximum Cubemap layered surface dimensions %d\n",hipDeviceProp_t.maxSurfaceCubemapLayered[1]);
		printf("Maximum 1D texture size %d\n",hipDeviceProp_t.maxTexture1D);
		printf("Maximum 1D layered texture dimensions %d\n",hipDeviceProp_t.maxTexture1DLayered[0]);
		printf("Maximum 1D layered texture dimensions %d\n",hipDeviceProp_t.maxTexture1DLayered[1]);
		printf("Maximum size for 1D textures bound to linear memory %d\n",hipDeviceProp_t.maxTexture1DLinear);
		printf("Maximum 1D mipmapped texture size %d\n",hipDeviceProp_t.maxTexture1DMipmap);
		printf("Maximum 2D texture dimensions %d\n",hipDeviceProp_t.maxTexture2D[0]);
		printf("Maximum 2D texture dimensions %d\n",hipDeviceProp_t.maxTexture2D[1]);
		printf("Maximum 2D texture dimensions if texture gather operations have to be performed %d\n",hipDeviceProp_t.maxTexture2DGather[0]);
		printf("Maximum 2D texture dimensions if texture gather operations have to be performed %d\n",hipDeviceProp_t.maxTexture2DGather[1]);
		printf("Maximum 2D layered texture dimensions %d\n",hipDeviceProp_t.maxTexture2DLayered[0]);
		printf("Maximum 2D layered texture dimensions %d\n",hipDeviceProp_t.maxTexture2DLayered[1]);
		printf("Maximum 2D layered texture dimensions %d\n",hipDeviceProp_t.maxTexture2DLayered[2]);
		printf("Maximum dimensions (width, height, pitch) for 2D textures bound to pitched memory %d\n",hipDeviceProp_t.maxTexture2DLinear[0]);
		printf("Maximum dimensions (width, height, pitch) for 2D textures bound to pitched memory %d\n",hipDeviceProp_t.maxTexture2DLinear[1]);
		printf("Maximum dimensions (width, height, pitch) for 2D textures bound to pitched memory %d\n",hipDeviceProp_t.maxTexture2DLinear[2]);
		printf("Maximum 2D mipmapped texture dimensions %d\n",hipDeviceProp_t.maxTexture2DMipmap[0]);
		printf("Maximum 2D mipmapped texture dimensions %d\n",hipDeviceProp_t.maxTexture2DMipmap[1]);
		printf("Maximum 3D texture dimensions %d\n",hipDeviceProp_t.maxTexture3D[0]);
		printf("Maximum 3D texture dimensions %d\n",hipDeviceProp_t.maxTexture3D[1]);
		printf("Maximum 3D texture dimensions %d\n",hipDeviceProp_t.maxTexture3D[2]);
		printf("Maximum alternate 3D texture dimensions %d\n",hipDeviceProp_t.maxTexture3DAlt[0]);
		printf("Maximum alternate 3D texture dimensions %d\n",hipDeviceProp_t.maxTexture3DAlt[1]);
		printf("Maximum alternate 3D texture dimensions %d\n",hipDeviceProp_t.maxTexture3DAlt[2]);
		printf("Maximum Cubemap texture dimensions %d\n",hipDeviceProp_t.maxTextureCubemap);
		printf("Maximum Cubemap layered texture dimensions %d\n",hipDeviceProp_t.maxTextureCubemapLayered[0]);
		printf("Maximum Cubemap layered texture dimensions %d\n",hipDeviceProp_t.maxTextureCubemapLayered[1]);
		printf("Maximum size of each dimension of a block %d\n",hipDeviceProp_t.maxThreadsDim[0]);
		printf("Maximum size of each dimension of a block %d\n",hipDeviceProp_t.maxThreadsDim[1]);
		printf("Maximum size of each dimension of a block %d\n",hipDeviceProp_t.maxThreadsDim[2]);
		printf("Maximum number of threads per block %d\n",hipDeviceProp_t.maxThreadsPerBlock);
		printf("Maximum resident threads per multiprocessor %d\n",hipDeviceProp_t.maxThreadsPerMultiProcessor);
		printf("Maximum pitch in bytes allowed by memory copies %d\n",hipDeviceProp_t.memPitch);
		printf("Global memory bus width in bits %d\n",hipDeviceProp_t.memoryBusWidth);
		printf("Peak memory clock frequency in kilohertz %d\n",hipDeviceProp_t.memoryClockRate);
		printf("Minor compute capability %d\n",hipDeviceProp_t.minor);
		printf("Number of multiprocessors on device %d\n",hipDeviceProp_t.multiProcessorCount);
		printf("PCI bus ID of the device %d\n",hipDeviceProp_t.pciBusID);
		printf("PCI device ID of the device %d\n",hipDeviceProp_t.pciDeviceID);
		printf("PCI domain ID of the device %d\n",hipDeviceProp_t.pciDomainID);
		printf("32-bit registers available per block %d\n",hipDeviceProp_t.regsPerBlock);
		printf("Shared memory available per block in bytes %d\n",hipDeviceProp_t.sharedMemPerBlock);
		printf("Device supports stream priorities %d\n",hipDeviceProp_t.streamPrioritiesSupported);
		printf("Alignment requirements for surfaces %d\n",hipDeviceProp_t.surfaceAlignment);
		printf("1 if device is a Tesla device using TCC driver, 0 otherwise %d\n",hipDeviceProp_t.tccDriver);
		printf("Alignment requirement for textures %d\n",hipDeviceProp_t.textureAlignment);
		printf("Pitch alignment requirement for texture references bound to pitched memory %d\n",hipDeviceProp_t.texturePitchAlignment);
		printf("Constant memory available on device in bytes %d\n",hipDeviceProp_t.totalConstMem);
		printf("Global memory available on device in bytes %d\n",hipDeviceProp_t.totalGlobalMem);
		printf("Device shares a unified address space with the host %d\n",hipDeviceProp_t.unifiedAddressing);
		printf("Warp size in threads %d\n",hipDeviceProp_t.warpSize);

		fflush(stdout);
	}

	host_floyd(gridSize, blockSize, matrix, intermedian, n);

	hipDeviceReset();

	/* Bыводим результаты */
	fs = fopen(outputFileName, "w");
	if (fs == NULL) {
		fprintf(stderr, "File open error (%s)\n", outputFileName); fflush(stderr);
		exit(-1);
	}
	for (i = 0; i < n; i++){
		for (j = 0; j < n; j++){
			if (matrix[i*n + j] != LONG_MAX)
				fprintf(fs, "%ld%s", matrix[i*n + j], ((j == n - 1) ? "\n" : ";"));
			else
				fprintf(fs, "%s", ((j == n - 1) ? "\n" : ";"));
		}
	}
	fclose(fs);

	/* Bыводим результаты */
	fs = fopen(intermedianFileName, "w");
	if (fs == NULL) {
		fprintf(stderr, "File open error (%s)\n", intermedianFileName); fflush(stderr);
		exit(-1);
	}
	for (i = 0; i < n; i++){
		for (j = 0; j < n; j++){
			if (intermedian[i*n + j] >= 0)
				fprintf(fs, "%d%s", intermedian[i*n + j], ((j == n - 1) ? "\n" : ";"));
			else
				fprintf(fs, "%s", ((j == n - 1) ? "\n" : ";"));
		}
	}
	fclose(fs);

	free(matrix);
	free(intermedian);

	exit(0);
}
	.text
	.file	"floyd.hip"
                                        # Start of file scope inline assembly
	.globl	_ZSt21ios_base_library_initv

                                        # End of file scope inline assembly
	.globl	_Z8stremptyPKc                  # -- Begin function _Z8stremptyPKc
	.type	_Z8stremptyPKc,@function
_Z8stremptyPKc:                         # @_Z8stremptyPKc
	.cfi_startproc
# %bb.0:
	pushq	%r14
	.cfi_def_cfa_offset 16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	pushq	%rax
	.cfi_def_cfa_offset 32
	.cfi_offset %rbx, -24
	.cfi_offset %r14, -16
	movl	$1, %ebx
	testq	%rdi, %rdi
	je	.LBB0_6
# %bb.1:                                # %.preheader
	movq	%rdi, %r14
	movb	(%rdi), %al
	testb	%al, %al
	je	.LBB0_6
# %bb.2:                                # %.lr.ph.preheader
	incq	%r14
.LBB0_4:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movsbl	%al, %edi
	callq	isspace
	testl	%eax, %eax
	je	.LBB0_5
# %bb.3:                                #   in Loop: Header=BB0_4 Depth=1
	movb	(%r14), %al
	incq	%r14
	testb	%al, %al
	jne	.LBB0_4
	jmp	.LBB0_6
.LBB0_5:
	xorl	%ebx, %ebx
.LBB0_6:                                # %.loopexit
	movl	%ebx, %eax
	addq	$8, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end0:
	.size	_Z8stremptyPKc, .Lfunc_end0-_Z8stremptyPKc
	.cfi_endproc
                                        # -- End function
	.globl	_Z8mystrtokPPcS_c               # -- Begin function _Z8mystrtokPPcS_c
	.type	_Z8mystrtokPPcS_c,@function
_Z8mystrtokPPcS_c:                      # @_Z8mystrtokPPcS_c
	.cfi_startproc
# %bb.0:
	pushq	%r14
	.cfi_def_cfa_offset 16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	pushq	%rax
	.cfi_def_cfa_offset 32
	.cfi_offset %rbx, -24
	.cfi_offset %r14, -16
	movq	%rsi, %rbx
	movq	%rdi, %r14
	testq	%rsi, %rsi
	jne	.LBB1_2
# %bb.1:
	movq	(%r14), %rbx
.LBB1_2:
	cmpb	$0, (%rbx)
	je	.LBB1_3
# %bb.4:
	movsbl	%dl, %esi
	movq	%rbx, %rdi
	callq	strchr
	movq	%rax, (%r14)
	testq	%rax, %rax
	je	.LBB1_6
# %bb.5:
	movq	%rax, %rcx
	incq	%rcx
	movq	%rcx, (%r14)
	movb	$0, (%rax)
	jmp	.LBB1_7
.LBB1_3:
	xorl	%ebx, %ebx
	jmp	.LBB1_7
.LBB1_6:
	movq	%rbx, %rdi
	callq	strlen
	addq	%rbx, %rax
	movq	%rax, (%r14)
.LBB1_7:
	movq	%rbx, %rax
	addq	$8, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end1:
	.size	_Z8mystrtokPPcS_c, .Lfunc_end1-_Z8mystrtokPPcS_c
	.cfi_endproc
                                        # -- End function
	.globl	_Z26__device_stub__global_initPlS_Piii # -- Begin function _Z26__device_stub__global_initPlS_Piii
	.type	_Z26__device_stub__global_initPlS_Piii,@function
_Z26__device_stub__global_initPlS_Piii: # @_Z26__device_stub__global_initPlS_Piii
	.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	24(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	4(%rsp), %rdx
	movl	%ecx, (%rdx)
	movq	%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	16(%rsp), %r12
	leaq	8(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z11global_initPlS_Piii, %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_end2:
	.size	_Z26__device_stub__global_initPlS_Piii, .Lfunc_end2-_Z26__device_stub__global_initPlS_Piii
	.cfi_endproc
                                        # -- End function
	.globl	_Z27__device_stub__global_floydPlS_Piii # -- Begin function _Z27__device_stub__global_floydPlS_Piii
	.type	_Z27__device_stub__global_floydPlS_Piii,@function
_Z27__device_stub__global_floydPlS_Piii: # @_Z27__device_stub__global_floydPlS_Piii
	.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	24(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	4(%rsp), %rdx
	movl	%ecx, (%rdx)
	movq	%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	16(%rsp), %r12
	leaq	8(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z12global_floydPlS_Piii, %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_end3:
	.size	_Z27__device_stub__global_floydPlS_Piii, .Lfunc_end3-_Z27__device_stub__global_floydPlS_Piii
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function _Z10host_floydiiPlPii
.LCPI4_0:
	.quad	0x3fd555555555554f              # double 0.33333333333333298
	.text
	.globl	_Z10host_floydiiPlPii
	.type	_Z10host_floydiiPlPii,@function
_Z10host_floydiiPlPii:                  # @_Z10host_floydiiPlPii
	.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	$104, %rsp
	.cfi_def_cfa_offset 160
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movl	%r8d, %r14d
	movq	%rcx, 72(%rsp)                  # 8-byte Spill
	movq	%rdx, %r12
	movl	%esi, 8(%rsp)                   # 4-byte Spill
	movl	%edi, %ebp
	movl	%r8d, %r15d
	imull	%r15d, %r15d
	leaq	(,%r15,8), %r13
	leaq	24(%rsp), %rbx
	movq	%rbx, %rdi
	movq	%r13, %rsi
	callq	hipMalloc
	leaq	16(%rsp), %rdi
	movq	%r13, %rsi
	callq	hipMalloc
	shlq	$2, %r15
	leaq	40(%rsp), %rdi
	movq	%r15, 64(%rsp)                  # 8-byte Spill
	movq	%r15, %rsi
	callq	hipMalloc
	movq	(%rbx), %rdi
	movq	%r12, 48(%rsp)                  # 8-byte Spill
	movq	%r12, %rsi
	movq	%r13, 56(%rsp)                  # 8-byte Spill
	movq	%r13, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	testl	%ebp, %ebp
	jg	.LBB4_6
# %bb.1:
	cvtsi2sd	%r14d, %xmm0
	mulsd	%xmm0, %xmm0
	movsd	.LCPI4_0(%rip), %xmm1           # xmm1 = mem[0],zero
	movsd	%xmm0, 32(%rsp)                 # 8-byte Spill
	callq	pow
	cvttsd2si	%xmm0, %eax
	testl	%eax, %eax
	jle	.LBB4_3
# %bb.2:
	movsd	.LCPI4_0(%rip), %xmm1           # xmm1 = mem[0],zero
	movsd	32(%rsp), %xmm0                 # 8-byte Reload
                                        # xmm0 = mem[0],zero
	callq	pow
	cvttsd2si	%xmm0, %eax
	movl	$15, %ebp
	cmpl	$14, %eax
	jg	.LBB4_6
.LBB4_3:                                # %.critedge
	movsd	.LCPI4_0(%rip), %xmm1           # xmm1 = mem[0],zero
	movsd	32(%rsp), %xmm0                 # 8-byte Reload
                                        # xmm0 = mem[0],zero
	callq	pow
	cvttsd2si	%xmm0, %eax
	testl	%eax, %eax
	jle	.LBB4_4
# %bb.5:
	movsd	.LCPI4_0(%rip), %xmm1           # xmm1 = mem[0],zero
	movsd	32(%rsp), %xmm0                 # 8-byte Reload
                                        # xmm0 = mem[0],zero
	callq	pow
	cvttsd2si	%xmm0, %ebp
	jmp	.LBB4_6
.LBB4_4:
	movl	$1, %ebp
.LBB4_6:
	movl	8(%rsp), %ecx                   # 4-byte Reload
	testl	%ecx, %ecx
	jg	.LBB4_12
# %bb.7:
	xorps	%xmm0, %xmm0
	cvtsi2sd	%r14d, %xmm0
	mulsd	%xmm0, %xmm0
	movsd	.LCPI4_0(%rip), %xmm1           # xmm1 = mem[0],zero
	movsd	%xmm0, 8(%rsp)                  # 8-byte Spill
	callq	pow
	cvttsd2si	%xmm0, %eax
	testl	%eax, %eax
	jle	.LBB4_9
# %bb.8:
	movsd	.LCPI4_0(%rip), %xmm1           # xmm1 = mem[0],zero
	movsd	8(%rsp), %xmm0                  # 8-byte Reload
                                        # xmm0 = mem[0],zero
	callq	pow
	cvttsd2si	%xmm0, %eax
	movl	$15, %ecx
	cmpl	$14, %eax
	jg	.LBB4_12
.LBB4_9:                                # %.critedge74
	movsd	.LCPI4_0(%rip), %xmm1           # xmm1 = mem[0],zero
	movsd	8(%rsp), %xmm0                  # 8-byte Reload
                                        # xmm0 = mem[0],zero
	callq	pow
	cvttsd2si	%xmm0, %eax
	testl	%eax, %eax
	jle	.LBB4_10
# %bb.11:
	movsd	.LCPI4_0(%rip), %xmm1           # xmm1 = mem[0],zero
	movsd	8(%rsp), %xmm0                  # 8-byte Reload
                                        # xmm0 = mem[0],zero
	callq	pow
	cvttsd2si	%xmm0, %ecx
	jmp	.LBB4_12
.LBB4_10:
	movl	$1, %ecx
.LBB4_12:
	movl	%ebp, %edi
	btsq	$32, %rdi
	movl	%ecx, %edx
	btsq	$32, %rdx
	movq	%rdi, 88(%rsp)                  # 8-byte Spill
	movl	$1, %esi
	movq	%rdx, 80(%rsp)                  # 8-byte Spill
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB4_14
# %bb.13:
	movq	24(%rsp), %rdi
	movq	16(%rsp), %rsi
	movq	40(%rsp), %rdx
	movl	$-1, %ecx
	movl	%r14d, %r8d
	callq	_Z26__device_stub__global_initPlS_Piii
.LBB4_14:
	testl	%r14d, %r14d
	jle	.LBB4_15
# %bb.17:                               # %.lr.ph89
	leal	-1(%r14), %r12d
	movl	%r14d, %ebp
	movl	$.L.str.3, %ebx
	xorl	%r15d, %r15d
	movq	%r14, 96(%rsp)                  # 8-byte Spill
.LBB4_18:                               # =>This Loop Header: Depth=1
                                        #     Child Loop BB4_21 Depth 2
                                        #       Child Loop BB4_22 Depth 3
	movq	88(%rsp), %rdi                  # 8-byte Reload
	movl	$1, %esi
	movq	80(%rsp), %rdx                  # 8-byte Reload
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB4_20
# %bb.19:                               #   in Loop: Header=BB4_18 Depth=1
	movq	24(%rsp), %rdi
	movq	16(%rsp), %rsi
	movq	40(%rsp), %rdx
	movl	%r15d, %ecx
	movl	%r14d, %r8d
	callq	_Z27__device_stub__global_floydPlS_Piii
.LBB4_20:                               # %.preheader.preheader
                                        #   in Loop: Header=BB4_18 Depth=1
	movl	%r15d, 32(%rsp)                 # 4-byte Spill
	movq	16(%rsp), %rsi
	movq	48(%rsp), %rdi                  # 8-byte Reload
	movq	56(%rsp), %rdx                  # 8-byte Reload
	movl	$2, %ecx
	callq	hipMemcpy
	xorl	%eax, %eax
	xorl	%r15d, %r15d
.LBB4_21:                               # %.preheader
                                        #   Parent Loop BB4_18 Depth=1
                                        # =>  This Loop Header: Depth=2
                                        #       Child Loop BB4_22 Depth 3
	movq	%rax, 8(%rsp)                   # 8-byte Spill
	movl	%eax, %eax
	movq	48(%rsp), %rcx                  # 8-byte Reload
	leaq	(%rcx,%rax,8), %r13
	xorl	%r14d, %r14d
.LBB4_22:                               #   Parent Loop BB4_18 Depth=1
                                        #     Parent Loop BB4_21 Depth=2
                                        # =>    This Inner Loop Header: Depth=3
	movq	(%r13,%r14,8), %rsi
	cmpq	%r14, %r12
	movl	$.L.str.4, %edx
	cmoveq	%rbx, %rdx
	movl	$.L.str.2, %edi
	xorl	%eax, %eax
	callq	printf
	incq	%r14
	cmpq	%r14, %rbp
	jne	.LBB4_22
# %bb.23:                               # %._crit_edge
                                        #   in Loop: Header=BB4_21 Depth=2
	incq	%r15
	movq	96(%rsp), %r14                  # 8-byte Reload
	movq	8(%rsp), %rax                   # 8-byte Reload
	addl	%r14d, %eax
	cmpq	%rbp, %r15
	jne	.LBB4_21
# %bb.24:                               # %._crit_edge86
                                        #   in Loop: Header=BB4_18 Depth=1
	movl	$10, %edi
	callq	putchar@PLT
	movq	24(%rsp), %rax
	movq	16(%rsp), %rsi
	movq	%rsi, 24(%rsp)
	movq	%rax, 16(%rsp)
	movl	32(%rsp), %r15d                 # 4-byte Reload
	incl	%r15d
	cmpl	%r14d, %r15d
	jne	.LBB4_18
	jmp	.LBB4_16
.LBB4_15:                               # %.._crit_edge90_crit_edge
	movq	24(%rsp), %rsi
.LBB4_16:                               # %._crit_edge90
	movq	48(%rsp), %rdi                  # 8-byte Reload
	movq	56(%rsp), %rdx                  # 8-byte Reload
	movl	$2, %ecx
	callq	hipMemcpy
	movq	40(%rsp), %rsi
	movq	72(%rsp), %rdi                  # 8-byte Reload
	movq	64(%rsp), %rdx                  # 8-byte Reload
	movl	$2, %ecx
	callq	hipMemcpy
	movq	24(%rsp), %rdi
	callq	hipFree
	movq	16(%rsp), %rdi
	callq	hipFree
	movq	40(%rsp), %rdi
	callq	hipFree
	addq	$104, %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_end4:
	.size	_Z10host_floydiiPlPii, .Lfunc_end4-_Z10host_floydiiPlPii
	.cfi_endproc
                                        # -- End function
	.globl	main                            # -- Begin function main
	.type	main,@function
main:                                   # @main
	.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	$5672, %rsp                     # imm = 0x1628
	.cfi_def_cfa_offset 5728
	.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, %rbx
	movl	%edi, %ebp
	movq	title(%rip), %rsi
	xorl	%r12d, %r12d
	movl	$.L.str.5, %edi
	xorl	%eax, %eax
	callq	printf
	movq	stdout(%rip), %rdi
	callq	fflush
	cmpl	$4, %ebp
	jge	.LBB5_1
# %bb.69:
	movq	(%rbx), %rsi
	movl	$.L.str.6, %edi
	xorl	%eax, %eax
	callq	printf
	movq	stdout(%rip), %rdi
	callq	fflush
	movl	$.Lstr, %edi
	callq	puts@PLT
	movq	stdout(%rip), %rdi
	callq	fflush
	movl	$.Lstr.1, %edi
	callq	puts@PLT
	movq	stdout(%rip), %rdi
	callq	fflush
	movl	$.Lstr.2, %edi
	callq	puts@PLT
	movq	stdout(%rip), %rdi
	jmp	.LBB5_25
.LBB5_1:                                # %.preheader196.preheader
	movl	$1, %r15d
	xorl	%r13d, %r13d
.LBB5_2:                                # %.preheader196
                                        # =>This Inner Loop Header: Depth=1
	movslq	%r15d, %rax
	movq	(%rbx,%rax,8), %r14
	cmpb	$45, (%r14)
	jne	.LBB5_9
# %bb.3:                                #   in Loop: Header=BB5_2 Depth=1
	movzbl	1(%r14), %eax
	cmpl	$98, %eax
	je	.LBB5_6
# %bb.4:                                #   in Loop: Header=BB5_2 Depth=1
	cmpl	$103, %eax
	jne	.LBB5_7
# %bb.5:                                #   in Loop: Header=BB5_2 Depth=1
	movslq	%r15d, %rax
	incl	%r15d
	movq	8(%rbx,%rax,8), %rdi
	callq	atoi
	movl	%eax, %r12d
	jmp	.LBB5_7
.LBB5_6:                                #   in Loop: Header=BB5_2 Depth=1
	movslq	%r15d, %rax
	incl	%r15d
	movq	8(%rbx,%rax,8), %rdi
	callq	atoi
	movl	%eax, %r13d
.LBB5_7:                                #   in Loop: Header=BB5_2 Depth=1
	incl	%r15d
	cmpl	%ebp, %r15d
	jl	.LBB5_2
# %bb.8:                                # %..critedge_crit_edge
	movslq	%r15d, %rax
	movq	(%rbx,%rax,8), %r14
.LBB5_9:                                # %.critedge
	movq	8(%rbx,%rax,8), %rbp
	movq	16(%rbx,%rax,8), %r15
	xorl	%ebx, %ebx
	movl	$.L.str.10, %edi
	movq	%r14, %rsi
	xorl	%eax, %eax
	callq	printf
	movq	stdout(%rip), %rdi
	callq	fflush
	movl	$.L.str.11, %edi
	movq	%rbp, 80(%rsp)                  # 8-byte Spill
	movq	%rbp, %rsi
	xorl	%eax, %eax
	callq	printf
	movq	stdout(%rip), %rdi
	callq	fflush
	movl	$.L.str.12, %edi
	movq	%r15, %rsi
	xorl	%eax, %eax
	callq	printf
	movq	stdout(%rip), %rdi
	callq	fflush
	movl	$.L.str.13, %esi
	movq	%r14, %rdi
	callq	fopen
	testq	%rax, %rax
	je	.LBB5_24
# %bb.10:                               # %.preheader195
	movl	%r13d, 60(%rsp)                 # 4-byte Spill
	movl	%r12d, 64(%rsp)                 # 4-byte Spill
	leaq	1568(%rsp), %rdi
	movl	$4096, %esi                     # imm = 0x1000
	movq	%rax, 32(%rsp)                  # 8-byte Spill
	movq	%rax, %rdx
	callq	fgets
	movl	$0, %ebp
	testq	%rax, %rax
	je	.LBB5_16
# %bb.11:                               # %.lr.ph205.preheader
	xorl	%ebp, %ebp
	leaq	88(%rsp), %r12
	leaq	1568(%rsp), %r13
	xorl	%ebx, %ebx
.LBB5_12:                               # %.lr.ph205
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB5_14 Depth 2
	movq	%rbp, %r14
	movq	%r12, %rdi
	movq	%rax, %rsi
	movl	$59, %edx
	callq	_Z8mystrtokPPcS_c
	movl	$0, %ebp
	testq	%rax, %rax
	je	.LBB5_15
# %bb.13:                               # %.lr.ph.preheader
                                        #   in Loop: Header=BB5_12 Depth=1
	xorl	%ebp, %ebp
.LBB5_14:                               # %.lr.ph
                                        #   Parent Loop BB5_12 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	incl	%ebp
	movq	%r12, %rdi
	xorl	%esi, %esi
	movl	$59, %edx
	callq	_Z8mystrtokPPcS_c
	testq	%rax, %rax
	jne	.LBB5_14
.LBB5_15:                               # %._crit_edge
                                        #   in Loop: Header=BB5_12 Depth=1
	cmpl	%ebp, %ebx
	cmovlel	%ebp, %ebx
	movq	%r14, %rbp
	incl	%ebp
	movq	%r13, %rdi
	movl	$4096, %esi                     # imm = 0x1000
	movq	32(%rsp), %rdx                  # 8-byte Reload
	callq	fgets
	testq	%rax, %rax
	jne	.LBB5_12
.LBB5_16:                               # %._crit_edge206
	movq	%r15, 72(%rsp)                  # 8-byte Spill
	cmpl	%ebp, %ebx
	cmovgl	%ebx, %ebp
	movl	%ebp, %r15d
	imull	%r15d, %r15d
	leaq	(,%r15,8), %rdi
	callq	malloc
	movq	%rax, %r14
	shlq	$2, %r15
	movq	%r15, %rdi
	callq	malloc
	movq	%rax, 48(%rsp)                  # 8-byte Spill
	xorl	%r15d, %r15d
	movq	32(%rsp), %rbx                  # 8-byte Reload
	movq	%rbx, %rdi
	xorl	%esi, %esi
	xorl	%edx, %edx
	callq	fseek
	leaq	1568(%rsp), %rdi
	movl	$4096, %esi                     # imm = 0x1000
	movq	%rbx, %rdx
	callq	fgets
	testq	%rax, %rax
	movq	%rbp, 40(%rsp)                  # 8-byte Spill
	movq	%r14, 16(%rsp)                  # 8-byte Spill
	je	.LBB5_27
# %bb.17:                               # %.lr.ph219.preheader
	movl	%ebp, %ebx
	xorl	%ecx, %ecx
	movq	%rcx, 24(%rsp)                  # 8-byte Spill
	leaq	88(%rsp), %r13
	xorl	%r15d, %r15d
	movabsq	$9223372036854775807, %r14      # imm = 0x7FFFFFFFFFFFFFFF
.LBB5_18:                               # %.lr.ph219
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB5_20 Depth 2
                                        #       Child Loop BB5_22 Depth 3
                                        #     Child Loop BB5_41 Depth 2
	movq	%r15, 8(%rsp)                   # 8-byte Spill
	movq	%r13, %rdi
	movq	%rax, %rsi
	movl	$59, %edx
	callq	_Z8mystrtokPPcS_c
	movl	$0, %r12d
	testq	%rax, %rax
	je	.LBB5_39
# %bb.19:                               # %.preheader.i.lr.ph
                                        #   in Loop: Header=BB5_18 Depth=1
	movq	%rax, %r15
	movl	%ebp, %eax
	imull	8(%rsp), %eax                   # 4-byte Folded Reload
	movq	16(%rsp), %rcx                  # 8-byte Reload
	leaq	(%rcx,%rax,8), %rbp
	xorl	%r12d, %r12d
.LBB5_20:                               # %.preheader.i
                                        #   Parent Loop BB5_18 Depth=1
                                        # =>  This Loop Header: Depth=2
                                        #       Child Loop BB5_22 Depth 3
	movb	(%r15), %cl
	movq	%r14, %rax
	testb	%cl, %cl
	je	.LBB5_45
# %bb.21:                               # %.lr.ph.i.preheader
                                        #   in Loop: Header=BB5_20 Depth=2
	leaq	1(%r15), %r14
.LBB5_22:                               # %.lr.ph.i
                                        #   Parent Loop BB5_18 Depth=1
                                        #     Parent Loop BB5_20 Depth=2
                                        # =>    This Inner Loop Header: Depth=3
	movsbl	%cl, %edi
	callq	isspace
	testl	%eax, %eax
	je	.LBB5_23
# %bb.43:                               #   in Loop: Header=BB5_22 Depth=3
	movb	(%r14), %cl
	incq	%r14
	testb	%cl, %cl
	jne	.LBB5_22
# %bb.44:                               #   in Loop: Header=BB5_20 Depth=2
	movabsq	$9223372036854775807, %r14      # imm = 0x7FFFFFFFFFFFFFFF
	movq	%r14, %rax
	jmp	.LBB5_45
.LBB5_23:                               # %_Z8stremptyPKc.exit
                                        #   in Loop: Header=BB5_20 Depth=2
	movq	%r15, %rdi
	callq	atol
	movabsq	$9223372036854775807, %r14      # imm = 0x7FFFFFFFFFFFFFFF
.LBB5_45:                               # %_Z8stremptyPKc.exit.thread
                                        #   in Loop: Header=BB5_20 Depth=2
	movq	%rax, (%rbp,%r12,8)
	incq	%r12
	movq	%r13, %rdi
	xorl	%esi, %esi
	movl	$59, %edx
	callq	_Z8mystrtokPPcS_c
	movq	%rax, %r15
	testq	%rax, %rax
	jne	.LBB5_20
# %bb.38:                               # %.preheader194.loopexit
                                        #   in Loop: Header=BB5_18 Depth=1
	movq	40(%rsp), %rbp                  # 8-byte Reload
.LBB5_39:                               # %.preheader194
                                        #   in Loop: Header=BB5_18 Depth=1
	cmpl	%ebp, %r12d
	jae	.LBB5_42
# %bb.40:                               # %.lr.ph214
                                        #   in Loop: Header=BB5_18 Depth=1
	movl	24(%rsp), %eax                  # 4-byte Reload
	movq	16(%rsp), %rcx                  # 8-byte Reload
	leaq	(%rcx,%rax,8), %rax
	movl	%r12d, %ecx
.LBB5_41:                               #   Parent Loop BB5_18 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movq	%r14, (%rax,%rcx,8)
	incq	%rcx
	cmpq	%rcx, %rbx
	jne	.LBB5_41
.LBB5_42:                               # %._crit_edge215
                                        #   in Loop: Header=BB5_18 Depth=1
	movq	8(%rsp), %r15                   # 8-byte Reload
	incq	%r15
	leaq	1568(%rsp), %rdi
	movl	$4096, %esi                     # imm = 0x1000
	movq	32(%rsp), %rdx                  # 8-byte Reload
	callq	fgets
	movq	24(%rsp), %rcx                  # 8-byte Reload
	addl	%ebp, %ecx
	movq	%rcx, 24(%rsp)                  # 8-byte Spill
	testq	%rax, %rax
	jne	.LBB5_18
# %bb.26:                               # %.preheader193.loopexit
	movq	16(%rsp), %r14                  # 8-byte Reload
.LBB5_27:                               # %.preheader193
	movl	%ebp, %eax
	subl	%r15d, %eax
	imull	%ebp, %eax
	testl	%eax, %eax
	movabsq	$9223372036854775807, %rsi      # imm = 0x7FFFFFFFFFFFFFFF
	jle	.LBB5_30
# %bb.28:                               # %.lr.ph222
	imull	%ebp, %r15d
	movl	%eax, %eax
	leaq	(%r14,%r15,8), %rcx
	xorl	%edx, %edx
.LBB5_29:                               # =>This Inner Loop Header: Depth=1
	movq	%rsi, (%rcx,%rdx,8)
	incq	%rdx
	cmpq	%rdx, %rax
	jne	.LBB5_29
.LBB5_30:                               # %.preheader192
	testl	%ebp, %ebp
	jne	.LBB5_31
# %bb.46:                               # %._crit_edge228.critedge
	movq	32(%rsp), %rdi                  # 8-byte Reload
	callq	fclose
	movl	$.L.str.15, %edi
	movl	%ebp, %esi
	xorl	%eax, %eax
	callq	printf
	jmp	.LBB5_47
.LBB5_31:                               # %.lr.ph224
	leal	1(%rbp), %eax
	movl	%ebp, %ebx
	xorl	%ecx, %ecx
	movq	%rbx, %rdx
	movabsq	$9223372036854775807, %rsi      # imm = 0x7FFFFFFFFFFFFFFF
.LBB5_32:                               # =>This Inner Loop Header: Depth=1
	movslq	%ecx, %rcx
	movq	%rsi, (%r14,%rcx,8)
	addl	%eax, %ecx
	decq	%rdx
	jne	.LBB5_32
# %bb.33:                               # %._crit_edge225
	movq	32(%rsp), %rdi                  # 8-byte Reload
	callq	fclose
	movl	$.L.str.15, %edi
	movl	%ebp, %esi
	xorl	%eax, %eax
	callq	printf
	testl	%ebp, %ebp
	je	.LBB5_47
# %bb.34:                               # %.preheader191.lr.ph
	leal	-1(%rbp), %r14d
	movl	$.L.str.3, %r15d
	xorl	%eax, %eax
	xorl	%r12d, %r12d
.LBB5_35:                               # %.preheader191
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB5_36 Depth 2
	movl	%eax, 8(%rsp)                   # 4-byte Spill
	movl	%eax, %eax
	movq	16(%rsp), %rcx                  # 8-byte Reload
	leaq	(%rcx,%rax,8), %r13
	xorl	%ebp, %ebp
.LBB5_36:                               #   Parent Loop BB5_35 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movq	(%r13,%rbp,8), %rsi
	cmpq	%rbp, %r14
	movl	$.L.str.16, %edx
	cmoveq	%r15, %rdx
	movl	$.L.str.2, %edi
	xorl	%eax, %eax
	callq	printf
	incq	%rbp
	cmpq	%rbp, %rbx
	jne	.LBB5_36
# %bb.37:                               #   in Loop: Header=BB5_35 Depth=1
	incq	%r12
	movq	40(%rsp), %rbp                  # 8-byte Reload
	movl	8(%rsp), %eax                   # 4-byte Reload
	addl	%ebp, %eax
	cmpq	%rbx, %r12
	jne	.LBB5_35
.LBB5_47:                               # %._crit_edge228
	movq	stdout(%rip), %rdi
	callq	fflush
	leaq	68(%rsp), %rbx
	movl	$0, (%rbx)
	movq	%rbx, %rdi
	callq	hipGetDeviceCount
	cmpl	$0, (%rbx)
	jle	.LBB5_50
# %bb.48:                               # %.lr.ph231
	leaq	96(%rsp), %r13
	xorl	%ebx, %ebx
.LBB5_49:                               # =>This Inner Loop Header: Depth=1
	movq	%r13, %rdi
	movl	%ebx, %esi
	callq	hipGetDevicePropertiesR0600
	movl	$.L.str.17, %edi
	movl	%ebx, %esi
	movq	%r13, %rdx
	xorl	%eax, %eax
	callq	printf
	movl	676(%rsp), %esi
	movl	$.L.str.18, %edi
	xorl	%eax, %eax
	callq	printf
	movl	696(%rsp), %esi
	movl	$.L.str.19, %edi
	xorl	%eax, %eax
	callq	printf
	movl	496(%rsp), %esi
	movl	$.L.str.20, %edi
	xorl	%eax, %eax
	callq	printf
	movl	444(%rsp), %esi
	movl	$.L.str.21, %edi
	xorl	%eax, %eax
	callq	printf
	movl	500(%rsp), %esi
	movl	$.L.str.22, %edi
	xorl	%eax, %eax
	callq	printf
	movl	672(%rsp), %esi
	movl	$.L.str.23, %edi
	xorl	%eax, %eax
	callq	printf
	movl	480(%rsp), %esi
	movl	$.L.str.24, %edi
	xorl	%eax, %eax
	callq	printf
	movl	492(%rsp), %esi
	movl	$.L.str.25, %edi
	xorl	%eax, %eax
	callq	printf
	movl	488(%rsp), %esi
	movl	$.L.str.26, %edi
	xorl	%eax, %eax
	callq	printf
	movl	712(%rsp), %esi
	movl	$.L.str.27, %edi
	xorl	%eax, %eax
	callq	printf
	movl	456(%rsp), %esi
	movl	$.L.str.28, %edi
	xorl	%eax, %eax
	callq	printf
	movl	432(%rsp), %esi
	movl	$.L.str.29, %edi
	xorl	%eax, %eax
	callq	printf
	movl	436(%rsp), %esi
	movl	$.L.str.29, %edi
	xorl	%eax, %eax
	callq	printf
	movl	440(%rsp), %esi
	movl	$.L.str.29, %edi
	xorl	%eax, %eax
	callq	printf
	movl	608(%rsp), %esi
	movl	$.L.str.30, %edi
	xorl	%eax, %eax
	callq	printf
	movl	632(%rsp), %esi
	movl	$.L.str.31, %edi
	xorl	%eax, %eax
	callq	printf
	movl	636(%rsp), %esi
	movl	$.L.str.31, %edi
	xorl	%eax, %eax
	callq	printf
	movl	612(%rsp), %esi
	movl	$.L.str.32, %edi
	xorl	%eax, %eax
	callq	printf
	movl	616(%rsp), %esi
	movl	$.L.str.32, %edi
	xorl	%eax, %eax
	callq	printf
	movl	640(%rsp), %esi
	movl	$.L.str.33, %edi
	xorl	%eax, %eax
	callq	printf
	movl	644(%rsp), %esi
	movl	$.L.str.33, %edi
	xorl	%eax, %eax
	callq	printf
	movl	648(%rsp), %esi
	movl	$.L.str.33, %edi
	xorl	%eax, %eax
	callq	printf
	movl	620(%rsp), %esi
	movl	$.L.str.34, %edi
	xorl	%eax, %eax
	callq	printf
	movl	624(%rsp), %esi
	movl	$.L.str.34, %edi
	xorl	%eax, %eax
	callq	printf
	movl	628(%rsp), %esi
	movl	$.L.str.34, %edi
	xorl	%eax, %eax
	callq	printf
	movl	652(%rsp), %esi
	movl	$.L.str.35, %edi
	xorl	%eax, %eax
	callq	printf
	movl	656(%rsp), %esi
	movl	$.L.str.36, %edi
	xorl	%eax, %eax
	callq	printf
	movl	660(%rsp), %esi
	movl	$.L.str.36, %edi
	xorl	%eax, %eax
	callq	printf
	movl	504(%rsp), %esi
	movl	$.L.str.37, %edi
	xorl	%eax, %eax
	callq	printf
	movl	580(%rsp), %esi
	movl	$.L.str.38, %edi
	xorl	%eax, %eax
	callq	printf
	movl	584(%rsp), %esi
	movl	$.L.str.38, %edi
	xorl	%eax, %eax
	callq	printf
	movl	512(%rsp), %esi
	movl	$.L.str.39, %edi
	xorl	%eax, %eax
	callq	printf
	movl	508(%rsp), %esi
	movl	$.L.str.40, %edi
	xorl	%eax, %eax
	callq	printf
	movl	516(%rsp), %esi
	movl	$.L.str.41, %edi
	xorl	%eax, %eax
	callq	printf
	movl	520(%rsp), %esi
	movl	$.L.str.41, %edi
	xorl	%eax, %eax
	callq	printf
	movl	544(%rsp), %esi
	movl	$.L.str.42, %edi
	xorl	%eax, %eax
	callq	printf
	movl	548(%rsp), %esi
	movl	$.L.str.42, %edi
	xorl	%eax, %eax
	callq	printf
	movl	588(%rsp), %esi
	movl	$.L.str.43, %edi
	xorl	%eax, %eax
	callq	printf
	movl	592(%rsp), %esi
	movl	$.L.str.43, %edi
	xorl	%eax, %eax
	callq	printf
	movl	596(%rsp), %esi
	movl	$.L.str.43, %edi
	xorl	%eax, %eax
	callq	printf
	movl	532(%rsp), %esi
	movl	$.L.str.44, %edi
	xorl	%eax, %eax
	callq	printf
	movl	536(%rsp), %esi
	movl	$.L.str.44, %edi
	xorl	%eax, %eax
	callq	printf
	movl	540(%rsp), %esi
	movl	$.L.str.44, %edi
	xorl	%eax, %eax
	callq	printf
	movl	524(%rsp), %esi
	movl	$.L.str.45, %edi
	xorl	%eax, %eax
	callq	printf
	movl	528(%rsp), %esi
	movl	$.L.str.45, %edi
	xorl	%eax, %eax
	callq	printf
	movl	552(%rsp), %esi
	movl	$.L.str.46, %edi
	xorl	%eax, %eax
	callq	printf
	movl	556(%rsp), %esi
	movl	$.L.str.46, %edi
	xorl	%eax, %eax
	callq	printf
	movl	560(%rsp), %esi
	movl	$.L.str.46, %edi
	xorl	%eax, %eax
	callq	printf
	movl	564(%rsp), %esi
	movl	$.L.str.47, %edi
	xorl	%eax, %eax
	callq	printf
	movl	568(%rsp), %esi
	movl	$.L.str.47, %edi
	xorl	%eax, %eax
	callq	printf
	movl	572(%rsp), %esi
	movl	$.L.str.47, %edi
	xorl	%eax, %eax
	callq	printf
	movl	576(%rsp), %esi
	movl	$.L.str.48, %edi
	xorl	%eax, %eax
	callq	printf
	movl	600(%rsp), %esi
	movl	$.L.str.49, %edi
	xorl	%eax, %eax
	callq	printf
	movl	604(%rsp), %esi
	movl	$.L.str.49, %edi
	xorl	%eax, %eax
	callq	printf
	movl	420(%rsp), %esi
	movl	$.L.str.50, %edi
	xorl	%eax, %eax
	callq	printf
	movl	424(%rsp), %esi
	movl	$.L.str.50, %edi
	xorl	%eax, %eax
	callq	printf
	movl	428(%rsp), %esi
	movl	$.L.str.50, %edi
	xorl	%eax, %eax
	callq	printf
	movl	416(%rsp), %esi
	movl	$.L.str.51, %edi
	xorl	%eax, %eax
	callq	printf
	movl	720(%rsp), %esi
	movl	$.L.str.52, %edi
	xorl	%eax, %eax
	callq	printf
	movq	408(%rsp), %rsi
	movl	$.L.str.53, %edi
	xorl	%eax, %eax
	callq	printf
	movl	708(%rsp), %esi
	movl	$.L.str.54, %edi
	xorl	%eax, %eax
	callq	printf
	movl	704(%rsp), %esi
	movl	$.L.str.55, %edi
	xorl	%eax, %eax
	callq	printf
	movl	460(%rsp), %esi
	movl	$.L.str.56, %edi
	xorl	%eax, %eax
	callq	printf
	movl	484(%rsp), %esi
	movl	$.L.str.57, %edi
	xorl	%eax, %eax
	callq	printf
	movl	680(%rsp), %esi
	movl	$.L.str.58, %edi
	xorl	%eax, %eax
	callq	printf
	movl	684(%rsp), %esi
	movl	$.L.str.59, %edi
	xorl	%eax, %eax
	callq	printf
	movl	688(%rsp), %esi
	movl	$.L.str.60, %edi
	xorl	%eax, %eax
	callq	printf
	movl	400(%rsp), %esi
	movl	$.L.str.61, %edi
	xorl	%eax, %eax
	callq	printf
	movq	392(%rsp), %rsi
	movl	$.L.str.62, %edi
	xorl	%eax, %eax
	callq	printf
	movl	724(%rsp), %esi
	movl	$.L.str.63, %edi
	xorl	%eax, %eax
	callq	printf
	movq	664(%rsp), %rsi
	movl	$.L.str.64, %edi
	xorl	%eax, %eax
	callq	printf
	movl	692(%rsp), %esi
	movl	$.L.str.65, %edi
	xorl	%eax, %eax
	callq	printf
	movq	464(%rsp), %rsi
	movl	$.L.str.66, %edi
	xorl	%eax, %eax
	callq	printf
	movq	472(%rsp), %rsi
	movl	$.L.str.67, %edi
	xorl	%eax, %eax
	callq	printf
	movq	448(%rsp), %rsi
	movl	$.L.str.68, %edi
	xorl	%eax, %eax
	callq	printf
	movq	384(%rsp), %rsi
	movl	$.L.str.69, %edi
	xorl	%eax, %eax
	callq	printf
	movl	700(%rsp), %esi
	movl	$.L.str.70, %edi
	xorl	%eax, %eax
	callq	printf
	movl	404(%rsp), %esi
	movl	$.L.str.71, %edi
	xorl	%eax, %eax
	callq	printf
	movq	stdout(%rip), %rdi
	callq	fflush
	incl	%ebx
	cmpl	68(%rsp), %ebx
	jl	.LBB5_49
.LBB5_50:                               # %._crit_edge232
	movl	64(%rsp), %edi                  # 4-byte Reload
	movl	60(%rsp), %esi                  # 4-byte Reload
	movq	16(%rsp), %r13                  # 8-byte Reload
	movq	%r13, %rdx
	movq	48(%rsp), %rcx                  # 8-byte Reload
	movl	%ebp, %r8d
	callq	_Z10host_floydiiPlPii
	callq	hipDeviceReset
	movl	$.L.str.72, %esi
	movq	80(%rsp), %r14                  # 8-byte Reload
	movq	%r14, %rdi
	callq	fopen
	testq	%rax, %rax
	je	.LBB5_24
# %bb.51:                               # %.preheader190
	movq	%rax, %rbx
	testl	%ebp, %ebp
	je	.LBB5_59
# %bb.52:                               # %.preheader189.lr.ph
	leal	-1(%rbp), %r14d
	movl	%ebp, %r15d
	movl	$.L.str.3, %r12d
	xorl	%eax, %eax
	xorl	%ecx, %ecx
.LBB5_53:                               # %.preheader189
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB5_54 Depth 2
	movq	%rcx, 24(%rsp)                  # 8-byte Spill
	movl	%eax, 8(%rsp)                   # 4-byte Spill
	movl	%eax, %eax
	leaq	(,%rax,8), %rbp
	addq	%r13, %rbp
	xorl	%r13d, %r13d
.LBB5_54:                               #   Parent Loop BB5_53 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movq	(%rbp,%r13,8), %rdx
	cmpq	%r13, %r14
	movl	$.L.str.4, %ecx
	cmoveq	%r12, %rcx
	movabsq	$9223372036854775807, %rax      # imm = 0x7FFFFFFFFFFFFFFF
	cmpq	%rax, %rdx
	jne	.LBB5_55
# %bb.56:                               #   in Loop: Header=BB5_54 Depth=2
	movq	%rcx, %rdi
	movq	%rbx, %rsi
	callq	fputs@PLT
	jmp	.LBB5_57
.LBB5_55:                               #   in Loop: Header=BB5_54 Depth=2
	movl	$.L.str.2, %esi
	movq	%rbx, %rdi
	xorl	%eax, %eax
	callq	fprintf
.LBB5_57:                               #   in Loop: Header=BB5_54 Depth=2
	incq	%r13
	cmpq	%r13, %r15
	jne	.LBB5_54
# %bb.58:                               #   in Loop: Header=BB5_53 Depth=1
	movq	24(%rsp), %rcx                  # 8-byte Reload
	incq	%rcx
	movq	40(%rsp), %rbp                  # 8-byte Reload
	movl	8(%rsp), %eax                   # 4-byte Reload
	addl	%ebp, %eax
	cmpq	%r15, %rcx
	movq	16(%rsp), %r13                  # 8-byte Reload
	jne	.LBB5_53
.LBB5_59:                               # %._crit_edge235
	movq	%rbx, %rdi
	callq	fclose
	movl	$.L.str.72, %esi
	movq	72(%rsp), %r14                  # 8-byte Reload
	movq	%r14, %rdi
	callq	fopen
	testq	%rax, %rax
	jne	.LBB5_60
.LBB5_24:
	movq	stderr(%rip), %rdi
	movl	$.L.str.14, %esi
	movq	%r14, %rdx
	xorl	%eax, %eax
	callq	fprintf
	movq	stderr(%rip), %rdi
.LBB5_25:
	callq	fflush
	movl	$-1, %edi
	callq	exit
.LBB5_60:                               # %.preheader188
	movq	%rax, %rbx
	testl	%ebp, %ebp
	je	.LBB5_68
# %bb.61:                               # %.preheader.lr.ph
	leal	-1(%rbp), %r14d
	movl	%ebp, %r15d
	movl	$.L.str.3, %r12d
	xorl	%eax, %eax
	xorl	%ecx, %ecx
.LBB5_62:                               # %.preheader
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB5_63 Depth 2
	movq	%rcx, 24(%rsp)                  # 8-byte Spill
	movl	%eax, 8(%rsp)                   # 4-byte Spill
	movl	%eax, %eax
	movq	48(%rsp), %rcx                  # 8-byte Reload
	leaq	(%rcx,%rax,4), %rbp
	xorl	%r13d, %r13d
.LBB5_63:                               #   Parent Loop BB5_62 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movl	(%rbp,%r13,4), %edx
	cmpq	%r13, %r14
	movl	$.L.str.4, %ecx
	cmoveq	%r12, %rcx
	testl	%edx, %edx
	js	.LBB5_65
# %bb.64:                               #   in Loop: Header=BB5_63 Depth=2
	movl	$.L.str.74, %esi
	movq	%rbx, %rdi
	xorl	%eax, %eax
	callq	fprintf
	jmp	.LBB5_66
.LBB5_65:                               #   in Loop: Header=BB5_63 Depth=2
	movq	%rcx, %rdi
	movq	%rbx, %rsi
	callq	fputs@PLT
.LBB5_66:                               #   in Loop: Header=BB5_63 Depth=2
	incq	%r13
	cmpq	%r13, %r15
	jne	.LBB5_63
# %bb.67:                               #   in Loop: Header=BB5_62 Depth=1
	movq	24(%rsp), %rcx                  # 8-byte Reload
	incq	%rcx
	movl	8(%rsp), %eax                   # 4-byte Reload
	addl	40(%rsp), %eax                  # 4-byte Folded Reload
	cmpq	%r15, %rcx
	jne	.LBB5_62
.LBB5_68:                               # %._crit_edge238
	movq	%rbx, %rdi
	callq	fclose
	movq	16(%rsp), %rdi                  # 8-byte Reload
	callq	free
	movq	48(%rsp), %rdi                  # 8-byte Reload
	callq	free
	xorl	%edi, %edi
	callq	exit
.Lfunc_end5:
	.size	main, .Lfunc_end5-main
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_ctor
	.type	__hip_module_ctor,@function
__hip_module_ctor:                      # @__hip_module_ctor
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	subq	$32, %rsp
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -16
	movq	__hip_gpubin_handle(%rip), %rbx
	testq	%rbx, %rbx
	jne	.LBB6_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB6_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z11global_initPlS_Piii, %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	$_Z12global_floydPlS_Piii, %esi
	movl	$.L__unnamed_2, %edx
	movl	$.L__unnamed_2, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end6:
	.size	__hip_module_ctor, .Lfunc_end6-__hip_module_ctor
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_dtor
	.type	__hip_module_dtor,@function
__hip_module_dtor:                      # @__hip_module_dtor
	.cfi_startproc
# %bb.0:
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	je	.LBB7_2
# %bb.1:
	pushq	%rax
	.cfi_def_cfa_offset 16
	callq	__hipUnregisterFatBinary
	movq	$0, __hip_gpubin_handle(%rip)
	addq	$8, %rsp
	.cfi_def_cfa_offset 8
.LBB7_2:
	retq
.Lfunc_end7:
	.size	__hip_module_dtor, .Lfunc_end7-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"Floyd's algorithm"
	.size	.L.str, 18

	.type	title,@object                   # @title
	.data
	.globl	title
	.p2align	3, 0x0
title:
	.quad	.L.str
	.size	title, 8

	.type	.L.str.1,@object                # @.str.1
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str.1:
	.asciz	"\320\220\320\273\320\263\320\276\321\200\320\270\321\202\320\274 \320\244\320\273\320\276\320\271\320\264\320\260 - \320\277\320\276\320\270\321\201\320\272 \320\262\321\201\320\265\321\205 \320\272\321\200\320\260\321\202\321\207\320\260\320\271\321\210\320\270\321\205 \320\277\321\203\321\202\320\265\320\271 \320\262 \320\263\321\200\320\260\321\204\320\265"
	.size	.L.str.1, 98

	.type	description,@object             # @description
	.data
	.globl	description
	.p2align	3, 0x0
description:
	.quad	.L.str.1
	.size	description, 8

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

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

	.type	.L.str.2,@object                # @.str.2
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str.2:
	.asciz	"%ld%s"
	.size	.L.str.2, 6

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	"\n"
	.size	.L.str.3, 2

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

	.type	.L.str.5,@object                # @.str.5
.L.str.5:
	.asciz	"Title :\t%s\n"
	.size	.L.str.5, 12

	.type	.L.str.6,@object                # @.str.6
.L.str.6:
	.asciz	"Usage :\t%s [-g <gridSize>] [-b <blockSize>] <inputfilename> <outputfilename> <intermedianfilename>\n"
	.size	.L.str.6, 100

	.type	.L.str.10,@object               # @.str.10
.L.str.10:
	.asciz	"Input File Name :\t%s\n"
	.size	.L.str.10, 22

	.type	.L.str.11,@object               # @.str.11
.L.str.11:
	.asciz	"Output File Name :\t%s\n"
	.size	.L.str.11, 23

	.type	.L.str.12,@object               # @.str.12
.L.str.12:
	.asciz	"Intermedian File Name :\t%s\n"
	.size	.L.str.12, 28

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

	.type	.L.str.14,@object               # @.str.14
.L.str.14:
	.asciz	"File open error (%s)\n"
	.size	.L.str.14, 22

	.type	.L.str.15,@object               # @.str.15
.L.str.15:
	.asciz	"Matrix rank :\t%d\n"
	.size	.L.str.15, 18

	.type	.L.str.16,@object               # @.str.16
.L.str.16:
	.asciz	"\t"
	.size	.L.str.16, 2

	.type	.L.str.17,@object               # @.str.17
.L.str.17:
	.asciz	"Running on GPU %d (%s)\n"
	.size	.L.str.17, 24

	.type	.L.str.18,@object               # @.str.18
.L.str.18:
	.asciz	"Device has ECC support enabled %d\n"
	.size	.L.str.18, 35

	.type	.L.str.19,@object               # @.str.19
.L.str.19:
	.asciz	"Number of asynchronous engines %d\n"
	.size	.L.str.19, 35

	.type	.L.str.20,@object               # @.str.20
.L.str.20:
	.asciz	"Device can map host memory with cudaHostAlloc/cudaHostGetDevicePointer %d\n"
	.size	.L.str.20, 75

	.type	.L.str.21,@object               # @.str.21
.L.str.21:
	.asciz	"Clock frequency in kilohertz %d\n"
	.size	.L.str.21, 33

	.type	.L.str.22,@object               # @.str.22
.L.str.22:
	.asciz	"Compute mode (See cudaComputeMode) %d\n"
	.size	.L.str.22, 39

	.type	.L.str.23,@object               # @.str.23
.L.str.23:
	.asciz	"Device can possibly execute multiple kernels concurrently %d\n"
	.size	.L.str.23, 62

	.type	.L.str.24,@object               # @.str.24
.L.str.24:
	.asciz	"Device can concurrently copy memory and execute a kernel. Deprecated. Use instead asyncEngineCount. %d\n"
	.size	.L.str.24, 104

	.type	.L.str.25,@object               # @.str.25
.L.str.25:
	.asciz	"Device is integrated as opposed to discrete %d\n"
	.size	.L.str.25, 48

	.type	.L.str.26,@object               # @.str.26
.L.str.26:
	.asciz	"Specified whether there is a run time limit on kernels %d\n"
	.size	.L.str.26, 59

	.type	.L.str.27,@object               # @.str.27
.L.str.27:
	.asciz	"Size of L2 cache in bytes %d\n"
	.size	.L.str.27, 30

	.type	.L.str.28,@object               # @.str.28
.L.str.28:
	.asciz	"Major compute capability %d\n"
	.size	.L.str.28, 29

	.type	.L.str.29,@object               # @.str.29
.L.str.29:
	.asciz	"Maximum size of each dimension of a grid %d\n"
	.size	.L.str.29, 45

	.type	.L.str.30,@object               # @.str.30
.L.str.30:
	.asciz	"Maximum 1D surface size %d\n"
	.size	.L.str.30, 28

	.type	.L.str.31,@object               # @.str.31
.L.str.31:
	.asciz	"Maximum 1D layered surface dimensions %d\n"
	.size	.L.str.31, 42

	.type	.L.str.32,@object               # @.str.32
.L.str.32:
	.asciz	"Maximum 2D surface dimensions %d\n"
	.size	.L.str.32, 34

	.type	.L.str.33,@object               # @.str.33
.L.str.33:
	.asciz	"Maximum 2D layered surface dimensions %d\n"
	.size	.L.str.33, 42

	.type	.L.str.34,@object               # @.str.34
.L.str.34:
	.asciz	"Maximum 3D surface dimensions %d\n"
	.size	.L.str.34, 34

	.type	.L.str.35,@object               # @.str.35
.L.str.35:
	.asciz	"Maximum Cubemap surface dimensions %d\n"
	.size	.L.str.35, 39

	.type	.L.str.36,@object               # @.str.36
.L.str.36:
	.asciz	"Maximum Cubemap layered surface dimensions %d\n"
	.size	.L.str.36, 47

	.type	.L.str.37,@object               # @.str.37
.L.str.37:
	.asciz	"Maximum 1D texture size %d\n"
	.size	.L.str.37, 28

	.type	.L.str.38,@object               # @.str.38
.L.str.38:
	.asciz	"Maximum 1D layered texture dimensions %d\n"
	.size	.L.str.38, 42

	.type	.L.str.39,@object               # @.str.39
.L.str.39:
	.asciz	"Maximum size for 1D textures bound to linear memory %d\n"
	.size	.L.str.39, 56

	.type	.L.str.40,@object               # @.str.40
.L.str.40:
	.asciz	"Maximum 1D mipmapped texture size %d\n"
	.size	.L.str.40, 38

	.type	.L.str.41,@object               # @.str.41
.L.str.41:
	.asciz	"Maximum 2D texture dimensions %d\n"
	.size	.L.str.41, 34

	.type	.L.str.42,@object               # @.str.42
.L.str.42:
	.asciz	"Maximum 2D texture dimensions if texture gather operations have to be performed %d\n"
	.size	.L.str.42, 84

	.type	.L.str.43,@object               # @.str.43
.L.str.43:
	.asciz	"Maximum 2D layered texture dimensions %d\n"
	.size	.L.str.43, 42

	.type	.L.str.44,@object               # @.str.44
.L.str.44:
	.asciz	"Maximum dimensions (width, height, pitch) for 2D textures bound to pitched memory %d\n"
	.size	.L.str.44, 86

	.type	.L.str.45,@object               # @.str.45
.L.str.45:
	.asciz	"Maximum 2D mipmapped texture dimensions %d\n"
	.size	.L.str.45, 44

	.type	.L.str.46,@object               # @.str.46
.L.str.46:
	.asciz	"Maximum 3D texture dimensions %d\n"
	.size	.L.str.46, 34

	.type	.L.str.47,@object               # @.str.47
.L.str.47:
	.asciz	"Maximum alternate 3D texture dimensions %d\n"
	.size	.L.str.47, 44

	.type	.L.str.48,@object               # @.str.48
.L.str.48:
	.asciz	"Maximum Cubemap texture dimensions %d\n"
	.size	.L.str.48, 39

	.type	.L.str.49,@object               # @.str.49
.L.str.49:
	.asciz	"Maximum Cubemap layered texture dimensions %d\n"
	.size	.L.str.49, 47

	.type	.L.str.50,@object               # @.str.50
.L.str.50:
	.asciz	"Maximum size of each dimension of a block %d\n"
	.size	.L.str.50, 46

	.type	.L.str.51,@object               # @.str.51
.L.str.51:
	.asciz	"Maximum number of threads per block %d\n"
	.size	.L.str.51, 40

	.type	.L.str.52,@object               # @.str.52
.L.str.52:
	.asciz	"Maximum resident threads per multiprocessor %d\n"
	.size	.L.str.52, 48

	.type	.L.str.53,@object               # @.str.53
.L.str.53:
	.asciz	"Maximum pitch in bytes allowed by memory copies %d\n"
	.size	.L.str.53, 52

	.type	.L.str.54,@object               # @.str.54
.L.str.54:
	.asciz	"Global memory bus width in bits %d\n"
	.size	.L.str.54, 36

	.type	.L.str.55,@object               # @.str.55
.L.str.55:
	.asciz	"Peak memory clock frequency in kilohertz %d\n"
	.size	.L.str.55, 45

	.type	.L.str.56,@object               # @.str.56
.L.str.56:
	.asciz	"Minor compute capability %d\n"
	.size	.L.str.56, 29

	.type	.L.str.57,@object               # @.str.57
.L.str.57:
	.asciz	"Number of multiprocessors on device %d\n"
	.size	.L.str.57, 40

	.type	.L.str.58,@object               # @.str.58
.L.str.58:
	.asciz	"PCI bus ID of the device %d\n"
	.size	.L.str.58, 29

	.type	.L.str.59,@object               # @.str.59
.L.str.59:
	.asciz	"PCI device ID of the device %d\n"
	.size	.L.str.59, 32

	.type	.L.str.60,@object               # @.str.60
.L.str.60:
	.asciz	"PCI domain ID of the device %d\n"
	.size	.L.str.60, 32

	.type	.L.str.61,@object               # @.str.61
.L.str.61:
	.asciz	"32-bit registers available per block %d\n"
	.size	.L.str.61, 41

	.type	.L.str.62,@object               # @.str.62
.L.str.62:
	.asciz	"Shared memory available per block in bytes %d\n"
	.size	.L.str.62, 47

	.type	.L.str.63,@object               # @.str.63
.L.str.63:
	.asciz	"Device supports stream priorities %d\n"
	.size	.L.str.63, 38

	.type	.L.str.64,@object               # @.str.64
.L.str.64:
	.asciz	"Alignment requirements for surfaces %d\n"
	.size	.L.str.64, 40

	.type	.L.str.65,@object               # @.str.65
.L.str.65:
	.asciz	"1 if device is a Tesla device using TCC driver, 0 otherwise %d\n"
	.size	.L.str.65, 64

	.type	.L.str.66,@object               # @.str.66
.L.str.66:
	.asciz	"Alignment requirement for textures %d\n"
	.size	.L.str.66, 39

	.type	.L.str.67,@object               # @.str.67
.L.str.67:
	.asciz	"Pitch alignment requirement for texture references bound to pitched memory %d\n"
	.size	.L.str.67, 79

	.type	.L.str.68,@object               # @.str.68
.L.str.68:
	.asciz	"Constant memory available on device in bytes %d\n"
	.size	.L.str.68, 49

	.type	.L.str.69,@object               # @.str.69
.L.str.69:
	.asciz	"Global memory available on device in bytes %d\n"
	.size	.L.str.69, 47

	.type	.L.str.70,@object               # @.str.70
.L.str.70:
	.asciz	"Device shares a unified address space with the host %d\n"
	.size	.L.str.70, 56

	.type	.L.str.71,@object               # @.str.71
.L.str.71:
	.asciz	"Warp size in threads %d\n"
	.size	.L.str.71, 25

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

	.type	.L.str.74,@object               # @.str.74
.L.str.74:
	.asciz	"%d%s"
	.size	.L.str.74, 5

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

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"_Z12global_floydPlS_Piii"
	.size	.L__unnamed_2, 25

	.type	__hip_fatbin_wrapper,@object    # @__hip_fatbin_wrapper
	.section	.hipFatBinSegment,"a",@progbits
	.p2align	3, 0x0
__hip_fatbin_wrapper:
	.long	1212764230                      # 0x48495046
	.long	1                               # 0x1
	.quad	__hip_fatbin
	.quad	0
	.size	__hip_fatbin_wrapper, 24

	.type	__hip_gpubin_handle,@object     # @__hip_gpubin_handle
	.local	__hip_gpubin_handle
	.comm	__hip_gpubin_handle,8,8
	.section	.init_array,"aw",@init_array
	.p2align	3, 0x0
	.quad	__hip_module_ctor
	.type	__hip_cuid_,@object             # @__hip_cuid_
	.bss
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               # 0x0
	.size	__hip_cuid_, 1

	.type	.Lstr,@object                   # @str
	.section	.rodata.str1.1,"aMS",@progbits,1
.Lstr:
	.asciz	"\tinputfilename - source matrix of path prices or empty"
	.size	.Lstr, 55

	.type	.Lstr.1,@object                 # @str.1
.Lstr.1:
	.asciz	"\toutputfilename - output floyd's matrix of path prices"
	.size	.Lstr.1, 55

	.type	.Lstr.2,@object                 # @str.2
.Lstr.2:
	.asciz	"\tintermedianfilename - output matrix of intermedian points or empty"
	.size	.Lstr.2, 68

	.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 _Z26__device_stub__global_initPlS_Piii
	.addrsig_sym _Z27__device_stub__global_floydPlS_Piii
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z11global_initPlS_Piii
	.addrsig_sym _Z12global_floydPlS_Piii
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z11global_initPlS_Piii ; -- Begin function _Z11global_initPlS_Piii
	.globl	_Z11global_initPlS_Piii
	.p2align	8
	.type	_Z11global_initPlS_Piii,@function
_Z11global_initPlS_Piii:                ; @_Z11global_initPlS_Piii
; %bb.0:
	s_clause 0x1
	s_load_b32 s4, s[0:1], 0x2c
	s_load_b32 s6, s[0:1], 0x1c
	s_add_u32 s2, s0, 32
	s_addc_u32 s3, s1, 0
	s_waitcnt lgkmcnt(0)
	s_and_b32 s5, s4, 0xffff
	s_mul_i32 s4, s6, s6
	v_mad_u64_u32 v[1:2], null, s15, s5, v[0:1]
	s_mov_b32 s6, exec_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_gt_i32_e64 s4, v1
	s_cbranch_execz .LBB0_3
; %bb.1:                                ; %.lr.ph
	s_load_b32 s6, s[2:3], 0x0
	s_load_b64 s[2:3], s[0:1], 0x10
	v_mov_b32_e32 v0, -1
	s_waitcnt lgkmcnt(0)
	s_mul_i32 s1, s6, s5
	s_mov_b32 s5, 0
.LBB0_2:                                ; =>This Inner Loop Header: Depth=1
	v_ashrrev_i32_e32 v2, 31, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[2:3], 2, v[1:2]
	v_add_nc_u32_e32 v1, s1, v1
	v_cmp_le_i32_e32 vcc_lo, s4, v1
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_co_u32 v2, s0, s2, v2
	v_add_co_ci_u32_e64 v3, s0, s3, v3, s0
	s_or_b32 s5, vcc_lo, s5
	global_store_b32 v[2:3], v0, off
	s_and_not1_b32 exec_lo, exec_lo, s5
	s_cbranch_execnz .LBB0_2
.LBB0_3:                                ; %Flow11
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z11global_initPlS_Piii
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 288
		.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	_Z11global_initPlS_Piii, .Lfunc_end0-_Z11global_initPlS_Piii
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 180
; 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
	.text
	.protected	_Z12global_floydPlS_Piii ; -- Begin function _Z12global_floydPlS_Piii
	.globl	_Z12global_floydPlS_Piii
	.p2align	8
	.type	_Z12global_floydPlS_Piii,@function
_Z12global_floydPlS_Piii:               ; @_Z12global_floydPlS_Piii
; %bb.0:
	s_clause 0x1
	s_load_b32 s6, s[0:1], 0x2c
	s_load_b64 s[2:3], s[0:1], 0x18
	s_add_u32 s4, s0, 32
	s_addc_u32 s5, s1, 0
	s_waitcnt lgkmcnt(0)
	s_and_b32 s16, s6, 0xffff
	s_mul_i32 s12, s3, s3
	v_mad_u64_u32 v[1:2], null, s15, s16, v[0:1]
	s_mov_b32 s6, exec_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_gt_i32_e64 s12, v1
	s_cbranch_execz .LBB1_13
; %bb.1:                                ; %.lr.ph
	s_ashr_i32 s13, s3, 31
	s_load_b32 s15, s[4:5], 0x0
	s_add_i32 s6, s3, s13
	s_mov_b32 s10, -1
	s_xor_b32 s14, s6, s13
	s_brev_b32 s11, -2
	v_cvt_f32_u32_e32 v0, s14
	s_sub_i32 s4, 0, s14
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_iflag_f32_e32 v0, v0
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v0, 0x4f7ffffe, v0
	v_cvt_u32_f32_e32 v0, v0
	s_delay_alu instid0(VALU_DEP_1)
	v_mul_lo_u32 v2, s4, v0
	s_clause 0x1
	s_load_b128 s[4:7], s[0:1], 0x0
	s_load_b64 s[8:9], s[0:1], 0x10
	s_waitcnt lgkmcnt(0)
	s_mul_i32 s1, s15, s16
	s_mov_b32 s15, 0
	s_sub_i32 s16, 0, s3
	v_mul_hi_u32 v2, v0, v2
	s_delay_alu instid0(VALU_DEP_1)
	v_add_nc_u32_e32 v0, v0, v2
.LBB1_2:                                ; =>This Inner Loop Header: Depth=1
	v_ashrrev_i32_e32 v2, 31, v1
	s_mov_b32 s18, exec_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v3, v1, v2
	v_xor_b32_e32 v3, v3, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_hi_u32 v4, v3, v0
	v_mul_lo_u32 v5, v4, s14
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_sub_nc_u32_e32 v3, v3, v5
	v_add_nc_u32_e32 v5, 1, v4
	v_subrev_nc_u32_e32 v6, s14, v3
	v_cmp_le_u32_e32 vcc_lo, s14, v3
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_dual_cndmask_b32 v4, v4, v5 :: v_dual_cndmask_b32 v3, v3, v6
	v_xor_b32_e32 v5, s13, v2
	v_add_nc_u32_e32 v6, 1, v4
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cmp_le_u32_e32 vcc_lo, s14, v3
	v_cndmask_b32_e32 v3, v4, v6, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v6, v3, v5
	v_sub_nc_u32_e32 v3, v6, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[7:8], null, s16, v3, v[1:2]
	v_cmp_eq_u32_e64 s17, v7, v3
	v_cmpx_ne_u32_e64 v7, v3
	s_cbranch_execz .LBB1_10
; %bb.3:                                ;   in Loop: Header=BB1_2 Depth=1
	v_mad_u64_u32 v[7:8], null, v3, s3, s[2:3]
	s_mov_b32 s19, exec_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v8, 31, v7
	v_lshlrev_b64 v[3:4], 3, v[7:8]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v3, vcc_lo, s4, v3
	v_add_co_ci_u32_e32 v4, vcc_lo, s5, v4, vcc_lo
	global_load_b64 v[3:4], v[3:4], off
	s_waitcnt vmcnt(0)
	v_cmp_eq_u64_e32 vcc_lo, s[10:11], v[3:4]
	v_cmpx_ne_u64_e64 s[10:11], v[3:4]
	s_cbranch_execz .LBB1_9
; %bb.4:                                ;   in Loop: Header=BB1_2 Depth=1
	v_add_nc_u32_e32 v5, s2, v5
	s_mov_b32 s21, exec_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_nc_u32_e32 v7, v5, v6
	v_mad_u64_u32 v[5:6], null, s3, v7, v[1:2]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v6, 31, v5
	v_lshlrev_b64 v[5:6], 3, v[5:6]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_co_u32 v5, s0, s4, v5
	v_add_co_ci_u32_e64 v6, s0, s5, v6, s0
	global_load_b64 v[7:8], v[5:6], off
	s_waitcnt vmcnt(0)
	v_cmp_eq_u64_e64 s20, s[10:11], v[7:8]
	v_cmpx_ne_u64_e64 s[10:11], v[7:8]
	s_cbranch_execz .LBB1_8
; %bb.5:                                ;   in Loop: Header=BB1_2 Depth=1
	v_lshlrev_b64 v[5:6], 3, v[1:2]
	s_mov_b32 s23, -1
	s_mov_b32 s22, exec_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_co_u32 v9, s0, s4, v5
	v_add_co_ci_u32_e64 v10, s0, s5, v6, s0
	v_add_co_u32 v3, s0, v7, v3
	s_delay_alu instid0(VALU_DEP_1)
	v_add_co_ci_u32_e64 v4, s0, v8, v4, s0
	global_load_b64 v[9:10], v[9:10], off
	s_waitcnt vmcnt(0)
	v_cmpx_lt_i64_e64 v[3:4], v[9:10]
	s_cbranch_execz .LBB1_7
; %bb.6:                                ;   in Loop: Header=BB1_2 Depth=1
	v_lshlrev_b64 v[7:8], 2, v[1:2]
	v_add_co_u32 v5, s0, s6, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_add_co_ci_u32_e64 v6, s0, s7, v6, s0
	v_mov_b32_e32 v9, s2
	v_add_co_u32 v7, s0, s8, v7
	s_delay_alu instid0(VALU_DEP_1)
	v_add_co_ci_u32_e64 v8, s0, s9, v8, s0
	s_xor_b32 s23, exec_lo, -1
	global_store_b64 v[5:6], v[3:4], off
	global_store_b32 v[7:8], v9, off
.LBB1_7:                                ; %Flow75
                                        ;   in Loop: Header=BB1_2 Depth=1
	s_or_b32 exec_lo, exec_lo, s22
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	s_and_not1_b32 s0, s20, exec_lo
	s_and_b32 s20, s23, exec_lo
	s_or_b32 s20, s0, s20
.LBB1_8:                                ; %Flow74
                                        ;   in Loop: Header=BB1_2 Depth=1
	s_or_b32 exec_lo, exec_lo, s21
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	s_and_not1_b32 s0, vcc_lo, exec_lo
	s_and_b32 s20, s20, exec_lo
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 vcc_lo, s0, s20
.LBB1_9:                                ; %Flow73
                                        ;   in Loop: Header=BB1_2 Depth=1
	s_or_b32 exec_lo, exec_lo, s19
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	s_and_not1_b32 s0, s17, exec_lo
	s_and_b32 s17, vcc_lo, exec_lo
	s_or_b32 s17, s0, s17
.LBB1_10:                               ; %Flow
                                        ;   in Loop: Header=BB1_2 Depth=1
	s_or_b32 exec_lo, exec_lo, s18
	s_delay_alu instid0(VALU_DEP_2)
	s_and_saveexec_b32 s0, s17
	s_cbranch_execz .LBB1_12
; %bb.11:                               ;   in Loop: Header=BB1_2 Depth=1
	v_lshlrev_b64 v[2:3], 3, v[1:2]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v4, vcc_lo, s4, v2
	v_add_co_ci_u32_e32 v5, vcc_lo, s5, v3, vcc_lo
	v_add_co_u32 v2, vcc_lo, s6, v2
	v_add_co_ci_u32_e32 v3, vcc_lo, s7, v3, vcc_lo
	global_load_b64 v[4:5], v[4:5], off
	s_waitcnt vmcnt(0)
	global_store_b64 v[2:3], v[4:5], off
.LBB1_12:                               ;   in Loop: Header=BB1_2 Depth=1
	s_or_b32 exec_lo, exec_lo, s0
	v_add_nc_u32_e32 v1, s1, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	v_cmp_le_i32_e32 vcc_lo, s12, v1
	s_or_b32 s15, vcc_lo, s15
	s_and_not1_b32 exec_lo, exec_lo, s15
	s_cbranch_execnz .LBB1_2
.LBB1_13:                               ; %Flow77
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z12global_floydPlS_Piii
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 288
		.amdhsa_user_sgpr_count 15
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 11
		.amdhsa_next_free_sgpr 24
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end1:
	.size	_Z12global_floydPlS_Piii, .Lfunc_end1-_Z12global_floydPlS_Piii
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 824
; NumSgprs: 26
; NumVgprs: 11
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 3
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 26
; NumVGPRsForWavesPerEU: 11
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_
	.amdgpu_metadata
---
amdhsa.kernels:
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .offset:         24
        .size:           4
        .value_kind:     by_value
      - .offset:         28
        .size:           4
        .value_kind:     by_value
      - .offset:         32
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         36
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         40
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         44
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         46
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         48
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         50
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         52
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         54
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         88
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         96
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 288
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z11global_initPlS_Piii
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z11global_initPlS_Piii.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
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .offset:         24
        .size:           4
        .value_kind:     by_value
      - .offset:         28
        .size:           4
        .value_kind:     by_value
      - .offset:         32
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         36
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         40
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         44
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         46
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         48
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         50
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         52
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         54
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         88
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         96
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 288
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z12global_floydPlS_Piii
    .private_segment_fixed_size: 0
    .sgpr_count:     26
    .sgpr_spill_count: 0
    .symbol:         _Z12global_floydPlS_Piii.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     11
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
