#include <stdio.h>
#include <stdlib.h>
#include <math.h>
#include <time.h>
 
#define DEFAULT_WIDTH		1024
#define DEFAULT_HEIGHT		1024
#define DEFAULT_INTWNDW		32
#define DEFAULT_INTWNDH		32
#define DEFAULT_SRHWNDL		12
#define DEFAULT_SRHWNDR		12
#define DEFAULT_SRHWNDA		12
#define DEFAULT_SRHWNDB		12
 
 
__host__	void getData(int size, unsigned char *data0, unsigned char *data1, char *path0, char *path1);
 
__global__	void vectorComputation(unsigned char *f, unsigned char *g, int width, int totalVectorsX, int totalVectorsY, char m, char n, char l, char r, char a, char b, char ww, char wh, char fw, char fh, char stepX, char stepY, float *R_posSubX, float *R_posSubY, float *startX, float *startY, float *endX, float *endY);
 
__device__	void getData0intoSharedMemory(unsigned char *f, unsigned char *sh_f, int width, char m, char n, char l, char a, int startPixelX, int startPixelY, char stepX, char stepY);
 
__device__	void getData1intoSharedMemory(unsigned char *g, unsigned char *sh_g, int width, char fw, char fh, int startPixelX, int startPixelY, char stepX, char stepY);
 
__device__	void calculateR(unsigned char *sh_f, unsigned char *sh_g, float *R, char *R_posX, char *R_posY, char m, char n, char ww, char wh, char fw, char fh);
 
__device__	void bitonicSort(float *R, char *R_posX, char *R_posY);
 
__device__	void swapFloat(float *x, float *y);
 
__device__	void swapChar(char *x, char *y);
 
__device__	void subPixelAccuracy(int vectorID, float *R, char *R_posX, char *R_posY, float *subXN, float *subXD, float *subYN, float *subYD, float *R_posSubX, float *R_posSubY, char l, char a, float *endX, float *endY);
 
__global__ void vectorCorrection(int totalVectorsX, float *R_posSubX, float *R_posSubY, float *endX, float *endY);
 
__host__	void printSummary(int width, int height, char intWndW, char intWndH, char srhWndL, char srhWndR, char srhWndA, char srhWndB, char stepX, char stepY, int totalVectorsX, int totalVectorsY, int totalVectors, int Dbx, int Dby, int Dgx, int Dgy, float h_time, float d_time);
 
__host__	void putData(float *startX, float *startY, float *endX, float *endY, int totalVectors, char *path);
 
__host__	void putDataVTK(int height, float *startX, float *startY, float *endX, float *endY, int number, int numberX, int numberY, char *pathVTK);
 
int main(int argc, char *argv[]) {
	int width = DEFAULT_WIDTH, height = DEFAULT_HEIGHT;
	int size = width * height;
 
	unsigned char *h_data0, *h_data1, *d_data0, *d_data1;
	char intWndW = (char)DEFAULT_INTWNDW,
		 intWndH = (char)DEFAULT_INTWNDH,
		 srhWndL = (char)DEFAULT_SRHWNDL,
		 srhWndR = (char)DEFAULT_SRHWNDR,
		 srhWndA = (char)DEFAULT_SRHWNDA,
		 srhWndB = (char)DEFAULT_SRHWNDB,
		 srhWndW = srhWndL + srhWndR + 1,
		 srhWndH = srhWndA + srhWndB + 1,
		 srhFldW = srhWndL + srhWndR + intWndW,
		 srhFldH = srhWndL + srhWndR + intWndH;
 
	char stepX, stepY;
	float *h_startX, *h_startY, *d_startX, *d_startY;
	float *h_endX, *h_endY, *d_endX, *d_endY;
	float *R_posSubX, *R_posSubY;
	int totalVectorsX, totalVectorsY, totalVectors;
 
	char *inputFileName0 = "2ms2000018",
		 *inputFileName1 = "2ms2000019",
		 *inputFileExtension = "raw",
		 *outputFileName = "result",
		 *outputFileExtension = "dat";
 
	char inputFilepath0[64], inputFilepath1[64], outputFilepath[64], outputFilePathVTK[64];
 
	clock_t h_start, h_end, d_start, d_end;
	float h_time, d_time;
 
	cudaError_t cudaError;
 
	dim3 Dg, Db;
	int Dgx, Dgy, Dbx, Dby;
 
	h_data0 = (unsigned char *)malloc(sizeof(unsigned char) * size);
	h_data1 = (unsigned char *)malloc(sizeof(unsigned char) * size);
	cudaMalloc((void **)&d_data0, sizeof(unsigned char) * size);
	cudaMalloc((void **)&d_data1, sizeof(unsigned char) * size);
 
	stepX = (char)(intWndW / 2);
	stepY = (char)(intWndH / 2);
 
	totalVectorsX = (int)((width - srhFldW) / stepX) + 1;
	totalVectorsY = (int)((height - srhFldH) / stepY) + 1;
	totalVectors = totalVectorsX * totalVectorsY;
 
	h_startX = (float *)malloc(sizeof(float) * totalVectors);
	h_startY = (float *)malloc(sizeof(float) * totalVectors);
	h_endX = (float *)malloc(sizeof(float) * totalVectors);
	h_endY = (float *)malloc(sizeof(float) * totalVectors);
	cudaMalloc((void **)&d_startX, sizeof(float) * totalVectors);
	cudaMalloc((void **)&d_startY, sizeof(float) * totalVectors);
	cudaMalloc((void **)&d_endX, sizeof(float) * totalVectors);
	cudaMalloc((void **)&d_endY, sizeof(float) * totalVectors);
 
	cudaMalloc((void **)&R_posSubX, sizeof(float) * totalVectors * 5);
	cudaMalloc((void **)&R_posSubY, sizeof(float) * totalVectors * 5);
 
	sprintf(inputFilepath0, "%s.%s", inputFileName0, inputFileExtension);
	sprintf(inputFilepath1, "%s.%s", inputFileName1, inputFileExtension);
	sprintf(outputFilepath, "%s.%s", outputFileName, outputFileExtension);
	sprintf(outputFilePathVTK, "%s-vtk.vtk", outputFileName);
 
	Db.x = 32;
	Db.y = 32;
	Dg.x = totalVectorsX;
	Dg.y = totalVectorsY;
	Dbx = Db.x; Dby = Db.y;
	Dgx = Dg.x; Dgy = Dg.y;
 
	getData(size, h_data0, h_data1, inputFilepath0, inputFilepath1);
 
	h_start = clock();
	cudaMemcpy(d_data0, h_data0, sizeof(unsigned char) * size, cudaMemcpyHostToDevice);
	cudaMemcpy(d_data1, h_data1, sizeof(unsigned char) * size, cudaMemcpyHostToDevice);
 
	d_start = clock();
 
	vectorComputation<<<Dg, Db>>>(d_data0, d_data1, width, totalVectorsX, totalVectorsY, intWndW, intWndH, srhWndL, srhWndR, srhWndA, srhWndB, srhWndW, srhWndH, srhFldW, srhFldH, stepX, stepY, R_posSubX, R_posSubY, d_startX, d_startY, d_endX, d_endY);
	cudaThreadSynchronize();
 
	vectorCorrection<<<Dg, Db>>>(totalVectorsX, R_posSubX, R_posSubY, d_endX, d_endY);
	cudaThreadSynchronize();
 
	d_end = clock();
 
	cudaMemcpy(h_startX, d_startX, sizeof(float) * totalVectors, cudaMemcpyDeviceToHost);
	cudaMemcpy(h_startY, d_startY, sizeof(float) * totalVectors, cudaMemcpyDeviceToHost);
	cudaMemcpy(h_endX, d_endX, sizeof(float) * totalVectors, cudaMemcpyDeviceToHost);
	cudaMemcpy(h_endY, d_endY, sizeof(float) * totalVectors, cudaMemcpyDeviceToHost);	
	h_end = clock();
 
	h_time = (float)(h_end - h_start) / CLOCKS_PER_SEC;
	d_time = (float)(d_end - d_start) / CLOCKS_PER_SEC;
	printSummary(width, height, intWndW, intWndH, srhWndL, srhWndR, srhWndA, srhWndB, stepX, stepY, totalVectorsX, totalVectorsY, totalVectors, Dbx, Dby, Dgx, Dgy, h_time, d_time);
 
	putData(h_startX, h_startY, h_endX, h_endY, totalVectors, outputFilepath);
	putDataVTK(height, h_startX, h_startY, h_endX, h_endY, totalVectors, totalVectorsX, totalVectorsY, outputFilePathVTK);
 
	free(h_data0);
	free(h_data1);
	cudaFree(d_data0);
	cudaFree(d_data1);
 
	free(h_startX);
	free(h_startY);
	free(h_endX);
	free(h_endY);
	cudaFree(d_startX);
	cudaFree(d_startY);
	cudaFree(d_endX);
	cudaFree(d_endY);
 
	cudaFree(R_posSubX);
	cudaFree(R_posSubY);
 
	cudaError = cudaGetLastError();
 
	if(cudaError != cudaSuccess) {
		// print the CUDA error message and exit
		printf("CUDA error: %s\n\n", cudaGetErrorString(cudaError));
		exit(-1);
	} else {
		printf("CUDA success\n\n");
	}
 
	return 0;
}
 
__host__		void getData(int size, unsigned char *data0, unsigned char *data1, char *path0, char *path1) {
	int i;
	FILE *fp;
 
	fp = fopen(path0, "rb");
	if (fp == NULL) {
		printf("%s file open error\n", path0);
	}
 
	for (i = 0; i < size; i++) {
		data0[i] = fgetc(fp);
	}
 
	fclose(fp);
 
	fp = fopen(path1, "rb");
	if (fp == NULL) {
		printf("%s file open error\n", path1);
	}
 
	for (i = 0; i < size; i++) {
		data1[i] = fgetc(fp);
	}
 
	fclose(fp);
}
 
__global__	void vectorComputation(unsigned char *f, unsigned char *g, int width, int totalVectorsX, int totalVectorsY, char m, char n, char l, char r, char a, char b, char ww, char wh, char fw, char fh, char stepX, char stepY, float *R_posSubX, float *R_posSubY, float *startX, float *startY, float *endX, float *endY) {
	int vectorID;
	int startPixelX, startPixelY;
 
	__shared__ unsigned char
		sh_f[2304], sh_g[12100];
 
	__shared__ float
		R[4096];
 
	__shared__ char
		R_posX[4096], R_posY[4096];
 
	__shared__ float
		subXN[5], subXD[5], subYN[5], subYD[5];
 
	vectorID = blockIdx.y * gridDim.x + blockIdx.x;
 
	startPixelX = stepX * blockIdx.x;
	startPixelY = stepY * blockIdx.y;
 
	startX[vectorID] = startPixelX + l + (m / 2) + (float)0.5;
	startY[vectorID] = startPixelY + a + (n / 2) + (float)0.5;
	endX[vectorID] = 0;
	endY[vectorID] = 0;
 
	getData0intoSharedMemory(f, sh_f, width, m, n, l, a, startPixelX, startPixelY, stepX, stepY);
 
	getData1intoSharedMemory(g, sh_g, width, fw, fh, startPixelX, startPixelY, stepX, stepY);
 
		__syncthreads();
 
	calculateR(sh_f, sh_g, R, R_posX, R_posY, m, n, ww, wh, fw, fh);
 
		__syncthreads();
 
	bitonicSort(R, R_posX, R_posY);
 
		__syncthreads();
 
	subPixelAccuracy(vectorID, R, R_posX, R_posY, subXN, subXD, subYN, subYD, R_posSubX, R_posSubY, l, a, endX, endY);
 
		__syncthreads();
 
}
 
__device__ void getData0intoSharedMemory(unsigned char *f, unsigned char *sh_f, int width, char m, char n, char l, char a, int startPixelX, int startPixelY, char stepX, char stepY) {
	char j, i;
 
	for (j = threadIdx.y; j < n; j += blockDim.y) {
		for (i = threadIdx.x; i < m; i += blockDim.x) {
			sh_f[j * m + i] = f[(startPixelY + a + j) * width + (startPixelX + l + i)];
		}
	}
 
}
 
__device__ void getData1intoSharedMemory(unsigned char *g, unsigned char *sh_g, int width, char fw, char fh, int startPixelX, int startPixelY, char stepX, char stepY) {
	char j, i;
 
	for (j = threadIdx.y; j < fh; j += blockDim.y) {
		for (i = threadIdx.x; i < fw; i += blockDim.x) {
			sh_g[j * fw + i] = g[(startPixelY + j) * width + (startPixelX + i)];
		}
	}
 
}
 
__device__ void calculateR(unsigned char *sh_f, unsigned char *sh_g, float *R, char *R_posX, char *R_posY, char m, char n, char ww, char wh, char fw, char fh) {
	char i, j, zeta, eta;
	float fValue, gValue;
	float A, B, C;
	int theta;
 
	for (eta = threadIdx.y; eta < wh; eta += blockDim.y) {
		for (zeta = threadIdx.x; zeta < ww; zeta += blockDim.x) {
			theta = eta * ww + zeta;
			A = 0; B = 0; C = 0;
 
			for (j = 0; j < n; j++) {
				for (i = 0; i < m; i++) {
 
					fValue = (float)sh_f[j * m + i];
					gValue = (float)sh_g[(eta + j) * fw + (zeta + i)];
 
					A += fValue * gValue;
					B += fValue * fValue;
					C += gValue * gValue;
				}
			}
			R[theta] = sqrt((A * A) / (B * C));
			R_posX[theta] = zeta;
			R_posY[theta] = eta;
		}
	}
 
	for (theta = (threadIdx.y * blockDim.x + threadIdx.x) + (ww * wh); theta < 4096; theta += blockDim.x * blockDim.y) {
		R[theta] = 0;
		R_posX[theta] = -128;
		R_posY[theta] = -128;
	}
 
}
 
__device__ void bitonicSort(float *R, char *R_posX, char *R_posY) {
 
	int i, j, k, frac, xtemp, x, group, el0, el1;
	for (i = 0; 1 << i <= 2048; i++) {
		for (j = i; j >= 0; j--) {
 
			if (j == i) {
				for (k = threadIdx.y * blockDim.x + threadIdx.x; k < 2048; k += blockDim.x * blockDim.y) {
					frac = (1 << (j + 1)) - 1;
					xtemp = k >> j;
					x = xtemp << j;
 
					group = (x << 2) + frac;
					el0 = (x << 1) + (k & ((1 << j) - 1));
					el1 = group - el0;
 
					if (R[el0] < R[el1]) {
						swapFloat(&R[el0], &R[el1]);
						swapChar(&R_posX[el0], &R_posX[el1]);
						swapChar(&R_posY[el0], &R_posY[el1]);
					}
				}
 
			} else {
				for (k = threadIdx.y * blockDim.x + threadIdx.x; k < 2048; k += blockDim.x * blockDim.y) {
					xtemp = k >> j;
					x = xtemp << j;
 
					el0 = (x << 1) + (k & ((1 << j) - 1));
					el1 = el0 + (1 << j);
 
					if (R[el0] < R[el1]) {
						swapFloat(&R[el0], &R[el1]);
						swapChar(&R_posX[el0], &R_posX[el1]);
						swapChar(&R_posY[el0], &R_posY[el1]);
					}
				}
 
			}
			__syncthreads();
 
		}
	}
 
}
 
__device__ void swapFloat(float *x, float *y) {
 
	float t;
	 t = *x;
	*x = *y;
	*y =  t;
}
 
__device__ void swapChar(char *x, char *y) {
 
	char t;
	 t = *x;
	*x = *y;
	*y =  t;
}
 
__device__	void subPixelAccuracy(int vectorID, float *R, char *R_posX, char *R_posY, float *subXN, float *subXD, float *subYN, float *subYD, float *R_posSubX, float *R_posSubY, char l, char a, float *endX, float *endY) {
 
	int i, goal;
 
	if (threadIdx.y == 0 && threadIdx.x < 5) {
		subXN[threadIdx.x] = 0;
	} else if (threadIdx.y == 1 && threadIdx.x < 5) {
		subXD[threadIdx.x] = 0;
	} else if (threadIdx.y == 2 && threadIdx.x < 5) {
		subYN[threadIdx.x] = 0;
	} else if (threadIdx.y == 3 && threadIdx.x < 5) {
		subYD[threadIdx.x] = 0;
	}
 
	__syncthreads();
 
	for (goal = 0; goal < 5; goal++) {
		for (i = threadIdx.y * blockDim.x + threadIdx.x; i < 4096; i += blockDim.x * blockDim.y) {
			if ((R_posX[i] == R_posX[goal]) && (R_posY[i] == R_posY[goal])) {
				subXD[goal] -= (R[i] + R[i]);
				subYD[goal] -= (R[i] + R[i]);
			} else if ((R_posX[i] == R_posX[goal] - 1) && (R_posY[i] == R_posY[goal])) {
				subXN[goal] -= R[i];
				subXD[goal] += R[i];
			} else if ((R_posX[i] == R_posX[goal] + 1) && (R_posY[i] == R_posY[goal])) {
				subXN[goal] += R[i];
				subXD[goal] += R[i];
			} else if ((R_posX[i] == R_posX[goal]) && (R_posY[i] == R_posY[goal] - 1)) {
				subYN[goal] -= R[i];
				subYD[goal] += R[i];
			} else if ((R_posX[i] == R_posX[goal]) && (R_posY[i] == R_posY[goal] + 1)) {
				subYN[goal] += R[i];
				subYD[goal] += R[i];
			}
		}
	}
 
	__syncthreads();
 
	if (threadIdx.y == 0 && threadIdx.x < 5) {
 
		float subX;
		subX = subXN[threadIdx.x] / (subXD[threadIdx.x] + subXD[threadIdx.x]);
		if (subX > 1 || subX < -1) {
			subX = 0;
		}
		R_posSubX[vectorID * 5 + threadIdx.x] = (float)R_posX[threadIdx.x] + subX - (float)l;
 
	} else if (threadIdx.y == 1 && threadIdx.x < 5) {
		float subY;
		subY = subYN[threadIdx.x] / (subYD[threadIdx.x] + subYD[threadIdx.x]);
		if (subY > 1 || subY < -1) {
			subY = 0;
		}
		R_posSubY[vectorID * 5 + threadIdx.x] = (float)R_posY[threadIdx.x] + subY - (float)a;
 
	}
 
	__syncthreads();
 
	if (threadIdx.y == 0 && threadIdx.x == 0) {
		endX[vectorID] = R_posSubX[vectorID * 5];
 
	} else if (threadIdx.y == 1 && threadIdx.x == 0) {
		endY[vectorID] = R_posSubY[vectorID * 5];
 
	}
	__syncthreads();
}
 
__global__	void vectorCorrection(int totalVectorsX, float *R_posSubX, float *R_posSubY, float *endX, float *endY) {
 
	if ((blockIdx.x != 0) && (blockIdx.x != gridDim.x - 1) && (blockIdx.y != 0) && (blockIdx.y != gridDim.y - 1) && (threadIdx.y == 0) && (threadIdx.x == 0)) {
		if (threadIdx.x == 0 && threadIdx.y == 0) {
			int i, j, k, vectorID;
			float averageX, averageY, dX, dY, X, Y;
 
			vectorID = blockIdx.y * gridDim.x + blockIdx.x;
			averageX = 0;
			averageY = 0;
			X = R_posSubX[vectorID * 5];
			Y = R_posSubY[vectorID * 5];
 
			for (j = -1; j <= 1; j++) {
				for (i = -1; i <= 1; i++) {
					if ((j != 0) || (i != 0)) {
						k = j * totalVectorsX + i;
						averageX += R_posSubX[(vectorID + k) * 5];
						averageY += R_posSubY[(vectorID + k) * 5];
					}
				}
			}
			averageX /= 8;
			averageY /= 8;
 
			for (i = 1; i < 5; i++) {
				dX = R_posSubX[vectorID * 5 + i] - averageX;
				dY = R_posSubY[vectorID * 5 + i] - averageY;
				if (sqrt(dX * dX + dY * dY) < sqrt((X - averageX) * (X - averageX) + (Y - averageY) * (Y - averageY))) {
					X = R_posSubX[vectorID * 5 + i];
					Y = R_posSubY[vectorID * 5 + i];
				}
 
			}
 
			if (sqrt((X - averageX) * (X - averageX) + (Y - averageY) * (Y - averageY)) > 4) {
				X = averageX;
				Y = averageY;
			}
 
			endX[vectorID] = X;
			endY[vectorID] = Y;
 
		}
 
	}
}
 
__host__		void printSummary(int width, int height, char intWndW, char intWndH, char srhWndL, char srhWndR, char srhWndA, char srhWndB, char stepX, char stepY, int totalVectorsX, int totalVectorsY, int totalVectors, int Dbx, int Dby, int Dgx, int Dgy, float h_time, float d_time) {
	int srhWndW, srhWndH, threadsX, threadsY, threads;
	float overlap;
 
	srhWndW = 1 + srhWndL + srhWndR;
	srhWndH = 1 + srhWndA + srhWndB;
	threadsX = Dgx * Dbx;
	threadsY = Dgy * Dby;
	threads = threadsX * threadsY;
	overlap = 1 - ((float)stepX / (float)intWndW);
 
	printf(
		"Summary:\n"
		"Image size:                  %d * %d\n"
		"Interrogation window wize:   %d * %d\n"
		"Search window size:          %d * %d\n"
		"\n"
		"Steps:                       %d * %d\n"
		"Overlap:                     %.2f%% \n"
		"\n"
		"Vector size:                 %d * %d\n"
		"Total vectors                %d\n"
		"\n"
		"Grid size:                   %d * %d\n"
		"Block size:                  %d * %d\n"
		"Thread size:                 %d * %d\n"
		"Total threads:               %d\n"
		"\n"
		"Processing time (with memcpy):	%f sec\n"
		"Processing time (without memcpy): %f sec\n"
		"\n"
		, width, height
		, (int)intWndW, (int)intWndH
		, (int)srhWndW, (int)srhWndH
 
		, (int)stepX, (int)stepY
		, (overlap * 100)
 
		, totalVectorsX, totalVectorsY
		, totalVectors
 
		, Dgx, Dgy
		, Dbx, Dby
		, threadsX, threadsY
		, threads
 
		, h_time
		, d_time
	);
}
 
__host__	void putData(float *startX, float *startY, float *endX, float *endY, int number, char *path) {
	int i;
 
	FILE *fp;
 
	fp = fopen(path, "w");
	if (fp == NULL) {
		fprintf(stderr, "%s file open error\n", path);
	}
 
	for (i = 0; i < number; i++) {	
 
		fprintf(fp, "%.1f	%.1f	%.1f	%.1f\n", startX[i], startY[i], endX[i] * 2, endY[i] * 2);
	}
 
	fclose(fp);
 
}
 
/*vtk output*/
__host__	void putDataVTK(int height, float *startX, float *startY, float *endX, float *endY, int number, int numberX, int numberY, char *pathVTK) {
	int i;
	FILE *fp;
 
	fp = fopen(pathVTK, "w");
	if (fp == NULL) {
		fprintf(stderr, "%s file open error.\n", pathVTK);
	}
 
	fprintf(fp, "# vtk DataFile Version 3.0\nvtk output\nASCII\nDATASET STRUCTURED_GRID\nDIMENSIONS %d %d 1\nPOINTS %d float\n", numberX, numberY, number);
	for (i = 0; i < number; i++) {
		fprintf(fp,"%.1f %.1f %d\n", startX[i], height - startY[i], 0);
	}
 
	fprintf(fp, "POINT_DATA %d\nSCALARS velocity float\nLOOKUP_TABLE default\n", number);
	for (i = 0; i < number; i++) {
		fprintf(fp,"%.10f\n", sqrt((endX[i] * endX[i]) + (endY[i] * endY[i])));
	}
 
	fprintf(fp, "VECTORS velocity float\n");
	for (i = 0; i < number; i++) {
		fprintf(fp,"%.10f %.10f %.10f\n", endX[i], -endY[i], .0);
	}
 
}
	.file	"tmpxft_00383143_00000000-6_main.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2041:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2041:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"rb"
.LC1:
	.string	"%s file open error\n"
	.text
	.globl	_Z7getDataiPhS_PcS0_
	.type	_Z7getDataiPhS_PcS0_, @function
_Z7getDataiPhS_PcS0_:
.LFB2028:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	leaq	.LC0(%rip), %r15
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	movq	%rsi, %r14
	movq	%r15, %rsi
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	movq	%rdx, %r13
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movl	%edi, %r12d
	movq	%rcx, %rdi
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movq	%rcx, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	subq	$24, %rsp
	.cfi_def_cfa_offset 80
	movq	%r8, 8(%rsp)
	call	fopen@PLT
	movq	%rax, %rbx
	testq	%rax, %rax
	jne	.L3
	movq	%rbp, %rdx
	leaq	.LC1(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
.L3:
	xorl	%ebp, %ebp
.L4:
	movq	%rbx, %rdi
	cmpl	%ebp, %r12d
	jle	.L11
	call	fgetc@PLT
	movb	%al, (%r14,%rbp)
	incq	%rbp
	jmp	.L4
.L11:
	call	fclose@PLT
	movq	8(%rsp), %rdi
	movq	%r15, %rsi
	call	fopen@PLT
	movq	%rax, %rbx
	testq	%rax, %rax
	jne	.L6
	movq	8(%rsp), %rdx
	leaq	.LC1(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
.L6:
	xorl	%ebp, %ebp
.L7:
	movq	%rbx, %rdi
	cmpl	%ebp, %r12d
	jle	.L12
	call	fgetc@PLT
	movb	%al, 0(%r13,%rbp)
	incq	%rbp
	jmp	.L7
.L12:
	addq	$24, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	jmp	fclose@PLT
	.cfi_endproc
.LFE2028:
	.size	_Z7getDataiPhS_PcS0_, .-_Z7getDataiPhS_PcS0_
	.globl	_Z24getData0intoSharedMemoryPhS_icccciicc
	.type	_Z24getData0intoSharedMemoryPhS_icccciicc, @function
_Z24getData0intoSharedMemoryPhS_icccciicc:
.LFB2029:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2029:
	.size	_Z24getData0intoSharedMemoryPhS_icccciicc, .-_Z24getData0intoSharedMemoryPhS_icccciicc
	.globl	_Z24getData1intoSharedMemoryPhS_icciicc
	.type	_Z24getData1intoSharedMemoryPhS_icciicc, @function
_Z24getData1intoSharedMemoryPhS_icciicc:
.LFB2030:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2030:
	.size	_Z24getData1intoSharedMemoryPhS_icciicc, .-_Z24getData1intoSharedMemoryPhS_icciicc
	.globl	_Z10calculateRPhS_PfPcS1_cccccc
	.type	_Z10calculateRPhS_PfPcS1_cccccc, @function
_Z10calculateRPhS_PfPcS1_cccccc:
.LFB2031:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2031:
	.size	_Z10calculateRPhS_PfPcS1_cccccc, .-_Z10calculateRPhS_PfPcS1_cccccc
	.globl	_Z11bitonicSortPfPcS0_
	.type	_Z11bitonicSortPfPcS0_, @function
_Z11bitonicSortPfPcS0_:
.LFB2032:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2032:
	.size	_Z11bitonicSortPfPcS0_, .-_Z11bitonicSortPfPcS0_
	.globl	_Z9swapFloatPfS_
	.type	_Z9swapFloatPfS_, @function
_Z9swapFloatPfS_:
.LFB2033:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2033:
	.size	_Z9swapFloatPfS_, .-_Z9swapFloatPfS_
	.globl	_Z8swapCharPcS_
	.type	_Z8swapCharPcS_, @function
_Z8swapCharPcS_:
.LFB2185:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	pushq	%rax
	.cfi_def_cfa_offset 16
	call	_Z9swapFloatPfS_
	.cfi_endproc
.LFE2185:
	.size	_Z8swapCharPcS_, .-_Z8swapCharPcS_
	.globl	_Z16subPixelAccuracyiPfPcS0_S_S_S_S_S_S_ccS_S_
	.type	_Z16subPixelAccuracyiPfPcS0_S_S_S_S_S_S_ccS_S_, @function
_Z16subPixelAccuracyiPfPcS0_S_S_S_S_S_S_ccS_S_:
.LFB2035:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2035:
	.size	_Z16subPixelAccuracyiPfPcS0_S_S_S_S_S_S_ccS_S_, .-_Z16subPixelAccuracyiPfPcS0_S_S_S_S_S_S_ccS_S_
	.section	.rodata.str1.1
.LC4:
	.ascii	"Summary:\nImage size:                  %d * %d\nInterrogatio"
	.ascii	"n window wize:   %d * %d\nSearch window size:          %d * "
	.ascii	"%d\n\nSteps:                       %d * %d\nOverlap:        "
	.ascii	"             %.2f%% \n\nVector size:                 %d * %d"
	.ascii	"\nT"
	.string	"otal vectors                %d\n\nGrid size:                   %d * %d\nBlock size:                  %d * %d\nThread size:                 %d * %d\nTotal threads:               %d\n\nProcessing time (with memcpy):\t%f sec\nProcessing time (without memcpy): %f sec\n\n"
	.text
	.globl	_Z12printSummaryiicccccccciiiiiiiff
	.type	_Z12printSummaryiicccccccciiiiiiiff, @function
_Z12printSummaryiicccccccciiiiiiiff:
.LFB2036:
	.cfi_startproc
	endbr64
	pushq	%r14
	.cfi_def_cfa_offset 16
	.cfi_offset 14, -16
	movl	%r8d, %eax
	movsbl	%dl, %r8d
	movaps	%xmm0, %xmm5
	pushq	%r13
	.cfi_def_cfa_offset 24
	.cfi_offset 13, -24
	cvtsi2ssl	%r8d, %xmm0
	movsbl	%r9b, %r9d
	movsbl	%al, %eax
	pushq	%r12
	.cfi_def_cfa_offset 32
	.cfi_offset 12, -32
	movss	.LC2(%rip), %xmm3
	leal	1(%rax,%r9), %eax
	movsbl	%cl, %r9d
	pushq	%rbp
	.cfi_def_cfa_offset 40
	.cfi_offset 6, -40
	movl	%esi, %ecx
	leaq	.LC4(%rip), %rsi
	cvtss2sd	%xmm1, %xmm2
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset 3, -48
	movsbl	64(%rsp), %r10d
	cvtss2sd	%xmm5, %xmm1
	movl	120(%rsp), %r11d
	movl	104(%rsp), %ebp
	cvtsi2ssl	%r10d, %xmm4
	movl	128(%rsp), %ebx
	movl	112(%rsp), %r12d
	movl	%r11d, %r13d
	imull	%ebp, %r13d
	movl	%ebx, %r14d
	imull	%r12d, %r14d
	divss	%xmm0, %xmm4
	movl	%r13d, %edx
	imull	%r14d, %edx
	pushq	%rdx
	.cfi_def_cfa_offset 56
	pushq	%r14
	.cfi_def_cfa_offset 64
	pushq	%r13
	.cfi_def_cfa_offset 72
	pushq	%r12
	.cfi_def_cfa_offset 80
	pushq	%rbp
	.cfi_def_cfa_offset 88
	pushq	%rbx
	.cfi_def_cfa_offset 96
	pushq	%r11
	.cfi_def_cfa_offset 104
	movl	152(%rsp), %edx
	subss	%xmm4, %xmm3
	mulss	.LC3(%rip), %xmm3
	pushq	%rdx
	.cfi_def_cfa_offset 112
	movl	152(%rsp), %edx
	pushq	%rdx
	.cfi_def_cfa_offset 120
	movl	152(%rsp), %edx
	pushq	%rdx
	.cfi_def_cfa_offset 128
	movsbl	152(%rsp), %edx
	cvtss2sd	%xmm3, %xmm0
	pushq	%rdx
	.cfi_def_cfa_offset 136
	pushq	%r10
	.cfi_def_cfa_offset 144
	movsbl	152(%rsp), %edx
	movsbl	144(%rsp), %r10d
	leal	1(%r10,%rdx), %edx
	pushq	%rdx
	.cfi_def_cfa_offset 152
	movl	%edi, %edx
	movl	$2, %edi
	pushq	%rax
	.cfi_def_cfa_offset 160
	movb	$3, %al
	call	__printf_chk@PLT
	addq	$112, %rsp
	.cfi_def_cfa_offset 48
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%rbp
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r13
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2036:
	.size	_Z12printSummaryiicccccccciiiiiiiff, .-_Z12printSummaryiicccccccciiiiiiiff
	.section	.rodata.str1.1
.LC5:
	.string	"w"
.LC6:
	.string	"%.1f\t%.1f\t%.1f\t%.1f\n"
	.text
	.globl	_Z7putDataPfS_S_S_iPc
	.type	_Z7putDataPfS_S_S_iPc, @function
_Z7putDataPfS_S_S_iPc:
.LFB2037:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movq	%rcx, %r15
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	movq	%rdx, %r14
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	movq	%rsi, %r13
	leaq	.LC5(%rip), %rsi
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movq	%rdi, %r12
	movq	%r9, %rdi
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movq	%r9, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	subq	$24, %rsp
	.cfi_def_cfa_offset 80
	movl	%r8d, 12(%rsp)
	call	fopen@PLT
	movq	%rax, %rbx
	testq	%rax, %rax
	jne	.L30
	movq	stderr(%rip), %rdi
	movq	%rbp, %rcx
	movl	$2, %esi
	xorl	%eax, %eax
	leaq	.LC1(%rip), %rdx
	call	__fprintf_chk@PLT
.L30:
	xorl	%ebp, %ebp
.L31:
	cmpl	%ebp, 12(%rsp)
	jle	.L34
	movss	(%r15,%rbp,4), %xmm3
	movl	$2, %esi
	movq	%rbx, %rdi
	movb	$4, %al
	movss	(%r14,%rbp,4), %xmm2
	leaq	.LC6(%rip), %rdx
	cvtss2sd	(%r12,%rbp,4), %xmm0
	cvtss2sd	0(%r13,%rbp,4), %xmm1
	addss	%xmm3, %xmm3
	incq	%rbp
	addss	%xmm2, %xmm2
	cvtss2sd	%xmm3, %xmm3
	cvtss2sd	%xmm2, %xmm2
	call	__fprintf_chk@PLT
	jmp	.L31
.L34:
	addq	$24, %rsp
	.cfi_def_cfa_offset 56
	movq	%rbx, %rdi
	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
	jmp	fclose@PLT
	.cfi_endproc
.LFE2037:
	.size	_Z7putDataPfS_S_S_iPc, .-_Z7putDataPfS_S_S_iPc
	.section	.rodata.str1.1
.LC7:
	.string	"%s file open error.\n"
.LC8:
	.string	"# vtk DataFile Version 3.0\nvtk output\nASCII\nDATASET STRUCTURED_GRID\nDIMENSIONS %d %d 1\nPOINTS %d float\n"
.LC9:
	.string	"%.1f %.1f %d\n"
.LC10:
	.string	"POINT_DATA %d\nSCALARS velocity float\nLOOKUP_TABLE default\n"
.LC11:
	.string	"%.10f\n"
.LC12:
	.string	"VECTORS velocity float\n"
.LC14:
	.string	"%.10f %.10f %.10f\n"
	.text
	.globl	_Z10putDataVTKiPfS_S_S_iiiPc
	.type	_Z10putDataVTKiPfS_S_S_iiiPc, @function
_Z10putDataVTKiPfS_S_S_iiiPc:
.LFB2038:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movq	%rdx, %r15
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	movq	%rsi, %r14
	leaq	.LC5(%rip), %rsi
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	movl	%edi, %r13d
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movq	%r8, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movq	%rcx, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movl	%r9d, %ebx
	subq	$24, %rsp
	.cfi_def_cfa_offset 80
	movq	96(%rsp), %rdi
	call	fopen@PLT
	movq	%rax, (%rsp)
	testq	%rax, %rax
	jne	.L36
	movq	96(%rsp), %rcx
	movq	stderr(%rip), %rdi
	movl	$2, %esi
	leaq	.LC7(%rip), %rdx
	call	__fprintf_chk@PLT
.L36:
	movl	88(%rsp), %r8d
	movl	80(%rsp), %ecx
	movl	%ebx, %r9d
	xorl	%eax, %eax
	movq	(%rsp), %rdi
	leaq	.LC8(%rip), %rdx
	movl	$2, %esi
	call	__fprintf_chk@PLT
	xorl	%r8d, %r8d
.L37:
	cmpl	%r8d, %ebx
	jle	.L44
	cvtsi2ssl	%r13d, %xmm1
	subss	(%r15,%r8,4), %xmm1
	xorl	%ecx, %ecx
	movb	$2, %al
	movq	(%rsp), %rdi
	movl	$2, %esi
	movq	%r8, 8(%rsp)
	leaq	.LC9(%rip), %rdx
	cvtss2sd	(%r14,%r8,4), %xmm0
	cvtss2sd	%xmm1, %xmm1
	call	__fprintf_chk@PLT
	movq	8(%rsp), %r8
	incq	%r8
	jmp	.L37
.L44:
	movq	(%rsp), %rdi
	movl	%ebx, %ecx
	leaq	.LC10(%rip), %rdx
	xorl	%eax, %eax
	movl	$2, %esi
	xorl	%r14d, %r14d
	leaq	.LC11(%rip), %r13
	call	__fprintf_chk@PLT
.L39:
	cmpl	%r14d, %ebx
	jle	.L45
	movss	0(%rbp,%r14,4), %xmm0
	movss	(%r12,%r14,4), %xmm1
	incq	%r14
	mulss	%xmm1, %xmm1
	mulss	%xmm0, %xmm0
	addss	%xmm1, %xmm0
	call	sqrtf@PLT
	movq	(%rsp), %rdi
	movq	%r13, %rdx
	movb	$1, %al
	movl	$2, %esi
	cvtss2sd	%xmm0, %xmm0
	call	__fprintf_chk@PLT
	jmp	.L39
.L45:
	movq	(%rsp), %rdi
	leaq	.LC12(%rip), %rdx
	xorl	%eax, %eax
	xorl	%r13d, %r13d
	movl	$2, %esi
	call	__fprintf_chk@PLT
.L41:
	cmpl	%r13d, %ebx
	jle	.L46
	movss	(%r12,%r13,4), %xmm1
	movq	(%rsp), %rdi
	xorps	%xmm2, %xmm2
	movb	$3, %al
	xorps	.LC13(%rip), %xmm1
	leaq	.LC14(%rip), %rdx
	movl	$2, %esi
	cvtss2sd	0(%rbp,%r13,4), %xmm0
	incq	%r13
	cvtss2sd	%xmm1, %xmm1
	call	__fprintf_chk@PLT
	jmp	.L41
.L46:
	addq	$24, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2038:
	.size	_Z10putDataVTKiPfS_S_S_iiiPc, .-_Z10putDataVTKiPfS_S_S_iiiPc
	.globl	_Z71__device_stub__Z17vectorComputationPhS_iiiccccccccccccPfS0_S0_S0_S0_S0_PhS_iiiccccccccccccPfS0_S0_S0_S0_S0_
	.type	_Z71__device_stub__Z17vectorComputationPhS_iiiccccccccccccPfS0_S0_S0_S0_S0_PhS_iiiccccccccccccPfS0_S0_S0_S0_S0_, @function
_Z71__device_stub__Z17vectorComputationPhS_iiiccccccccccccPfS0_S0_S0_S0_S0_PhS_iiiccccccccccccPfS0_S0_S0_S0_S0_:
.LFB2063:
	.cfi_startproc
	endbr64
	subq	$392, %rsp
	.cfi_def_cfa_offset 400
	movl	400(%rsp), %eax
	movq	%rdi, 120(%rsp)
	leaq	144(%rsp), %rdi
	movq	%rsi, 112(%rsp)
	leaq	156(%rsp), %rsi
	movb	%al, 92(%rsp)
	movl	408(%rsp), %eax
	movl	%edx, 108(%rsp)
	leaq	128(%rsp), %rdx
	movb	%al, 88(%rsp)
	movl	416(%rsp), %eax
	movl	%ecx, 104(%rsp)
	leaq	136(%rsp), %rcx
	movb	%al, 84(%rsp)
	movl	424(%rsp), %eax
	movl	%r8d, 100(%rsp)
	movb	%al, 80(%rsp)
	movl	432(%rsp), %eax
	movb	%r9b, 96(%rsp)
	movb	%al, 76(%rsp)
	movl	440(%rsp), %eax
	movb	%al, 72(%rsp)
	movl	448(%rsp), %eax
	movb	%al, 68(%rsp)
	movl	456(%rsp), %eax
	movb	%al, 64(%rsp)
	movl	464(%rsp), %eax
	movb	%al, 60(%rsp)
	movl	472(%rsp), %eax
	movb	%al, 56(%rsp)
	movl	480(%rsp), %eax
	movb	%al, 52(%rsp)
	movq	488(%rsp), %rax
	movq	%rax, 40(%rsp)
	movq	496(%rsp), %rax
	movq	%rax, 32(%rsp)
	movq	504(%rsp), %rax
	movq	%rax, 24(%rsp)
	movq	512(%rsp), %rax
	movq	%rax, 16(%rsp)
	movq	520(%rsp), %rax
	movq	%rax, 8(%rsp)
	movq	528(%rsp), %rax
	movq	%rax, (%rsp)
	movq	%fs:40, %rax
	movq	%rax, 376(%rsp)
	xorl	%eax, %eax
	leaq	120(%rsp), %rax
	movq	%rax, 192(%rsp)
	leaq	112(%rsp), %rax
	movq	%rax, 200(%rsp)
	leaq	108(%rsp), %rax
	movq	%rax, 208(%rsp)
	leaq	104(%rsp), %rax
	movq	%rax, 216(%rsp)
	leaq	100(%rsp), %rax
	movq	%rax, 224(%rsp)
	leaq	96(%rsp), %rax
	movq	%rax, 232(%rsp)
	leaq	92(%rsp), %rax
	movq	%rax, 240(%rsp)
	leaq	88(%rsp), %rax
	movq	%rax, 248(%rsp)
	leaq	84(%rsp), %rax
	movq	%rax, 256(%rsp)
	leaq	80(%rsp), %rax
	movq	%rax, 264(%rsp)
	leaq	76(%rsp), %rax
	movq	%rax, 272(%rsp)
	leaq	72(%rsp), %rax
	movq	%rax, 280(%rsp)
	leaq	68(%rsp), %rax
	movq	%rax, 288(%rsp)
	leaq	64(%rsp), %rax
	movq	%rax, 296(%rsp)
	leaq	60(%rsp), %rax
	movq	%rax, 304(%rsp)
	leaq	56(%rsp), %rax
	movq	%rax, 312(%rsp)
	leaq	52(%rsp), %rax
	movq	%rax, 320(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 328(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 336(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 344(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 352(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 360(%rsp)
	movq	%rsp, %rax
	movq	%rax, 368(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 144(%rsp)
	movl	$1, 152(%rsp)
	movq	%rax, 156(%rsp)
	movl	$1, 164(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L47
	pushq	136(%rsp)
	.cfi_def_cfa_offset 408
	leaq	_Z17vectorComputationPhS_iiiccccccccccccPfS0_S0_S0_S0_S0_(%rip), %rdi
	pushq	136(%rsp)
	.cfi_def_cfa_offset 416
	movq	172(%rsp), %rcx
	movl	180(%rsp), %r8d
	movq	160(%rsp), %rsi
	movl	168(%rsp), %edx
	leaq	208(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 408
	popq	%rdx
	.cfi_def_cfa_offset 400
.L47:
	movq	376(%rsp), %rax
	subq	%fs:40, %rax
	je	.L49
	call	__stack_chk_fail@PLT
.L49:
	addq	$392, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2063:
	.size	_Z71__device_stub__Z17vectorComputationPhS_iiiccccccccccccPfS0_S0_S0_S0_S0_PhS_iiiccccccccccccPfS0_S0_S0_S0_S0_, .-_Z71__device_stub__Z17vectorComputationPhS_iiiccccccccccccPfS0_S0_S0_S0_S0_PhS_iiiccccccccccccPfS0_S0_S0_S0_S0_
	.globl	_Z17vectorComputationPhS_iiiccccccccccccPfS0_S0_S0_S0_S0_
	.type	_Z17vectorComputationPhS_iiiccccccccccccPfS0_S0_S0_S0_S0_, @function
_Z17vectorComputationPhS_iiiccccccccccccPfS0_S0_S0_S0_S0_:
.LFB2064:
	.cfi_startproc
	endbr64
	movsbl	88(%rsp), %eax
	movsbl	%r9b, %r9d
	movl	%eax, 88(%rsp)
	movsbl	80(%rsp), %eax
	movl	%eax, 80(%rsp)
	movsbl	72(%rsp), %eax
	movl	%eax, 72(%rsp)
	movsbl	64(%rsp), %eax
	movl	%eax, 64(%rsp)
	movsbl	56(%rsp), %eax
	movl	%eax, 56(%rsp)
	movsbl	48(%rsp), %eax
	movl	%eax, 48(%rsp)
	movsbl	40(%rsp), %eax
	movl	%eax, 40(%rsp)
	movsbl	32(%rsp), %eax
	movl	%eax, 32(%rsp)
	movsbl	24(%rsp), %eax
	movl	%eax, 24(%rsp)
	movsbl	16(%rsp), %eax
	movl	%eax, 16(%rsp)
	movsbl	8(%rsp), %eax
	movl	%eax, 8(%rsp)
	jmp	_Z71__device_stub__Z17vectorComputationPhS_iiiccccccccccccPfS0_S0_S0_S0_S0_PhS_iiiccccccccccccPfS0_S0_S0_S0_S0_
	.cfi_endproc
.LFE2064:
	.size	_Z17vectorComputationPhS_iiiccccccccccccPfS0_S0_S0_S0_S0_, .-_Z17vectorComputationPhS_iiiccccccccccccPfS0_S0_S0_S0_S0_
	.globl	_Z43__device_stub__Z16vectorCorrectioniPfS_S_S_iPfS_S_S_
	.type	_Z43__device_stub__Z16vectorCorrectioniPfS_S_S_iPfS_S_S_, @function
_Z43__device_stub__Z16vectorCorrectioniPfS_S_S_iPfS_S_S_:
.LFB2065:
	.cfi_startproc
	endbr64
	subq	$168, %rsp
	.cfi_def_cfa_offset 176
	movl	%edi, 44(%rsp)
	leaq	64(%rsp), %rdi
	movq	%rsi, 32(%rsp)
	leaq	76(%rsp), %rsi
	movq	%rdx, 24(%rsp)
	leaq	48(%rsp), %rdx
	movq	%rcx, 16(%rsp)
	leaq	56(%rsp), %rcx
	movq	%r8, 8(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 152(%rsp)
	xorl	%eax, %eax
	leaq	44(%rsp), %rax
	movl	$1, 72(%rsp)
	movq	%rax, 112(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 144(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 64(%rsp)
	movq	%rax, 76(%rsp)
	movl	$1, 84(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L52
	pushq	56(%rsp)
	.cfi_def_cfa_offset 184
	leaq	_Z16vectorCorrectioniPfS_S_S_(%rip), %rdi
	pushq	56(%rsp)
	.cfi_def_cfa_offset 192
	movq	92(%rsp), %rcx
	movl	100(%rsp), %r8d
	movq	80(%rsp), %rsi
	movl	88(%rsp), %edx
	leaq	128(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 184
	popq	%rdx
	.cfi_def_cfa_offset 176
.L52:
	movq	152(%rsp), %rax
	subq	%fs:40, %rax
	je	.L54
	call	__stack_chk_fail@PLT
.L54:
	addq	$168, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2065:
	.size	_Z43__device_stub__Z16vectorCorrectioniPfS_S_S_iPfS_S_S_, .-_Z43__device_stub__Z16vectorCorrectioniPfS_S_S_iPfS_S_S_
	.globl	_Z16vectorCorrectioniPfS_S_S_
	.type	_Z16vectorCorrectioniPfS_S_S_, @function
_Z16vectorCorrectioniPfS_S_S_:
.LFB2066:
	.cfi_startproc
	endbr64
	jmp	_Z43__device_stub__Z16vectorCorrectioniPfS_S_S_iPfS_S_S_
	.cfi_endproc
.LFE2066:
	.size	_Z16vectorCorrectioniPfS_S_S_, .-_Z16vectorCorrectioniPfS_S_S_
	.section	.rodata.str1.1
.LC16:
	.string	"raw"
.LC17:
	.string	"2ms2000018"
.LC18:
	.string	"%s.%s"
.LC19:
	.string	"2ms2000019"
.LC20:
	.string	"dat"
.LC21:
	.string	"result"
.LC22:
	.string	"%s-vtk.vtk"
.LC24:
	.string	"CUDA error: %s\n\n"
.LC25:
	.string	"CUDA success\n\n"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2027:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movl	$1048576, %edi
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	leaq	.LC18(%rip), %r13
	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	$408, %rsp
	.cfi_def_cfa_offset 464
	movq	%fs:40, %rax
	movq	%rax, 392(%rsp)
	xorl	%eax, %eax
	leaq	136(%rsp), %rbp
	leaq	200(%rsp), %r12
	call	malloc@PLT
	movl	$1048576, %edi
	leaq	264(%rsp), %r15
	leaq	328(%rsp), %r14
	movq	%rax, %rbx
	call	malloc@PLT
	leaq	48(%rsp), %rdi
	movl	$1048576, %esi
	movq	%rax, 40(%rsp)
	call	cudaMalloc@PLT
	movl	$1048576, %esi
	leaq	56(%rsp), %rdi
	call	cudaMalloc@PLT
	movl	$14884, %edi
	call	malloc@PLT
	movl	$14884, %edi
	movq	%rax, 8(%rsp)
	call	malloc@PLT
	movl	$14884, %edi
	movq	%rax, 16(%rsp)
	call	malloc@PLT
	movl	$14884, %edi
	movq	%rax, 24(%rsp)
	call	malloc@PLT
	leaq	64(%rsp), %rdi
	movl	$14884, %esi
	movq	%rax, 32(%rsp)
	call	cudaMalloc@PLT
	leaq	72(%rsp), %rdi
	movl	$14884, %esi
	call	cudaMalloc@PLT
	leaq	80(%rsp), %rdi
	movl	$14884, %esi
	call	cudaMalloc@PLT
	leaq	88(%rsp), %rdi
	movl	$14884, %esi
	call	cudaMalloc@PLT
	leaq	96(%rsp), %rdi
	movl	$74420, %esi
	call	cudaMalloc@PLT
	leaq	104(%rsp), %rdi
	movl	$74420, %esi
	call	cudaMalloc@PLT
	movq	%r13, %rcx
	movl	$64, %edx
	movq	%rbp, %rdi
	leaq	.LC16(%rip), %r9
	leaq	.LC17(%rip), %r8
	movl	$2, %esi
	xorl	%eax, %eax
	call	__sprintf_chk@PLT
	movq	%r13, %rcx
	movq	%r12, %rdi
	movl	$64, %edx
	leaq	.LC16(%rip), %r9
	leaq	.LC19(%rip), %r8
	movl	$2, %esi
	xorl	%eax, %eax
	call	__sprintf_chk@PLT
	movq	%r13, %rcx
	movl	$64, %edx
	movq	%r15, %rdi
	leaq	.LC20(%rip), %r9
	leaq	.LC21(%rip), %r8
	movl	$2, %esi
	xorl	%eax, %eax
	call	__sprintf_chk@PLT
	movl	$64, %edx
	movq	%r14, %rdi
	xorl	%eax, %eax
	leaq	.LC21(%rip), %r8
	leaq	.LC22(%rip), %rcx
	movl	$2, %esi
	call	__sprintf_chk@PLT
	movq	40(%rsp), %rdx
	movq	%r12, %r8
	movq	%rbp, %rcx
	movq	%rbx, %rsi
	movl	$1048576, %edi
	call	_Z7getDataiPhS_PcS0_
	call	clock@PLT
	movq	48(%rsp), %rdi
	movl	$1, %ecx
	movq	%rbx, %rsi
	movl	$1048576, %edx
	movq	%rax, %r12
	call	cudaMemcpy@PLT
	movq	40(%rsp), %rsi
	movq	56(%rsp), %rdi
	movl	$1, %ecx
	movl	$1048576, %edx
	call	cudaMemcpy@PLT
	call	clock@PLT
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$1, %ecx
	movl	$1, %esi
	movl	$1, 120(%rsp)
	movq	%rax, %r13
	movabsq	$261993005117, %rdi
	movabsq	$137438953504, %rdx
	movq	%rdi, 112(%rsp)
	movq	%rdx, 124(%rsp)
	movl	$1, 132(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L58
	pushq	%rcx
	.cfi_def_cfa_offset 472
	movl	$32, %r9d
	movl	$61, %r8d
	movl	$61, %ecx
	pushq	96(%rsp)
	.cfi_def_cfa_offset 480
	movl	$1024, %edx
	pushq	96(%rsp)
	.cfi_def_cfa_offset 488
	pushq	96(%rsp)
	.cfi_def_cfa_offset 496
	pushq	96(%rsp)
	.cfi_def_cfa_offset 504
	pushq	144(%rsp)
	.cfi_def_cfa_offset 512
	pushq	144(%rsp)
	.cfi_def_cfa_offset 520
	pushq	$16
	.cfi_def_cfa_offset 528
	pushq	$16
	.cfi_def_cfa_offset 536
	pushq	$56
	.cfi_def_cfa_offset 544
	pushq	$56
	.cfi_def_cfa_offset 552
	pushq	$25
	.cfi_def_cfa_offset 560
	pushq	$25
	.cfi_def_cfa_offset 568
	pushq	$12
	.cfi_def_cfa_offset 576
	pushq	$12
	.cfi_def_cfa_offset 584
	pushq	$12
	.cfi_def_cfa_offset 592
	pushq	$12
	.cfi_def_cfa_offset 600
	pushq	$32
	.cfi_def_cfa_offset 608
	movq	200(%rsp), %rsi
	movq	192(%rsp), %rdi
	call	_Z71__device_stub__Z17vectorComputationPhS_iiiccccccccccccPfS0_S0_S0_S0_S0_PhS_iiiccccccccccccPfS0_S0_S0_S0_S0_
	addq	$144, %rsp
	.cfi_def_cfa_offset 464
.L58:
	call	cudaThreadSynchronize@PLT
	movq	124(%rsp), %rdx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	132(%rsp), %ecx
	movq	112(%rsp), %rdi
	movl	120(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L59
	movq	88(%rsp), %r8
	movq	80(%rsp), %rcx
	movl	$61, %edi
	movq	104(%rsp), %rdx
	movq	96(%rsp), %rsi
	call	_Z43__device_stub__Z16vectorCorrectioniPfS_S_S_iPfS_S_S_
.L59:
	call	cudaThreadSynchronize@PLT
	call	clock@PLT
	movq	64(%rsp), %rsi
	movq	8(%rsp), %rdi
	movl	$2, %ecx
	movl	$14884, %edx
	movq	%rax, %rbp
	call	cudaMemcpy@PLT
	movq	72(%rsp), %rsi
	movq	16(%rsp), %rdi
	subq	%r13, %rbp
	movl	$2, %ecx
	movl	$14884, %edx
	call	cudaMemcpy@PLT
	movq	80(%rsp), %rsi
	movq	24(%rsp), %rdi
	movl	$2, %ecx
	movl	$14884, %edx
	call	cudaMemcpy@PLT
	movq	88(%rsp), %rsi
	movq	32(%rsp), %rdi
	movl	$2, %ecx
	movl	$14884, %edx
	call	cudaMemcpy@PLT
	call	clock@PLT
	cvtsi2ssq	%rbp, %xmm1
	movl	$32, %ecx
	movss	.LC23(%rip), %xmm2
	subq	%r12, %rax
	movl	$12, %r9d
	movl	$12, %r8d
	movl	$32, %edx
	cvtsi2ssq	%rax, %xmm0
	pushq	%rax
	.cfi_def_cfa_offset 472
	movl	$1024, %esi
	movl	$1024, %edi
	pushq	$61
	.cfi_def_cfa_offset 480
	divss	%xmm2, %xmm1
	pushq	$61
	.cfi_def_cfa_offset 488
	pushq	$32
	.cfi_def_cfa_offset 496
	pushq	$32
	.cfi_def_cfa_offset 504
	pushq	$3721
	.cfi_def_cfa_offset 512
	pushq	$61
	.cfi_def_cfa_offset 520
	pushq	$61
	.cfi_def_cfa_offset 528
	pushq	$16
	.cfi_def_cfa_offset 536
	pushq	$16
	.cfi_def_cfa_offset 544
	pushq	$12
	.cfi_def_cfa_offset 552
	pushq	$12
	.cfi_def_cfa_offset 560
	divss	%xmm2, %xmm0
	call	_Z12printSummaryiicccccccciiiiiiiff
	movq	120(%rsp), %rdx
	movq	112(%rsp), %rsi
	movq	%r15, %r9
	movq	104(%rsp), %rdi
	addq	$96, %rsp
	.cfi_def_cfa_offset 464
	movq	32(%rsp), %rcx
	movl	$3721, %r8d
	call	_Z7putDataPfS_S_S_iPc
	movl	$3721, %r9d
	movl	$1024, %edi
	pushq	%rdx
	.cfi_def_cfa_offset 472
	pushq	%r14
	.cfi_def_cfa_offset 480
	pushq	$61
	.cfi_def_cfa_offset 488
	pushq	$61
	.cfi_def_cfa_offset 496
	movq	64(%rsp), %r8
	movq	56(%rsp), %rcx
	movq	48(%rsp), %rdx
	movq	40(%rsp), %rsi
	call	_Z10putDataVTKiPfS_S_S_iiiPc
	addq	$32, %rsp
	.cfi_def_cfa_offset 464
	movq	%rbx, %rdi
	call	free@PLT
	movq	40(%rsp), %rdi
	call	free@PLT
	movq	48(%rsp), %rdi
	call	cudaFree@PLT
	movq	56(%rsp), %rdi
	call	cudaFree@PLT
	movq	8(%rsp), %rdi
	call	free@PLT
	movq	16(%rsp), %rdi
	call	free@PLT
	movq	24(%rsp), %rdi
	call	free@PLT
	movq	32(%rsp), %rdi
	call	free@PLT
	movq	64(%rsp), %rdi
	call	cudaFree@PLT
	movq	72(%rsp), %rdi
	call	cudaFree@PLT
	movq	80(%rsp), %rdi
	call	cudaFree@PLT
	movq	88(%rsp), %rdi
	call	cudaFree@PLT
	movq	96(%rsp), %rdi
	call	cudaFree@PLT
	movq	104(%rsp), %rdi
	call	cudaFree@PLT
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L60
	call	cudaGetErrorString@PLT
	movl	$2, %edi
	leaq	.LC24(%rip), %rsi
	movq	%rax, %rdx
	xorl	%eax, %eax
	call	__printf_chk@PLT
	orl	$-1, %edi
	call	exit@PLT
.L60:
	xorl	%eax, %eax
	leaq	.LC25(%rip), %rsi
	movl	$2, %edi
	call	__printf_chk@PLT
	movq	392(%rsp), %rax
	subq	%fs:40, %rax
	je	.L61
	call	__stack_chk_fail@PLT
.L61:
	addq	$408, %rsp
	.cfi_def_cfa_offset 56
	xorl	%eax, %eax
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2027:
	.size	main, .-main
	.section	.rodata.str1.1
.LC26:
	.string	"_Z16vectorCorrectioniPfS_S_S_"
.LC27:
	.string	"_Z17vectorComputationPhS_iiiccccccccccccPfS0_S0_S0_S0_S0_"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2068:
	.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	.LC26(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z16vectorCorrectioniPfS_S_S_(%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	.LC27(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z17vectorComputationPhS_iiiccccccccccccPfS0_S0_S0_S0_S0_(%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
.LFE2068:
	.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
	.section	.rodata.cst4,"aM",@progbits,4
	.align 4
.LC2:
	.long	1065353216
	.align 4
.LC3:
	.long	1120403456
	.section	.rodata.cst16,"aM",@progbits,16
	.align 16
.LC13:
	.long	-2147483648
	.long	0
	.long	0
	.long	0
	.section	.rodata.cst4
	.align 4
.LC23:
	.long	1232348160
	.ident	"GCC: (Ubuntu 13.3.0-6ubuntu2~24.04) 13.3.0"
	.section	.note.GNU-stack,"",@progbits
	.section	.note.gnu.property,"a"
	.align 8
	.long	1f - 0f
	.long	4f - 1f
	.long	5
0:
	.string	"GNU"
1:
	.align 8
	.long	0xc0000002
	.long	3f - 2f
2:
	.long	0x3
3:
	.align 8
4:


#include <hip/hip_runtime.h>
#include <stdio.h>
#include <stdlib.h>
#include <math.h>
#include <time.h>
 
#define DEFAULT_WIDTH		1024
#define DEFAULT_HEIGHT		1024
#define DEFAULT_INTWNDW		32
#define DEFAULT_INTWNDH		32
#define DEFAULT_SRHWNDL		12
#define DEFAULT_SRHWNDR		12
#define DEFAULT_SRHWNDA		12
#define DEFAULT_SRHWNDB		12
 
 
__host__	void getData(int size, unsigned char *data0, unsigned char *data1, char *path0, char *path1);
 
__global__	void vectorComputation(unsigned char *f, unsigned char *g, int width, int totalVectorsX, int totalVectorsY, char m, char n, char l, char r, char a, char b, char ww, char wh, char fw, char fh, char stepX, char stepY, float *R_posSubX, float *R_posSubY, float *startX, float *startY, float *endX, float *endY);
 
__device__	void getData0intoSharedMemory(unsigned char *f, unsigned char *sh_f, int width, char m, char n, char l, char a, int startPixelX, int startPixelY, char stepX, char stepY);
 
__device__	void getData1intoSharedMemory(unsigned char *g, unsigned char *sh_g, int width, char fw, char fh, int startPixelX, int startPixelY, char stepX, char stepY);
 
__device__	void calculateR(unsigned char *sh_f, unsigned char *sh_g, float *R, char *R_posX, char *R_posY, char m, char n, char ww, char wh, char fw, char fh);
 
__device__	void bitonicSort(float *R, char *R_posX, char *R_posY);
 
__device__	void swapFloat(float *x, float *y);
 
__device__	void swapChar(char *x, char *y);
 
__device__	void subPixelAccuracy(int vectorID, float *R, char *R_posX, char *R_posY, float *subXN, float *subXD, float *subYN, float *subYD, float *R_posSubX, float *R_posSubY, char l, char a, float *endX, float *endY);
 
__global__ void vectorCorrection(int totalVectorsX, float *R_posSubX, float *R_posSubY, float *endX, float *endY);
 
__host__	void printSummary(int width, int height, char intWndW, char intWndH, char srhWndL, char srhWndR, char srhWndA, char srhWndB, char stepX, char stepY, int totalVectorsX, int totalVectorsY, int totalVectors, int Dbx, int Dby, int Dgx, int Dgy, float h_time, float d_time);
 
__host__	void putData(float *startX, float *startY, float *endX, float *endY, int totalVectors, char *path);
 
__host__	void putDataVTK(int height, float *startX, float *startY, float *endX, float *endY, int number, int numberX, int numberY, char *pathVTK);
 
int main(int argc, char *argv[]) {
	int width = DEFAULT_WIDTH, height = DEFAULT_HEIGHT;
	int size = width * height;
 
	unsigned char *h_data0, *h_data1, *d_data0, *d_data1;
	char intWndW = (char)DEFAULT_INTWNDW,
		 intWndH = (char)DEFAULT_INTWNDH,
		 srhWndL = (char)DEFAULT_SRHWNDL,
		 srhWndR = (char)DEFAULT_SRHWNDR,
		 srhWndA = (char)DEFAULT_SRHWNDA,
		 srhWndB = (char)DEFAULT_SRHWNDB,
		 srhWndW = srhWndL + srhWndR + 1,
		 srhWndH = srhWndA + srhWndB + 1,
		 srhFldW = srhWndL + srhWndR + intWndW,
		 srhFldH = srhWndL + srhWndR + intWndH;
 
	char stepX, stepY;
	float *h_startX, *h_startY, *d_startX, *d_startY;
	float *h_endX, *h_endY, *d_endX, *d_endY;
	float *R_posSubX, *R_posSubY;
	int totalVectorsX, totalVectorsY, totalVectors;
 
	char *inputFileName0 = "2ms2000018",
		 *inputFileName1 = "2ms2000019",
		 *inputFileExtension = "raw",
		 *outputFileName = "result",
		 *outputFileExtension = "dat";
 
	char inputFilepath0[64], inputFilepath1[64], outputFilepath[64], outputFilePathVTK[64];
 
	clock_t h_start, h_end, d_start, d_end;
	float h_time, d_time;
 
	hipError_t hipError_t;
 
	dim3 Dg, Db;
	int Dgx, Dgy, Dbx, Dby;
 
	h_data0 = (unsigned char *)malloc(sizeof(unsigned char) * size);
	h_data1 = (unsigned char *)malloc(sizeof(unsigned char) * size);
	hipMalloc((void **)&d_data0, sizeof(unsigned char) * size);
	hipMalloc((void **)&d_data1, sizeof(unsigned char) * size);
 
	stepX = (char)(intWndW / 2);
	stepY = (char)(intWndH / 2);
 
	totalVectorsX = (int)((width - srhFldW) / stepX) + 1;
	totalVectorsY = (int)((height - srhFldH) / stepY) + 1;
	totalVectors = totalVectorsX * totalVectorsY;
 
	h_startX = (float *)malloc(sizeof(float) * totalVectors);
	h_startY = (float *)malloc(sizeof(float) * totalVectors);
	h_endX = (float *)malloc(sizeof(float) * totalVectors);
	h_endY = (float *)malloc(sizeof(float) * totalVectors);
	hipMalloc((void **)&d_startX, sizeof(float) * totalVectors);
	hipMalloc((void **)&d_startY, sizeof(float) * totalVectors);
	hipMalloc((void **)&d_endX, sizeof(float) * totalVectors);
	hipMalloc((void **)&d_endY, sizeof(float) * totalVectors);
 
	hipMalloc((void **)&R_posSubX, sizeof(float) * totalVectors * 5);
	hipMalloc((void **)&R_posSubY, sizeof(float) * totalVectors * 5);
 
	sprintf(inputFilepath0, "%s.%s", inputFileName0, inputFileExtension);
	sprintf(inputFilepath1, "%s.%s", inputFileName1, inputFileExtension);
	sprintf(outputFilepath, "%s.%s", outputFileName, outputFileExtension);
	sprintf(outputFilePathVTK, "%s-vtk.vtk", outputFileName);
 
	Db.x = 32;
	Db.y = 32;
	Dg.x = totalVectorsX;
	Dg.y = totalVectorsY;
	Dbx = Db.x; Dby = Db.y;
	Dgx = Dg.x; Dgy = Dg.y;
 
	getData(size, h_data0, h_data1, inputFilepath0, inputFilepath1);
 
	h_start = clock();
	hipMemcpy(d_data0, h_data0, sizeof(unsigned char) * size, hipMemcpyHostToDevice);
	hipMemcpy(d_data1, h_data1, sizeof(unsigned char) * size, hipMemcpyHostToDevice);
 
	d_start = clock();
 
	vectorComputation<<<Dg, Db>>>(d_data0, d_data1, width, totalVectorsX, totalVectorsY, intWndW, intWndH, srhWndL, srhWndR, srhWndA, srhWndB, srhWndW, srhWndH, srhFldW, srhFldH, stepX, stepY, R_posSubX, R_posSubY, d_startX, d_startY, d_endX, d_endY);
	hipDeviceSynchronize();
 
	vectorCorrection<<<Dg, Db>>>(totalVectorsX, R_posSubX, R_posSubY, d_endX, d_endY);
	hipDeviceSynchronize();
 
	d_end = clock();
 
	hipMemcpy(h_startX, d_startX, sizeof(float) * totalVectors, hipMemcpyDeviceToHost);
	hipMemcpy(h_startY, d_startY, sizeof(float) * totalVectors, hipMemcpyDeviceToHost);
	hipMemcpy(h_endX, d_endX, sizeof(float) * totalVectors, hipMemcpyDeviceToHost);
	hipMemcpy(h_endY, d_endY, sizeof(float) * totalVectors, hipMemcpyDeviceToHost);	
	h_end = clock();
 
	h_time = (float)(h_end - h_start) / CLOCKS_PER_SEC;
	d_time = (float)(d_end - d_start) / CLOCKS_PER_SEC;
	printSummary(width, height, intWndW, intWndH, srhWndL, srhWndR, srhWndA, srhWndB, stepX, stepY, totalVectorsX, totalVectorsY, totalVectors, Dbx, Dby, Dgx, Dgy, h_time, d_time);
 
	putData(h_startX, h_startY, h_endX, h_endY, totalVectors, outputFilepath);
	putDataVTK(height, h_startX, h_startY, h_endX, h_endY, totalVectors, totalVectorsX, totalVectorsY, outputFilePathVTK);
 
	free(h_data0);
	free(h_data1);
	hipFree(d_data0);
	hipFree(d_data1);
 
	free(h_startX);
	free(h_startY);
	free(h_endX);
	free(h_endY);
	hipFree(d_startX);
	hipFree(d_startY);
	hipFree(d_endX);
	hipFree(d_endY);
 
	hipFree(R_posSubX);
	hipFree(R_posSubY);
 
	hipError_t = hipGetLastError();
 
	if(hipError_t != hipSuccess) {
		// print the CUDA error message and exit
		printf("CUDA error: %s\n\n", hipGetErrorString(hipError_t));
		exit(-1);
	} else {
		printf("CUDA success\n\n");
	}
 
	return 0;
}
 
__host__		void getData(int size, unsigned char *data0, unsigned char *data1, char *path0, char *path1) {
	int i;
	FILE *fp;
 
	fp = fopen(path0, "rb");
	if (fp == NULL) {
		printf("%s file open error\n", path0);
	}
 
	for (i = 0; i < size; i++) {
		data0[i] = fgetc(fp);
	}
 
	fclose(fp);
 
	fp = fopen(path1, "rb");
	if (fp == NULL) {
		printf("%s file open error\n", path1);
	}
 
	for (i = 0; i < size; i++) {
		data1[i] = fgetc(fp);
	}
 
	fclose(fp);
}
 
__global__	void vectorComputation(unsigned char *f, unsigned char *g, int width, int totalVectorsX, int totalVectorsY, char m, char n, char l, char r, char a, char b, char ww, char wh, char fw, char fh, char stepX, char stepY, float *R_posSubX, float *R_posSubY, float *startX, float *startY, float *endX, float *endY) {
	int vectorID;
	int startPixelX, startPixelY;
 
	__shared__ unsigned char
		sh_f[2304], sh_g[12100];
 
	__shared__ float
		R[4096];
 
	__shared__ char
		R_posX[4096], R_posY[4096];
 
	__shared__ float
		subXN[5], subXD[5], subYN[5], subYD[5];
 
	vectorID = blockIdx.y * gridDim.x + blockIdx.x;
 
	startPixelX = stepX * blockIdx.x;
	startPixelY = stepY * blockIdx.y;
 
	startX[vectorID] = startPixelX + l + (m / 2) + (float)0.5;
	startY[vectorID] = startPixelY + a + (n / 2) + (float)0.5;
	endX[vectorID] = 0;
	endY[vectorID] = 0;
 
	getData0intoSharedMemory(f, sh_f, width, m, n, l, a, startPixelX, startPixelY, stepX, stepY);
 
	getData1intoSharedMemory(g, sh_g, width, fw, fh, startPixelX, startPixelY, stepX, stepY);
 
		__syncthreads();
 
	calculateR(sh_f, sh_g, R, R_posX, R_posY, m, n, ww, wh, fw, fh);
 
		__syncthreads();
 
	bitonicSort(R, R_posX, R_posY);
 
		__syncthreads();
 
	subPixelAccuracy(vectorID, R, R_posX, R_posY, subXN, subXD, subYN, subYD, R_posSubX, R_posSubY, l, a, endX, endY);
 
		__syncthreads();
 
}
 
__device__ void getData0intoSharedMemory(unsigned char *f, unsigned char *sh_f, int width, char m, char n, char l, char a, int startPixelX, int startPixelY, char stepX, char stepY) {
	char j, i;
 
	for (j = threadIdx.y; j < n; j += blockDim.y) {
		for (i = threadIdx.x; i < m; i += blockDim.x) {
			sh_f[j * m + i] = f[(startPixelY + a + j) * width + (startPixelX + l + i)];
		}
	}
 
}
 
__device__ void getData1intoSharedMemory(unsigned char *g, unsigned char *sh_g, int width, char fw, char fh, int startPixelX, int startPixelY, char stepX, char stepY) {
	char j, i;
 
	for (j = threadIdx.y; j < fh; j += blockDim.y) {
		for (i = threadIdx.x; i < fw; i += blockDim.x) {
			sh_g[j * fw + i] = g[(startPixelY + j) * width + (startPixelX + i)];
		}
	}
 
}
 
__device__ void calculateR(unsigned char *sh_f, unsigned char *sh_g, float *R, char *R_posX, char *R_posY, char m, char n, char ww, char wh, char fw, char fh) {
	char i, j, zeta, eta;
	float fValue, gValue;
	float A, B, C;
	int theta;
 
	for (eta = threadIdx.y; eta < wh; eta += blockDim.y) {
		for (zeta = threadIdx.x; zeta < ww; zeta += blockDim.x) {
			theta = eta * ww + zeta;
			A = 0; B = 0; C = 0;
 
			for (j = 0; j < n; j++) {
				for (i = 0; i < m; i++) {
 
					fValue = (float)sh_f[j * m + i];
					gValue = (float)sh_g[(eta + j) * fw + (zeta + i)];
 
					A += fValue * gValue;
					B += fValue * fValue;
					C += gValue * gValue;
				}
			}
			R[theta] = sqrt((A * A) / (B * C));
			R_posX[theta] = zeta;
			R_posY[theta] = eta;
		}
	}
 
	for (theta = (threadIdx.y * blockDim.x + threadIdx.x) + (ww * wh); theta < 4096; theta += blockDim.x * blockDim.y) {
		R[theta] = 0;
		R_posX[theta] = -128;
		R_posY[theta] = -128;
	}
 
}
 
__device__ void bitonicSort(float *R, char *R_posX, char *R_posY) {
 
	int i, j, k, frac, xtemp, x, group, el0, el1;
	for (i = 0; 1 << i <= 2048; i++) {
		for (j = i; j >= 0; j--) {
 
			if (j == i) {
				for (k = threadIdx.y * blockDim.x + threadIdx.x; k < 2048; k += blockDim.x * blockDim.y) {
					frac = (1 << (j + 1)) - 1;
					xtemp = k >> j;
					x = xtemp << j;
 
					group = (x << 2) + frac;
					el0 = (x << 1) + (k & ((1 << j) - 1));
					el1 = group - el0;
 
					if (R[el0] < R[el1]) {
						swapFloat(&R[el0], &R[el1]);
						swapChar(&R_posX[el0], &R_posX[el1]);
						swapChar(&R_posY[el0], &R_posY[el1]);
					}
				}
 
			} else {
				for (k = threadIdx.y * blockDim.x + threadIdx.x; k < 2048; k += blockDim.x * blockDim.y) {
					xtemp = k >> j;
					x = xtemp << j;
 
					el0 = (x << 1) + (k & ((1 << j) - 1));
					el1 = el0 + (1 << j);
 
					if (R[el0] < R[el1]) {
						swapFloat(&R[el0], &R[el1]);
						swapChar(&R_posX[el0], &R_posX[el1]);
						swapChar(&R_posY[el0], &R_posY[el1]);
					}
				}
 
			}
			__syncthreads();
 
		}
	}
 
}
 
__device__ void swapFloat(float *x, float *y) {
 
	float t;
	 t = *x;
	*x = *y;
	*y =  t;
}
 
__device__ void swapChar(char *x, char *y) {
 
	char t;
	 t = *x;
	*x = *y;
	*y =  t;
}
 
__device__	void subPixelAccuracy(int vectorID, float *R, char *R_posX, char *R_posY, float *subXN, float *subXD, float *subYN, float *subYD, float *R_posSubX, float *R_posSubY, char l, char a, float *endX, float *endY) {
 
	int i, goal;
 
	if (threadIdx.y == 0 && threadIdx.x < 5) {
		subXN[threadIdx.x] = 0;
	} else if (threadIdx.y == 1 && threadIdx.x < 5) {
		subXD[threadIdx.x] = 0;
	} else if (threadIdx.y == 2 && threadIdx.x < 5) {
		subYN[threadIdx.x] = 0;
	} else if (threadIdx.y == 3 && threadIdx.x < 5) {
		subYD[threadIdx.x] = 0;
	}
 
	__syncthreads();
 
	for (goal = 0; goal < 5; goal++) {
		for (i = threadIdx.y * blockDim.x + threadIdx.x; i < 4096; i += blockDim.x * blockDim.y) {
			if ((R_posX[i] == R_posX[goal]) && (R_posY[i] == R_posY[goal])) {
				subXD[goal] -= (R[i] + R[i]);
				subYD[goal] -= (R[i] + R[i]);
			} else if ((R_posX[i] == R_posX[goal] - 1) && (R_posY[i] == R_posY[goal])) {
				subXN[goal] -= R[i];
				subXD[goal] += R[i];
			} else if ((R_posX[i] == R_posX[goal] + 1) && (R_posY[i] == R_posY[goal])) {
				subXN[goal] += R[i];
				subXD[goal] += R[i];
			} else if ((R_posX[i] == R_posX[goal]) && (R_posY[i] == R_posY[goal] - 1)) {
				subYN[goal] -= R[i];
				subYD[goal] += R[i];
			} else if ((R_posX[i] == R_posX[goal]) && (R_posY[i] == R_posY[goal] + 1)) {
				subYN[goal] += R[i];
				subYD[goal] += R[i];
			}
		}
	}
 
	__syncthreads();
 
	if (threadIdx.y == 0 && threadIdx.x < 5) {
 
		float subX;
		subX = subXN[threadIdx.x] / (subXD[threadIdx.x] + subXD[threadIdx.x]);
		if (subX > 1 || subX < -1) {
			subX = 0;
		}
		R_posSubX[vectorID * 5 + threadIdx.x] = (float)R_posX[threadIdx.x] + subX - (float)l;
 
	} else if (threadIdx.y == 1 && threadIdx.x < 5) {
		float subY;
		subY = subYN[threadIdx.x] / (subYD[threadIdx.x] + subYD[threadIdx.x]);
		if (subY > 1 || subY < -1) {
			subY = 0;
		}
		R_posSubY[vectorID * 5 + threadIdx.x] = (float)R_posY[threadIdx.x] + subY - (float)a;
 
	}
 
	__syncthreads();
 
	if (threadIdx.y == 0 && threadIdx.x == 0) {
		endX[vectorID] = R_posSubX[vectorID * 5];
 
	} else if (threadIdx.y == 1 && threadIdx.x == 0) {
		endY[vectorID] = R_posSubY[vectorID * 5];
 
	}
	__syncthreads();
}
 
__global__	void vectorCorrection(int totalVectorsX, float *R_posSubX, float *R_posSubY, float *endX, float *endY) {
 
	if ((blockIdx.x != 0) && (blockIdx.x != gridDim.x - 1) && (blockIdx.y != 0) && (blockIdx.y != gridDim.y - 1) && (threadIdx.y == 0) && (threadIdx.x == 0)) {
		if (threadIdx.x == 0 && threadIdx.y == 0) {
			int i, j, k, vectorID;
			float averageX, averageY, dX, dY, X, Y;
 
			vectorID = blockIdx.y * gridDim.x + blockIdx.x;
			averageX = 0;
			averageY = 0;
			X = R_posSubX[vectorID * 5];
			Y = R_posSubY[vectorID * 5];
 
			for (j = -1; j <= 1; j++) {
				for (i = -1; i <= 1; i++) {
					if ((j != 0) || (i != 0)) {
						k = j * totalVectorsX + i;
						averageX += R_posSubX[(vectorID + k) * 5];
						averageY += R_posSubY[(vectorID + k) * 5];
					}
				}
			}
			averageX /= 8;
			averageY /= 8;
 
			for (i = 1; i < 5; i++) {
				dX = R_posSubX[vectorID * 5 + i] - averageX;
				dY = R_posSubY[vectorID * 5 + i] - averageY;
				if (sqrt(dX * dX + dY * dY) < sqrt((X - averageX) * (X - averageX) + (Y - averageY) * (Y - averageY))) {
					X = R_posSubX[vectorID * 5 + i];
					Y = R_posSubY[vectorID * 5 + i];
				}
 
			}
 
			if (sqrt((X - averageX) * (X - averageX) + (Y - averageY) * (Y - averageY)) > 4) {
				X = averageX;
				Y = averageY;
			}
 
			endX[vectorID] = X;
			endY[vectorID] = Y;
 
		}
 
	}
}
 
__host__		void printSummary(int width, int height, char intWndW, char intWndH, char srhWndL, char srhWndR, char srhWndA, char srhWndB, char stepX, char stepY, int totalVectorsX, int totalVectorsY, int totalVectors, int Dbx, int Dby, int Dgx, int Dgy, float h_time, float d_time) {
	int srhWndW, srhWndH, threadsX, threadsY, threads;
	float overlap;
 
	srhWndW = 1 + srhWndL + srhWndR;
	srhWndH = 1 + srhWndA + srhWndB;
	threadsX = Dgx * Dbx;
	threadsY = Dgy * Dby;
	threads = threadsX * threadsY;
	overlap = 1 - ((float)stepX / (float)intWndW);
 
	printf(
		"Summary:\n"
		"Image size:                  %d * %d\n"
		"Interrogation window wize:   %d * %d\n"
		"Search window size:          %d * %d\n"
		"\n"
		"Steps:                       %d * %d\n"
		"Overlap:                     %.2f%% \n"
		"\n"
		"Vector size:                 %d * %d\n"
		"Total vectors                %d\n"
		"\n"
		"Grid size:                   %d * %d\n"
		"Block size:                  %d * %d\n"
		"Thread size:                 %d * %d\n"
		"Total threads:               %d\n"
		"\n"
		"Processing time (with memcpy):	%f sec\n"
		"Processing time (without memcpy): %f sec\n"
		"\n"
		, width, height
		, (int)intWndW, (int)intWndH
		, (int)srhWndW, (int)srhWndH
 
		, (int)stepX, (int)stepY
		, (overlap * 100)
 
		, totalVectorsX, totalVectorsY
		, totalVectors
 
		, Dgx, Dgy
		, Dbx, Dby
		, threadsX, threadsY
		, threads
 
		, h_time
		, d_time
	);
}
 
__host__	void putData(float *startX, float *startY, float *endX, float *endY, int number, char *path) {
	int i;
 
	FILE *fp;
 
	fp = fopen(path, "w");
	if (fp == NULL) {
		fprintf(stderr, "%s file open error\n", path);
	}
 
	for (i = 0; i < number; i++) {	
 
		fprintf(fp, "%.1f	%.1f	%.1f	%.1f\n", startX[i], startY[i], endX[i] * 2, endY[i] * 2);
	}
 
	fclose(fp);
 
}
 
/*vtk output*/
__host__	void putDataVTK(int height, float *startX, float *startY, float *endX, float *endY, int number, int numberX, int numberY, char *pathVTK) {
	int i;
	FILE *fp;
 
	fp = fopen(pathVTK, "w");
	if (fp == NULL) {
		fprintf(stderr, "%s file open error.\n", pathVTK);
	}
 
	fprintf(fp, "# vtk DataFile Version 3.0\nvtk output\nASCII\nDATASET STRUCTURED_GRID\nDIMENSIONS %d %d 1\nPOINTS %d float\n", numberX, numberY, number);
	for (i = 0; i < number; i++) {
		fprintf(fp,"%.1f %.1f %d\n", startX[i], height - startY[i], 0);
	}
 
	fprintf(fp, "POINT_DATA %d\nSCALARS velocity float\nLOOKUP_TABLE default\n", number);
	for (i = 0; i < number; i++) {
		fprintf(fp,"%.10f\n", sqrt((endX[i] * endX[i]) + (endY[i] * endY[i])));
	}
 
	fprintf(fp, "VECTORS velocity float\n");
	for (i = 0; i < number; i++) {
		fprintf(fp,"%.10f %.10f %.10f\n", endX[i], -endY[i], .0);
	}
 
}
	.text
	.file	"main.hip"
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0                          # -- Begin function main
.LCPI0_0:
	.long	0x49742400                      # float 1.0E+6
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0
.LCPI0_1:
	.quad	0x4049000000000000              # double 50
	.text
	.globl	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	$376, %rsp                      # imm = 0x178
	.cfi_def_cfa_offset 432
	.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	$1048576, %edi                  # imm = 0x100000
	callq	malloc
	movq	%rax, %r15
	movl	$1048576, %edi                  # imm = 0x100000
	callq	malloc
	movq	%rax, %r12
	leaq	64(%rsp), %rdi
	movl	$1048576, %esi                  # imm = 0x100000
	callq	hipMalloc
	leaq	56(%rsp), %rdi
	movl	$1048576, %esi                  # imm = 0x100000
	callq	hipMalloc
	movl	$14884, %edi                    # imm = 0x3A24
	callq	malloc
	movq	%rax, 96(%rsp)                  # 8-byte Spill
	movl	$14884, %edi                    # imm = 0x3A24
	callq	malloc
	movq	%rax, 88(%rsp)                  # 8-byte Spill
	movl	$14884, %edi                    # imm = 0x3A24
	callq	malloc
	movq	%rax, %r13
	movl	$14884, %edi                    # imm = 0x3A24
	callq	malloc
	movq	%rax, %rbp
	leaq	48(%rsp), %rdi
	movl	$14884, %esi                    # imm = 0x3A24
	callq	hipMalloc
	leaq	40(%rsp), %rdi
	movl	$14884, %esi                    # imm = 0x3A24
	callq	hipMalloc
	leaq	16(%rsp), %rdi
	movl	$14884, %esi                    # imm = 0x3A24
	callq	hipMalloc
	leaq	8(%rsp), %rdi
	movl	$14884, %esi                    # imm = 0x3A24
	callq	hipMalloc
	leaq	32(%rsp), %rdi
	movl	$74420, %esi                    # imm = 0x122B4
	callq	hipMalloc
	leaq	24(%rsp), %rdi
	movl	$74420, %esi                    # imm = 0x122B4
	callq	hipMalloc
	leaq	304(%rsp), %rbx
	movl	$.L.str.5, %esi
	movl	$.L.str, %edx
	movl	$.L.str.2, %ecx
	movq	%rbx, %rdi
	xorl	%eax, %eax
	callq	sprintf
	leaq	240(%rsp), %r14
	movl	$.L.str.5, %esi
	movl	$.L.str.1, %edx
	movl	$.L.str.2, %ecx
	movq	%r14, %rdi
	xorl	%eax, %eax
	callq	sprintf
	leaq	176(%rsp), %rdi
	movl	$.L.str.5, %esi
	movl	$.L.str.3, %edx
	movl	$.L.str.4, %ecx
	xorl	%eax, %eax
	callq	sprintf
	leaq	112(%rsp), %rdi
	movl	$.L.str.6, %esi
	movl	$.L.str.3, %edx
	xorl	%eax, %eax
	callq	sprintf
	movl	$1048576, %edi                  # imm = 0x100000
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%rbx, %rcx
	movq	%r14, %r8
	callq	_Z7getDataiPhS_PcS0_
	callq	clock
	movq	%rax, 72(%rsp)                  # 8-byte Spill
	leaq	64(%rsp), %rax
	movq	(%rax), %rdi
	movl	$1048576, %edx                  # imm = 0x100000
	movq	%r15, 104(%rsp)                 # 8-byte Spill
	movq	%r15, %rsi
	movabsq	$137438953504, %r15             # imm = 0x2000000020
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	56(%rsp), %rax
	movq	(%rax), %rdi
	movl	$1048576, %edx                  # imm = 0x100000
	movq	%r12, 80(%rsp)                  # 8-byte Spill
	movq	%r12, %rsi
	movabsq	$261993005117, %r12             # imm = 0x3D0000003D
	movl	$1, %ecx
	callq	hipMemcpy
	callq	clock
	movq	%rax, %r14
	movq	%r12, %rdi
	movl	$1, %esi
	movq	%r15, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_2
# %bb.1:
	movq	64(%rsp), %rdi
	movq	56(%rsp), %rsi
	subq	$8, %rsp
	.cfi_adjust_cfa_offset 8
	movl	$16, %ebx
	movl	$56, %r11d
	movl	$25, %r10d
	movl	$12, %eax
	movl	$1024, %edx                     # imm = 0x400
	movl	$61, %ecx
	movl	$61, %r8d
	movl	$32, %r9d
	pushq	16(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	32(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	64(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	80(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	64(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	80(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	%rbx
	.cfi_adjust_cfa_offset 8
	pushq	%rbx
	.cfi_adjust_cfa_offset 8
	pushq	%r11
	.cfi_adjust_cfa_offset 8
	pushq	%r11
	.cfi_adjust_cfa_offset 8
	pushq	%r10
	.cfi_adjust_cfa_offset 8
	pushq	%r10
	.cfi_adjust_cfa_offset 8
	pushq	%rax
	.cfi_adjust_cfa_offset 8
	pushq	%rax
	.cfi_adjust_cfa_offset 8
	pushq	%rax
	.cfi_adjust_cfa_offset 8
	pushq	%rax
	.cfi_adjust_cfa_offset 8
	pushq	$32
	.cfi_adjust_cfa_offset 8
	callq	_Z32__device_stub__vectorComputationPhS_iiiccccccccccccPfS0_S0_S0_S0_S0_
	addq	$144, %rsp
	.cfi_adjust_cfa_offset -144
.LBB0_2:
	callq	hipDeviceSynchronize
	movq	%r12, %rdi
	movl	$1, %esi
	movq	%r15, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_4
# %bb.3:
	movq	32(%rsp), %rsi
	movq	24(%rsp), %rdx
	movq	16(%rsp), %rcx
	movq	8(%rsp), %r8
	movl	$61, %edi
	callq	_Z31__device_stub__vectorCorrectioniPfS_S_S_
.LBB0_4:
	callq	hipDeviceSynchronize
	callq	clock
	movq	%rax, %rbx
	movq	48(%rsp), %rsi
	movl	$14884, %edx                    # imm = 0x3A24
	movq	96(%rsp), %r12                  # 8-byte Reload
	movq	%r12, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movq	40(%rsp), %rsi
	movl	$14884, %edx                    # imm = 0x3A24
	movq	88(%rsp), %r15                  # 8-byte Reload
	movq	%r15, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movq	16(%rsp), %rsi
	movl	$14884, %edx                    # imm = 0x3A24
	movq	%r13, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movq	8(%rsp), %rsi
	movl	$14884, %edx                    # imm = 0x3A24
	movq	%rbp, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	callq	clock
	subq	72(%rsp), %rax                  # 8-byte Folded Reload
	cvtsi2ss	%rax, %xmm0
	movss	.LCPI0_0(%rip), %xmm1           # xmm1 = mem[0],zero,zero,zero
	subq	%r14, %rbx
	cvtsi2ss	%rbx, %xmm2
	divss	%xmm1, %xmm0
	divss	%xmm1, %xmm2
	xorps	%xmm1, %xmm1
	cvtss2sd	%xmm0, %xmm1
	cvtss2sd	%xmm2, %xmm2
	subq	$8, %rsp
	.cfi_adjust_cfa_offset 8
	movl	$1952, %r11d                    # imm = 0x7A0
	movl	$32, %r14d
	movl	$61, %ebx
	movl	$16, %r10d
	movsd	.LCPI0_1(%rip), %xmm0           # xmm0 = mem[0],zero
	movl	$.L.str.11, %edi
	movl	$1024, %esi                     # imm = 0x400
	movl	$1024, %edx                     # imm = 0x400
	movl	$32, %ecx
	movl	$32, %r8d
	movl	$25, %r9d
	movb	$3, %al
	pushq	$3810304                        # imm = 0x3A2400
	.cfi_adjust_cfa_offset 8
	pushq	%r11
	.cfi_adjust_cfa_offset 8
	pushq	%r11
	.cfi_adjust_cfa_offset 8
	pushq	%r14
	.cfi_adjust_cfa_offset 8
	pushq	%r14
	.cfi_adjust_cfa_offset 8
	pushq	%rbx
	.cfi_adjust_cfa_offset 8
	pushq	%rbx
	.cfi_adjust_cfa_offset 8
	pushq	$3721                           # imm = 0xE89
	.cfi_adjust_cfa_offset 8
	pushq	%rbx
	.cfi_adjust_cfa_offset 8
	pushq	%rbx
	.cfi_adjust_cfa_offset 8
	pushq	%r10
	.cfi_adjust_cfa_offset 8
	pushq	%r10
	.cfi_adjust_cfa_offset 8
	pushq	$25
	.cfi_adjust_cfa_offset 8
	callq	printf
	addq	$112, %rsp
	.cfi_adjust_cfa_offset -112
	leaq	176(%rsp), %r9
	movq	%r12, %rdi
	movq	%r15, %rsi
	movq	%r13, %rdx
	movq	%rbp, %rcx
	movl	$3721, %r8d                     # imm = 0xE89
	callq	_Z7putDataPfS_S_S_iPc
	subq	$8, %rsp
	.cfi_adjust_cfa_offset 8
	movl	$1024, %edi                     # imm = 0x400
	movq	%r12, %rsi
	movq	%r15, %rdx
	movq	%r13, %rcx
	movq	%rbp, %r8
	movl	$3721, %r9d                     # imm = 0xE89
	leaq	120(%rsp), %rax
	pushq	%rax
	.cfi_adjust_cfa_offset 8
	pushq	%rbx
	.cfi_adjust_cfa_offset 8
	pushq	%rbx
	.cfi_adjust_cfa_offset 8
	callq	_Z10putDataVTKiPfS_S_S_iiiPc
	addq	$32, %rsp
	.cfi_adjust_cfa_offset -32
	movq	104(%rsp), %rdi                 # 8-byte Reload
	callq	free
	movq	80(%rsp), %rdi                  # 8-byte Reload
	callq	free
	movq	64(%rsp), %rdi
	callq	hipFree
	movq	56(%rsp), %rdi
	callq	hipFree
	movq	%r12, %rdi
	callq	free
	movq	%r15, %rdi
	callq	free
	movq	%r13, %rdi
	callq	free
	movq	%rbp, %rdi
	callq	free
	movq	48(%rsp), %rdi
	callq	hipFree
	movq	40(%rsp), %rdi
	callq	hipFree
	movq	16(%rsp), %rdi
	callq	hipFree
	movq	8(%rsp), %rdi
	callq	hipFree
	movq	32(%rsp), %rdi
	callq	hipFree
	movq	24(%rsp), %rdi
	callq	hipFree
	callq	hipGetLastError
	testl	%eax, %eax
	jne	.LBB0_6
# %bb.5:
	movl	$.Lstr, %edi
	callq	puts@PLT
	xorl	%eax, %eax
	addq	$376, %rsp                      # imm = 0x178
	.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
.LBB0_6:
	.cfi_def_cfa_offset 432
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.7, %edi
	movq	%rax, %rsi
	xorl	%eax, %eax
	callq	printf
	movl	$-1, %edi
	callq	exit
.Lfunc_end0:
	.size	main, .Lfunc_end0-main
	.cfi_endproc
                                        # -- End function
	.globl	_Z7getDataiPhS_PcS0_            # -- Begin function _Z7getDataiPhS_PcS0_
	.type	_Z7getDataiPhS_PcS0_,@function
_Z7getDataiPhS_PcS0_:                   # @_Z7getDataiPhS_PcS0_
	.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
	pushq	%rax
	.cfi_def_cfa_offset 64
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movq	%r8, (%rsp)                     # 8-byte Spill
	movq	%rcx, %r13
	movq	%rdx, %rbx
	movq	%rsi, %r15
	movl	%edi, %ebp
	movl	$.L.str.9, %esi
	movq	%rcx, %rdi
	callq	fopen
	movq	%rax, %r12
	testq	%rax, %rax
	jne	.LBB1_2
# %bb.1:
	movl	$.L.str.10, %edi
	movq	%r13, %rsi
	xorl	%eax, %eax
	callq	printf
.LBB1_2:
	movl	%ebp, %r13d
	testl	%ebp, %ebp
	jle	.LBB1_5
# %bb.3:                                # %.lr.ph.preheader
	xorl	%r14d, %r14d
.LBB1_4:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movq	%r12, %rdi
	callq	fgetc
	movb	%al, (%r15,%r14)
	incq	%r14
	cmpq	%r14, %r13
	jne	.LBB1_4
.LBB1_5:                                # %._crit_edge
	movq	%r12, %rdi
	callq	fclose
	movl	$.L.str.9, %esi
	movq	(%rsp), %r14                    # 8-byte Reload
	movq	%r14, %rdi
	callq	fopen
	movq	%rax, %r15
	testq	%rax, %rax
	jne	.LBB1_7
# %bb.6:
	movl	$.L.str.10, %edi
	movq	%r14, %rsi
	xorl	%eax, %eax
	callq	printf
.LBB1_7:
	testl	%ebp, %ebp
	jle	.LBB1_10
# %bb.8:                                # %.lr.ph23.preheader
	xorl	%r14d, %r14d
.LBB1_9:                                # %.lr.ph23
                                        # =>This Inner Loop Header: Depth=1
	movq	%r15, %rdi
	callq	fgetc
	movb	%al, (%rbx,%r14)
	incq	%r14
	cmpq	%r14, %r13
	jne	.LBB1_9
.LBB1_10:                               # %._crit_edge24
	movq	%r15, %rdi
	addq	$8, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	jmp	fclose                          # TAILCALL
.Lfunc_end1:
	.size	_Z7getDataiPhS_PcS0_, .Lfunc_end1-_Z7getDataiPhS_PcS0_
	.cfi_endproc
                                        # -- End function
	.globl	_Z32__device_stub__vectorComputationPhS_iiiccccccccccccPfS0_S0_S0_S0_S0_ # -- Begin function _Z32__device_stub__vectorComputationPhS_iiiccccccccccccPfS0_S0_S0_S0_S0_
	.type	_Z32__device_stub__vectorComputationPhS_iiiccccccccccccPfS0_S0_S0_S0_S0_,@function
_Z32__device_stub__vectorComputationPhS_iiiccccccccccccPfS0_S0_S0_S0_S0_: # @_Z32__device_stub__vectorComputationPhS_iiiccccccccccccPfS0_S0_S0_S0_S0_
	.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	$272, %rsp                      # imm = 0x110
	.cfi_def_cfa_offset 320
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	40(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	32(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	12(%rsp), %rsi
	movl	%edx, (%rsi)
	leaq	8(%rsp), %rdx
	movl	%ecx, (%rdx)
	leaq	4(%rsp), %rcx
	movl	%r8d, (%rcx)
	leaq	3(%rsp), %r8
	movb	%r9b, (%r8)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%rbx)
	movq	%r8, 40(%rbx)
	leaq	320(%rsp), %rax
	movq	%rax, 48(%rbx)
	leaq	328(%rsp), %rax
	movq	%rax, 56(%rbx)
	leaq	336(%rsp), %rax
	movq	%rax, 64(%rbx)
	leaq	344(%rsp), %rax
	movq	%rax, 72(%rbx)
	leaq	352(%rsp), %rax
	movq	%rax, 80(%rbx)
	leaq	360(%rsp), %rax
	movq	%rax, 88(%rbx)
	leaq	368(%rsp), %rax
	movq	%rax, 96(%rbx)
	leaq	376(%rsp), %rax
	movq	%rax, 104(%rbx)
	leaq	384(%rsp), %rax
	movq	%rax, 112(%rbx)
	leaq	392(%rsp), %rax
	movq	%rax, 120(%rbx)
	leaq	400(%rsp), %rax
	movq	%rax, 128(%rbx)
	leaq	408(%rsp), %rax
	movq	%rax, 136(%rbx)
	leaq	416(%rsp), %rax
	movq	%rax, 144(%rbx)
	leaq	424(%rsp), %rax
	movq	%rax, 152(%rbx)
	leaq	432(%rsp), %rax
	movq	%rax, 160(%rbx)
	leaq	440(%rsp), %rax
	movq	%rax, 168(%rbx)
	leaq	448(%rsp), %rax
	movq	%rax, 176(%rbx)
	leaq	64(%rsp), %r14
	leaq	48(%rsp), %r15
	leaq	24(%rsp), %r12
	leaq	16(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z17vectorComputationPhS_iiiccccccccccccPfS0_S0_S0_S0_S0_, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$288, %rsp                      # imm = 0x120
	.cfi_adjust_cfa_offset -288
	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	_Z32__device_stub__vectorComputationPhS_iiiccccccccccccPfS0_S0_S0_S0_S0_, .Lfunc_end2-_Z32__device_stub__vectorComputationPhS_iiiccccccccccccPfS0_S0_S0_S0_S0_
	.cfi_endproc
                                        # -- End function
	.globl	_Z31__device_stub__vectorCorrectioniPfS_S_S_ # -- Begin function _Z31__device_stub__vectorCorrectioniPfS_S_S_
	.type	_Z31__device_stub__vectorCorrectioniPfS_S_S_,@function
_Z31__device_stub__vectorCorrectioniPfS_S_S_: # @_Z31__device_stub__vectorCorrectioniPfS_S_S_
	.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	$144, %rsp
	.cfi_def_cfa_offset 192
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	12(%rsp), %rax
	movl	%edi, (%rax)
	leaq	56(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	48(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	40(%rsp), %rdx
	movq	%rcx, (%rdx)
	leaq	32(%rsp), %rcx
	movq	%r8, (%rcx)
	leaq	96(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%rbx)
	leaq	80(%rsp), %r14
	leaq	64(%rsp), %r15
	leaq	24(%rsp), %r12
	leaq	16(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z16vectorCorrectioniPfS_S_S_, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$160, %rsp
	.cfi_adjust_cfa_offset -160
	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	_Z31__device_stub__vectorCorrectioniPfS_S_S_, .Lfunc_end3-_Z31__device_stub__vectorCorrectioniPfS_S_S_
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0                          # -- Begin function _Z12printSummaryiicccccccciiiiiiiff
.LCPI4_0:
	.long	0x3f800000                      # float 1
.LCPI4_1:
	.long	0x42c80000                      # float 100
	.text
	.globl	_Z12printSummaryiicccccccciiiiiiiff
	.type	_Z12printSummaryiicccccccciiiiiiiff,@function
_Z12printSummaryiicccccccciiiiiiiff:    # @_Z12printSummaryiicccccccciiiiiiiff
	.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
	pushq	%rax
	.cfi_def_cfa_offset 64
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
                                        # kill: def $r9d killed $r9d def $r9
                                        # kill: def $r8d killed $r8d def $r8
	movl	%ecx, %ebp
	movl	%edx, %ecx
	movl	%esi, %edx
	movl	%edi, %esi
	movsbl	88(%rsp), %r12d
	movsbl	80(%rsp), %r11d
	movl	128(%rsp), %r13d
	movl	144(%rsp), %eax
	movl	120(%rsp), %r10d
	movl	136(%rsp), %r14d
	movsbl	64(%rsp), %edi
	movsbl	72(%rsp), %ebx
	addl	%r8d, %r9d
	incl	%r9d
	addl	%edi, %ebx
	incl	%ebx
                                        # kill: def $r14d killed $r14d killed $r14 def $r14
	imull	%r10d, %r14d
	movl	%eax, %r15d
	imull	%r13d, %r15d
	cvtsi2ss	%r11d, %xmm2
	movl	%r15d, %r10d
	cvtsi2ss	%ecx, %xmm3
	imull	%r14d, %r10d
	divss	%xmm3, %xmm2
	movss	.LCPI4_0(%rip), %xmm3           # xmm3 = mem[0],zero,zero,zero
	subss	%xmm2, %xmm3
	mulss	.LCPI4_1(%rip), %xmm3
	cvtss2sd	%xmm3, %xmm4
	xorps	%xmm3, %xmm3
	cvtss2sd	%xmm0, %xmm3
	xorps	%xmm2, %xmm2
	cvtss2sd	%xmm1, %xmm2
	subq	$8, %rsp
	.cfi_adjust_cfa_offset 8
	movl	$.L.str.11, %edi
	movl	%ebp, %r8d
	movaps	%xmm4, %xmm0
	movaps	%xmm3, %xmm1
	movb	$3, %al
	pushq	%r10
	.cfi_adjust_cfa_offset 8
	pushq	%r15
	.cfi_adjust_cfa_offset 8
	pushq	%r14
	.cfi_adjust_cfa_offset 8
	pushq	%r13
	.cfi_adjust_cfa_offset 8
	movl	160(%rsp), %r10d
	pushq	%r10
	.cfi_adjust_cfa_offset 8
	movl	192(%rsp), %r10d
	pushq	%r10
	.cfi_adjust_cfa_offset 8
	movl	192(%rsp), %r10d
	pushq	%r10
	.cfi_adjust_cfa_offset 8
	movl	176(%rsp), %r10d
	pushq	%r10
	.cfi_adjust_cfa_offset 8
	movl	176(%rsp), %r10d
	pushq	%r10
	.cfi_adjust_cfa_offset 8
	movl	176(%rsp), %r10d
	pushq	%r10
	.cfi_adjust_cfa_offset 8
	pushq	%r12
	.cfi_adjust_cfa_offset 8
	pushq	%r11
	.cfi_adjust_cfa_offset 8
	pushq	%rbx
	.cfi_adjust_cfa_offset 8
	callq	printf
	addq	$120, %rsp
	.cfi_adjust_cfa_offset -120
	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	_Z12printSummaryiicccccccciiiiiiiff, .Lfunc_end4-_Z12printSummaryiicccccccciiiiiiiff
	.cfi_endproc
                                        # -- End function
	.globl	_Z7putDataPfS_S_S_iPc           # -- Begin function _Z7putDataPfS_S_S_iPc
	.type	_Z7putDataPfS_S_S_iPc,@function
_Z7putDataPfS_S_S_iPc:                  # @_Z7putDataPfS_S_S_iPc
	.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
	pushq	%rax
	.cfi_def_cfa_offset 64
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movq	%r9, %r13
	movl	%r8d, %ebp
	movq	%rcx, %r14
	movq	%rdx, %r15
	movq	%rsi, %r12
	movq	%rdi, %rbx
	movl	$.L.str.12, %esi
	movq	%r9, %rdi
	callq	fopen
	testq	%rax, %rax
	je	.LBB5_1
.LBB5_2:
	testl	%ebp, %ebp
	jle	.LBB5_5
# %bb.3:                                # %.lr.ph.preheader
	movq	%rbx, %rdi
	movl	%ebp, %r8d
	movq	%r8, (%rsp)                     # 8-byte Spill
	xorl	%ebx, %ebx
.LBB5_4:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	xorps	%xmm0, %xmm0
	cvtss2sd	(%rdi,%rbx,4), %xmm0
	xorps	%xmm1, %xmm1
	cvtss2sd	(%r12,%rbx,4), %xmm1
	movss	(%r15,%rbx,4), %xmm2            # xmm2 = mem[0],zero,zero,zero
	addss	%xmm2, %xmm2
	cvtss2sd	%xmm2, %xmm2
	movss	(%r14,%rbx,4), %xmm3            # xmm3 = mem[0],zero,zero,zero
	addss	%xmm3, %xmm3
	cvtss2sd	%xmm3, %xmm3
	movl	$.L.str.13, %esi
	movq	%rax, %r13
	movq	%rdi, %rbp
	movq	%rax, %rdi
	movb	$4, %al
	callq	fprintf
	movq	%rbp, %rdi
	movq	%r13, %rax
	incq	%rbx
	cmpq	%rbx, (%rsp)                    # 8-byte Folded Reload
	jne	.LBB5_4
.LBB5_5:                                # %._crit_edge
	movq	%rax, %rdi
	addq	$8, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	jmp	fclose                          # TAILCALL
.LBB5_1:
	.cfi_def_cfa_offset 64
	movq	%r13, %rdx
	movq	%rax, %r13
	movq	stderr(%rip), %rdi
	movl	$.L.str.10, %esi
	xorl	%eax, %eax
	callq	fprintf
	movq	%r13, %rax
	jmp	.LBB5_2
.Lfunc_end5:
	.size	_Z7putDataPfS_S_S_iPc, .Lfunc_end5-_Z7putDataPfS_S_S_iPc
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst16,"aM",@progbits,16
	.p2align	4, 0x0                          # -- Begin function _Z10putDataVTKiPfS_S_S_iiiPc
.LCPI6_0:
	.long	0x80000000                      # float -0
	.long	0x80000000                      # float -0
	.long	0x80000000                      # float -0
	.long	0x80000000                      # float -0
	.text
	.globl	_Z10putDataVTKiPfS_S_S_iiiPc
	.type	_Z10putDataVTKiPfS_S_S_iiiPc,@function
_Z10putDataVTKiPfS_S_S_iiiPc:           # @_Z10putDataVTKiPfS_S_S_iiiPc
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	subq	$24, %rsp
	.cfi_def_cfa_offset 80
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movl	%r9d, %r14d
	movq	%r8, 16(%rsp)                   # 8-byte Spill
	movq	%rcx, %r15
	movq	%rdx, %r12
	movq	%rsi, %r13
	movl	%edi, 12(%rsp)                  # 4-byte Spill
	movq	96(%rsp), %rbp
	movl	88(%rsp), %eax
	movl	%eax, 8(%rsp)                   # 4-byte Spill
	movl	80(%rsp), %eax
	movl	%eax, 4(%rsp)                   # 4-byte Spill
	movl	$.L.str.12, %esi
	movq	%rbp, %rdi
	callq	fopen
	movq	%rax, %rbx
	testq	%rax, %rax
	je	.LBB6_1
.LBB6_2:
	movl	$.L.str.15, %esi
	movq	%rbx, %rdi
	movl	4(%rsp), %edx                   # 4-byte Reload
	movl	8(%rsp), %ecx                   # 4-byte Reload
	movl	%r14d, %r8d
	xorl	%eax, %eax
	callq	fprintf
	testl	%r14d, %r14d
	jle	.LBB6_3
# %bb.5:                                # %.lr.ph
	cvtsi2ssl	12(%rsp), %xmm0         # 4-byte Folded Reload
	movss	%xmm0, 8(%rsp)                  # 4-byte Spill
	movl	%r14d, 4(%rsp)                  # 4-byte Spill
	movl	%r14d, %ebp
	xorl	%r14d, %r14d
.LBB6_6:                                # =>This Inner Loop Header: Depth=1
	xorps	%xmm0, %xmm0
	cvtss2sd	(%r13,%r14,4), %xmm0
	movss	8(%rsp), %xmm1                  # 4-byte Reload
                                        # xmm1 = mem[0],zero,zero,zero
	subss	(%r12,%r14,4), %xmm1
	cvtss2sd	%xmm1, %xmm1
	movl	$.L.str.16, %esi
	movq	%rbx, %rdi
	xorl	%edx, %edx
	movb	$2, %al
	callq	fprintf
	incq	%r14
	cmpq	%r14, %rbp
	jne	.LBB6_6
# %bb.7:                                # %._crit_edge
	movl	$.L.str.17, %esi
	movq	%rbx, %rdi
	movl	4(%rsp), %r14d                  # 4-byte Reload
	movl	%r14d, %edx
	xorl	%eax, %eax
	callq	fprintf
	testl	%r14d, %r14d
	movq	16(%rsp), %r14                  # 8-byte Reload
	jle	.LBB6_4
# %bb.8:                                # %.lr.ph42.preheader
	xorl	%r12d, %r12d
.LBB6_9:                                # %.lr.ph42
                                        # =>This Inner Loop Header: Depth=1
	movss	(%r15,%r12,4), %xmm0            # xmm0 = mem[0],zero,zero,zero
	mulss	%xmm0, %xmm0
	movss	(%r14,%r12,4), %xmm1            # xmm1 = mem[0],zero,zero,zero
	mulss	%xmm1, %xmm1
	addss	%xmm0, %xmm1
	xorps	%xmm0, %xmm0
	sqrtss	%xmm1, %xmm0
	cvtss2sd	%xmm0, %xmm0
	movl	$.L.str.18, %esi
	movq	%rbx, %rdi
	movb	$1, %al
	callq	fprintf
	incq	%r12
	cmpq	%r12, %rbp
	jne	.LBB6_9
# %bb.10:                               # %._crit_edge43
	movl	$.L.str.19, %edi
	movl	$23, %esi
	movl	$1, %edx
	movq	%rbx, %rcx
	callq	fwrite@PLT
	cmpl	$0, 4(%rsp)                     # 4-byte Folded Reload
	jle	.LBB6_13
# %bb.11:                               # %.lr.ph46.preheader
	xorl	%r12d, %r12d
.LBB6_12:                               # %.lr.ph46
                                        # =>This Inner Loop Header: Depth=1
	xorps	%xmm0, %xmm0
	cvtss2sd	(%r15,%r12,4), %xmm0
	movss	(%r14,%r12,4), %xmm1            # xmm1 = mem[0],zero,zero,zero
	xorps	.LCPI6_0(%rip), %xmm1
	cvtss2sd	%xmm1, %xmm1
	movl	$.L.str.20, %esi
	xorps	%xmm2, %xmm2
	movq	%rbx, %rdi
	movb	$3, %al
	callq	fprintf
	incq	%r12
	cmpq	%r12, %rbp
	jne	.LBB6_12
.LBB6_13:                               # %._crit_edge47
	addq	$24, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.LBB6_3:                                # %._crit_edge.thread
	.cfi_def_cfa_offset 80
	movl	$.L.str.17, %esi
	movq	%rbx, %rdi
	movl	%r14d, %edx
	xorl	%eax, %eax
	callq	fprintf
.LBB6_4:                                # %._crit_edge43.thread
	movl	$.L.str.19, %edi
	movl	$23, %esi
	movl	$1, %edx
	movq	%rbx, %rcx
	addq	$24, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	jmp	fwrite@PLT                      # TAILCALL
.LBB6_1:
	.cfi_def_cfa_offset 80
	movq	stderr(%rip), %rdi
	movl	$.L.str.14, %esi
	movq	%rbp, %rdx
	xorl	%eax, %eax
	callq	fprintf
	jmp	.LBB6_2
.Lfunc_end6:
	.size	_Z10putDataVTKiPfS_S_S_iiiPc, .Lfunc_end6-_Z10putDataVTKiPfS_S_S_iiiPc
	.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	.LBB7_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB7_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z17vectorComputationPhS_iiiccccccccccccPfS0_S0_S0_S0_S0_, %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	$_Z16vectorCorrectioniPfS_S_S_, %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_end7:
	.size	__hip_module_ctor, .Lfunc_end7-__hip_module_ctor
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_dtor
	.type	__hip_module_dtor,@function
__hip_module_dtor:                      # @__hip_module_dtor
	.cfi_startproc
# %bb.0:
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	je	.LBB8_2
# %bb.1:
	pushq	%rax
	.cfi_def_cfa_offset 16
	callq	__hipUnregisterFatBinary
	movq	$0, __hip_gpubin_handle(%rip)
	addq	$8, %rsp
	.cfi_def_cfa_offset 8
.LBB8_2:
	retq
.Lfunc_end8:
	.size	__hip_module_dtor, .Lfunc_end8-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"2ms2000018"
	.size	.L.str, 11

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

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

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	"result"
	.size	.L.str.3, 7

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

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

	.type	.L.str.6,@object                # @.str.6
.L.str.6:
	.asciz	"%s-vtk.vtk"
	.size	.L.str.6, 11

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

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

	.type	.L.str.7,@object                # @.str.7
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str.7:
	.asciz	"CUDA error: %s\n\n"
	.size	.L.str.7, 17

	.type	.L.str.9,@object                # @.str.9
.L.str.9:
	.asciz	"rb"
	.size	.L.str.9, 3

	.type	.L.str.10,@object               # @.str.10
.L.str.10:
	.asciz	"%s file open error\n"
	.size	.L.str.10, 20

	.type	.L.str.11,@object               # @.str.11
.L.str.11:
	.asciz	"Summary:\nImage size:                  %d * %d\nInterrogation window wize:   %d * %d\nSearch window size:          %d * %d\n\nSteps:                       %d * %d\nOverlap:                     %.2f%% \n\nVector size:                 %d * %d\nTotal vectors                %d\n\nGrid size:                   %d * %d\nBlock size:                  %d * %d\nThread size:                 %d * %d\nTotal threads:               %d\n\nProcessing time (with memcpy):\t%f sec\nProcessing time (without memcpy): %f sec\n\n"
	.size	.L.str.11, 491

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

	.type	.L.str.13,@object               # @.str.13
.L.str.13:
	.asciz	"%.1f\t%.1f\t%.1f\t%.1f\n"
	.size	.L.str.13, 21

	.type	.L.str.14,@object               # @.str.14
.L.str.14:
	.asciz	"%s file open error.\n"
	.size	.L.str.14, 21

	.type	.L.str.15,@object               # @.str.15
.L.str.15:
	.asciz	"# vtk DataFile Version 3.0\nvtk output\nASCII\nDATASET STRUCTURED_GRID\nDIMENSIONS %d %d 1\nPOINTS %d float\n"
	.size	.L.str.15, 104

	.type	.L.str.16,@object               # @.str.16
.L.str.16:
	.asciz	"%.1f %.1f %d\n"
	.size	.L.str.16, 14

	.type	.L.str.17,@object               # @.str.17
.L.str.17:
	.asciz	"POINT_DATA %d\nSCALARS velocity float\nLOOKUP_TABLE default\n"
	.size	.L.str.17, 59

	.type	.L.str.18,@object               # @.str.18
.L.str.18:
	.asciz	"%.10f\n"
	.size	.L.str.18, 7

	.type	.L.str.19,@object               # @.str.19
.L.str.19:
	.asciz	"VECTORS velocity float\n"
	.size	.L.str.19, 24

	.type	.L.str.20,@object               # @.str.20
.L.str.20:
	.asciz	"%.10f %.10f %.10f\n"
	.size	.L.str.20, 19

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z17vectorComputationPhS_iiiccccccccccccPfS0_S0_S0_S0_S0_"
	.size	.L__unnamed_1, 58

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"_Z16vectorCorrectioniPfS_S_S_"
	.size	.L__unnamed_2, 30

	.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	"CUDA success\n"
	.size	.Lstr, 14

	.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 _Z32__device_stub__vectorComputationPhS_iiiccccccccccccPfS0_S0_S0_S0_S0_
	.addrsig_sym _Z31__device_stub__vectorCorrectioniPfS_S_S_
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z17vectorComputationPhS_iiiccccccccccccPfS0_S0_S0_S0_S0_
	.addrsig_sym _Z16vectorCorrectioniPfS_S_S_
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z17vectorComputationPhS_iiiccccccccccccPfS0_S0_S0_S0_S0_ ; -- Begin function _Z17vectorComputationPhS_iiiccccccccccccPfS0_S0_S0_S0_S0_
	.globl	_Z17vectorComputationPhS_iiiccccccccccccPfS0_S0_S0_S0_S0_
	.p2align	8
	.type	_Z17vectorComputationPhS_iiiccccccccccccPfS0_S0_S0_S0_S0_,@function
_Z17vectorComputationPhS_iiiccccccccccccPfS0_S0_S0_S0_S0_: ; @_Z17vectorComputationPhS_iiiccccccccccccPfS0_S0_S0_S0_S0_
; %bb.0:
	s_clause 0x7
	s_load_b32 s33, s[0:1], 0x1c
	s_load_b32 s31, s[0:1], 0x58
	s_load_b64 s[26:27], s[0:1], 0x1c
	s_load_b32 s38, s[0:1], 0x24
	s_load_b128 s[16:19], s[0:1], 0x48
	s_load_b128 s[20:23], s[0:1], 0x0
	s_load_b32 s36, s[0:1], 0x10
	s_load_b256 s[4:11], s[0:1], 0x28
	v_bfe_u32 v5, v0, 10, 10
	v_bfe_i32 v4, v0, 0, 8
	s_delay_alu instid0(VALU_DEP_2)
	v_bfe_i32 v3, v5, 0, 8
	s_waitcnt lgkmcnt(0)
	s_lshr_b32 s35, s33, 8
	v_lshrrev_b16 v2, 7, s33
	v_lshrrev_b16 v6, 7, s35
	v_mov_b32_e32 v1, 0
	s_mul_i32 s2, s31, s15
	s_add_u32 s28, s0, 0x58
	v_and_b32_e32 v2, 1, v2
	v_and_b32_e32 v6, 1, v6
	s_addc_u32 s29, s1, 0
	s_add_i32 s12, s2, s14
	s_bfe_i32 s2, s38, 0x80010
	v_add_nc_u16 v2, s33, v2
	v_add_nc_u16 v6, s35, v6
	s_ashr_i32 s13, s38, 24
	s_bfe_i32 s24, s26, 0x80010
	s_mul_i32 s30, s14, s2
	v_bfe_i32 v2, v2, 0, 8
	v_bfe_i32 v6, v6, 0, 8
	s_mul_i32 s39, s15, s13
	s_ashr_i32 s13, s12, 31
	s_add_i32 s34, s30, s24
	v_ashrrev_i16 v2, 1, v2
	v_ashrrev_i16 v6, 1, v6
	s_lshl_b64 s[24:25], s[12:13], 2
	s_sext_i32_i8 s3, s27
	s_add_u32 s8, s8, s24
	v_bfe_i32 v2, v2, 0, 16
	v_bfe_i32 v6, v6, 0, 16
	s_addc_u32 s9, s9, s25
	s_add_i32 s3, s39, s3
	s_add_u32 s40, s10, s24
	v_add_nc_u32_e32 v2, s34, v2
	v_add_nc_u32_e32 v6, s3, v6
	s_addc_u32 s41, s11, s25
	s_add_u32 s42, s16, s24
	s_addc_u32 s43, s17, s25
	v_cvt_f32_i32_e32 v2, v2
	v_cvt_f32_i32_e32 v6, v6
	s_add_u32 s44, s18, s24
	s_addc_u32 s45, s19, s25
	s_bfe_i32 s37, s33, 0x80008
	v_add_f32_e32 v2, 0.5, v2
	v_add_f32_e32 v6, 0.5, v6
	s_sext_i32_i8 s10, s33
	s_clause 0x3
	global_store_b32 v1, v2, s[8:9]
	global_store_b32 v1, v6, s[40:41]
	global_store_b32 v1, v1, s[42:43]
	global_store_b32 v1, v1, s[44:45]
	s_mov_b32 s8, exec_lo
	v_cmpx_gt_i32_e64 s37, v3
	s_cbranch_execz .LBB0_6
; %bb.1:                                ; %.lr.ph26.i
	s_load_b32 s2, s[28:29], 0xc
	v_bfe_i32 v6, v0, 0, 8
	v_mov_b32_e32 v7, v3
	s_mov_b32 s11, 0
	s_delay_alu instid0(VALU_DEP_2)
	v_cmp_gt_i32_e32 vcc_lo, s10, v6
	s_waitcnt lgkmcnt(0)
	s_lshr_b32 s9, s2, 16
.LBB0_2:                                ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB0_4 Depth 2
	s_and_saveexec_b32 s13, vcc_lo
	s_cbranch_execz .LBB0_5
; %bb.3:                                ; %.lr.ph.i
                                        ;   in Loop: Header=BB0_2 Depth=1
	s_load_b32 s2, s[28:29], 0xc
	v_add_nc_u32_e32 v8, s3, v7
	s_mov_b32 s41, 0
	s_delay_alu instid0(VALU_DEP_1)
	v_mad_u64_u32 v[1:2], null, v8, s36, s[34:35]
	v_mad_i32_i24 v2, v7, s10, 0x8f50
	v_mov_b32_e32 v8, v6
	s_waitcnt lgkmcnt(0)
	s_and_b32 s40, s2, 0xffff
.LBB0_4:                                ;   Parent Loop BB0_2 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v9, v1, v8
	v_add_nc_u32_e32 v11, v2, v8
	v_ashrrev_i32_e32 v10, 31, v9
	v_add_co_u32 v9, s2, s20, v9
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_add_co_ci_u32_e64 v10, s2, s21, v10, s2
	global_load_u8 v9, v[9:10], off
	v_add_nc_u32_e32 v10, s40, v8
	v_bfe_i32 v8, v10, 0, 8
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cmp_le_i32_e64 s2, s10, v8
	s_or_b32 s41, s2, s41
	s_waitcnt vmcnt(0)
	ds_store_b8 v11, v9
	s_and_not1_b32 exec_lo, exec_lo, s41
	s_cbranch_execnz .LBB0_4
.LBB0_5:                                ; %Flow226
                                        ;   in Loop: Header=BB0_2 Depth=1
	s_or_b32 exec_lo, exec_lo, s13
	v_add_nc_u32_e32 v1, s9, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_bfe_i32 v7, v1, 0, 8
	v_cmp_le_i32_e64 s2, s37, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_or_b32 s11, s2, s11
	s_and_not1_b32 exec_lo, exec_lo, s11
	s_cbranch_execnz .LBB0_2
.LBB0_6:                                ; %Flow228
	s_or_b32 exec_lo, exec_lo, s8
	s_load_b64 s[8:9], s[0:1], 0x20
	s_lshr_b32 s3, s26, 16
	s_bfe_i32 s11, s38, 0x80008
	s_mov_b32 s13, exec_lo
	s_waitcnt lgkmcnt(0)
	s_sext_i32_i8 s9, s9
	v_cmpx_gt_i32_e64 s11, v3
	s_cbranch_execz .LBB0_12
; %bb.7:                                ; %.lr.ph24.i
	s_load_b32 s2, s[28:29], 0xc
	v_cmp_gt_i32_e32 vcc_lo, s9, v4
	v_mov_b32_e32 v6, v3
	s_mov_b32 s21, 0
	s_waitcnt lgkmcnt(0)
	s_lshr_b32 s20, s2, 16
.LBB0_8:                                ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB0_10 Depth 2
	s_and_saveexec_b32 s26, vcc_lo
	s_cbranch_execz .LBB0_11
; %bb.9:                                ; %.lr.ph.i56
                                        ;   in Loop: Header=BB0_8 Depth=1
	s_load_b32 s2, s[28:29], 0xc
	v_add_nc_u32_e32 v7, s39, v6
	s_mov_b32 s38, 0
	s_delay_alu instid0(VALU_DEP_1)
	v_mad_u64_u32 v[1:2], null, v7, s36, s[30:31]
	v_mad_i32_i24 v2, v6, s9, 0x4000
	v_mov_b32_e32 v7, v4
	s_waitcnt lgkmcnt(0)
	s_and_b32 s34, s2, 0xffff
.LBB0_10:                               ;   Parent Loop BB0_8 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v8, v1, v7
	v_add_nc_u32_e32 v10, v2, v7
	v_ashrrev_i32_e32 v9, 31, v8
	v_add_co_u32 v8, s2, s22, v8
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_add_co_ci_u32_e64 v9, s2, s23, v9, s2
	global_load_u8 v8, v[8:9], off
	v_add_nc_u32_e32 v9, s34, v7
	v_bfe_i32 v7, v9, 0, 8
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cmp_le_i32_e64 s2, s9, v7
	s_or_b32 s38, s2, s38
	s_waitcnt vmcnt(0)
	ds_store_b8 v10, v8
	s_and_not1_b32 exec_lo, exec_lo, s38
	s_cbranch_execnz .LBB0_10
.LBB0_11:                               ; %Flow222
                                        ;   in Loop: Header=BB0_8 Depth=1
	s_or_b32 exec_lo, exec_lo, s26
	v_add_nc_u32_e32 v1, s20, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_bfe_i32 v6, v1, 0, 8
	v_cmp_le_i32_e64 s2, s11, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_or_b32 s21, s2, s21
	s_and_not1_b32 exec_lo, exec_lo, s21
	s_cbranch_execnz .LBB0_8
.LBB0_12:                               ; %Flow224
	s_or_b32 exec_lo, exec_lo, s13
	s_load_b32 s0, s[0:1], 0x20
	v_and_b32_e32 v4, 0x3ff, v0
	s_ashr_i32 s2, s8, 24
	s_mov_b32 s11, exec_lo
	s_waitcnt lgkmcnt(0)
	s_waitcnt_vscnt null, 0x0
	s_barrier
	buffer_gl0_inv
	s_bfe_i32 s8, s0, 0x80010
	v_cmpx_gt_i32_e64 s2, v3
	s_cbranch_execz .LBB0_25
; %bb.13:                               ; %.lr.ph100.i
	s_load_b32 s1, s[28:29], 0xc
	v_bfe_i32 v6, v4, 0, 8
	s_cmp_gt_i32 s37, 0
	v_mov_b32_e32 v7, v5
	s_cselect_b32 s20, -1, 0
	s_cmp_gt_i32 s10, 0
	v_cmp_gt_i32_e64 s0, s8, v6
	s_mov_b32 s13, 0
	s_cselect_b32 s21, -1, 0
	s_and_b32 s22, s35, 0xff
	s_and_b32 s26, s33, 0xff
	s_waitcnt lgkmcnt(0)
	s_lshr_b32 s23, s1, 16
.LBB0_14:                               ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB0_16 Depth 2
                                        ;       Child Loop BB0_18 Depth 3
                                        ;         Child Loop BB0_20 Depth 4
	s_and_saveexec_b32 s30, s0
	s_cbranch_execz .LBB0_24
; %bb.15:                               ; %.lr.ph94.i
                                        ;   in Loop: Header=BB0_14 Depth=1
	s_load_b32 s1, s[28:29], 0xc
	v_mad_u64_u32 v[1:2], null, s9, v3, 0x4000
	v_mul_i32_i24_e32 v2, s8, v3
	v_dual_mov_b32 v8, v6 :: v_dual_mov_b32 v9, v4
	s_mov_b32 s34, 0
	s_waitcnt lgkmcnt(0)
	s_and_b32 s33, s1, 0xffff
.LBB0_16:                               ;   Parent Loop BB0_14 Depth=1
                                        ; =>  This Loop Header: Depth=2
                                        ;       Child Loop BB0_18 Depth 3
                                        ;         Child Loop BB0_20 Depth 4
	v_dual_mov_b32 v11, 0 :: v_dual_mov_b32 v10, 0
	s_and_not1_b32 vcc_lo, exec_lo, s20
	s_cbranch_vccnz .LBB0_23
; %bb.17:                               ; %.preheader.lr.ph.i
                                        ;   in Loop: Header=BB0_16 Depth=2
	v_dual_mov_b32 v12, 0 :: v_dual_add_nc_u32 v11, v1, v8
	v_dual_mov_b32 v13, 0 :: v_dual_mov_b32 v10, 0
	s_mov_b32 s1, 0
	s_mov_b32 s35, 0x8f50
.LBB0_18:                               ; %.preheader.i
                                        ;   Parent Loop BB0_14 Depth=1
                                        ;     Parent Loop BB0_16 Depth=2
                                        ; =>    This Loop Header: Depth=3
                                        ;         Child Loop BB0_20 Depth 4
	s_and_not1_b32 vcc_lo, exec_lo, s21
	s_cbranch_vccnz .LBB0_21
; %bb.19:                               ; %.lr.ph.i61
                                        ;   in Loop: Header=BB0_18 Depth=3
	v_mov_b32_e32 v14, v11
	s_mov_b32 s36, s35
	s_mov_b32 s37, s26
.LBB0_20:                               ;   Parent Loop BB0_14 Depth=1
                                        ;     Parent Loop BB0_16 Depth=2
                                        ;       Parent Loop BB0_18 Depth=3
                                        ; =>      This Inner Loop Header: Depth=4
	v_mov_b32_e32 v15, s36
	s_add_i32 s37, s37, -1
	s_add_i32 s36, s36, 1
	s_cmp_lg_u32 s37, 0
	ds_load_u8 v16, v14
	ds_load_u8 v15, v15
	v_add_nc_u32_e32 v14, 1, v14
	s_waitcnt lgkmcnt(1)
	v_cvt_f32_ubyte0_e32 v16, v16
	s_waitcnt lgkmcnt(0)
	v_cvt_f32_ubyte0_e32 v15, v15
	s_delay_alu instid0(VALU_DEP_1)
	v_dual_fmac_f32 v12, v16, v16 :: v_dual_fmac_f32 v13, v15, v15
	v_fmac_f32_e32 v10, v15, v16
	s_cbranch_scc1 .LBB0_20
.LBB0_21:                               ; %._crit_edge.i60
                                        ;   in Loop: Header=BB0_18 Depth=3
	v_add_nc_u32_e32 v11, s9, v11
	s_add_i32 s1, s1, 1
	s_add_i32 s35, s35, s10
	s_cmp_lg_u32 s1, s22
	s_cbranch_scc1 .LBB0_18
; %bb.22:                               ; %._crit_edge86.loopexit.i
                                        ;   in Loop: Header=BB0_16 Depth=2
	v_mul_f32_e32 v11, v13, v12
.LBB0_23:                               ; %._crit_edge86.i
                                        ;   in Loop: Header=BB0_16 Depth=2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v10, v10, v10
	v_div_scale_f32 v12, null, v11, v11, v10
	v_div_scale_f32 v15, vcc_lo, v10, v11, v10
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_f32_e32 v13, v12
	s_waitcnt_depctr 0xfff
	v_fma_f32 v14, -v12, v13, 1.0
	v_fmac_f32_e32 v13, v14, v13
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v14, v15, v13
	v_fma_f32 v16, -v12, v14, v15
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v14, v16, v13
	v_fma_f32 v12, -v12, v14, v15
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_div_fmas_f32 v12, v12, v13, v14
	v_div_fixup_f32 v10, v12, v11, v10
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_mul_f32_e32 v11, 0x4f800000, v10
	v_cmp_gt_f32_e32 vcc_lo, 0xf800000, v10
	v_cndmask_b32_e32 v10, v10, v11, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_2)
	v_sqrt_f32_e32 v11, v10
	s_waitcnt_depctr 0xfff
	v_add_nc_u32_e32 v12, -1, v11
	v_add_nc_u32_e32 v13, 1, v11
	v_fma_f32 v14, -v12, v11, v10
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f32 v15, -v13, v11, v10
	v_cmp_ge_f32_e64 s1, 0, v14
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_cndmask_b32_e64 v11, v11, v12, s1
	v_cmp_lt_f32_e64 s1, 0, v15
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e64 v11, v11, v13, s1
	v_dual_mul_f32 v13, 0x37800000, v11 :: v_dual_add_nc_u32 v12, v8, v2
	ds_store_b8 v12, v9 offset:28496
	v_add_nc_u32_e32 v9, s33, v8
	v_cndmask_b32_e32 v11, v11, v13, vcc_lo
	v_cmp_class_f32_e64 s1, v10, 0x260
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_bfe_i32 v8, v9, 0, 8
	v_lshlrev_b32_e32 v13, 2, v12
	v_cndmask_b32_e64 v10, v11, v10, s1
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_3) | instid1(SALU_CYCLE_1)
	v_cmp_le_i32_e32 vcc_lo, s8, v8
	ds_store_b32 v13, v10
	ds_store_b8 v12, v7 offset:32592
	s_or_b32 s34, vcc_lo, s34
	s_and_not1_b32 exec_lo, exec_lo, s34
	s_cbranch_execnz .LBB0_16
.LBB0_24:                               ; %Flow218
                                        ;   in Loop: Header=BB0_14 Depth=1
	s_or_b32 exec_lo, exec_lo, s30
	v_add_nc_u32_e32 v7, s23, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_bfe_i32 v3, v7, 0, 8
	v_cmp_le_i32_e32 vcc_lo, s2, v3
	s_or_b32 s13, vcc_lo, s13
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s13
	s_cbranch_execnz .LBB0_14
.LBB0_25:                               ; %Flow220
	s_or_b32 exec_lo, exec_lo, s11
	s_cmp_lt_u32 s14, s31
	s_mul_i32 s8, s8, s2
	s_cselect_b32 s0, 12, 18
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_4) | instid1(VALU_DEP_1)
	v_mov_b32_e32 v1, s0
	s_mov_b32 s0, exec_lo
	global_load_u16 v2, v1, s[28:29]
	s_waitcnt vmcnt(0)
	v_mul_u32_u24_e32 v1, v5, v2
	v_add3_u32 v1, s8, v4, v1
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_gt_i32_e32 0x1000, v1
	s_cbranch_execz .LBB0_28
; %bb.26:                               ; %.lr.ph105.i
	s_load_b32 s1, s[28:29], 0x4
	v_dual_mov_b32 v4, 0 :: v_dual_mov_b32 v7, 0x80
	s_waitcnt lgkmcnt(0)
	s_cmp_lt_u32 s15, s1
	s_cselect_b32 s1, 14, 20
	s_delay_alu instid0(SALU_CYCLE_1)
	v_mov_b32_e32 v3, s1
	s_mov_b32 s1, 0
	global_load_u16 v3, v3, s[28:29]
	s_waitcnt vmcnt(0)
	v_mul_lo_u32 v2, v3, v2
	v_lshlrev_b32_e32 v3, 2, v1
	s_delay_alu instid0(VALU_DEP_2)
	v_lshlrev_b32_e32 v6, 2, v2
.LBB0_27:                               ; =>This Inner Loop Header: Depth=1
	ds_store_b8 v1, v7 offset:28496
	ds_store_b8 v1, v7 offset:32592
	v_add_nc_u32_e32 v1, v1, v2
	ds_store_b32 v3, v4
	v_add_nc_u32_e32 v3, v3, v6
	v_cmp_lt_i32_e32 vcc_lo, 0xfff, v1
	s_or_b32 s1, vcc_lo, s1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s1
	s_cbranch_execnz .LBB0_27
.LBB0_28:                               ; %Flow213
	s_or_b32 exec_lo, exec_lo, s0
	s_waitcnt lgkmcnt(0)
	s_barrier
	buffer_gl0_inv
	s_load_b32 s0, s[28:29], 0xc
	v_and_b32_e32 v0, 0x3ff, v0
	s_mov_b32 s1, 0
	s_waitcnt lgkmcnt(0)
	s_and_b32 s8, s0, 0xffff
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_mad_u32_u24 v1, v5, s8, v0
	s_delay_alu instid0(VALU_DEP_1)
	v_cmp_gt_u32_e64 s0, 0x800, v1
.LBB0_29:                               ; %.preheader.i62
                                        ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB0_30 Depth 2
                                        ;       Child Loop BB0_33 Depth 3
                                        ;       Child Loop BB0_40 Depth 3
	s_lshl_b32 s9, 2, s1
	s_lshl_b32 s2, -1, s1
	s_add_i32 s9, s9, -1
	s_not_b32 s10, s2
	s_mov_b32 s11, s1
.LBB0_30:                               ; %.lr.ph76.i
                                        ;   Parent Loop BB0_29 Depth=1
                                        ; =>  This Loop Header: Depth=2
                                        ;       Child Loop BB0_33 Depth 3
                                        ;       Child Loop BB0_40 Depth 3
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmp_lg_u32 s11, s1
	s_mov_b32 s13, -1
	s_cbranch_scc0 .LBB0_37
; %bb.31:                               ;   in Loop: Header=BB0_30 Depth=2
	s_and_saveexec_b32 s13, s0
	s_cbranch_execz .LBB0_36
; %bb.32:                               ; %.lr.ph.i66
                                        ;   in Loop: Header=BB0_30 Depth=2
	s_load_b32 s21, s[28:29], 0x4
	s_lshl_b32 s14, -1, s11
	s_lshl_b32 s20, 1, s11
	s_mov_b32 s22, 0
	v_mov_b32_e32 v3, v1
	s_waitcnt lgkmcnt(0)
	s_cmp_lt_u32 s15, s21
	s_cselect_b32 s21, 14, 20
	s_delay_alu instid0(SALU_CYCLE_1)
	v_mov_b32_e32 v2, s21
	s_bfm_b32 s21, s11, 0
	global_load_u16 v2, v2, s[28:29]
	s_waitcnt vmcnt(0)
	v_mul_lo_u32 v2, v2, s8
.LBB0_33:                               ;   Parent Loop BB0_29 Depth=1
                                        ;     Parent Loop BB0_30 Depth=2
                                        ; =>    This Inner Loop Header: Depth=3
	v_and_b32_e32 v4, s14, v3
	v_and_b32_e32 v6, s21, v3
	s_mov_b32 s23, exec_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshl_add_u32 v4, v4, 1, v6
	v_lshlrev_b32_e32 v6, 2, v4
	s_delay_alu instid0(VALU_DEP_1)
	v_lshl_add_u32 v7, s20, 2, v6
	ds_load_b32 v8, v6
	ds_load_b32 v9, v7
	s_waitcnt lgkmcnt(0)
	v_cmpx_lt_f32_e32 v8, v9
	s_cbranch_execz .LBB0_35
; %bb.34:                               ;   in Loop: Header=BB0_33 Depth=3
	v_add_nc_u32_e32 v10, s20, v4
	ds_load_u8 v11, v10 offset:28496
	ds_load_u8 v12, v10 offset:32592
	ds_load_u8 v13, v4 offset:28496
	ds_load_u8 v14, v4 offset:32592
	s_waitcnt lgkmcnt(3)
	ds_store_b8 v4, v11 offset:28496
	s_waitcnt lgkmcnt(3)
	ds_store_b8 v4, v12 offset:32592
	ds_store_b32 v6, v9
	ds_store_b32 v7, v8
	s_waitcnt lgkmcnt(5)
	ds_store_b8 v10, v13 offset:28496
	s_waitcnt lgkmcnt(5)
	ds_store_b8 v10, v14 offset:32592
.LBB0_35:                               ;   in Loop: Header=BB0_33 Depth=3
	s_or_b32 exec_lo, exec_lo, s23
	v_add_nc_u32_e32 v3, v3, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	v_cmp_lt_i32_e32 vcc_lo, 0x7ff, v3
	s_or_b32 s22, vcc_lo, s22
	s_and_not1_b32 exec_lo, exec_lo, s22
	s_cbranch_execnz .LBB0_33
.LBB0_36:                               ; %Flow208
                                        ;   in Loop: Header=BB0_30 Depth=2
	s_or_b32 exec_lo, exec_lo, s13
	s_mov_b32 s13, 0
.LBB0_37:                               ; %Flow211
                                        ;   in Loop: Header=BB0_30 Depth=2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 vcc_lo, exec_lo, s13
	s_cbranch_vccnz .LBB0_44
; %bb.38:                               ;   in Loop: Header=BB0_30 Depth=2
	s_and_saveexec_b32 s13, s0
	s_cbranch_execz .LBB0_43
; %bb.39:                               ; %.lr.ph73.i
                                        ;   in Loop: Header=BB0_30 Depth=2
	s_load_b32 s14, s[28:29], 0x4
	v_mov_b32_e32 v3, v1
	s_waitcnt lgkmcnt(0)
	s_cmp_lt_u32 s15, s14
	s_cselect_b32 s14, 14, 20
	s_delay_alu instid0(SALU_CYCLE_1)
	v_mov_b32_e32 v2, s14
	s_mov_b32 s14, 0
	global_load_u16 v2, v2, s[28:29]
	s_waitcnt vmcnt(0)
	v_mul_lo_u32 v2, v2, s8
.LBB0_40:                               ;   Parent Loop BB0_29 Depth=1
                                        ;     Parent Loop BB0_30 Depth=2
                                        ; =>    This Inner Loop Header: Depth=3
	v_and_b32_e32 v4, s2, v3
	v_and_b32_e32 v6, s10, v3
	s_mov_b32 s20, exec_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_lshl_add_u32 v7, v4, 2, s9
	v_lshl_add_u32 v6, v4, 1, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_sub_nc_u32_e32 v4, v7, v6
	v_lshlrev_b32_e32 v8, 2, v6
	v_lshlrev_b32_e32 v7, 2, v4
	ds_load_b32 v9, v8
	ds_load_b32 v10, v7
	s_waitcnt lgkmcnt(0)
	v_cmpx_lt_f32_e32 v9, v10
	s_cbranch_execz .LBB0_42
; %bb.41:                               ;   in Loop: Header=BB0_40 Depth=3
	ds_load_u8 v11, v4 offset:28496
	ds_load_u8 v12, v6 offset:28496
	ds_load_u8 v13, v4 offset:32592
	ds_load_u8 v14, v6 offset:32592
	s_waitcnt lgkmcnt(3)
	ds_store_b8 v6, v11 offset:28496
	ds_store_b32 v8, v10
	s_waitcnt lgkmcnt(4)
	ds_store_b8 v4, v12 offset:28496
	s_waitcnt lgkmcnt(4)
	ds_store_b8 v6, v13 offset:32592
	ds_store_b32 v7, v9
	s_waitcnt lgkmcnt(5)
	ds_store_b8 v4, v14 offset:32592
.LBB0_42:                               ;   in Loop: Header=BB0_40 Depth=3
	s_or_b32 exec_lo, exec_lo, s20
	v_add_nc_u32_e32 v3, v3, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	v_cmp_lt_i32_e32 vcc_lo, 0x7ff, v3
	s_or_b32 s14, vcc_lo, s14
	s_and_not1_b32 exec_lo, exec_lo, s14
	s_cbranch_execnz .LBB0_40
.LBB0_43:                               ; %Flow210
                                        ;   in Loop: Header=BB0_30 Depth=2
	s_or_b32 exec_lo, exec_lo, s13
.LBB0_44:                               ; %.loopexit.i
                                        ;   in Loop: Header=BB0_30 Depth=2
	s_add_i32 s13, s11, -1
	s_cmp_gt_i32 s11, 0
	s_mov_b32 s11, s13
	s_waitcnt lgkmcnt(0)
	s_barrier
	buffer_gl0_inv
	s_cbranch_scc1 .LBB0_30
; %bb.45:                               ; %._crit_edge.i64
                                        ;   in Loop: Header=BB0_29 Depth=1
	s_add_i32 s1, s1, 1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmp_lg_u32 s1, 12
	s_cbranch_scc1 .LBB0_29
; %bb.46:                               ; %_Z11bitonicSortPfPcS0_.exit
	v_cmp_eq_u32_e64 s0, 0, v5
	v_cmp_gt_u32_e32 vcc_lo, 5, v0
	v_mov_b32_e32 v2, 0x9870
	s_barrier
	buffer_gl0_inv
	s_and_b32 s1, s0, vcc_lo
	s_mov_b32 s0, -1
	s_xor_b32 s9, s1, -1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_saveexec_b32 s1, s9
	s_cbranch_execz .LBB0_60
; %bb.47:                               ; %NodeBlock178
	v_mov_b32_e32 v2, 0x9870
	s_mov_b32 s10, 0
	s_mov_b32 s2, exec_lo
	v_cmpx_lt_i32_e32 1, v5
	s_xor_b32 s2, exec_lo, s2
	s_cbranch_execz .LBB0_55
; %bb.48:                               ; %NodeBlock
	v_mov_b32_e32 v2, 0x9870
	s_mov_b32 s11, exec_lo
	v_cmpx_lt_i32_e32 2, v5
	s_xor_b32 s11, exec_lo, s11
	s_cbranch_execz .LBB0_52
; %bb.49:                               ; %LeafBlock176
	v_mov_b32_e32 v2, 0x9870
	s_mov_b32 s13, exec_lo
	v_cmpx_eq_u32_e32 3, v5
; %bb.50:                               ; %._crit_edge132.i
	v_mov_b32_e32 v2, 0x9890
	s_and_b32 s10, vcc_lo, exec_lo
; %bb.51:                               ; %Flow203
	s_or_b32 exec_lo, exec_lo, s13
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_b32 s10, s10, exec_lo
.LBB0_52:                               ; %Flow202
	s_and_not1_saveexec_b32 s0, s11
; %bb.53:                               ; %._crit_edge133.i
	s_and_not1_b32 s10, s10, exec_lo
	s_and_b32 s11, vcc_lo, exec_lo
	v_mov_b32_e32 v2, 0x98b0
	s_or_b32 s10, s10, s11
; %bb.54:                               ; %Flow204
	s_or_b32 exec_lo, exec_lo, s0
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_b32 s10, s10, exec_lo
.LBB0_55:                               ; %Flow201
	s_and_not1_saveexec_b32 s2, s2
	s_cbranch_execz .LBB0_59
; %bb.56:                               ; %LeafBlock
	s_mov_b32 s11, s10
	s_mov_b32 s13, exec_lo
	v_cmpx_eq_u32_e32 1, v5
; %bb.57:
	v_mov_b32_e32 v2, 0x9850
	s_and_not1_b32 s0, s10, exec_lo
	s_and_b32 s11, vcc_lo, exec_lo
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 s11, s0, s11
; %bb.58:                               ; %Flow206
	s_or_b32 exec_lo, exec_lo, s13
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	s_and_not1_b32 s0, s10, exec_lo
	s_and_b32 s10, s11, exec_lo
	s_or_b32 s10, s0, s10
.LBB0_59:                               ; %Flow205
	s_or_b32 exec_lo, exec_lo, s2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_not1_b32 s0, s10, exec_lo
.LBB0_60:                               ; %Flow200
	s_or_b32 exec_lo, exec_lo, s1
	s_and_saveexec_b32 s1, s0
	s_cbranch_execz .LBB0_62
; %bb.61:                               ; %._crit_edge134.sink.split.i
	v_lshl_add_u32 v2, v0, 2, v2
	v_mov_b32_e32 v3, 0
	ds_store_b32 v2, v3
.LBB0_62:                               ; %._crit_edge134.i
	s_or_b32 exec_lo, exec_lo, s1
	v_cmp_gt_u32_e64 s0, 0x1000, v1
	v_lshlrev_b32_e32 v2, 2, v1
	s_lshl_b32 s10, s8, 2
	s_mov_b32 s11, 0
	s_waitcnt lgkmcnt(0)
	s_barrier
	buffer_gl0_inv
.LBB0_63:                               ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB0_65 Depth 2
	s_and_saveexec_b32 s13, s0
	s_cbranch_execz .LBB0_91
; %bb.64:                               ; %.lr.ph.i72
                                        ;   in Loop: Header=BB0_63 Depth=1
	v_mov_b32_e32 v3, s11
	s_load_b32 s1, s[28:29], 0xc
	s_lshl_b32 s21, s11, 2
	s_mov_b32 s26, 0
	s_add_i32 s14, s21, 0x98b0
	ds_load_u8 v4, v3 offset:28496
	s_add_i32 s15, s21, 0x9890
	s_add_i32 s20, s21, 0x9870
	s_add_i32 s21, s21, 0x9850
	v_mov_b32_e32 v8, v1
	s_waitcnt lgkmcnt(0)
	s_lshr_b32 s1, s1, 16
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_3) | instid1(VALU_DEP_2)
	s_mul_i32 s22, s1, s8
	s_mul_i32 s23, s10, s1
	v_bfe_i32 v7, v4, 0, 8
	v_dual_mov_b32 v3, v2 :: v_dual_and_b32 v4, 0xff, v4
	v_add_nc_u32_e32 v6, -1, v7
	v_add_nc_u32_e32 v7, 1, v7
.LBB0_65:                               ;   Parent Loop BB0_63 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	ds_load_u8 v9, v8 offset:28496
	s_mov_b32 s30, 0
	s_waitcnt lgkmcnt(0)
	v_cmp_ne_u16_e64 s31, v9, v4
	v_cmp_eq_u16_e64 s1, v9, v4
	s_delay_alu instid0(VALU_DEP_1)
	s_and_saveexec_b32 s33, s1
	s_cbranch_execz .LBB0_67
; %bb.66:                               ;   in Loop: Header=BB0_65 Depth=2
	v_mov_b32_e32 v10, s11
	s_and_not1_b32 s31, s31, exec_lo
	s_mov_b32 s30, exec_lo
	ds_load_u8 v11, v8 offset:32592
	ds_load_u8 v10, v10 offset:32592
	s_waitcnt lgkmcnt(0)
	v_cmp_ne_u16_e64 s2, v11, v10
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b32 s2, s2, exec_lo
	s_or_b32 s31, s31, s2
.LBB0_67:                               ; %Flow196
                                        ;   in Loop: Header=BB0_65 Depth=2
	s_or_b32 exec_lo, exec_lo, s33
	s_and_saveexec_b32 s2, s31
	s_delay_alu instid0(SALU_CYCLE_1)
	s_xor_b32 s31, exec_lo, s2
	s_cbranch_execz .LBB0_88
; %bb.68:                               ;   in Loop: Header=BB0_65 Depth=2
	v_bfe_i32 v9, v9, 0, 8
	s_mov_b32 s33, 0
	s_mov_b32 s35, exec_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_cmp_ne_u32_e64 s34, v6, v9
	v_cmpx_eq_u32_e64 v6, v9
	s_cbranch_execz .LBB0_70
; %bb.69:                               ;   in Loop: Header=BB0_65 Depth=2
	v_mov_b32_e32 v10, s11
	s_delay_alu instid0(VALU_DEP_3)
	s_and_not1_b32 s34, s34, exec_lo
	s_mov_b32 s33, exec_lo
	ds_load_u8 v11, v8 offset:32592
	ds_load_u8 v10, v10 offset:32592
	s_waitcnt lgkmcnt(0)
	v_cmp_ne_u16_e64 s2, v11, v10
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b32 s2, s2, exec_lo
	s_or_b32 s34, s34, s2
.LBB0_70:                               ; %Flow192
                                        ;   in Loop: Header=BB0_65 Depth=2
	s_or_b32 exec_lo, exec_lo, s35
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_saveexec_b32 s2, s34
	s_xor_b32 s34, exec_lo, s2
	s_cbranch_execz .LBB0_85
; %bb.71:                               ;   in Loop: Header=BB0_65 Depth=2
	v_cmp_ne_u32_e64 s36, v7, v9
	s_mov_b32 s35, 0
	s_mov_b32 s37, exec_lo
	v_cmpx_eq_u32_e64 v7, v9
	s_cbranch_execz .LBB0_73
; %bb.72:                               ;   in Loop: Header=BB0_65 Depth=2
	v_mov_b32_e32 v9, s11
	s_and_not1_b32 s36, s36, exec_lo
	s_mov_b32 s35, exec_lo
	ds_load_u8 v10, v8 offset:32592
	ds_load_u8 v9, v9 offset:32592
	s_waitcnt lgkmcnt(0)
	v_cmp_ne_u16_e64 s2, v10, v9
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b32 s2, s2, exec_lo
	s_or_b32 s36, s36, s2
.LBB0_73:                               ; %Flow188
                                        ;   in Loop: Header=BB0_65 Depth=2
	s_or_b32 exec_lo, exec_lo, s37
	s_and_saveexec_b32 s2, s36
	s_delay_alu instid0(SALU_CYCLE_1)
	s_xor_b32 s2, exec_lo, s2
	s_cbranch_execz .LBB0_82
; %bb.74:                               ;   in Loop: Header=BB0_65 Depth=2
	s_and_saveexec_b32 s36, s1
	s_cbranch_execz .LBB0_81
; %bb.75:                               ;   in Loop: Header=BB0_65 Depth=2
	v_mov_b32_e32 v9, s11
	s_mov_b32 s37, exec_lo
	ds_load_i8 v10, v9 offset:32592
	ds_load_i8 v9, v8 offset:32592
	s_waitcnt lgkmcnt(1)
	v_add_nc_u32_e32 v11, -1, v10
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_ne_u32_e64 v11, v9
	s_xor_b32 s37, exec_lo, s37
	s_cbranch_execz .LBB0_79
; %bb.76:                               ;   in Loop: Header=BB0_65 Depth=2
	v_add_nc_u32_e32 v10, 1, v10
	s_mov_b32 s38, exec_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_eq_u32_e64 v10, v9
	s_cbranch_execz .LBB0_78
; %bb.77:                               ;   in Loop: Header=BB0_65 Depth=2
	v_dual_mov_b32 v9, s14 :: v_dual_mov_b32 v10, s15
	ds_load_b32 v11, v3
	ds_load_b32 v12, v9
	ds_load_b32 v13, v10
	s_waitcnt lgkmcnt(1)
	v_add_f32_e32 v12, v11, v12
	s_waitcnt lgkmcnt(0)
	v_add_f32_e32 v11, v11, v13
	ds_store_b32 v9, v12
	ds_store_b32 v10, v11
.LBB0_78:                               ; %Flow183
                                        ;   in Loop: Header=BB0_65 Depth=2
	s_or_b32 exec_lo, exec_lo, s38
.LBB0_79:                               ; %Flow184
                                        ;   in Loop: Header=BB0_65 Depth=2
	s_and_not1_saveexec_b32 s1, s37
	s_cbranch_execz .LBB0_81
; %bb.80:                               ;   in Loop: Header=BB0_65 Depth=2
	v_dual_mov_b32 v9, s14 :: v_dual_mov_b32 v10, s15
	ds_load_b32 v11, v3
	ds_load_b32 v12, v9
	ds_load_b32 v13, v10
	s_waitcnt lgkmcnt(0)
	v_dual_sub_f32 v12, v12, v11 :: v_dual_add_f32 v11, v11, v13
	ds_store_b32 v9, v12
	ds_store_b32 v10, v11
.LBB0_81:                               ; %Flow186
                                        ;   in Loop: Header=BB0_65 Depth=2
	s_or_b32 exec_lo, exec_lo, s36
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 s35, s35, exec_lo
.LBB0_82:                               ; %Flow189
                                        ;   in Loop: Header=BB0_65 Depth=2
	s_or_b32 exec_lo, exec_lo, s2
	s_and_saveexec_b32 s1, s35
	s_cbranch_execz .LBB0_84
; %bb.83:                               ;   in Loop: Header=BB0_65 Depth=2
	v_dual_mov_b32 v9, s20 :: v_dual_mov_b32 v10, s21
	ds_load_b32 v11, v3
	ds_load_b32 v12, v9
	ds_load_b32 v13, v10
	s_waitcnt lgkmcnt(1)
	v_add_f32_e32 v12, v11, v12
	s_waitcnt lgkmcnt(0)
	v_add_f32_e32 v11, v11, v13
	ds_store_b32 v9, v12
	ds_store_b32 v10, v11
.LBB0_84:                               ; %Flow190
                                        ;   in Loop: Header=BB0_65 Depth=2
	s_or_b32 exec_lo, exec_lo, s1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 s33, s33, exec_lo
.LBB0_85:                               ; %Flow193
                                        ;   in Loop: Header=BB0_65 Depth=2
	s_or_b32 exec_lo, exec_lo, s34
	s_and_saveexec_b32 s1, s33
	s_cbranch_execz .LBB0_87
; %bb.86:                               ;   in Loop: Header=BB0_65 Depth=2
	v_dual_mov_b32 v9, s20 :: v_dual_mov_b32 v10, s21
	ds_load_b32 v11, v3
	ds_load_b32 v12, v9
	ds_load_b32 v13, v10
	s_waitcnt lgkmcnt(0)
	v_dual_sub_f32 v12, v12, v11 :: v_dual_add_f32 v11, v11, v13
	ds_store_b32 v9, v12
	ds_store_b32 v10, v11
.LBB0_87:                               ; %Flow194
                                        ;   in Loop: Header=BB0_65 Depth=2
	s_or_b32 exec_lo, exec_lo, s1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 s30, s30, exec_lo
.LBB0_88:                               ; %Flow197
                                        ;   in Loop: Header=BB0_65 Depth=2
	s_or_b32 exec_lo, exec_lo, s31
	s_and_saveexec_b32 s1, s30
	s_cbranch_execz .LBB0_90
; %bb.89:                               ;   in Loop: Header=BB0_65 Depth=2
	v_dual_mov_b32 v9, s21 :: v_dual_mov_b32 v10, s15
	ds_load_b32 v11, v3
	ds_load_b32 v12, v9
	ds_load_b32 v13, v10
	s_waitcnt lgkmcnt(1)
	v_fmac_f32_e32 v12, -2.0, v11
	s_waitcnt lgkmcnt(0)
	v_fmac_f32_e32 v13, -2.0, v11
	ds_store_b32 v9, v12
	ds_store_b32 v10, v13
.LBB0_90:                               ; %.critedge.i
                                        ;   in Loop: Header=BB0_65 Depth=2
	s_or_b32 exec_lo, exec_lo, s1
	v_add_nc_u32_e32 v8, s22, v8
	v_add_nc_u32_e32 v3, s23, v3
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cmp_lt_i32_e64 s1, 0xfff, v8
	s_or_b32 s26, s1, s26
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s26
	s_cbranch_execnz .LBB0_65
.LBB0_91:                               ; %Flow199
                                        ;   in Loop: Header=BB0_63 Depth=1
	s_or_b32 exec_lo, exec_lo, s13
	s_add_i32 s11, s11, 1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmp_lg_u32 s11, 5
	s_cbranch_scc1 .LBB0_63
; %bb.92:
	s_mov_b32 s0, 0
	s_waitcnt lgkmcnt(0)
	s_barrier
	buffer_gl0_inv
                                        ; implicit-def: $sgpr2
                                        ; implicit-def: $vgpr3
	s_and_saveexec_b32 s1, s9
	s_delay_alu instid0(SALU_CYCLE_1)
	s_xor_b32 s1, exec_lo, s1
	s_cbranch_execz .LBB0_96
; %bb.93:
	v_cmp_eq_u32_e64 s0, 1, v5
                                        ; implicit-def: $sgpr2
                                        ; implicit-def: $vgpr3
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(SALU_CYCLE_1)
	s_and_b32 s8, s0, vcc_lo
	s_mov_b32 s0, 0
	s_and_saveexec_b32 s9, s8
	s_xor_b32 s8, exec_lo, s9
	s_cbranch_execz .LBB0_95
; %bb.94:
	v_lshlrev_b32_e32 v1, 2, v0
	s_mov_b32 s0, exec_lo
	s_movk_i32 s2, 0x7f50
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v1, 0x9800, v1
	ds_load_2addr_b32 v[1:2], v1 offset0:36 offset1:44
	s_waitcnt lgkmcnt(0)
	v_add_f32_e32 v1, v1, v1
	v_div_scale_f32 v3, null, v1, v1, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_f32_e32 v4, v3
	s_waitcnt_depctr 0xfff
	v_fma_f32 v6, -v3, v4, 1.0
	v_fmac_f32_e32 v4, v6, v4
	v_div_scale_f32 v6, vcc_lo, v2, v1, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v7, v6, v4
	v_fma_f32 v8, -v3, v7, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v7, v8, v4
	v_fma_f32 v3, -v3, v7, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_div_fmas_f32 v3, v3, v4, v7
	v_div_fixup_f32 v3, v3, v1, v2
.LBB0_95:                               ; %Flow181
	s_or_b32 exec_lo, exec_lo, s8
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_b32 s0, s0, exec_lo
.LBB0_96:                               ; %Flow180
	s_or_saveexec_b32 s1, s1
	v_dual_mov_b32 v1, s6 :: v_dual_mov_b32 v2, s7
	v_mov_b32_e32 v4, s27
	v_mov_b32_e32 v6, s2
	s_xor_b32 exec_lo, exec_lo, s1
	s_cbranch_execz .LBB0_98
; %bb.97:
	v_lshlrev_b32_e32 v1, 2, v0
	s_or_b32 s0, s0, exec_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v1, 0x9800, v1
	ds_load_2addr_b32 v[2:3], v1 offset0:20 offset1:28
	s_waitcnt lgkmcnt(0)
	v_add_f32_e32 v6, v2, v2
	v_div_scale_f32 v1, null, v6, v6, v3
	v_div_scale_f32 v7, vcc_lo, v3, v6, v3
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_f32_e32 v2, v1
	s_waitcnt_depctr 0xfff
	v_fma_f32 v4, -v1, v2, 1.0
	v_fmac_f32_e32 v2, v4, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v4, v7, v2
	v_fma_f32 v8, -v1, v4, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v4, v8, v2
	v_fma_f32 v1, -v1, v4, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_div_fmas_f32 v7, v1, v2, v4
	v_dual_mov_b32 v1, s4 :: v_dual_mov_b32 v2, s5
	v_mov_b32_e32 v4, s3
	v_div_fixup_f32 v3, v7, v6, v3
	v_mov_b32_e32 v6, 0x6f50
.LBB0_98:                               ; %Flow182
	s_or_b32 exec_lo, exec_lo, s1
	s_and_saveexec_b32 s1, s0
	s_cbranch_execz .LBB0_100
; %bb.99:                               ; %.sink.split.i
	s_delay_alu instid0(VALU_DEP_1)
	v_add_nc_u32_e32 v6, v6, v0
	v_cmp_lt_f32_e32 vcc_lo, 1.0, v3
	v_cmp_gt_f32_e64 s0, -1.0, v3
	v_bfe_i32 v4, v4, 0, 8
	ds_load_i8 v8, v6
	v_mad_u64_u32 v[6:7], null, s12, 5, v[0:1]
	v_mov_b32_e32 v7, 0
	s_or_b32 s0, vcc_lo, s0
	v_bfe_i32 v4, v4, 0, 16
	v_cndmask_b32_e64 v3, v3, 0, s0
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_cvt_f32_i32_e32 v9, v4
	s_waitcnt lgkmcnt(0)
	v_cvt_f32_i32_e32 v8, v8
	v_add_f32_e32 v8, v3, v8
	v_lshlrev_b64 v[3:4], 2, v[6:7]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_sub_f32_e32 v6, v8, v9
	v_add_co_u32 v1, vcc_lo, v1, v3
	s_delay_alu instid0(VALU_DEP_3)
	v_add_co_ci_u32_e32 v2, vcc_lo, v2, v4, vcc_lo
	global_store_b32 v[1:2], v6, off
.LBB0_100:
	s_or_b32 exec_lo, exec_lo, s1
	v_or_b32_e32 v3, v0, v5
	v_dual_mov_b32 v1, s16 :: v_dual_mov_b32 v2, s17
	s_waitcnt_vscnt null, 0x0
	s_barrier
	s_delay_alu instid0(VALU_DEP_2)
	v_cmp_eq_u32_e64 s1, 0, v3
	v_cmp_ne_u32_e32 vcc_lo, 0, v3
	v_dual_mov_b32 v3, s4 :: v_dual_mov_b32 v4, s5
	buffer_gl0_inv
	s_and_saveexec_b32 s2, vcc_lo
; %bb.101:
	v_cmp_eq_u32_e32 vcc_lo, 0, v0
	v_cmp_eq_u32_e64 s0, 1, v5
	v_dual_mov_b32 v1, s18 :: v_dual_mov_b32 v2, s19
	v_dual_mov_b32 v3, s6 :: v_dual_mov_b32 v4, s7
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_2) | instid1(SALU_CYCLE_1)
	s_and_b32 s0, s0, vcc_lo
	s_and_not1_b32 s1, s1, exec_lo
	s_and_b32 s0, s0, exec_lo
	s_or_b32 s1, s1, s0
; %bb.102:                              ; %Flow
	s_or_b32 exec_lo, exec_lo, s2
	s_and_saveexec_b32 s0, s1
	s_cbranch_execz .LBB0_104
; %bb.103:                              ; %.sink.split158.i
	s_mul_i32 s2, s12, 5
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_ashr_i32 s3, s2, 31
	s_lshl_b64 s[2:3], s[2:3], 2
	s_delay_alu instid0(SALU_CYCLE_1)
	v_add_co_u32 v3, vcc_lo, v3, s2
	v_add_co_ci_u32_e32 v4, vcc_lo, s3, v4, vcc_lo
	v_add_co_u32 v0, vcc_lo, v1, s24
	v_add_co_ci_u32_e32 v1, vcc_lo, s25, v2, vcc_lo
	global_load_b32 v3, v[3:4], off
	s_waitcnt vmcnt(0)
	global_store_b32 v[0:1], v3, off
.LBB0_104:                              ; %_Z16subPixelAccuracyiPfPcS0_S_S_S_S_S_S_ccS_S_.exit
	s_or_b32 exec_lo, exec_lo, s0
	s_waitcnt_vscnt null, 0x0
	s_barrier
	buffer_gl0_inv
	s_barrier
	buffer_gl0_inv
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z17vectorComputationPhS_iiiccccccccccccPfS0_S0_S0_S0_S0_
		.amdhsa_group_segment_fixed_size 39108
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 344
		.amdhsa_user_sgpr_count 14
		.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 1
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 1
		.amdhsa_next_free_vgpr 17
		.amdhsa_next_free_sgpr 46
		.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	_Z17vectorComputationPhS_iiiccccccccccccPfS0_S0_S0_S0_S0_, .Lfunc_end0-_Z17vectorComputationPhS_iiiccccccccccccPfS0_S0_S0_S0_S0_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 4616
; NumSgprs: 48
; NumVgprs: 17
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 39108 bytes/workgroup (compile time only)
; SGPRBlocks: 5
; VGPRBlocks: 2
; NumSGPRsForWavesPerEU: 48
; NumVGPRsForWavesPerEU: 17
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 14
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 1
	.text
	.protected	_Z16vectorCorrectioniPfS_S_S_ ; -- Begin function _Z16vectorCorrectioniPfS_S_S_
	.globl	_Z16vectorCorrectioniPfS_S_S_
	.p2align	8
	.type	_Z16vectorCorrectioniPfS_S_S_,@function
_Z16vectorCorrectioniPfS_S_S_:          ; @_Z16vectorCorrectioniPfS_S_S_
; %bb.0:
	s_cmp_eq_u32 s14, 0
	s_cbranch_scc1 .LBB1_13
; %bb.1:
	s_load_b32 s12, s[0:1], 0x28
	s_add_u32 s2, s0, 40
	s_addc_u32 s3, s1, 0
	s_waitcnt lgkmcnt(0)
	s_add_i32 s4, s12, -1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmp_eq_u32 s14, s4
	s_cbranch_scc1 .LBB1_13
; %bb.2:
	s_cmp_eq_u32 s15, 0
	s_cbranch_scc1 .LBB1_13
; %bb.3:
	s_load_b32 s2, s[2:3], 0x4
	v_lshrrev_b32_e32 v1, 10, v0
	s_mov_b32 s3, -1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_or_b32_e32 v0, v1, v0
	v_and_b32_e32 v0, 0x3ff, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(SALU_CYCLE_1)
	v_cmp_eq_u32_e32 vcc_lo, 0, v0
	s_waitcnt lgkmcnt(0)
	s_add_i32 s2, s2, -1
	s_cmp_lg_u32 s15, s2
	s_cselect_b32 s2, -1, 0
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b32 s2, vcc_lo, s2
	s_and_saveexec_b32 s4, s2
	s_cbranch_execz .LBB1_13
; %bb.4:
	s_load_b256 s[4:11], s[0:1], 0x8
	s_mul_i32 s2, s12, s15
	s_load_b32 s18, s[0:1], 0x0
	s_add_i32 s2, s2, s14
	v_dual_mov_b32 v1, 0 :: v_dual_mov_b32 v0, 0
	s_mul_i32 s12, s2, 5
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_ashr_i32 s13, s12, 31
	s_lshl_b64 s[0:1], s[12:13], 2
	s_waitcnt lgkmcnt(0)
	s_add_u32 s12, s4, s0
	s_addc_u32 s13, s5, s1
	s_add_u32 s16, s6, s0
	s_addc_u32 s17, s7, s1
	s_load_b32 s14, s[12:13], 0x0
	s_load_b32 s15, s[16:17], 0x0
	s_not_b32 s12, s18
	s_mul_i32 s17, s18, 5
	s_add_i32 s12, s12, s2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_mul_i32 s16, s12, 5
.LBB1_5:                                ; %.preheader
                                        ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB1_6 Depth 2
	s_mov_b32 s18, -1
	s_mov_b32 s12, s16
.LBB1_6:                                ;   Parent Loop BB1_5 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_or_b32 s13, s18, s3
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmp_eq_u32 s13, 0
	s_cbranch_scc1 .LBB1_8
; %bb.7:                                ;   in Loop: Header=BB1_6 Depth=2
	s_ashr_i32 s13, s12, 31
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_lshl_b64 s[20:21], s[12:13], 2
	s_add_u32 s22, s4, s20
	s_addc_u32 s23, s5, s21
	s_add_u32 s20, s6, s20
	s_addc_u32 s21, s7, s21
	s_load_b32 s13, s[22:23], 0x0
	s_load_b32 s19, s[20:21], 0x0
	s_waitcnt lgkmcnt(0)
	v_dual_add_f32 v0, s13, v0 :: v_dual_add_f32 v1, s19, v1
.LBB1_8:                                ;   in Loop: Header=BB1_6 Depth=2
	s_add_i32 s18, s18, 1
	s_add_i32 s12, s12, 5
	s_cmp_lg_u32 s18, 2
	s_cbranch_scc1 .LBB1_6
; %bb.9:                                ;   in Loop: Header=BB1_5 Depth=1
	s_add_i32 s3, s3, 1
	s_add_i32 s16, s16, s17
	s_cmp_lg_u32 s3, 2
	s_cbranch_scc1 .LBB1_5
; %bb.10:
	s_add_u32 s0, s0, 4
	s_addc_u32 s1, s1, 0
	v_dual_mul_f32 v0, 0x3e000000, v0 :: v_dual_mul_f32 v1, 0x3e000000, v1
	s_add_u32 s3, s6, s0
	s_waitcnt lgkmcnt(0)
	v_dual_mov_b32 v2, s15 :: v_dual_mov_b32 v3, s14
	s_addc_u32 s6, s7, s1
	s_add_u32 s7, s4, s0
	s_addc_u32 s12, s5, s1
	s_mov_b64 s[4:5], 0
.LBB1_11:                               ; =>This Inner Loop Header: Depth=1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_add_u32 s0, s7, s4
	s_addc_u32 s1, s12, s5
	s_add_u32 s14, s3, s4
	s_addc_u32 s15, s6, s5
	v_dual_sub_f32 v4, v2, v1 :: v_dual_sub_f32 v5, v3, v0
	s_load_b32 s13, s[14:15], 0x0
	s_load_b32 s14, s[0:1], 0x0
	s_add_u32 s4, s4, 4
	s_addc_u32 s5, s5, 0
	s_cmp_lg_u32 s4, 16
	s_waitcnt lgkmcnt(0)
	v_dual_sub_f32 v6, s13, v1 :: v_dual_sub_f32 v7, s14, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v6, v6, v6
	v_fmac_f32_e32 v6, v7, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_mul_f32 v4, v4, v4 :: v_dual_mul_f32 v7, 0x4f800000, v6
	v_fmac_f32_e32 v4, v5, v5
	v_cmp_gt_f32_e64 s0, 0xf800000, v6
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_mul_f32_e32 v5, 0x4f800000, v4
	v_cmp_gt_f32_e32 vcc_lo, 0xf800000, v4
	v_cndmask_b32_e32 v4, v4, v5, vcc_lo
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e64 v5, v6, v7, s0
	v_sqrt_f32_e32 v7, v5
	s_waitcnt_depctr 0xfff
	v_add_nc_u32_e32 v13, 1, v7
	v_sqrt_f32_e32 v6, v4
	v_add_nc_u32_e32 v10, -1, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_2)
	v_fma_f32 v14, -v10, v7, v5
	s_waitcnt_depctr 0xfff
	v_add_nc_u32_e32 v8, -1, v6
	v_add_nc_u32_e32 v9, 1, v6
	v_fma_f32 v11, -v8, v6, v4
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f32 v12, -v9, v6, v4
	v_cmp_ge_f32_e64 s1, 0, v11
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_cndmask_b32_e64 v6, v6, v8, s1
	v_cmp_ge_f32_e64 s1, 0, v14
	v_fma_f32 v8, -v13, v7, v5
	v_cndmask_b32_e64 v7, v7, v10, s1
	v_cmp_lt_f32_e64 s1, 0, v12
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_cndmask_b32_e64 v6, v6, v9, s1
	v_cmp_lt_f32_e64 s1, 0, v8
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_f32_e32 v8, 0x37800000, v6
	v_cndmask_b32_e64 v7, v7, v13, s1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_dual_cndmask_b32 v6, v6, v8 :: v_dual_mul_f32 v9, 0x37800000, v7
	v_cmp_class_f32_e64 vcc_lo, v4, 0x260
	v_cndmask_b32_e64 v7, v7, v9, s0
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_cndmask_b32_e32 v4, v6, v4, vcc_lo
	v_cmp_class_f32_e64 vcc_lo, v5, 0x260
	v_cndmask_b32_e32 v5, v7, v5, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_cmp_lt_f32_e32 vcc_lo, v5, v4
	v_cndmask_b32_e64 v3, v3, s14, vcc_lo
	v_cndmask_b32_e64 v2, v2, s13, vcc_lo
	s_cbranch_scc1 .LBB1_11
; %bb.12:
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_dual_sub_f32 v4, v2, v1 :: v_dual_sub_f32 v5, v3, v0
	s_ashr_i32 s3, s2, 31
	v_mul_f32_e32 v4, v4, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v4, v5, v5
	v_mul_f32_e32 v5, 0x4f800000, v4
	v_cmp_gt_f32_e32 vcc_lo, 0xf800000, v4
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v4, v4, v5, vcc_lo
	v_sqrt_f32_e32 v5, v4
	s_waitcnt_depctr 0xfff
	v_add_nc_u32_e32 v6, -1, v5
	v_add_nc_u32_e32 v7, 1, v5
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f32 v8, -v6, v5, v4
	v_fma_f32 v9, -v7, v5, v4
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cmp_ge_f32_e64 s0, 0, v8
	v_cndmask_b32_e64 v5, v5, v6, s0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cmp_lt_f32_e64 s0, 0, v9
	v_cndmask_b32_e64 v5, v5, v7, s0
	s_lshl_b64 s[0:1], s[2:3], 2
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	s_add_u32 s2, s8, s0
	s_addc_u32 s3, s9, s1
	v_mul_f32_e32 v6, 0x37800000, v5
	s_add_u32 s0, s10, s0
	s_addc_u32 s1, s11, s1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_cndmask_b32_e32 v5, v5, v6, vcc_lo
	v_cmp_class_f32_e64 vcc_lo, v4, 0x260
	v_cndmask_b32_e32 v4, v5, v4, vcc_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_cmp_lt_f32_e32 vcc_lo, 4.0, v4
	v_dual_mov_b32 v4, 0 :: v_dual_cndmask_b32 v1, v2, v1
	v_cndmask_b32_e32 v0, v3, v0, vcc_lo
	s_clause 0x1
	global_store_b32 v4, v0, s[2:3]
	global_store_b32 v4, v1, s[0:1]
.LBB1_13:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z16vectorCorrectioniPfS_S_S_
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 296
		.amdhsa_user_sgpr_count 14
		.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 1
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 1
		.amdhsa_next_free_vgpr 15
		.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	_Z16vectorCorrectioniPfS_S_S_, .Lfunc_end1-_Z16vectorCorrectioniPfS_S_S_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 1020
; NumSgprs: 26
; NumVgprs: 15
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 3
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 26
; NumVGPRsForWavesPerEU: 15
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 14
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 1
	.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
      - .offset:         16
        .size:           4
        .value_kind:     by_value
      - .offset:         20
        .size:           4
        .value_kind:     by_value
      - .offset:         24
        .size:           4
        .value_kind:     by_value
      - .offset:         28
        .size:           1
        .value_kind:     by_value
      - .offset:         29
        .size:           1
        .value_kind:     by_value
      - .offset:         30
        .size:           1
        .value_kind:     by_value
      - .offset:         31
        .size:           1
        .value_kind:     by_value
      - .offset:         32
        .size:           1
        .value_kind:     by_value
      - .offset:         33
        .size:           1
        .value_kind:     by_value
      - .offset:         34
        .size:           1
        .value_kind:     by_value
      - .offset:         35
        .size:           1
        .value_kind:     by_value
      - .offset:         36
        .size:           1
        .value_kind:     by_value
      - .offset:         37
        .size:           1
        .value_kind:     by_value
      - .offset:         38
        .size:           1
        .value_kind:     by_value
      - .offset:         39
        .size:           1
        .value_kind:     by_value
      - .address_space:  global
        .offset:         40
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         48
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         56
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         64
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         72
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         80
        .size:           8
        .value_kind:     global_buffer
      - .offset:         88
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         92
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         96
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         100
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         102
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         104
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         106
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         108
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         110
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         128
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         136
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         144
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         152
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 39108
    .kernarg_segment_align: 8
    .kernarg_segment_size: 344
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z17vectorComputationPhS_iiiccccccccccccPfS0_S0_S0_S0_S0_
    .private_segment_fixed_size: 0
    .sgpr_count:     48
    .sgpr_spill_count: 0
    .symbol:         _Z17vectorComputationPhS_iiiccccccccccccPfS0_S0_S0_S0_S0_.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     17
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .offset:         0
        .size:           4
        .value_kind:     by_value
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         24
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         32
        .size:           8
        .value_kind:     global_buffer
      - .offset:         40
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         44
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         48
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         52
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         54
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         56
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         58
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         60
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         62
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         88
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         96
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         104
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 296
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z16vectorCorrectioniPfS_S_S_
    .private_segment_fixed_size: 0
    .sgpr_count:     26
    .sgpr_spill_count: 0
    .symbol:         _Z16vectorCorrectioniPfS_S_S_.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     15
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
