/*
	file: AES.cu
	BY: FU ERDONG
*/

#define PX 283
#define MODPX 27

#include <iostream>
#include <stdlib.h>
#include <sys/time.h>
#include <math.h>
#include <cuda_runtime.h>
using namespace std;

unsigned short int inversionMatrix[256] =
{
	0x0, 0x1, 0x8d, 0xf6, 0xcb, 0x52, 0x7b, 0xd1, 0xe8, 0x4f, 0x29, 0xc0, 0xb0, 0xe1, 0xe5, 0xc7, 
	0x74, 0xb4, 0xaa, 0x4b, 0x99, 0x2b, 0x60, 0x5f, 0x58, 0x3f, 0xfd, 0xcc, 0xff, 0x40, 0xee, 0xb2, 
	0x3a, 0x6e, 0x5a, 0xf1, 0x55, 0x4d, 0xa8, 0xc9, 0xc1, 0xa, 0x98, 0x15, 0x30, 0x44, 0xa2, 0xc2, 
	0x2c, 0x45, 0x92, 0x6c, 0xf3, 0x39, 0x66, 0x42, 0xf2, 0x35, 0x20, 0x6f, 0x77, 0xbb, 0x59, 0x19, 
	0x1d, 0xfe, 0x37, 0x67, 0x2d, 0x31, 0xf5, 0x69, 0xa7, 0x64, 0xab, 0x13, 0x54, 0x25, 0xe9, 0x9, 
	0xed, 0x5c, 0x5, 0xca, 0x4c, 0x24, 0x87, 0xbf, 0x18, 0x3e, 0x22, 0xf0, 0x51, 0xec, 0x61, 0x17, 
	0x16, 0x5e, 0xaf, 0xd3, 0x49, 0xa6, 0x36, 0x43, 0xf4, 0x47, 0x91, 0xdf, 0x33, 0x93, 0x21, 0x3b, 
	0x79, 0xb7, 0x97, 0x85, 0x10, 0xb5, 0xba, 0x3c, 0xb6, 0x70, 0xd0, 0x6, 0xa1, 0xfa, 0x81, 0x82, 
	0x83, 0x7e, 0x7f, 0x80, 0x96, 0x73, 0xbe, 0x56, 0x9b, 0x9e, 0x95, 0xd9, 0xf7, 0x2, 0xb9, 0xa4, 
	0xde, 0x6a, 0x32, 0x6d, 0xd8, 0x8a, 0x84, 0x72, 0x2a, 0x14, 0x9f, 0x88, 0xf9, 0xdc, 0x89, 0x9a, 
	0xfb, 0x7c, 0x2e, 0xc3, 0x8f, 0xb8, 0x65, 0x48, 0x26, 0xc8, 0x12, 0x4a, 0xce, 0xe7, 0xd2, 0x62, 
	0xc, 0xe0, 0x1f, 0xef, 0x11, 0x75, 0x78, 0x71, 0xa5, 0x8e, 0x76, 0x3d, 0xbd, 0xbc, 0x86, 0x57, 
	0xb, 0x28, 0x2f, 0xa3, 0xda, 0xd4, 0xe4, 0xf, 0xa9, 0x27, 0x53, 0x4, 0x1b, 0xfc, 0xac, 0xe6, 
	0x7a, 0x7, 0xae, 0x63, 0xc5, 0xdb, 0xe2, 0xea, 0x94, 0x8b, 0xc4, 0xd5, 0x9d, 0xf8, 0x90, 0x6b, 
	0xb1, 0xd, 0xd6, 0xeb, 0xc6, 0xe, 0xcf, 0xad, 0x8, 0x4e, 0xd7, 0xe3, 0x5d, 0x50, 0x1e, 0xb3, 
	0x5b, 0x23, 0x38, 0x34, 0x68, 0x46, 0x3, 0x8c, 0xdd, 0x9c, 0x7d, 0xa0, 0xcd, 0x1a, 0x41, 0x1c
};

unsigned short int modPXMatrix[256] = 
{
	0x0, 0x1b, 0x36, 0x2d, 0x6c, 0x77, 0x5a, 0x41, 0xd8, 0xc3, 0xee, 0xf5, 0xb4, 0xaf, 0x82, 0x99, 
	0xab, 0xb0, 0x9d, 0x86, 0xc7, 0xdc, 0xf1, 0xea, 0x73, 0x68, 0x45, 0x5e, 0x1f, 0x4, 0x29, 0x32, 
	0x4d, 0x56, 0x7b, 0x60, 0x21, 0x3a, 0x17, 0xc, 0x95, 0x8e, 0xa3, 0xb8, 0xf9, 0xe2, 0xcf, 0xd4, 
	0xe6, 0xfd, 0xd0, 0xcb, 0x8a, 0x91, 0xbc, 0xa7, 0x3e, 0x25, 0x8, 0x13, 0x52, 0x49, 0x64, 0x7f, 
	0x9a, 0x81, 0xac, 0xb7, 0xf6, 0xed, 0xc0, 0xdb, 0x42, 0x59, 0x74, 0x6f, 0x2e, 0x35, 0x18, 0x3, 
	0x31, 0x2a, 0x7, 0x1c, 0x5d, 0x46, 0x6b, 0x70, 0xe9, 0xf2, 0xdf, 0xc4, 0x85, 0x9e, 0xb3, 0xa8, 
	0xd7, 0xcc, 0xe1, 0xfa, 0xbb, 0xa0, 0x8d, 0x96, 0xf, 0x14, 0x39, 0x22, 0x63, 0x78, 0x55, 0x4e, 
	0x7c, 0x67, 0x4a, 0x51, 0x10, 0xb, 0x26, 0x3d, 0xa4, 0xbf, 0x92, 0x89, 0xc8, 0xd3, 0xfe, 0xe5, 
	0x2f, 0x34, 0x19, 0x2, 0x43, 0x58, 0x75, 0x6e, 0xf7, 0xec, 0xc1, 0xda, 0x9b, 0x80, 0xad, 0xb6, 
	0x84, 0x9f, 0xb2, 0xa9, 0xe8, 0xf3, 0xde, 0xc5, 0x5c, 0x47, 0x6a, 0x71, 0x30, 0x2b, 0x6, 0x1d, 
	0x62, 0x79, 0x54, 0x4f, 0xe, 0x15, 0x38, 0x23, 0xba, 0xa1, 0x8c, 0x97, 0xd6, 0xcd, 0xe0, 0xfb, 
	0xc9, 0xd2, 0xff, 0xe4, 0xa5, 0xbe, 0x93, 0x88, 0x11, 0xa, 0x27, 0x3c, 0x7d, 0x66, 0x4b, 0x50, 
	0xb5, 0xae, 0x83, 0x98, 0xd9, 0xc2, 0xef, 0xf4, 0x6d, 0x76, 0x5b, 0x40, 0x1, 0x1a, 0x37, 0x2c, 
	0x1e, 0x5, 0x28, 0x33, 0x72, 0x69, 0x44, 0x5f, 0xc6, 0xdd, 0xf0, 0xeb, 0xaa, 0xb1, 0x9c, 0x87, 
	0xf8, 0xe3, 0xce, 0xd5, 0x94, 0x8f, 0xa2, 0xb9, 0x20, 0x3b, 0x16, 0xd, 0x4c, 0x57, 0x7a, 0x61, 
	0x53, 0x48, 0x65, 0x7e, 0x3f, 0x24, 0x9, 0x12, 0x8b, 0x90, 0xbd, 0xa6, 0xe7, 0xfc, 0xd1, 0xca, 
};

unsigned char sBox[] =
{ /*  0    1    2    3    4    5    6    7    8    9    a    b    c    d    e    f */
    0x63,0x7c,0x77,0x7b,0xf2,0x6b,0x6f,0xc5,0x30,0x01,0x67,0x2b,0xfe,0xd7,0xab,0x76, /*0*/ 
    0xca,0x82,0xc9,0x7d,0xfa,0x59,0x47,0xf0,0xad,0xd4,0xa2,0xaf,0x9c,0xa4,0x72,0xc0, /*1*/
    0xb7,0xfd,0x93,0x26,0x36,0x3f,0xf7,0xcc,0x34,0xa5,0xe5,0xf1,0x71,0xd8,0x31,0x15, /*2*/
    0x04,0xc7,0x23,0xc3,0x18,0x96,0x05,0x9a,0x07,0x12,0x80,0xe2,0xeb,0x27,0xb2,0x75, /*3*/
    0x09,0x83,0x2c,0x1a,0x1b,0x6e,0x5a,0xa0,0x52,0x3b,0xd6,0xb3,0x29,0xe3,0x2f,0x84, /*4*/
    0x53,0xd1,0x00,0xed,0x20,0xfc,0xb1,0x5b,0x6a,0xcb,0xbe,0x39,0x4a,0x4c,0x58,0xcf, /*5*/
    0xd0,0xef,0xaa,0xfb,0x43,0x4d,0x33,0x85,0x45,0xf9,0x02,0x7f,0x50,0x3c,0x9f,0xa8, /*6*/ 
    0x51,0xa3,0x40,0x8f,0x92,0x9d,0x38,0xf5,0xbc,0xb6,0xda,0x21,0x10,0xff,0xf3,0xd2, /*7*/
    0xcd,0x0c,0x13,0xec,0x5f,0x97,0x44,0x17,0xc4,0xa7,0x7e,0x3d,0x64,0x5d,0x19,0x73, /*8*/
    0x60,0x81,0x4f,0xdc,0x22,0x2a,0x90,0x88,0x46,0xee,0xb8,0x14,0xde,0x5e,0x0b,0xdb, /*9*/
    0xe0,0x32,0x3a,0x0a,0x49,0x06,0x24,0x5c,0xc2,0xd3,0xac,0x62,0x91,0x95,0xe4,0x79, /*a*/
    0xe7,0xc8,0x37,0x6d,0x8d,0xd5,0x4e,0xa9,0x6c,0x56,0xf4,0xea,0x65,0x7a,0xae,0x08, /*b*/
    0xba,0x78,0x25,0x2e,0x1c,0xa6,0xb4,0xc6,0xe8,0xdd,0x74,0x1f,0x4b,0xbd,0x8b,0x8a, /*c*/
    0x70,0x3e,0xb5,0x66,0x48,0x03,0xf6,0x0e,0x61,0x35,0x57,0xb9,0x86,0xc1,0x1d,0x9e, /*d*/
    0xe1,0xf8,0x98,0x11,0x69,0xd9,0x8e,0x94,0x9b,0x1e,0x87,0xe9,0xce,0x55,0x28,0xdf, /*e*/
    0x8c,0xa1,0x89,0x0d,0xbf,0xe6,0x42,0x68,0x41,0x99,0x2d,0x0f,0xb0,0x54,0xbb,0x16  /*f*/
};

void printBits(unsigned short int num)
{
	for (int i = 0; i < 16; i++ ) {
		if (i == 8) cout << " ";
		if (num & (1 << (15 - i))) {
			cout << "1";
		} else {
			cout << "0";
		}
	}
}

void division(unsigned short int &q, unsigned short int &r, unsigned short int a, unsigned short int b)
{
	int la, lb, i;

	q = 0;
	for (i = 15; i > 0; i--) if (a & (1 << i)) break;
	la = i;
	for (i = 15; i > 0; i--) if (b & (1 << i)) break;
	lb = i;
	for (i = la; i >= lb; i--)
		if (a & (1 << i)) {
			a ^= b << (i - lb);
			q ^= 1 << (i - lb);
		}
	r = a;
}

unsigned short int modPX(unsigned short int num)
{
	unsigned short int px;

	px = PX << 7;
	for (int i = 15; i >= 8; i--) {
		if (num & (1 << i)) num = num ^ px;
		px >>= 1;
	}
	return num;
}

void printModPXMatrix()
{
	for (int i = 0; i < 16; i++) {
		cout << "\t";
		for (int j = 0; j < 16; j++) {
			cout << "0x" << hex << modPX((i * 16 + j) << 8) << ", ";
		}
		cout << endl;
	}
}

unsigned short int multiply(unsigned short int a, unsigned short int b)
{
	unsigned short int c = 0;
	for (int i = 0; i < 8; i++)
		if (a & (1 << i)) c ^= (b << i);
	return ((c & 0x00ff) ^ modPXMatrix[c >> 8]);
}

unsigned short int inversion(unsigned short int num)
{
	unsigned short int a, b, q, r, k11, k12, k21, k22, k11n, k12n, k21n, k22n;

	if (num == 0) return 0;
	k11 = 1;
	k22 = 1;
	k12 = 0;
	k21 = 0;
	a = PX;
	b = num;

	while (b != 1) {
		division(q, r, a, b);
		a = b;
		b = r;
		k11n = k21;
		k12n = k22;
		k21n = k11 ^ multiply(q, k21);
		k22n = k12 ^ multiply(q, k22);
		k11 = k11n;
		k12 = k12n;
		k21 = k21n;
		k22 = k22n;
	}
	return modPX(k22);
}

void printInversionMatrix()
{
	for (int i = 0; i < 16; i++) {
		cout << "\t";
		for (int j = 0; j < 16; j++) {
			cout << "0x" << hex << inversion(i * 16 + j) << ", ";
		}
		cout << endl;
	}
}

void testMultiply()
{
	struct timeval beginTime, endTime;

//multiply * 100,000
	gettimeofday(&beginTime, 0);
	cout << "Multiply for 100,000 times." << endl;
	for (int i = 0; i < 100000; i++) multiply(rand() % 256, rand() % 256);
	gettimeofday(&endTime, 0);
	cout << "uSeconds: "
		<< (endTime.tv_sec - beginTime.tv_sec) * 1000000 
		+ endTime.tv_usec - beginTime.tv_usec << endl << endl;

}

void testInversion()
{
	struct timeval beginTime, endTime;
	unsigned short int temp;


// inversion * 100,000
	gettimeofday(&beginTime, 0);
	cout << "Inversion for 100,000 times." << endl;
	for (int i = 0; i < 100000; i++) {
		while (!(temp = rand() % 256));
		if (multiply(inversionMatrix[temp], temp) - 1) cout << temp;
	}
	gettimeofday(&endTime, 0);
	cout << "uSeconds: "
		<< (endTime.tv_sec - beginTime.tv_sec) * 1000000 
		+ endTime.tv_usec - beginTime.tv_usec << endl << endl;

}

__global__ void count(unsigned short int *in, unsigned short int *out) {
	unsigned short int gindex = threadIdx.x + blockIdx.x * blockDim.x;
	unsigned short int baseIndex = blockIdx.x * blockDim.x;

	unsigned short int counter = 0;
	unsigned short int temp;
	for (int i = 0; i < 256; i++) {
		temp = (i << 8) ^ in[baseIndex + i];
		temp &= gindex;
		for (int j = 1; j < 16; j++)
			if (temp & (1 << j)) temp ^= 0x01;
		if (!(temp & 0x1)) counter++;
	}

	out[gindex] = counter;
}

__global__ void d_multiply(unsigned short int *d_a, unsigned short int *d_b) {

	int index = 320 * 32 * blockIdx.y + 320 * threadIdx.y + 32 * blockIdx.x + threadIdx.x;
	if (index >= 100000) return;

	unsigned short int d_c = 0;
	for (int i = 7; i >= 0; i--) {
		d_c <<= 1;
		if (d_c & 256) d_c ^= PX;
		if (d_b[index] & (1 << i)) d_c ^= d_a[index];
	}
	d_a[index] = d_c;
}

__global__ void d_inversion(unsigned short int *d_a) {

	int gindex = 320 * 32 * blockIdx.y + 320 * threadIdx.y + 32 * blockIdx.x + threadIdx.x;
	unsigned short int lindex = threadIdx.y * 32 + threadIdx.x;

	__shared__ unsigned short int matrix[256];

	if (threadIdx.y < 8) matrix[lindex] = d_a[100000 + lindex];

	__syncthreads();

	if (gindex >= 100000) return;
	d_a[gindex] = matrix[d_a[gindex]];
}


void sBoxAnalysis()
{
//	struct timeval begin, end;

/*

// sBox Linear Analysis in normal way
	cout << "sBox linear analysis in normal way" << endl;
	gettimeofday(&begin, 0);
	int check[256 * 256];
	int temp;
	int result;
	for (int i = 1; i < (256 * 256); i++) {
		result = 0;
		for (int j = 0; j < 256; j++) {
			temp = (j << 8) ^ sBox[j];
			temp &= i;
			for (int k = 1; k < 16; k++)
				if (temp & (1 << k)) temp ^= 0x01;
			if (!(temp & 0x1)) result++;
		}
		check[i] = result;
	}

	gettimeofday(&end, 0);
	cout << "uSeconds: "
		<< (end.tv_sec - begin.tv_sec) * 1000000 
		+ end.tv_usec - begin.tv_usec << endl << endl;

*/

// use CUDA to accelerate
	cout << "sBox linear analysis with CUDA" << endl;
	unsigned short int *in, *counter;
	unsigned short int *d_in, *d_counter;
	int size = 256 * 256 * sizeof(unsigned short int);

//	gettimeofday(&begin, 0);

	cudaMalloc((void**)&d_in, size);
	cudaMalloc((void**)&d_counter, size);

	in = (unsigned short int*)malloc(size);
	counter = (unsigned short int*)malloc(size);


	for (int i = 0; i < 256; i++)
		for (int j = 0; j < 256; j++)
			in[i * 256 + j] = sBox[j];

	cudaMemcpy(d_in, in, size, cudaMemcpyHostToDevice);

	count<<<256, 256>>>(d_in, d_counter);

	cudaMemcpy(counter, d_counter, size, cudaMemcpyDeviceToHost);

	int N = 10;
	int dmax[N], imax[N];
	int det;
	for (int i = 0; i < N; i++) dmax[i] = 0;
	for (int i = 1; i < 65536; i++) {
		int j;
		det = abs(counter[i] - 128);
		if (det <= dmax[N - 1]) continue;
		for (j = N - 2; j >= 0; j--) {
			if (det < dmax[j]) break;
			dmax[j + 1] = dmax[j];
			imax[j + 1] = imax[j];
		}
		dmax[j + 1] = det;
		imax[j + 1] = i;
	}
	for (int i = 0; i < N; i++) {
		printBits(imax[i]);
		cout << "\t" << dmax[i] << endl;
	}


//	gettimeofday(&end, 0);
//	cout << "Time in total" << endl;
//	cout << "uSeconds: " << (end.tv_sec - begin.tv_sec) * 1000000 + end.tv_usec - begin.tv_usec << endl << endl;

/*
// check
	for (int i = 1; i < 256 * 256; i++)
		if (check[i] != counter[i]) {
			cout << i << "\t" << check[i] << "\t" << hex << counter[i] << endl;
			//break;
		}
*/


	free(in);
	free(counter);
	cudaFree(d_in);
	cudaFree(d_counter);
}

void GPUMultiply()
{
	struct timeval begin, end;

	cout << "Multiply for 100,000 times with GPU." << endl;

	int memSize = 100000 * sizeof(unsigned short int);
	
	unsigned short int *a, *b, *c;

	a = (unsigned short int *)malloc(memSize);
	b = (unsigned short int *)malloc(memSize);
	c = (unsigned short int *)malloc(memSize);

	for (int i = 0; i < 100000; i++) {
		a[i] = rand() % 256;
		b[i] = rand() % 256;
	}


	gettimeofday(&begin, 0);

	unsigned short int *d_a, *d_b;

	cudaMalloc((void**) &d_a, memSize);
	cudaMalloc((void**) &d_b, memSize);

	cudaMemcpy(d_a, a, memSize, cudaMemcpyHostToDevice);
	cudaMemcpy(d_b, b, memSize, cudaMemcpyHostToDevice);

	dim3 threads(32, 32);
	dim3 grid(10, 10);

	d_multiply<<<grid, threads>>>(d_a, d_b);

	cudaDeviceSynchronize();

	cudaMemcpy(c, d_a, memSize, cudaMemcpyDeviceToHost);

	cudaFree(d_a);
	cudaFree(d_b);

	gettimeofday(&end, 0);

//check
	for (int i = 0; i < 100000; i++)
		if (c[i] != multiply(a[i], b[i])) {
			cout << "Error" << i << endl;
			break;
		}


	free(a);
	free(b);
	free(c);

	cout << "Time(uSeconds): " << (end.tv_sec - begin.tv_sec) * 1000000 
		+ end.tv_usec - begin.tv_usec << endl << endl;

}

void GPUInversion()
{
	struct timeval begin, end;

	cout << "Inversion for 100,000 times with GPU." << endl;

	int memSize = (100000 + 256) * sizeof(unsigned short int);
	
	unsigned short int *a, *b;

	a = (unsigned short int *)malloc(memSize);
	b = (unsigned short int *)malloc(memSize);

	for (int i = 0; i < 100000; i++) {
		a[i] = rand() % 256;
	}

	for (int i = 0; i < 256; i++) {
		a[i + 100000] = inversionMatrix[i];
	}

	gettimeofday(&begin, 0);

	unsigned short int *d_a;

	cudaMalloc((void**) &d_a, memSize);

	cudaMemcpy(d_a, a, memSize, cudaMemcpyHostToDevice);

	dim3 threads(32, 32);
	dim3 grid(10, 10);

	d_inversion<<<grid, threads>>>(d_a);

	cudaDeviceSynchronize();

	cudaMemcpy(b, d_a, memSize, cudaMemcpyDeviceToHost);

	cudaFree(d_a);

	gettimeofday(&end, 0);


//check
	for (int i = 0; i < 100000; i++)
		if ((a[i] == 0 && b[i] != 0) || (a[i] != 0 && multiply(a[i], b[i]) != 1)) {
			cout << "Error " << i << endl;
			break;
		}


	free(a);
	free(b);

	cout << "Time(uSeconds): " << (end.tv_sec - begin.tv_sec) * 1000000 
		+ end.tv_usec - begin.tv_usec << endl << endl;

}





int main()
{
	srand(time(0));

	unsigned short int *temp;
	cudaMalloc((void**)&temp, 100 * sizeof(unsigned short int));
	cudaFree(temp);

	GPUMultiply();
	GPUInversion();

//	testMultiply();
//	testInversion();

	sBoxAnalysis();

}







	.file	"tmpxft_00319086_00000000-6_AES.cudafe1.cpp"
	.text
#APP
	.globl _ZSt21ios_base_library_initv

#NO_APP
	.section	.text._ZSt3hexRSt8ios_base,"axG",@progbits,_ZSt3hexRSt8ios_base,comdat
	.weak	_ZSt3hexRSt8ios_base
	.type	_ZSt3hexRSt8ios_base, @function
_ZSt3hexRSt8ios_base:
.LFB3283:
	.cfi_startproc
	endbr64
	movl	24(%rdi), %edx
	movq	%rdi, %rax
	andl	$-75, %edx
	orl	$8, %edx
	movl	%edx, 24(%rdi)
	ret
	.cfi_endproc
.LFE3283:
	.size	_ZSt3hexRSt8ios_base, .-_ZSt3hexRSt8ios_base
	.text
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB3650:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE3650:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	" "
.LC1:
	.string	"1"
.LC2:
	.string	"0"
	.text
	.globl	_Z9printBitst
	.type	_Z9printBitst, @function
_Z9printBitst:
.LFB3635:
	.cfi_startproc
	endbr64
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	movzwl	%di, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	xorl	%ebx, %ebx
.L4:
	movl	$15, %eax
	leaq	_ZSt4cout(%rip), %r12
	leaq	.LC1(%rip), %rsi
	subl	%ebx, %eax
	btl	%eax, %ebp
	jc	.L10
	leaq	.LC2(%rip), %rsi
.L10:
	movq	%r12, %rdi
	incl	%ebx
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	cmpl	$16, %ebx
	je	.L12
	cmpl	$8, %ebx
	jne	.L4
	leaq	.LC0(%rip), %rsi
	movq	%r12, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	jmp	.L4
.L12:
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3635:
	.size	_Z9printBitst, .-_Z9printBitst
	.globl	_Z8divisionRtS_tt
	.type	_Z8divisionRtS_tt, @function
_Z8divisionRtS_tt:
.LFB3636:
	.cfi_startproc
	endbr64
	movw	$0, (%rdi)
	movq	%rsi, %r8
	movzwl	%dx, %r9d
	movl	%edx, %esi
	movl	$15, %eax
	btl	$15, %edx
.L30:
	jnc	.L15
.L14:
	movzwl	%cx, %r9d
	movl	$15, %edx
	btl	$15, %ecx
	jmp	.L32
.L15:
	decl	%eax
	je	.L14
	btl	%eax, %r9d
	jmp	.L30
.L33:
	decl	%edx
	je	.L16
	btl	%edx, %r9d
.L32:
	jnc	.L33
.L16:
	movl	$1, %r10d
.L21:
	cmpl	%edx, %eax
	jl	.L34
	movzwl	%si, %ecx
	btl	%eax, %ecx
	jnc	.L20
	movl	%eax, %ecx
	movl	%r9d, %r11d
	subl	%edx, %ecx
	sall	%cl, %r11d
	xorl	%r11d, %esi
	movl	%r10d, %r11d
	sall	%cl, %r11d
	xorw	%r11w, (%rdi)
.L20:
	decl	%eax
	jmp	.L21
.L34:
	movw	%si, (%r8)
	ret
	.cfi_endproc
.LFE3636:
	.size	_Z8divisionRtS_tt, .-_Z8divisionRtS_tt
	.globl	_Z5modPXt
	.type	_Z5modPXt, @function
_Z5modPXt:
.LFB3637:
	.cfi_startproc
	endbr64
	movl	%edi, %eax
	movl	$15, %ecx
	movl	$-29312, %esi
.L36:
	movzwl	%ax, %edx
	sarl	%cl, %edx
	decl	%ecx
	andl	$1, %edx
	negl	%edx
	andl	%esi, %edx
	shrw	%si
	xorl	%edx, %eax
	cmpl	$7, %ecx
	jne	.L36
	ret
	.cfi_endproc
.LFE3637:
	.size	_Z5modPXt, .-_Z5modPXt
	.section	.rodata.str1.1
.LC3:
	.string	"\t"
.LC4:
	.string	"0x"
.LC5:
	.string	", "
	.text
	.globl	_Z16printModPXMatrixv
	.type	_Z16printModPXMatrixv, @function
_Z16printModPXMatrixv:
.LFB3638:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	leaq	.LC3(%rip), %r14
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	xorl	%ebx, %ebx
	pushq	%rdx
	.cfi_def_cfa_offset 64
.L40:
	leaq	_ZSt4cout(%rip), %rbp
	movq	%r14, %rsi
	movl	%ebx, %r13d
	xorl	%r12d, %r12d
	movq	%rbp, %rdi
	sall	$8, %r13d
	leaq	.LC4(%rip), %r15
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
.L39:
	movq	%r15, %rsi
	movq	%rbp, %rdi
	incl	%r12d
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	leaq	_ZSt3hexRSt8ios_base(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSt8ios_baseS0_E@PLT
	movzwl	%r13w, %edi
	addw	$256, %r13w
	movq	%rax, %r8
	call	_Z5modPXt
	movq	%r8, %rdi
	movzwl	%ax, %esi
	call	_ZNSo9_M_insertImEERSoT_@PLT
	leaq	.LC5(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	cmpl	$16, %r12d
	jne	.L39
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rbp, %rdi
	addl	$16, %ebx
	call	_ZNSolsEPFRSoS_E@PLT
	cmpw	$256, %bx
	jne	.L40
	popq	%rax
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3638:
	.size	_Z16printModPXMatrixv, .-_Z16printModPXMatrixv
	.globl	_Z8multiplytt
	.type	_Z8multiplytt, @function
_Z8multiplytt:
.LFB3639:
	.cfi_startproc
	endbr64
	xorl	%ecx, %ecx
	xorl	%eax, %eax
	movzwl	%di, %edi
	movzwl	%si, %esi
.L46:
	btl	%ecx, %edi
	jnc	.L45
	movl	%esi, %edx
	sall	%cl, %edx
	xorl	%edx, %eax
.L45:
	incl	%ecx
	cmpl	$8, %ecx
	jne	.L46
	movzbl	%ah, %ecx
	leaq	modPXMatrix(%rip), %rdx
	movzbl	%al, %eax
	xorw	(%rdx,%rcx,2), %ax
	ret
	.cfi_endproc
.LFE3639:
	.size	_Z8multiplytt, .-_Z8multiplytt
	.globl	_Z9inversiont
	.type	_Z9inversiont, @function
_Z9inversiont:
.LFB3640:
	.cfi_startproc
	endbr64
	pushq	%r14
	.cfi_def_cfa_offset 16
	.cfi_offset 14, -16
	pushq	%r13
	.cfi_def_cfa_offset 24
	.cfi_offset 13, -24
	pushq	%r12
	.cfi_def_cfa_offset 32
	.cfi_offset 12, -32
	pushq	%rbp
	.cfi_def_cfa_offset 40
	.cfi_offset 6, -40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset 3, -48
	subq	$16, %rsp
	.cfi_def_cfa_offset 64
	movq	%fs:40, %rax
	movq	%rax, 8(%rsp)
	xorl	%eax, %eax
	testw	%di, %di
	je	.L52
	movl	%edi, %ebx
	movl	$1, %ebp
	xorl	%r14d, %r14d
	movl	$283, %edx
	leaq	6(%rsp), %r12
.L53:
	movzwl	%bp, %r13d
	cmpw	$1, %bx
	je	.L59
	movzwl	%bx, %ecx
	leaq	4(%rsp), %rdi
	movq	%r12, %rsi
	call	_Z8divisionRtS_tt
	movzwl	4(%rsp), %edi
	movl	%r13d, %esi
	call	_Z8multiplytt
	movzwl	%bx, %edx
	movw	6(%rsp), %bx
	xorl	%r14d, %eax
	movl	%ebp, %r14d
	movl	%eax, %ebp
	jmp	.L53
.L59:
	movq	8(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L58
	addq	$16, %rsp
	.cfi_remember_state
	.cfi_def_cfa_offset 48
	movl	%r13d, %edi
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%rbp
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r13
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	jmp	_Z5modPXt
.L52:
	.cfi_restore_state
	movq	8(%rsp), %rax
	subq	%fs:40, %rax
	je	.L56
.L58:
	call	__stack_chk_fail@PLT
.L56:
	addq	$16, %rsp
	.cfi_def_cfa_offset 48
	xorl	%eax, %eax
	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
.LFE3640:
	.size	_Z9inversiont, .-_Z9inversiont
	.globl	_Z20printInversionMatrixv
	.type	_Z20printInversionMatrixv, @function
_Z20printInversionMatrixv:
.LFB3641:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	leaq	.LC3(%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
	xorl	%ebx, %ebx
	pushq	%rdx
	.cfi_def_cfa_offset 64
.L62:
	leaq	_ZSt4cout(%rip), %r12
	movq	%r13, %rsi
	xorl	%ebp, %ebp
	movq	%r12, %rdi
	leaq	.LC4(%rip), %r14
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
.L61:
	movq	%r14, %rsi
	movq	%r12, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	leaq	_ZSt3hexRSt8ios_base(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSt8ios_baseS0_E@PLT
	leal	0(%rbp,%rbx), %edi
	incl	%ebp
	movq	%rax, %r15
	call	_Z9inversiont
	movq	%r15, %rdi
	movzwl	%ax, %esi
	call	_ZNSo9_M_insertImEERSoT_@PLT
	leaq	.LC5(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	cmpl	$16, %ebp
	jne	.L61
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%r12, %rdi
	addl	$16, %ebx
	call	_ZNSolsEPFRSoS_E@PLT
	cmpl	$256, %ebx
	jne	.L62
	popq	%rax
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3641:
	.size	_Z20printInversionMatrixv, .-_Z20printInversionMatrixv
	.section	.rodata.str1.1
.LC6:
	.string	"Multiply for 100,000 times."
.LC7:
	.string	"uSeconds: "
	.text
	.globl	_Z12testMultiplyv
	.type	_Z12testMultiplyv, @function
_Z12testMultiplyv:
.LFB3642:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	xorl	%esi, %esi
	leaq	_ZSt4cout(%rip), %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	movl	$100000, %ebx
	subq	$56, %rsp
	.cfi_def_cfa_offset 80
	movq	%fs:40, %rax
	movq	%rax, 40(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rdi
	call	gettimeofday@PLT
	leaq	.LC6(%rip), %rsi
	movq	%rbp, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
.L67:
	call	rand@PLT
	call	rand@PLT
	decl	%ebx
	jne	.L67
	xorl	%esi, %esi
	leaq	24(%rsp), %rdi
	call	gettimeofday@PLT
	leaq	.LC7(%rip), %rsi
	movq	%rbp, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	24(%rsp), %rsi
	subq	8(%rsp), %rsi
	imulq	$1000000, %rsi, %rsi
	movq	%rax, %rdi
	addq	32(%rsp), %rsi
	subq	16(%rsp), %rsi
	call	_ZNSo9_M_insertIlEERSoT_@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
	movq	40(%rsp), %rax
	subq	%fs:40, %rax
	je	.L68
	call	__stack_chk_fail@PLT
.L68:
	addq	$56, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3642:
	.size	_Z12testMultiplyv, .-_Z12testMultiplyv
	.section	.rodata.str1.1
.LC8:
	.string	"Inversion for 100,000 times."
	.text
	.globl	_Z13testInversionv
	.type	_Z13testInversionv, @function
_Z13testInversionv:
.LFB3643:
	.cfi_startproc
	endbr64
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	xorl	%esi, %esi
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	leaq	inversionMatrix(%rip), %r12
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	movl	$100000, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	leaq	_ZSt4cout(%rip), %rbx
	subq	$56, %rsp
	.cfi_def_cfa_offset 96
	movq	%fs:40, %rax
	movq	%rax, 40(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rdi
	call	gettimeofday@PLT
	leaq	.LC8(%rip), %rsi
	movq	%rbx, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
.L74:
	movl	$256, %r13d
.L72:
	call	rand@PLT
	cltd
	idivl	%r13d
	movl	%edx, %r8d
	testw	%dx, %dx
	je	.L72
	movzwl	%dx, %eax
	movzwl	%dx, %esi
	movzwl	(%r12,%rax,2), %edi
	call	_Z8multiplytt
	decw	%ax
	je	.L73
	movzwl	%r8w, %esi
	movq	%rbx, %rdi
	call	_ZNSo9_M_insertImEERSoT_@PLT
.L73:
	decl	%ebp
	jne	.L74
	xorl	%esi, %esi
	leaq	24(%rsp), %rdi
	call	gettimeofday@PLT
	leaq	.LC7(%rip), %rsi
	movq	%rbx, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	24(%rsp), %rsi
	subq	8(%rsp), %rsi
	imulq	$1000000, %rsi, %rsi
	movq	%rax, %rdi
	addq	32(%rsp), %rsi
	subq	16(%rsp), %rsi
	call	_ZNSo9_M_insertIlEERSoT_@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
	movq	40(%rsp), %rax
	subq	%fs:40, %rax
	je	.L75
	call	__stack_chk_fail@PLT
.L75:
	addq	$56, %rsp
	.cfi_def_cfa_offset 40
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%rbp
	.cfi_def_cfa_offset 24
	popq	%r12
	.cfi_def_cfa_offset 16
	popq	%r13
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3643:
	.size	_Z13testInversionv, .-_Z13testInversionv
	.globl	_Z26__device_stub__Z5countPtS_PtS_
	.type	_Z26__device_stub__Z5countPtS_PtS_, @function
_Z26__device_stub__Z5countPtS_PtS_:
.LFB3672:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movq	%rdi, 8(%rsp)
	leaq	32(%rsp), %rcx
	leaq	24(%rsp), %rdx
	movq	%rsi, (%rsp)
	leaq	40(%rsp), %rdi
	leaq	52(%rsp), %rsi
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rax
	movl	$1, 48(%rsp)
	movq	%rax, 88(%rsp)
	movq	%rsp, %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 40(%rsp)
	movq	%rax, 52(%rsp)
	movl	$1, 60(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L82
	pushq	32(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z5countPtS_(%rip), %rdi
	pushq	32(%rsp)
	.cfi_def_cfa_offset 144
	movq	68(%rsp), %rcx
	movl	76(%rsp), %r8d
	movq	56(%rsp), %rsi
	movl	64(%rsp), %edx
	leaq	104(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L82:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L84
	call	__stack_chk_fail@PLT
.L84:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3672:
	.size	_Z26__device_stub__Z5countPtS_PtS_, .-_Z26__device_stub__Z5countPtS_PtS_
	.globl	_Z5countPtS_
	.type	_Z5countPtS_, @function
_Z5countPtS_:
.LFB3673:
	.cfi_startproc
	endbr64
	jmp	_Z26__device_stub__Z5countPtS_PtS_
	.cfi_endproc
.LFE3673:
	.size	_Z5countPtS_, .-_Z5countPtS_
	.section	.rodata.str1.1
.LC9:
	.string	"sBox linear analysis with CUDA"
	.text
	.globl	_Z12sBoxAnalysisv
	.type	_Z12sBoxAnalysisv, @function
_Z12sBoxAnalysisv:
.LFB3644:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	leaq	.LC9(%rip), %rsi
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	leaq	_ZSt4cout(%rip), %r12
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movq	%r12, %rdi
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	subq	$136, %rsp
	.cfi_def_cfa_offset 192
	movq	%fs:40, %rax
	movq	%rax, 120(%rsp)
	xorl	%eax, %eax
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
	movq	%rsp, %rdi
	movl	$131072, %esi
	call	cudaMalloc@PLT
	movl	$131072, %esi
	leaq	8(%rsp), %rdi
	call	cudaMalloc@PLT
	movl	$131072, %edi
	call	malloc@PLT
	movl	$131072, %edi
	movq	%rax, %rbp
	call	malloc@PLT
	movq	%rbp, %rcx
	xorl	%edx, %edx
	leaq	sBox(%rip), %rsi
	movq	%rax, %rbx
.L88:
	xorl	%eax, %eax
.L89:
	movzbl	(%rsi,%rax), %edi
	movw	%di, (%rcx,%rax,2)
	incq	%rax
	cmpq	$256, %rax
	jne	.L89
	addl	$256, %edx
	addq	$512, %rcx
	cmpl	$65536, %edx
	jne	.L88
	movq	(%rsp), %rdi
	movl	$1, %ecx
	movl	$131072, %edx
	movq	%rbp, %rsi
	call	cudaMemcpy@PLT
	movl	$16777217, %edi
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	salq	$8, %rdi
	movl	$1, %ecx
	movl	$1, %esi
	movl	$1, 36(%rsp)
	movq	%rdi, %rdx
	movq	%rdi, 28(%rsp)
	movq	%rdi, 16(%rsp)
	movl	$1, 24(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L91
	movq	8(%rsp), %rsi
	movq	(%rsp), %rdi
	call	_Z26__device_stub__Z5countPtS_PtS_
.L91:
	movq	8(%rsp), %rsi
	movl	$2, %ecx
	movl	$131072, %edx
	movq	%rbx, %rdi
	leaq	40(%rsp), %r14
	call	cudaMemcpy@PLT
	movl	$10, %ecx
	xorl	%eax, %eax
	movq	%r14, %rdi
	rep stosl
	movl	$1, %edx
.L96:
	movzwl	(%rbx,%rdx,2), %eax
	addl	$-128, %eax
	movl	%eax, %ecx
	negl	%ecx
	cmovs	%eax, %ecx
	cmpl	%ecx, 76(%rsp)
	jge	.L92
	movl	$8, %eax
.L94:
	movl	(%r14,%rax,4), %esi
	cmpl	%ecx, %esi
	jg	.L95
	movl	%esi, 4(%r14,%rax,4)
	movl	80(%rsp,%rax,4), %esi
	movl	%esi, 84(%rsp,%rax,4)
	subq	$1, %rax
	jnb	.L94
.L95:
	incl	%eax
	cltq
	movl	%ecx, 40(%rsp,%rax,4)
	movl	%edx, 80(%rsp,%rax,4)
.L92:
	incq	%rdx
	cmpq	$65536, %rdx
	jne	.L96
	xorl	%r13d, %r13d
	leaq	.LC3(%rip), %r15
.L97:
	movzwl	80(%rsp,%r13), %edi
	call	_Z9printBitst
	movq	%r15, %rsi
	movq	%r12, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	(%r14,%r13), %esi
	addq	$4, %r13
	movq	%rax, %rdi
	call	_ZNSolsEi@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
	cmpq	$40, %r13
	jne	.L97
	movq	%rbp, %rdi
	call	free@PLT
	movq	%rbx, %rdi
	call	free@PLT
	movq	(%rsp), %rdi
	call	cudaFree@PLT
	movq	8(%rsp), %rdi
	call	cudaFree@PLT
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L98
	call	__stack_chk_fail@PLT
.L98:
	addq	$136, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3644:
	.size	_Z12sBoxAnalysisv, .-_Z12sBoxAnalysisv
	.globl	_Z32__device_stub__Z10d_multiplyPtS_PtS_
	.type	_Z32__device_stub__Z10d_multiplyPtS_PtS_, @function
_Z32__device_stub__Z10d_multiplyPtS_PtS_:
.LFB3674:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movq	%rdi, 8(%rsp)
	leaq	32(%rsp), %rcx
	leaq	24(%rsp), %rdx
	movq	%rsi, (%rsp)
	leaq	40(%rsp), %rdi
	leaq	52(%rsp), %rsi
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rax
	movl	$1, 48(%rsp)
	movq	%rax, 88(%rsp)
	movq	%rsp, %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 40(%rsp)
	movq	%rax, 52(%rsp)
	movl	$1, 60(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L104
	pushq	32(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z10d_multiplyPtS_(%rip), %rdi
	pushq	32(%rsp)
	.cfi_def_cfa_offset 144
	movq	68(%rsp), %rcx
	movl	76(%rsp), %r8d
	movq	56(%rsp), %rsi
	movl	64(%rsp), %edx
	leaq	104(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L104:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L106
	call	__stack_chk_fail@PLT
.L106:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3674:
	.size	_Z32__device_stub__Z10d_multiplyPtS_PtS_, .-_Z32__device_stub__Z10d_multiplyPtS_PtS_
	.globl	_Z10d_multiplyPtS_
	.type	_Z10d_multiplyPtS_, @function
_Z10d_multiplyPtS_:
.LFB3675:
	.cfi_startproc
	endbr64
	jmp	_Z32__device_stub__Z10d_multiplyPtS_PtS_
	.cfi_endproc
.LFE3675:
	.size	_Z10d_multiplyPtS_, .-_Z10d_multiplyPtS_
	.section	.rodata.str1.1
.LC10:
	.string	"Multiply for 100,000 times with GPU."
.LC11:
	.string	"Error"
.LC12:
	.string	"Time(uSeconds): "
	.text
	.globl	_Z11GPUMultiplyv
	.type	_Z11GPUMultiplyv, @function
_Z11GPUMultiplyv:
.LFB3645:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	leaq	.LC10(%rip), %rsi
	movl	$256, %r15d
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	leaq	_ZSt4cout(%rip), %r14
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	movq	%r14, %rdi
	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
	xorl	%ebx, %ebx
	subq	$88, %rsp
	.cfi_def_cfa_offset 144
	movq	%fs:40, %rax
	movq	%rax, 72(%rsp)
	xorl	%eax, %eax
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
	movl	$200000, %edi
	call	malloc@PLT
	movl	$200000, %edi
	movq	%rax, %r12
	call	malloc@PLT
	movl	$200000, %edi
	movq	%rax, %rbp
	call	malloc@PLT
	movq	%rax, %r13
.L110:
	call	rand@PLT
	cltd
	idivl	%r15d
	movw	%dx, (%r12,%rbx)
	call	rand@PLT
	cltd
	idivl	%r15d
	movw	%dx, 0(%rbp,%rbx)
	addq	$2, %rbx
	cmpq	$200000, %rbx
	jne	.L110
	xorl	%esi, %esi
	leaq	40(%rsp), %rdi
	call	gettimeofday@PLT
	movq	%rsp, %rdi
	movl	$200000, %esi
	call	cudaMalloc@PLT
	leaq	8(%rsp), %rdi
	movl	$200000, %esi
	call	cudaMalloc@PLT
	movq	(%rsp), %rdi
	movl	$1, %ecx
	movq	%r12, %rsi
	movl	$200000, %edx
	call	cudaMemcpy@PLT
	movq	8(%rsp), %rdi
	movl	$1, %ecx
	movq	%rbp, %rsi
	movl	$200000, %edx
	call	cudaMemcpy@PLT
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$1, %ecx
	movabsq	$137438953504, %rdx
	movl	$1, %esi
	movabsq	$42949672970, %rdi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L111
	movq	8(%rsp), %rsi
	movq	(%rsp), %rdi
	call	_Z32__device_stub__Z10d_multiplyPtS_PtS_
.L111:
	call	cudaDeviceSynchronize@PLT
	movq	(%rsp), %rsi
	movl	$2, %ecx
	xorl	%ebx, %ebx
	movl	$200000, %edx
	movq	%r13, %rdi
	call	cudaMemcpy@PLT
	movq	(%rsp), %rdi
	call	cudaFree@PLT
	movq	8(%rsp), %rdi
	call	cudaFree@PLT
	leaq	56(%rsp), %rdi
	xorl	%esi, %esi
	call	gettimeofday@PLT
.L114:
	movzwl	0(%rbp,%rbx,2), %esi
	movzwl	(%r12,%rbx,2), %edi
	movw	0(%r13,%rbx,2), %r8w
	call	_Z8multiplytt
	cmpw	%ax, %r8w
	je	.L112
	movq	%r14, %rdi
	leaq	.LC11(%rip), %rsi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	%ebx, %esi
	movq	%rax, %rdi
	call	_ZNSolsEi@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
	jmp	.L113
.L112:
	incq	%rbx
	cmpq	$100000, %rbx
	jne	.L114
.L113:
	movq	%r12, %rdi
	call	free@PLT
	movq	%rbp, %rdi
	call	free@PLT
	movq	%r13, %rdi
	call	free@PLT
	leaq	.LC12(%rip), %rsi
	movq	%r14, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	56(%rsp), %rsi
	subq	40(%rsp), %rsi
	imulq	$1000000, %rsi, %rsi
	movq	%rax, %rdi
	addq	64(%rsp), %rsi
	subq	48(%rsp), %rsi
	call	_ZNSo9_M_insertIlEERSoT_@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
	movq	72(%rsp), %rax
	subq	%fs:40, %rax
	je	.L115
	call	__stack_chk_fail@PLT
.L115:
	addq	$88, %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
.LFE3645:
	.size	_Z11GPUMultiplyv, .-_Z11GPUMultiplyv
	.globl	_Z31__device_stub__Z11d_inversionPtPt
	.type	_Z31__device_stub__Z11d_inversionPtPt, @function
_Z31__device_stub__Z11d_inversionPtPt:
.LFB3676:
	.cfi_startproc
	endbr64
	subq	$104, %rsp
	.cfi_def_cfa_offset 112
	movq	%rdi, 8(%rsp)
	leaq	24(%rsp), %rcx
	leaq	16(%rsp), %rdx
	leaq	44(%rsp), %rsi
	leaq	32(%rsp), %rdi
	movq	%fs:40, %rax
	movq	%rax, 88(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rax
	movl	$1, 40(%rsp)
	movq	%rax, 80(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 32(%rsp)
	movq	%rax, 44(%rsp)
	movl	$1, 52(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L119
	pushq	24(%rsp)
	.cfi_def_cfa_offset 120
	leaq	_Z11d_inversionPt(%rip), %rdi
	pushq	24(%rsp)
	.cfi_def_cfa_offset 128
	movq	60(%rsp), %rcx
	movl	68(%rsp), %r8d
	movq	48(%rsp), %rsi
	movl	56(%rsp), %edx
	leaq	96(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 120
	popq	%rdx
	.cfi_def_cfa_offset 112
.L119:
	movq	88(%rsp), %rax
	subq	%fs:40, %rax
	je	.L121
	call	__stack_chk_fail@PLT
.L121:
	addq	$104, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3676:
	.size	_Z31__device_stub__Z11d_inversionPtPt, .-_Z31__device_stub__Z11d_inversionPtPt
	.globl	_Z11d_inversionPt
	.type	_Z11d_inversionPt, @function
_Z11d_inversionPt:
.LFB3677:
	.cfi_startproc
	endbr64
	jmp	_Z31__device_stub__Z11d_inversionPtPt
	.cfi_endproc
.LFE3677:
	.size	_Z11d_inversionPt, .-_Z11d_inversionPt
	.section	.rodata.str1.1
.LC13:
	.string	"Inversion for 100,000 times with GPU."
.LC14:
	.string	"Error "
	.text
	.globl	_Z12GPUInversionv
	.type	_Z12GPUInversionv, @function
_Z12GPUInversionv:
.LFB3646:
	.cfi_startproc
	endbr64
	pushq	%r14
	.cfi_def_cfa_offset 16
	.cfi_offset 14, -16
	leaq	.LC13(%rip), %rsi
	movl	$256, %r14d
	pushq	%r13
	.cfi_def_cfa_offset 24
	.cfi_offset 13, -24
	leaq	_ZSt4cout(%rip), %r13
	pushq	%r12
	.cfi_def_cfa_offset 32
	.cfi_offset 12, -32
	movq	%r13, %rdi
	pushq	%rbp
	.cfi_def_cfa_offset 40
	.cfi_offset 6, -40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset 3, -48
	xorl	%ebx, %ebx
	subq	$80, %rsp
	.cfi_def_cfa_offset 128
	movq	%fs:40, %rax
	movq	%rax, 72(%rsp)
	xorl	%eax, %eax
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
	movl	$200512, %edi
	call	malloc@PLT
	movl	$200512, %edi
	movq	%rax, %rbp
	call	malloc@PLT
	movq	%rax, %r12
.L125:
	call	rand@PLT
	cltd
	idivl	%r14d
	movw	%dx, 0(%rbp,%rbx,2)
	incq	%rbx
	cmpq	$100000, %rbx
	jne	.L125
	leaq	inversionMatrix(%rip), %rsi
	leaq	200000(%rbp), %rdi
	movl	$128, %ecx
	rep movsl
	xorl	%esi, %esi
	leaq	40(%rsp), %rdi
	call	gettimeofday@PLT
	leaq	8(%rsp), %rdi
	movl	$200512, %esi
	call	cudaMalloc@PLT
	movq	8(%rsp), %rdi
	movl	$1, %ecx
	movq	%rbp, %rsi
	movl	$200512, %edx
	call	cudaMemcpy@PLT
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$1, %ecx
	movabsq	$137438953504, %rdx
	movl	$1, %esi
	movabsq	$42949672970, %rdi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L126
	movq	8(%rsp), %rdi
	call	_Z31__device_stub__Z11d_inversionPtPt
.L126:
	call	cudaDeviceSynchronize@PLT
	movq	8(%rsp), %rsi
	movq	%r12, %rdi
	xorl	%ebx, %ebx
	movl	$2, %ecx
	movl	$200512, %edx
	call	cudaMemcpy@PLT
	movq	8(%rsp), %rdi
	call	cudaFree@PLT
	leaq	56(%rsp), %rdi
	xorl	%esi, %esi
	call	gettimeofday@PLT
.L131:
	movzwl	0(%rbp,%rbx,2), %edi
	movzwl	(%r12,%rbx,2), %esi
	testw	%di, %di
	jne	.L127
	testw	%si, %si
	jmp	.L140
.L127:
	call	_Z8multiplytt
	decw	%ax
.L140:
	je	.L129
	movq	%r13, %rdi
	leaq	.LC14(%rip), %rsi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	%ebx, %esi
	movq	%rax, %rdi
	call	_ZNSolsEi@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
	jmp	.L130
.L129:
	incq	%rbx
	cmpq	$100000, %rbx
	jne	.L131
.L130:
	movq	%rbp, %rdi
	call	free@PLT
	movq	%r12, %rdi
	call	free@PLT
	leaq	.LC12(%rip), %rsi
	movq	%r13, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	56(%rsp), %rsi
	subq	40(%rsp), %rsi
	imulq	$1000000, %rsi, %rsi
	movq	%rax, %rdi
	addq	64(%rsp), %rsi
	subq	48(%rsp), %rsi
	call	_ZNSo9_M_insertIlEERSoT_@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
	movq	72(%rsp), %rax
	subq	%fs:40, %rax
	je	.L132
	call	__stack_chk_fail@PLT
.L132:
	addq	$80, %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
.LFE3646:
	.size	_Z12GPUInversionv, .-_Z12GPUInversionv
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB3647:
	.cfi_startproc
	endbr64
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	xorl	%edi, %edi
	movq	%fs:40, %rax
	movq	%rax, 8(%rsp)
	xorl	%eax, %eax
	call	time@PLT
	movl	%eax, %edi
	call	srand@PLT
	movl	$200, %esi
	movq	%rsp, %rdi
	call	cudaMalloc@PLT
	movq	(%rsp), %rdi
	call	cudaFree@PLT
	call	_Z11GPUMultiplyv
	call	_Z12GPUInversionv
	call	_Z12sBoxAnalysisv
	movq	8(%rsp), %rax
	subq	%fs:40, %rax
	je	.L142
	call	__stack_chk_fail@PLT
.L142:
	xorl	%eax, %eax
	addq	$24, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3647:
	.size	main, .-main
	.section	.rodata.str1.1
.LC15:
	.string	"_Z11d_inversionPt"
.LC16:
	.string	"_Z10d_multiplyPtS_"
.LC17:
	.string	"_Z5countPtS_"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB3679:
	.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	.LC15(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z11d_inversionPt(%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	.LC16(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z10d_multiplyPtS_(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC17(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z5countPtS_(%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
.LFE3679:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.nvFatBinSegment,"aw"
	.align 8
	.type	_ZL15__fatDeviceText, @object
	.size	_ZL15__fatDeviceText, 24
_ZL15__fatDeviceText:
	.long	1180844977
	.long	1
	.quad	fatbinData
	.quad	0
	.local	_ZL20__cudaFatCubinHandle
	.comm	_ZL20__cudaFatCubinHandle,8,8
	.globl	sBox
	.data
	.align 32
	.type	sBox, @object
	.size	sBox, 256
sBox:
	.string	"c|w{\362ko\3050\001g+\376\327\253v\312\202\311}\372YG\360\255\324\242\257\234\244r\300\267\375\223&6?\367\3144\245\345\361q\3301\025\004\307#\303\030\226\005\232\007\022\200\342\353'\262u\t\203,\032\033nZ\240R;\326\263)\343/\204S\321"
	.ascii	"\355 \374\261[j\313\2769JLX\317\320\357\252\373CM3\205E\371\002"
	.ascii	"\177P<\237\250Q\243@\217\222\2358\365\274\266\332!\020\377\363"
	.ascii	"\322\315\f\023\354_\227D\027\304\247~=d]\031s`\201O\334\"*\220"
	.ascii	"\210F\356\270\024\336^\013\333\3402:\nI\006$\\\302\323\254b\221"
	.ascii	"\225\344y\347\3107m\215\325N\251lV\364\352ez\256\b\272x%.\034"
	.ascii	"\246\264\306\350\335t\037K\275\213\212p>\265fH\003\366\016a5"
	.ascii	"W\271\206\301\035\236\341\370\230\021i\331\216\224\233\036\207"
	.ascii	"\351\316U(\337\214\241\211\r\277\346BhA\231-\017\260T\273\026"
	.globl	modPXMatrix
	.align 32
	.type	modPXMatrix, @object
	.size	modPXMatrix, 512
modPXMatrix:
	.value	0
	.value	27
	.value	54
	.value	45
	.value	108
	.value	119
	.value	90
	.value	65
	.value	216
	.value	195
	.value	238
	.value	245
	.value	180
	.value	175
	.value	130
	.value	153
	.value	171
	.value	176
	.value	157
	.value	134
	.value	199
	.value	220
	.value	241
	.value	234
	.value	115
	.value	104
	.value	69
	.value	94
	.value	31
	.value	4
	.value	41
	.value	50
	.value	77
	.value	86
	.value	123
	.value	96
	.value	33
	.value	58
	.value	23
	.value	12
	.value	149
	.value	142
	.value	163
	.value	184
	.value	249
	.value	226
	.value	207
	.value	212
	.value	230
	.value	253
	.value	208
	.value	203
	.value	138
	.value	145
	.value	188
	.value	167
	.value	62
	.value	37
	.value	8
	.value	19
	.value	82
	.value	73
	.value	100
	.value	127
	.value	154
	.value	129
	.value	172
	.value	183
	.value	246
	.value	237
	.value	192
	.value	219
	.value	66
	.value	89
	.value	116
	.value	111
	.value	46
	.value	53
	.value	24
	.value	3
	.value	49
	.value	42
	.value	7
	.value	28
	.value	93
	.value	70
	.value	107
	.value	112
	.value	233
	.value	242
	.value	223
	.value	196
	.value	133
	.value	158
	.value	179
	.value	168
	.value	215
	.value	204
	.value	225
	.value	250
	.value	187
	.value	160
	.value	141
	.value	150
	.value	15
	.value	20
	.value	57
	.value	34
	.value	99
	.value	120
	.value	85
	.value	78
	.value	124
	.value	103
	.value	74
	.value	81
	.value	16
	.value	11
	.value	38
	.value	61
	.value	164
	.value	191
	.value	146
	.value	137
	.value	200
	.value	211
	.value	254
	.value	229
	.value	47
	.value	52
	.value	25
	.value	2
	.value	67
	.value	88
	.value	117
	.value	110
	.value	247
	.value	236
	.value	193
	.value	218
	.value	155
	.value	128
	.value	173
	.value	182
	.value	132
	.value	159
	.value	178
	.value	169
	.value	232
	.value	243
	.value	222
	.value	197
	.value	92
	.value	71
	.value	106
	.value	113
	.value	48
	.value	43
	.value	6
	.value	29
	.value	98
	.value	121
	.value	84
	.value	79
	.value	14
	.value	21
	.value	56
	.value	35
	.value	186
	.value	161
	.value	140
	.value	151
	.value	214
	.value	205
	.value	224
	.value	251
	.value	201
	.value	210
	.value	255
	.value	228
	.value	165
	.value	190
	.value	147
	.value	136
	.value	17
	.value	10
	.value	39
	.value	60
	.value	125
	.value	102
	.value	75
	.value	80
	.value	181
	.value	174
	.value	131
	.value	152
	.value	217
	.value	194
	.value	239
	.value	244
	.value	109
	.value	118
	.value	91
	.value	64
	.value	1
	.value	26
	.value	55
	.value	44
	.value	30
	.value	5
	.value	40
	.value	51
	.value	114
	.value	105
	.value	68
	.value	95
	.value	198
	.value	221
	.value	240
	.value	235
	.value	170
	.value	177
	.value	156
	.value	135
	.value	248
	.value	227
	.value	206
	.value	213
	.value	148
	.value	143
	.value	162
	.value	185
	.value	32
	.value	59
	.value	22
	.value	13
	.value	76
	.value	87
	.value	122
	.value	97
	.value	83
	.value	72
	.value	101
	.value	126
	.value	63
	.value	36
	.value	9
	.value	18
	.value	139
	.value	144
	.value	189
	.value	166
	.value	231
	.value	252
	.value	209
	.value	202
	.globl	inversionMatrix
	.align 32
	.type	inversionMatrix, @object
	.size	inversionMatrix, 512
inversionMatrix:
	.value	0
	.value	1
	.value	141
	.value	246
	.value	203
	.value	82
	.value	123
	.value	209
	.value	232
	.value	79
	.value	41
	.value	192
	.value	176
	.value	225
	.value	229
	.value	199
	.value	116
	.value	180
	.value	170
	.value	75
	.value	153
	.value	43
	.value	96
	.value	95
	.value	88
	.value	63
	.value	253
	.value	204
	.value	255
	.value	64
	.value	238
	.value	178
	.value	58
	.value	110
	.value	90
	.value	241
	.value	85
	.value	77
	.value	168
	.value	201
	.value	193
	.value	10
	.value	152
	.value	21
	.value	48
	.value	68
	.value	162
	.value	194
	.value	44
	.value	69
	.value	146
	.value	108
	.value	243
	.value	57
	.value	102
	.value	66
	.value	242
	.value	53
	.value	32
	.value	111
	.value	119
	.value	187
	.value	89
	.value	25
	.value	29
	.value	254
	.value	55
	.value	103
	.value	45
	.value	49
	.value	245
	.value	105
	.value	167
	.value	100
	.value	171
	.value	19
	.value	84
	.value	37
	.value	233
	.value	9
	.value	237
	.value	92
	.value	5
	.value	202
	.value	76
	.value	36
	.value	135
	.value	191
	.value	24
	.value	62
	.value	34
	.value	240
	.value	81
	.value	236
	.value	97
	.value	23
	.value	22
	.value	94
	.value	175
	.value	211
	.value	73
	.value	166
	.value	54
	.value	67
	.value	244
	.value	71
	.value	145
	.value	223
	.value	51
	.value	147
	.value	33
	.value	59
	.value	121
	.value	183
	.value	151
	.value	133
	.value	16
	.value	181
	.value	186
	.value	60
	.value	182
	.value	112
	.value	208
	.value	6
	.value	161
	.value	250
	.value	129
	.value	130
	.value	131
	.value	126
	.value	127
	.value	128
	.value	150
	.value	115
	.value	190
	.value	86
	.value	155
	.value	158
	.value	149
	.value	217
	.value	247
	.value	2
	.value	185
	.value	164
	.value	222
	.value	106
	.value	50
	.value	109
	.value	216
	.value	138
	.value	132
	.value	114
	.value	42
	.value	20
	.value	159
	.value	136
	.value	249
	.value	220
	.value	137
	.value	154
	.value	251
	.value	124
	.value	46
	.value	195
	.value	143
	.value	184
	.value	101
	.value	72
	.value	38
	.value	200
	.value	18
	.value	74
	.value	206
	.value	231
	.value	210
	.value	98
	.value	12
	.value	224
	.value	31
	.value	239
	.value	17
	.value	117
	.value	120
	.value	113
	.value	165
	.value	142
	.value	118
	.value	61
	.value	189
	.value	188
	.value	134
	.value	87
	.value	11
	.value	40
	.value	47
	.value	163
	.value	218
	.value	212
	.value	228
	.value	15
	.value	169
	.value	39
	.value	83
	.value	4
	.value	27
	.value	252
	.value	172
	.value	230
	.value	122
	.value	7
	.value	174
	.value	99
	.value	197
	.value	219
	.value	226
	.value	234
	.value	148
	.value	139
	.value	196
	.value	213
	.value	157
	.value	248
	.value	144
	.value	107
	.value	177
	.value	13
	.value	214
	.value	235
	.value	198
	.value	14
	.value	207
	.value	173
	.value	8
	.value	78
	.value	215
	.value	227
	.value	93
	.value	80
	.value	30
	.value	179
	.value	91
	.value	35
	.value	56
	.value	52
	.value	104
	.value	70
	.value	3
	.value	140
	.value	221
	.value	156
	.value	125
	.value	160
	.value	205
	.value	26
	.value	65
	.value	28
	.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:

/*
	file: AES.cu
	BY: FU ERDONG
*/

#define PX 283
#define MODPX 27

#include <iostream>
#include <stdlib.h>
#include <sys/time.h>
#include <math.h>
#include <hip/hip_runtime.h>
using namespace std;

unsigned short int inversionMatrix[256] =
{
	0x0, 0x1, 0x8d, 0xf6, 0xcb, 0x52, 0x7b, 0xd1, 0xe8, 0x4f, 0x29, 0xc0, 0xb0, 0xe1, 0xe5, 0xc7, 
	0x74, 0xb4, 0xaa, 0x4b, 0x99, 0x2b, 0x60, 0x5f, 0x58, 0x3f, 0xfd, 0xcc, 0xff, 0x40, 0xee, 0xb2, 
	0x3a, 0x6e, 0x5a, 0xf1, 0x55, 0x4d, 0xa8, 0xc9, 0xc1, 0xa, 0x98, 0x15, 0x30, 0x44, 0xa2, 0xc2, 
	0x2c, 0x45, 0x92, 0x6c, 0xf3, 0x39, 0x66, 0x42, 0xf2, 0x35, 0x20, 0x6f, 0x77, 0xbb, 0x59, 0x19, 
	0x1d, 0xfe, 0x37, 0x67, 0x2d, 0x31, 0xf5, 0x69, 0xa7, 0x64, 0xab, 0x13, 0x54, 0x25, 0xe9, 0x9, 
	0xed, 0x5c, 0x5, 0xca, 0x4c, 0x24, 0x87, 0xbf, 0x18, 0x3e, 0x22, 0xf0, 0x51, 0xec, 0x61, 0x17, 
	0x16, 0x5e, 0xaf, 0xd3, 0x49, 0xa6, 0x36, 0x43, 0xf4, 0x47, 0x91, 0xdf, 0x33, 0x93, 0x21, 0x3b, 
	0x79, 0xb7, 0x97, 0x85, 0x10, 0xb5, 0xba, 0x3c, 0xb6, 0x70, 0xd0, 0x6, 0xa1, 0xfa, 0x81, 0x82, 
	0x83, 0x7e, 0x7f, 0x80, 0x96, 0x73, 0xbe, 0x56, 0x9b, 0x9e, 0x95, 0xd9, 0xf7, 0x2, 0xb9, 0xa4, 
	0xde, 0x6a, 0x32, 0x6d, 0xd8, 0x8a, 0x84, 0x72, 0x2a, 0x14, 0x9f, 0x88, 0xf9, 0xdc, 0x89, 0x9a, 
	0xfb, 0x7c, 0x2e, 0xc3, 0x8f, 0xb8, 0x65, 0x48, 0x26, 0xc8, 0x12, 0x4a, 0xce, 0xe7, 0xd2, 0x62, 
	0xc, 0xe0, 0x1f, 0xef, 0x11, 0x75, 0x78, 0x71, 0xa5, 0x8e, 0x76, 0x3d, 0xbd, 0xbc, 0x86, 0x57, 
	0xb, 0x28, 0x2f, 0xa3, 0xda, 0xd4, 0xe4, 0xf, 0xa9, 0x27, 0x53, 0x4, 0x1b, 0xfc, 0xac, 0xe6, 
	0x7a, 0x7, 0xae, 0x63, 0xc5, 0xdb, 0xe2, 0xea, 0x94, 0x8b, 0xc4, 0xd5, 0x9d, 0xf8, 0x90, 0x6b, 
	0xb1, 0xd, 0xd6, 0xeb, 0xc6, 0xe, 0xcf, 0xad, 0x8, 0x4e, 0xd7, 0xe3, 0x5d, 0x50, 0x1e, 0xb3, 
	0x5b, 0x23, 0x38, 0x34, 0x68, 0x46, 0x3, 0x8c, 0xdd, 0x9c, 0x7d, 0xa0, 0xcd, 0x1a, 0x41, 0x1c
};

unsigned short int modPXMatrix[256] = 
{
	0x0, 0x1b, 0x36, 0x2d, 0x6c, 0x77, 0x5a, 0x41, 0xd8, 0xc3, 0xee, 0xf5, 0xb4, 0xaf, 0x82, 0x99, 
	0xab, 0xb0, 0x9d, 0x86, 0xc7, 0xdc, 0xf1, 0xea, 0x73, 0x68, 0x45, 0x5e, 0x1f, 0x4, 0x29, 0x32, 
	0x4d, 0x56, 0x7b, 0x60, 0x21, 0x3a, 0x17, 0xc, 0x95, 0x8e, 0xa3, 0xb8, 0xf9, 0xe2, 0xcf, 0xd4, 
	0xe6, 0xfd, 0xd0, 0xcb, 0x8a, 0x91, 0xbc, 0xa7, 0x3e, 0x25, 0x8, 0x13, 0x52, 0x49, 0x64, 0x7f, 
	0x9a, 0x81, 0xac, 0xb7, 0xf6, 0xed, 0xc0, 0xdb, 0x42, 0x59, 0x74, 0x6f, 0x2e, 0x35, 0x18, 0x3, 
	0x31, 0x2a, 0x7, 0x1c, 0x5d, 0x46, 0x6b, 0x70, 0xe9, 0xf2, 0xdf, 0xc4, 0x85, 0x9e, 0xb3, 0xa8, 
	0xd7, 0xcc, 0xe1, 0xfa, 0xbb, 0xa0, 0x8d, 0x96, 0xf, 0x14, 0x39, 0x22, 0x63, 0x78, 0x55, 0x4e, 
	0x7c, 0x67, 0x4a, 0x51, 0x10, 0xb, 0x26, 0x3d, 0xa4, 0xbf, 0x92, 0x89, 0xc8, 0xd3, 0xfe, 0xe5, 
	0x2f, 0x34, 0x19, 0x2, 0x43, 0x58, 0x75, 0x6e, 0xf7, 0xec, 0xc1, 0xda, 0x9b, 0x80, 0xad, 0xb6, 
	0x84, 0x9f, 0xb2, 0xa9, 0xe8, 0xf3, 0xde, 0xc5, 0x5c, 0x47, 0x6a, 0x71, 0x30, 0x2b, 0x6, 0x1d, 
	0x62, 0x79, 0x54, 0x4f, 0xe, 0x15, 0x38, 0x23, 0xba, 0xa1, 0x8c, 0x97, 0xd6, 0xcd, 0xe0, 0xfb, 
	0xc9, 0xd2, 0xff, 0xe4, 0xa5, 0xbe, 0x93, 0x88, 0x11, 0xa, 0x27, 0x3c, 0x7d, 0x66, 0x4b, 0x50, 
	0xb5, 0xae, 0x83, 0x98, 0xd9, 0xc2, 0xef, 0xf4, 0x6d, 0x76, 0x5b, 0x40, 0x1, 0x1a, 0x37, 0x2c, 
	0x1e, 0x5, 0x28, 0x33, 0x72, 0x69, 0x44, 0x5f, 0xc6, 0xdd, 0xf0, 0xeb, 0xaa, 0xb1, 0x9c, 0x87, 
	0xf8, 0xe3, 0xce, 0xd5, 0x94, 0x8f, 0xa2, 0xb9, 0x20, 0x3b, 0x16, 0xd, 0x4c, 0x57, 0x7a, 0x61, 
	0x53, 0x48, 0x65, 0x7e, 0x3f, 0x24, 0x9, 0x12, 0x8b, 0x90, 0xbd, 0xa6, 0xe7, 0xfc, 0xd1, 0xca, 
};

unsigned char sBox[] =
{ /*  0    1    2    3    4    5    6    7    8    9    a    b    c    d    e    f */
    0x63,0x7c,0x77,0x7b,0xf2,0x6b,0x6f,0xc5,0x30,0x01,0x67,0x2b,0xfe,0xd7,0xab,0x76, /*0*/ 
    0xca,0x82,0xc9,0x7d,0xfa,0x59,0x47,0xf0,0xad,0xd4,0xa2,0xaf,0x9c,0xa4,0x72,0xc0, /*1*/
    0xb7,0xfd,0x93,0x26,0x36,0x3f,0xf7,0xcc,0x34,0xa5,0xe5,0xf1,0x71,0xd8,0x31,0x15, /*2*/
    0x04,0xc7,0x23,0xc3,0x18,0x96,0x05,0x9a,0x07,0x12,0x80,0xe2,0xeb,0x27,0xb2,0x75, /*3*/
    0x09,0x83,0x2c,0x1a,0x1b,0x6e,0x5a,0xa0,0x52,0x3b,0xd6,0xb3,0x29,0xe3,0x2f,0x84, /*4*/
    0x53,0xd1,0x00,0xed,0x20,0xfc,0xb1,0x5b,0x6a,0xcb,0xbe,0x39,0x4a,0x4c,0x58,0xcf, /*5*/
    0xd0,0xef,0xaa,0xfb,0x43,0x4d,0x33,0x85,0x45,0xf9,0x02,0x7f,0x50,0x3c,0x9f,0xa8, /*6*/ 
    0x51,0xa3,0x40,0x8f,0x92,0x9d,0x38,0xf5,0xbc,0xb6,0xda,0x21,0x10,0xff,0xf3,0xd2, /*7*/
    0xcd,0x0c,0x13,0xec,0x5f,0x97,0x44,0x17,0xc4,0xa7,0x7e,0x3d,0x64,0x5d,0x19,0x73, /*8*/
    0x60,0x81,0x4f,0xdc,0x22,0x2a,0x90,0x88,0x46,0xee,0xb8,0x14,0xde,0x5e,0x0b,0xdb, /*9*/
    0xe0,0x32,0x3a,0x0a,0x49,0x06,0x24,0x5c,0xc2,0xd3,0xac,0x62,0x91,0x95,0xe4,0x79, /*a*/
    0xe7,0xc8,0x37,0x6d,0x8d,0xd5,0x4e,0xa9,0x6c,0x56,0xf4,0xea,0x65,0x7a,0xae,0x08, /*b*/
    0xba,0x78,0x25,0x2e,0x1c,0xa6,0xb4,0xc6,0xe8,0xdd,0x74,0x1f,0x4b,0xbd,0x8b,0x8a, /*c*/
    0x70,0x3e,0xb5,0x66,0x48,0x03,0xf6,0x0e,0x61,0x35,0x57,0xb9,0x86,0xc1,0x1d,0x9e, /*d*/
    0xe1,0xf8,0x98,0x11,0x69,0xd9,0x8e,0x94,0x9b,0x1e,0x87,0xe9,0xce,0x55,0x28,0xdf, /*e*/
    0x8c,0xa1,0x89,0x0d,0xbf,0xe6,0x42,0x68,0x41,0x99,0x2d,0x0f,0xb0,0x54,0xbb,0x16  /*f*/
};

void printBits(unsigned short int num)
{
	for (int i = 0; i < 16; i++ ) {
		if (i == 8) cout << " ";
		if (num & (1 << (15 - i))) {
			cout << "1";
		} else {
			cout << "0";
		}
	}
}

void division(unsigned short int &q, unsigned short int &r, unsigned short int a, unsigned short int b)
{
	int la, lb, i;

	q = 0;
	for (i = 15; i > 0; i--) if (a & (1 << i)) break;
	la = i;
	for (i = 15; i > 0; i--) if (b & (1 << i)) break;
	lb = i;
	for (i = la; i >= lb; i--)
		if (a & (1 << i)) {
			a ^= b << (i - lb);
			q ^= 1 << (i - lb);
		}
	r = a;
}

unsigned short int modPX(unsigned short int num)
{
	unsigned short int px;

	px = PX << 7;
	for (int i = 15; i >= 8; i--) {
		if (num & (1 << i)) num = num ^ px;
		px >>= 1;
	}
	return num;
}

void printModPXMatrix()
{
	for (int i = 0; i < 16; i++) {
		cout << "\t";
		for (int j = 0; j < 16; j++) {
			cout << "0x" << hex << modPX((i * 16 + j) << 8) << ", ";
		}
		cout << endl;
	}
}

unsigned short int multiply(unsigned short int a, unsigned short int b)
{
	unsigned short int c = 0;
	for (int i = 0; i < 8; i++)
		if (a & (1 << i)) c ^= (b << i);
	return ((c & 0x00ff) ^ modPXMatrix[c >> 8]);
}

unsigned short int inversion(unsigned short int num)
{
	unsigned short int a, b, q, r, k11, k12, k21, k22, k11n, k12n, k21n, k22n;

	if (num == 0) return 0;
	k11 = 1;
	k22 = 1;
	k12 = 0;
	k21 = 0;
	a = PX;
	b = num;

	while (b != 1) {
		division(q, r, a, b);
		a = b;
		b = r;
		k11n = k21;
		k12n = k22;
		k21n = k11 ^ multiply(q, k21);
		k22n = k12 ^ multiply(q, k22);
		k11 = k11n;
		k12 = k12n;
		k21 = k21n;
		k22 = k22n;
	}
	return modPX(k22);
}

void printInversionMatrix()
{
	for (int i = 0; i < 16; i++) {
		cout << "\t";
		for (int j = 0; j < 16; j++) {
			cout << "0x" << hex << inversion(i * 16 + j) << ", ";
		}
		cout << endl;
	}
}

void testMultiply()
{
	struct timeval beginTime, endTime;

//multiply * 100,000
	gettimeofday(&beginTime, 0);
	cout << "Multiply for 100,000 times." << endl;
	for (int i = 0; i < 100000; i++) multiply(rand() % 256, rand() % 256);
	gettimeofday(&endTime, 0);
	cout << "uSeconds: "
		<< (endTime.tv_sec - beginTime.tv_sec) * 1000000 
		+ endTime.tv_usec - beginTime.tv_usec << endl << endl;

}

void testInversion()
{
	struct timeval beginTime, endTime;
	unsigned short int temp;


// inversion * 100,000
	gettimeofday(&beginTime, 0);
	cout << "Inversion for 100,000 times." << endl;
	for (int i = 0; i < 100000; i++) {
		while (!(temp = rand() % 256));
		if (multiply(inversionMatrix[temp], temp) - 1) cout << temp;
	}
	gettimeofday(&endTime, 0);
	cout << "uSeconds: "
		<< (endTime.tv_sec - beginTime.tv_sec) * 1000000 
		+ endTime.tv_usec - beginTime.tv_usec << endl << endl;

}

__global__ void count(unsigned short int *in, unsigned short int *out) {
	unsigned short int gindex = threadIdx.x + blockIdx.x * blockDim.x;
	unsigned short int baseIndex = blockIdx.x * blockDim.x;

	unsigned short int counter = 0;
	unsigned short int temp;
	for (int i = 0; i < 256; i++) {
		temp = (i << 8) ^ in[baseIndex + i];
		temp &= gindex;
		for (int j = 1; j < 16; j++)
			if (temp & (1 << j)) temp ^= 0x01;
		if (!(temp & 0x1)) counter++;
	}

	out[gindex] = counter;
}

__global__ void d_multiply(unsigned short int *d_a, unsigned short int *d_b) {

	int index = 320 * 32 * blockIdx.y + 320 * threadIdx.y + 32 * blockIdx.x + threadIdx.x;
	if (index >= 100000) return;

	unsigned short int d_c = 0;
	for (int i = 7; i >= 0; i--) {
		d_c <<= 1;
		if (d_c & 256) d_c ^= PX;
		if (d_b[index] & (1 << i)) d_c ^= d_a[index];
	}
	d_a[index] = d_c;
}

__global__ void d_inversion(unsigned short int *d_a) {

	int gindex = 320 * 32 * blockIdx.y + 320 * threadIdx.y + 32 * blockIdx.x + threadIdx.x;
	unsigned short int lindex = threadIdx.y * 32 + threadIdx.x;

	__shared__ unsigned short int matrix[256];

	if (threadIdx.y < 8) matrix[lindex] = d_a[100000 + lindex];

	__syncthreads();

	if (gindex >= 100000) return;
	d_a[gindex] = matrix[d_a[gindex]];
}


void sBoxAnalysis()
{
//	struct timeval begin, end;

/*

// sBox Linear Analysis in normal way
	cout << "sBox linear analysis in normal way" << endl;
	gettimeofday(&begin, 0);
	int check[256 * 256];
	int temp;
	int result;
	for (int i = 1; i < (256 * 256); i++) {
		result = 0;
		for (int j = 0; j < 256; j++) {
			temp = (j << 8) ^ sBox[j];
			temp &= i;
			for (int k = 1; k < 16; k++)
				if (temp & (1 << k)) temp ^= 0x01;
			if (!(temp & 0x1)) result++;
		}
		check[i] = result;
	}

	gettimeofday(&end, 0);
	cout << "uSeconds: "
		<< (end.tv_sec - begin.tv_sec) * 1000000 
		+ end.tv_usec - begin.tv_usec << endl << endl;

*/

// use CUDA to accelerate
	cout << "sBox linear analysis with CUDA" << endl;
	unsigned short int *in, *counter;
	unsigned short int *d_in, *d_counter;
	int size = 256 * 256 * sizeof(unsigned short int);

//	gettimeofday(&begin, 0);

	hipMalloc((void**)&d_in, size);
	hipMalloc((void**)&d_counter, size);

	in = (unsigned short int*)malloc(size);
	counter = (unsigned short int*)malloc(size);


	for (int i = 0; i < 256; i++)
		for (int j = 0; j < 256; j++)
			in[i * 256 + j] = sBox[j];

	hipMemcpy(d_in, in, size, hipMemcpyHostToDevice);

	count<<<256, 256>>>(d_in, d_counter);

	hipMemcpy(counter, d_counter, size, hipMemcpyDeviceToHost);

	int N = 10;
	int dmax[N], imax[N];
	int det;
	for (int i = 0; i < N; i++) dmax[i] = 0;
	for (int i = 1; i < 65536; i++) {
		int j;
		det = abs(counter[i] - 128);
		if (det <= dmax[N - 1]) continue;
		for (j = N - 2; j >= 0; j--) {
			if (det < dmax[j]) break;
			dmax[j + 1] = dmax[j];
			imax[j + 1] = imax[j];
		}
		dmax[j + 1] = det;
		imax[j + 1] = i;
	}
	for (int i = 0; i < N; i++) {
		printBits(imax[i]);
		cout << "\t" << dmax[i] << endl;
	}


//	gettimeofday(&end, 0);
//	cout << "Time in total" << endl;
//	cout << "uSeconds: " << (end.tv_sec - begin.tv_sec) * 1000000 + end.tv_usec - begin.tv_usec << endl << endl;

/*
// check
	for (int i = 1; i < 256 * 256; i++)
		if (check[i] != counter[i]) {
			cout << i << "\t" << check[i] << "\t" << hex << counter[i] << endl;
			//break;
		}
*/


	free(in);
	free(counter);
	hipFree(d_in);
	hipFree(d_counter);
}

void GPUMultiply()
{
	struct timeval begin, end;

	cout << "Multiply for 100,000 times with GPU." << endl;

	int memSize = 100000 * sizeof(unsigned short int);
	
	unsigned short int *a, *b, *c;

	a = (unsigned short int *)malloc(memSize);
	b = (unsigned short int *)malloc(memSize);
	c = (unsigned short int *)malloc(memSize);

	for (int i = 0; i < 100000; i++) {
		a[i] = rand() % 256;
		b[i] = rand() % 256;
	}


	gettimeofday(&begin, 0);

	unsigned short int *d_a, *d_b;

	hipMalloc((void**) &d_a, memSize);
	hipMalloc((void**) &d_b, memSize);

	hipMemcpy(d_a, a, memSize, hipMemcpyHostToDevice);
	hipMemcpy(d_b, b, memSize, hipMemcpyHostToDevice);

	dim3 threads(32, 32);
	dim3 grid(10, 10);

	d_multiply<<<grid, threads>>>(d_a, d_b);

	hipDeviceSynchronize();

	hipMemcpy(c, d_a, memSize, hipMemcpyDeviceToHost);

	hipFree(d_a);
	hipFree(d_b);

	gettimeofday(&end, 0);

//check
	for (int i = 0; i < 100000; i++)
		if (c[i] != multiply(a[i], b[i])) {
			cout << "Error" << i << endl;
			break;
		}


	free(a);
	free(b);
	free(c);

	cout << "Time(uSeconds): " << (end.tv_sec - begin.tv_sec) * 1000000 
		+ end.tv_usec - begin.tv_usec << endl << endl;

}

void GPUInversion()
{
	struct timeval begin, end;

	cout << "Inversion for 100,000 times with GPU." << endl;

	int memSize = (100000 + 256) * sizeof(unsigned short int);
	
	unsigned short int *a, *b;

	a = (unsigned short int *)malloc(memSize);
	b = (unsigned short int *)malloc(memSize);

	for (int i = 0; i < 100000; i++) {
		a[i] = rand() % 256;
	}

	for (int i = 0; i < 256; i++) {
		a[i + 100000] = inversionMatrix[i];
	}

	gettimeofday(&begin, 0);

	unsigned short int *d_a;

	hipMalloc((void**) &d_a, memSize);

	hipMemcpy(d_a, a, memSize, hipMemcpyHostToDevice);

	dim3 threads(32, 32);
	dim3 grid(10, 10);

	d_inversion<<<grid, threads>>>(d_a);

	hipDeviceSynchronize();

	hipMemcpy(b, d_a, memSize, hipMemcpyDeviceToHost);

	hipFree(d_a);

	gettimeofday(&end, 0);


//check
	for (int i = 0; i < 100000; i++)
		if ((a[i] == 0 && b[i] != 0) || (a[i] != 0 && multiply(a[i], b[i]) != 1)) {
			cout << "Error " << i << endl;
			break;
		}


	free(a);
	free(b);

	cout << "Time(uSeconds): " << (end.tv_sec - begin.tv_sec) * 1000000 
		+ end.tv_usec - begin.tv_usec << endl << endl;

}





int main()
{
	srand(time(0));

	unsigned short int *temp;
	hipMalloc((void**)&temp, 100 * sizeof(unsigned short int));
	hipFree(temp);

	GPUMultiply();
	GPUInversion();

//	testMultiply();
//	testInversion();

	sBoxAnalysis();

}







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

                                        # End of file scope inline assembly
	.globl	_Z9printBitst                   # -- Begin function _Z9printBitst
	.type	_Z9printBitst,@function
_Z9printBitst:                          # @_Z9printBitst
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	pushq	%rax
	.cfi_def_cfa_offset 32
	.cfi_offset %rbx, -24
	.cfi_offset %rbp, -16
	movl	%edi, %ebx
	xorl	%ebp, %ebp
.LBB0_1:                                # =>This Inner Loop Header: Depth=1
	cmpl	$8, %ebp
	jne	.LBB0_3
# %bb.2:                                #   in Loop: Header=BB0_1 Depth=1
	movl	$_ZSt4cout, %edi
	movl	$.L.str, %esi
	movl	$1, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
.LBB0_3:                                #   in Loop: Header=BB0_1 Depth=1
	movl	%ebp, %eax
	xorb	$15, %al
	movzbl	%al, %eax
	btl	%eax, %ebx
	movl	$.L.str.2, %esi
	jae	.LBB0_5
# %bb.4:                                #   in Loop: Header=BB0_1 Depth=1
	movl	$.L.str.1, %esi
.LBB0_5:                                #   in Loop: Header=BB0_1 Depth=1
	movl	$_ZSt4cout, %edi
	movl	$1, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	incl	%ebp
	cmpl	$16, %ebp
	jne	.LBB0_1
# %bb.6:
	addq	$8, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end0:
	.size	_Z9printBitst, .Lfunc_end0-_Z9printBitst
	.cfi_endproc
                                        # -- End function
	.globl	_Z8divisionRtS_tt               # -- Begin function _Z8divisionRtS_tt
	.type	_Z8divisionRtS_tt,@function
_Z8divisionRtS_tt:                      # @_Z8divisionRtS_tt
	.cfi_startproc
# %bb.0:
	movw	$0, (%rdi)
	movl	$16, %eax
.LBB1_1:                                # =>This Inner Loop Header: Depth=1
	leal	-1(%rax), %r8d
	movzbl	%r8b, %r8d
	decl	%eax
	btl	%r8d, %edx
	jb	.LBB1_4
# %bb.2:                                #   in Loop: Header=BB1_1 Depth=1
	cmpl	$1, %eax
	ja	.LBB1_1
# %bb.3:
	xorl	%eax, %eax
.LBB1_4:
	movzwl	%cx, %r8d
	movl	$16, %r9d
.LBB1_5:                                # =>This Inner Loop Header: Depth=1
	leal	-1(%r9), %ecx
	movzbl	%cl, %ecx
	decl	%r9d
	btl	%ecx, %r8d
	jb	.LBB1_8
# %bb.6:                                #   in Loop: Header=BB1_5 Depth=1
	cmpl	$1, %r9d
	ja	.LBB1_5
# %bb.7:
	xorl	%r9d, %r9d
.LBB1_8:
	cmpl	%r9d, %eax
	jl	.LBB1_14
# %bb.9:                                # %.lr.ph.preheader
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset %rbx, -16
	movl	%r9d, %r10d
	negl	%r10d
	xorl	%r11d, %r11d
.LBB1_10:                               # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movzwl	%dx, %ecx
	btl	%eax, %ecx
	jae	.LBB1_12
# %bb.11:                               #   in Loop: Header=BB1_10 Depth=1
	leal	(%r10,%rax), %ecx
	movl	%r8d, %ebx
	shll	%cl, %ebx
	xorl	%ebx, %edx
	btcl	%ecx, %r11d
	movw	%r11w, (%rdi)
.LBB1_12:                               #   in Loop: Header=BB1_10 Depth=1
	leal	-1(%rax), %ecx
	cmpl	%r9d, %eax
	movl	%ecx, %eax
	jg	.LBB1_10
# %bb.13:
	popq	%rbx
	.cfi_def_cfa_offset 8
	.cfi_restore %rbx
.LBB1_14:                               # %._crit_edge
	movw	%dx, (%rsi)
	retq
.Lfunc_end1:
	.size	_Z8divisionRtS_tt, .Lfunc_end1-_Z8divisionRtS_tt
	.cfi_endproc
                                        # -- End function
	.globl	_Z5modPXt                       # -- Begin function _Z5modPXt
	.type	_Z5modPXt,@function
_Z5modPXt:                              # @_Z5modPXt
	.cfi_startproc
# %bb.0:
	movl	%edi, %eax
	movw	$-29312, %cx                    # imm = 0x8D80
	movl	$16, %edx
.LBB2_1:                                # =>This Inner Loop Header: Depth=1
	decl	%edx
	movzwl	%ax, %edi
	xorl	%esi, %esi
	btl	%edx, %edi
	jae	.LBB2_3
# %bb.2:                                #   in Loop: Header=BB2_1 Depth=1
	movl	%ecx, %esi
.LBB2_3:                                #   in Loop: Header=BB2_1 Depth=1
	xorl	%esi, %eax
	movzwl	%cx, %ecx
	shrl	%ecx
	cmpl	$8, %edx
	ja	.LBB2_1
# %bb.4:
                                        # kill: def $ax killed $ax killed $eax
	retq
.Lfunc_end2:
	.size	_Z5modPXt, .Lfunc_end2-_Z5modPXt
	.cfi_endproc
                                        # -- End function
	.globl	_Z16printModPXMatrixv           # -- Begin function _Z16printModPXMatrixv
	.type	_Z16printModPXMatrixv,@function
_Z16printModPXMatrixv:                  # @_Z16printModPXMatrixv
	.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	%rbx
	.cfi_def_cfa_offset 40
	pushq	%rax
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	xorl	%ebx, %ebx
	movl	$_ZSt4cout, %r14d
.LBB3_1:                                # =>This Loop Header: Depth=1
                                        #     Child Loop BB3_2 Depth 2
                                        #       Child Loop BB3_3 Depth 3
	movl	$_ZSt4cout, %edi
	movl	$.L.str.3, %esi
	movl	$1, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	%ebx, %ebp
	shll	$12, %ebp
	xorl	%r15d, %r15d
.LBB3_2:                                #   Parent Loop BB3_1 Depth=1
                                        # =>  This Loop Header: Depth=2
                                        #       Child Loop BB3_3 Depth 3
	movl	$_ZSt4cout, %edi
	movl	$.L.str.4, %esi
	movl	$2, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rax
	movl	_ZSt4cout+24(%rax), %ecx
	andl	$-75, %ecx
	orl	$8, %ecx
	movl	%ecx, _ZSt4cout+24(%rax)
	movl	%r15d, %eax
	shll	$8, %eax
	addl	%ebp, %eax
	movw	$-29312, %cx                    # imm = 0x8D80
	movl	$16, %edx
.LBB3_3:                                #   Parent Loop BB3_1 Depth=1
                                        #     Parent Loop BB3_2 Depth=2
                                        # =>    This Inner Loop Header: Depth=3
	decl	%edx
	movzwl	%ax, %edi
	xorl	%esi, %esi
	btl	%edx, %edi
	jae	.LBB3_5
# %bb.4:                                #   in Loop: Header=BB3_3 Depth=3
	movl	%ecx, %esi
.LBB3_5:                                #   in Loop: Header=BB3_3 Depth=3
	xorl	%esi, %eax
	movzwl	%cx, %ecx
	shrl	%ecx
	cmpl	$8, %edx
	ja	.LBB3_3
# %bb.6:                                # %_Z5modPXt.exit
                                        #   in Loop: Header=BB3_2 Depth=2
	movzwl	%ax, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo9_M_insertImEERSoT_
	movl	$.L.str.5, %esi
	movl	$2, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	incl	%r15d
	cmpl	$16, %r15d
	jne	.LBB3_2
# %bb.7:                                #   in Loop: Header=BB3_1 Depth=1
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%r14, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	incl	%ebx
	cmpl	$16, %ebx
	jne	.LBB3_1
# %bb.8:
	addq	$8, %rsp
	.cfi_def_cfa_offset 40
	popq	%rbx
	.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_end3:
	.size	_Z16printModPXMatrixv, .Lfunc_end3-_Z16printModPXMatrixv
	.cfi_endproc
                                        # -- End function
	.globl	_Z8multiplytt                   # -- Begin function _Z8multiplytt
	.type	_Z8multiplytt,@function
_Z8multiplytt:                          # @_Z8multiplytt
	.cfi_startproc
# %bb.0:
	xorl	%eax, %eax
	xorl	%ecx, %ecx
	xorl	%edx, %edx
.LBB4_1:                                # =>This Inner Loop Header: Depth=1
	movl	%esi, %r8d
	shll	%cl, %r8d
	btl	%ecx, %edi
	cmovael	%eax, %r8d
	xorl	%r8d, %edx
	incl	%ecx
	cmpl	$8, %ecx
	jne	.LBB4_1
# %bb.2:
	movl	$255, %ecx
	movl	%edx, %eax
	andl	%ecx, %eax
	shrl	$8, %edx
	andl	%ecx, %edx
	xorw	modPXMatrix(%rdx,%rdx), %ax
                                        # kill: def $ax killed $ax killed $eax
	retq
.Lfunc_end4:
	.size	_Z8multiplytt, .Lfunc_end4-_Z8multiplytt
	.cfi_endproc
                                        # -- End function
	.globl	_Z9inversiont                   # -- Begin function _Z9inversiont
	.type	_Z9inversiont,@function
_Z9inversiont:                          # @_Z9inversiont
	.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
	movl	%edi, %ebx
	testl	%edi, %edi
	je	.LBB5_11
# %bb.1:
	cmpl	$1, %ebx
	jne	.LBB5_3
# %bb.2:
	movl	%ebx, %esi
	jmp	.LBB5_7
.LBB5_3:                                # %.lr.ph.preheader
	movw	$1, %si
	movw	$283, %cx                       # imm = 0x11B
	xorl	%ebp, %ebp
	leaq	4(%rsp), %r15
	movl	$255, %r12d
	xorl	%r13d, %r13d
.LBB5_4:                                # %.lr.ph
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB5_5 Depth 2
	movl	%esi, %r14d
	movzwl	%cx, %edx
	movzwl	%bx, %ecx
	leaq	6(%rsp), %rdi
	movq	%r15, %rsi
	callq	_Z8divisionRtS_tt
	movzwl	4(%rsp), %eax
	movzwl	6(%rsp), %esi
	movzwl	%r14w, %edi
	xorl	%ecx, %ecx
	xorl	%edx, %edx
.LBB5_5:                                #   Parent Loop BB5_4 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movl	%edi, %r8d
	shll	%cl, %r8d
	btl	%ecx, %esi
	cmovael	%ebp, %r8d
	xorl	%r8d, %edx
	incl	%ecx
	cmpl	$8, %ecx
	jne	.LBB5_5
# %bb.6:                                # %_Z8multiplytt.exit29
                                        #   in Loop: Header=BB5_4 Depth=1
	movl	%edx, %ecx
	andl	%r12d, %ecx
	shrl	$8, %edx
	andl	%r12d, %edx
	xorl	%ecx, %r13d
	movl	%r13d, %esi
	xorw	modPXMatrix(%rdx,%rdx), %si
	movl	%ebx, %ecx
	movl	%eax, %ebx
	movl	%r14d, %r13d
	cmpw	$1, %ax
	jne	.LBB5_4
.LBB5_7:                                # %.preheader.preheader
	movw	$-29312, %ax                    # imm = 0x8D80
	movl	$16, %ecx
	movl	%esi, %ebx
.LBB5_8:                                # %.preheader
                                        # =>This Inner Loop Header: Depth=1
	decl	%ecx
	movzwl	%bx, %esi
	xorl	%edx, %edx
	btl	%ecx, %esi
	jae	.LBB5_10
# %bb.9:                                # %.preheader
                                        #   in Loop: Header=BB5_8 Depth=1
	movl	%eax, %edx
.LBB5_10:                               # %.preheader
                                        #   in Loop: Header=BB5_8 Depth=1
	xorl	%edx, %ebx
	movzwl	%ax, %eax
	shrl	%eax
	cmpl	$8, %ecx
	ja	.LBB5_8
.LBB5_11:                               # %_Z5modPXt.exit
	movl	%ebx, %eax
	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
	retq
.Lfunc_end5:
	.size	_Z9inversiont, .Lfunc_end5-_Z9inversiont
	.cfi_endproc
                                        # -- End function
	.globl	_Z20printInversionMatrixv       # -- Begin function _Z20printInversionMatrixv
	.type	_Z20printInversionMatrixv,@function
_Z20printInversionMatrixv:              # @_Z20printInversionMatrixv
	.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	%rbx
	.cfi_def_cfa_offset 40
	pushq	%rax
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	xorl	%ebx, %ebx
	movl	$_ZSt4cout, %r14d
	xorl	%ebp, %ebp
.LBB6_1:                                # =>This Loop Header: Depth=1
                                        #     Child Loop BB6_2 Depth 2
	movl	$_ZSt4cout, %edi
	movl	$.L.str.3, %esi
	movl	$1, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	xorl	%r15d, %r15d
.LBB6_2:                                #   Parent Loop BB6_1 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movl	$_ZSt4cout, %edi
	movl	$.L.str.4, %esi
	movl	$2, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rax
	movl	_ZSt4cout+24(%rax), %ecx
	andl	$-75, %ecx
	orl	$8, %ecx
	movl	%ecx, _ZSt4cout+24(%rax)
	leal	(%rbx,%r15), %eax
	movzwl	%ax, %edi
	callq	_Z9inversiont
	movzwl	%ax, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo9_M_insertImEERSoT_
	movl	$.L.str.5, %esi
	movl	$2, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	incl	%r15d
	cmpl	$16, %r15d
	jne	.LBB6_2
# %bb.3:                                #   in Loop: Header=BB6_1 Depth=1
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%r14, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	incl	%ebp
	addl	$16, %ebx
	cmpl	$16, %ebp
	jne	.LBB6_1
# %bb.4:
	addq	$8, %rsp
	.cfi_def_cfa_offset 40
	popq	%rbx
	.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_end6:
	.size	_Z20printInversionMatrixv, .Lfunc_end6-_Z20printInversionMatrixv
	.cfi_endproc
                                        # -- End function
	.globl	_Z12testMultiplyv               # -- Begin function _Z12testMultiplyv
	.type	_Z12testMultiplyv,@function
_Z12testMultiplyv:                      # @_Z12testMultiplyv
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	subq	$32, %rsp
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -16
	movq	%rsp, %rdi
	xorl	%esi, %esi
	callq	gettimeofday
	movl	$_ZSt4cout, %ebx
	movl	$_ZSt4cout, %edi
	movl	$.L.str.6, %esi
	movl	$27, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	$100000, %ebx                   # imm = 0x186A0
.LBB7_1:                                # %_Z8multiplytt.exit
                                        # =>This Inner Loop Header: Depth=1
	callq	rand
	callq	rand
	decl	%ebx
	jne	.LBB7_1
# %bb.2:
	leaq	16(%rsp), %rbx
	movq	%rbx, %rdi
	xorl	%esi, %esi
	callq	gettimeofday
	movl	$_ZSt4cout, %edi
	movl	$.L.str.7, %esi
	movl	$10, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	(%rbx), %rax
	subq	(%rsp), %rax
	imulq	$1000000, %rax, %rsi            # imm = 0xF4240
	addq	8(%rbx), %rsi
	subq	8(%rsp), %rsi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo9_M_insertIlEERSoT_
	movq	%rax, %rbx
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%rbx, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movq	%rax, %rbx
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%rbx, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	_ZNSo5flushEv                   # TAILCALL
.Lfunc_end7:
	.size	_Z12testMultiplyv, .Lfunc_end7-_Z12testMultiplyv
	.cfi_endproc
                                        # -- End function
	.globl	_Z13testInversionv              # -- Begin function _Z13testInversionv
	.type	_Z13testInversionv,@function
_Z13testInversionv:                     # @_Z13testInversionv
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	subq	$32, %rsp
	.cfi_def_cfa_offset 64
	.cfi_offset %rbx, -32
	.cfi_offset %r14, -24
	.cfi_offset %rbp, -16
	xorl	%ebx, %ebx
	movq	%rsp, %rdi
	xorl	%esi, %esi
	callq	gettimeofday
	movl	$_ZSt4cout, %r14d
	movl	$_ZSt4cout, %edi
	movl	$.L.str.8, %esi
	movl	$28, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%r14, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	$255, %ebp
	xorl	%r14d, %r14d
.LBB8_1:                                # =>This Loop Header: Depth=1
                                        #     Child Loop BB8_3 Depth 2
	callq	rand
                                        # kill: def $eax killed $eax def $rax
	leal	255(%rax), %ecx
	testl	%eax, %eax
	cmovnsl	%eax, %ecx
	andl	$-256, %ecx
	subl	%ecx, %eax
	movl	%eax, %ecx
	andl	$65535, %ecx                    # imm = 0xFFFF
	je	.LBB8_1
# %bb.2:                                #   in Loop: Header=BB8_1 Depth=1
	movl	%ecx, %esi
	movzwl	inversionMatrix(%rsi,%rsi), %edi
	xorl	%ecx, %ecx
	xorl	%edx, %edx
.LBB8_3:                                #   Parent Loop BB8_1 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movl	%eax, %r8d
	shll	%cl, %r8d
	btl	%ecx, %edi
	cmovael	%ebx, %r8d
	xorl	%r8d, %edx
	incl	%ecx
	cmpl	$8, %ecx
	jne	.LBB8_3
# %bb.4:                                # %_Z8multiplytt.exit
                                        #   in Loop: Header=BB8_1 Depth=1
	movl	%edx, %eax
	andl	%ebp, %eax
	shrl	$8, %edx
	andl	%ebp, %edx
	xorw	modPXMatrix(%rdx,%rdx), %ax
	cmpw	$1, %ax
	je	.LBB8_6
# %bb.5:                                #   in Loop: Header=BB8_1 Depth=1
	movl	$_ZSt4cout, %edi
	callq	_ZNSo9_M_insertImEERSoT_
.LBB8_6:                                #   in Loop: Header=BB8_1 Depth=1
	incl	%r14d
	cmpl	$100000, %r14d                  # imm = 0x186A0
	jne	.LBB8_1
# %bb.7:
	leaq	16(%rsp), %rbx
	movq	%rbx, %rdi
	xorl	%esi, %esi
	callq	gettimeofday
	movl	$_ZSt4cout, %edi
	movl	$.L.str.7, %esi
	movl	$10, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	(%rbx), %rax
	subq	(%rsp), %rax
	imulq	$1000000, %rax, %rsi            # imm = 0xF4240
	addq	8(%rbx), %rsi
	subq	8(%rsp), %rsi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo9_M_insertIlEERSoT_
	movq	%rax, %rbx
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%rbx, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movq	%rax, %rbx
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%rbx, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	addq	$32, %rsp
	.cfi_def_cfa_offset 32
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	jmp	_ZNSo5flushEv                   # TAILCALL
.Lfunc_end8:
	.size	_Z13testInversionv, .Lfunc_end8-_Z13testInversionv
	.cfi_endproc
                                        # -- End function
	.globl	_Z20__device_stub__countPtS_    # -- Begin function _Z20__device_stub__countPtS_
	.type	_Z20__device_stub__countPtS_,@function
_Z20__device_stub__countPtS_:           # @_Z20__device_stub__countPtS_
	.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	$80, %rsp
	.cfi_def_cfa_offset 128
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	24(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	16(%rsp), %rcx
	movq	%rsi, (%rcx)
	leaq	64(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%rbx)
	leaq	48(%rsp), %r14
	leaq	32(%rsp), %r15
	leaq	8(%rsp), %r12
	movq	%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	$_Z5countPtS_, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$96, %rsp
	.cfi_adjust_cfa_offset -96
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end9:
	.size	_Z20__device_stub__countPtS_, .Lfunc_end9-_Z20__device_stub__countPtS_
	.cfi_endproc
                                        # -- End function
	.globl	_Z25__device_stub__d_multiplyPtS_ # -- Begin function _Z25__device_stub__d_multiplyPtS_
	.type	_Z25__device_stub__d_multiplyPtS_,@function
_Z25__device_stub__d_multiplyPtS_:      # @_Z25__device_stub__d_multiplyPtS_
	.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	$80, %rsp
	.cfi_def_cfa_offset 128
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	24(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	16(%rsp), %rcx
	movq	%rsi, (%rcx)
	leaq	64(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%rbx)
	leaq	48(%rsp), %r14
	leaq	32(%rsp), %r15
	leaq	8(%rsp), %r12
	movq	%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	$_Z10d_multiplyPtS_, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$96, %rsp
	.cfi_adjust_cfa_offset -96
	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_end10:
	.size	_Z25__device_stub__d_multiplyPtS_, .Lfunc_end10-_Z25__device_stub__d_multiplyPtS_
	.cfi_endproc
                                        # -- End function
	.globl	_Z26__device_stub__d_inversionPt # -- Begin function _Z26__device_stub__d_inversionPt
	.type	_Z26__device_stub__d_inversionPt,@function
_Z26__device_stub__d_inversionPt:       # @_Z26__device_stub__d_inversionPt
	.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	$64, %rsp
	.cfi_def_cfa_offset 112
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	24(%rsp), %rax
	movq	%rdi, (%rax)
	movq	%rsp, %rbx
	movq	%rax, (%rbx)
	leaq	48(%rsp), %r14
	leaq	32(%rsp), %r15
	leaq	16(%rsp), %r12
	leaq	8(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z11d_inversionPt, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$80, %rsp
	.cfi_adjust_cfa_offset -80
	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_end11:
	.size	_Z26__device_stub__d_inversionPt, .Lfunc_end11-_Z26__device_stub__d_inversionPt
	.cfi_endproc
                                        # -- End function
	.globl	_Z12sBoxAnalysisv               # -- Begin function _Z12sBoxAnalysisv
	.type	_Z12sBoxAnalysisv,@function
_Z12sBoxAnalysisv:                      # @_Z12sBoxAnalysisv
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset %rbp, -16
	movq	%rsp, %rbp
	.cfi_def_cfa_register %rbp
	pushq	%r15
	pushq	%r14
	pushq	%r13
	pushq	%r12
	pushq	%rbx
	subq	$40, %rsp
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	movl	$_ZSt4cout, %ebx
	movl	$_ZSt4cout, %edi
	movl	$.L.str.9, %esi
	movl	$30, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	leaq	-56(%rbp), %rdi
	movl	$131072, %esi                   # imm = 0x20000
	callq	hipMalloc
	leaq	-48(%rbp), %rdi
	movl	$131072, %esi                   # imm = 0x20000
	callq	hipMalloc
	movl	$131072, %edi                   # imm = 0x20000
	callq	malloc
	movq	%rax, %rbx
	movl	$131072, %edi                   # imm = 0x20000
	callq	malloc
	movq	%rax, %r14
	xorl	%eax, %eax
	movq	%rbx, %rcx
.LBB12_1:                               # %.preheader63
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB12_2 Depth 2
	xorl	%edx, %edx
.LBB12_2:                               #   Parent Loop BB12_1 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movzbl	sBox(%rdx), %esi
	movw	%si, (%rcx,%rdx,2)
	incq	%rdx
	cmpq	$256, %rdx                      # imm = 0x100
	jne	.LBB12_2
# %bb.3:                                #   in Loop: Header=BB12_1 Depth=1
	incq	%rax
	addq	$512, %rcx                      # imm = 0x200
	cmpq	$256, %rax                      # imm = 0x100
	jne	.LBB12_1
# %bb.4:
	movq	-56(%rbp), %rdi
	movl	$131072, %edx                   # imm = 0x20000
	movq	%rbx, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movabsq	$4294967552, %rdi               # imm = 0x100000100
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB12_6
# %bb.5:
	movq	-56(%rbp), %rdi
	movq	-48(%rbp), %rsi
	callq	_Z20__device_stub__countPtS_
.LBB12_6:                               # %.preheader62
	movq	%rbx, -72(%rbp)                 # 8-byte Spill
	movq	-48(%rbp), %rsi
	movl	$131072, %edx                   # imm = 0x20000
	movq	%r14, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movq	%rsp, -64(%rbp)                 # 8-byte Spill
	movq	%rsp, %rax
	leaq	-48(%rax), %r13
	movq	%r13, %rsp
	movq	%rsp, %r12
	addq	$-48, %r12
	movq	%r12, %rsp
	xorps	%xmm0, %xmm0
	movaps	%xmm0, -32(%rax)
	movaps	%xmm0, -48(%rax)
	movq	$0, -16(%rax)
	movl	$1, %eax
.LBB12_7:                               # =>This Loop Header: Depth=1
                                        #     Child Loop BB12_9 Depth 2
	movzwl	(%r14,%rax,2), %edx
	addl	$-128, %edx
	movl	%edx, %ecx
	negl	%ecx
	cmovsl	%edx, %ecx
	cmpl	36(%r13), %ecx
	jle	.LBB12_13
# %bb.8:                                # %.preheader61.preheader
                                        #   in Loop: Header=BB12_7 Depth=1
	movl	$8, %edx
	movl	$9, %esi
.LBB12_9:                               # %.preheader61
                                        #   Parent Loop BB12_7 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movl	%edx, %edi
	movl	(%r13,%rdi,4), %r8d
	cmpl	%r8d, %ecx
	jl	.LBB12_12
# %bb.10:                               #   in Loop: Header=BB12_9 Depth=2
	movl	%r8d, (%r13,%rsi,4)
	movl	(%r12,%rdi,4), %edi
	movl	%edi, (%r12,%rsi,4)
	decq	%rsi
	addl	$-1, %edx
	jb	.LBB12_9
# %bb.11:                               #   in Loop: Header=BB12_7 Depth=1
	movl	$-1, %edx
.LBB12_12:                              #   in Loop: Header=BB12_7 Depth=1
	movslq	%edx, %rdx
	movl	%ecx, 4(%r13,%rdx,4)
	movl	%eax, 4(%r12,%rdx,4)
.LBB12_13:                              #   in Loop: Header=BB12_7 Depth=1
	incq	%rax
	cmpq	$65536, %rax                    # imm = 0x10000
	jne	.LBB12_7
# %bb.14:                               # %.preheader.preheader
	xorl	%ebx, %ebx
.LBB12_15:                              # %.preheader
                                        # =>This Inner Loop Header: Depth=1
	movzwl	(%r12,%rbx,4), %edi
	callq	_Z9printBitst
	movl	$_ZSt4cout, %edi
	movl	$.L.str.3, %esi
	movl	$1, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	(%r13,%rbx,4), %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSolsEi
	movq	%rax, %r15
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%r15, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%r15, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	incq	%rbx
	cmpq	$10, %rbx
	jne	.LBB12_15
# %bb.16:
	movq	-72(%rbp), %rdi                 # 8-byte Reload
	callq	free
	movq	%r14, %rdi
	callq	free
	movq	-56(%rbp), %rdi
	callq	hipFree
	movq	-48(%rbp), %rdi
	callq	hipFree
	movq	-64(%rbp), %rsp                 # 8-byte Reload
	leaq	-40(%rbp), %rsp
	popq	%rbx
	popq	%r12
	popq	%r13
	popq	%r14
	popq	%r15
	popq	%rbp
	.cfi_def_cfa %rsp, 8
	retq
.Lfunc_end12:
	.size	_Z12sBoxAnalysisv, .Lfunc_end12-_Z12sBoxAnalysisv
	.cfi_endproc
                                        # -- End function
	.globl	_Z11GPUMultiplyv                # -- Begin function _Z11GPUMultiplyv
	.type	_Z11GPUMultiplyv,@function
_Z11GPUMultiplyv:                       # @_Z11GPUMultiplyv
	.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	$56, %rsp
	.cfi_def_cfa_offset 112
	.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	$_ZSt4cout, %ebx
	movl	$_ZSt4cout, %edi
	movl	$.L.str.10, %esi
	movl	$36, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	$200000, %edi                   # imm = 0x30D40
	callq	malloc
	movq	%rax, %rbx
	movl	$200000, %edi                   # imm = 0x30D40
	callq	malloc
	movq	%rax, %r14
	movl	$200000, %edi                   # imm = 0x30D40
	callq	malloc
	movq	%rax, %r15
	xorl	%r12d, %r12d
	movl	$255, %r13d
	movl	$65280, %ebp                    # imm = 0xFF00
.LBB13_1:                               # =>This Inner Loop Header: Depth=1
	callq	rand
                                        # kill: def $eax killed $eax def $rax
	leal	(%rax,%r13), %ecx
	testl	%eax, %eax
	cmovnsl	%eax, %ecx
	andl	%ebp, %ecx
	subl	%ecx, %eax
	movw	%ax, (%rbx,%r12,2)
	callq	rand
                                        # kill: def $eax killed $eax def $rax
	leal	(%rax,%r13), %ecx
	testl	%eax, %eax
	cmovnsl	%eax, %ecx
	andl	%ebp, %ecx
	subl	%ecx, %eax
	movw	%ax, (%r14,%r12,2)
	incq	%r12
	cmpq	$100000, %r12                   # imm = 0x186A0
	jne	.LBB13_1
# %bb.2:
	xorl	%ebp, %ebp
	leaq	40(%rsp), %rdi
	xorl	%esi, %esi
	callq	gettimeofday
	leaq	8(%rsp), %r12
	movl	$200000, %esi                   # imm = 0x30D40
	movq	%r12, %rdi
	callq	hipMalloc
	leaq	16(%rsp), %r13
	movl	$200000, %esi                   # imm = 0x30D40
	movq	%r13, %rdi
	callq	hipMalloc
	movq	(%r12), %rdi
	movl	$200000, %edx                   # imm = 0x30D40
	movq	%rbx, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movq	(%r13), %rdi
	movl	$200000, %edx                   # imm = 0x30D40
	movq	%r14, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movabsq	$42949672970, %rdi              # imm = 0xA0000000A
	movabsq	$137438953504, %rdx             # imm = 0x2000000020
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB13_4
# %bb.3:
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	callq	_Z25__device_stub__d_multiplyPtS_
.LBB13_4:
	callq	hipDeviceSynchronize
	movq	8(%rsp), %rsi
	movl	$200000, %edx                   # imm = 0x30D40
	movq	%r15, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movq	8(%rsp), %rdi
	callq	hipFree
	movq	16(%rsp), %rdi
	callq	hipFree
	leaq	24(%rsp), %rdi
	xorl	%esi, %esi
	callq	gettimeofday
	movl	$255, %eax
	xorl	%r12d, %r12d
.LBB13_5:                               # =>This Loop Header: Depth=1
                                        #     Child Loop BB13_6 Depth 2
	movzwl	(%r15,%r12,2), %edx
	movzwl	(%rbx,%r12,2), %edi
	movzwl	(%r14,%r12,2), %r8d
	xorl	%ecx, %ecx
	xorl	%esi, %esi
.LBB13_6:                               #   Parent Loop BB13_5 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movl	%r8d, %r9d
	shll	%cl, %r9d
	btl	%ecx, %edi
	cmovael	%ebp, %r9d
	xorl	%r9d, %esi
	incl	%ecx
	cmpl	$8, %ecx
	jne	.LBB13_6
# %bb.7:                                # %_Z8multiplytt.exit
                                        #   in Loop: Header=BB13_5 Depth=1
	movl	%esi, %ecx
	andl	%eax, %ecx
	shrl	$8, %esi
	andl	%eax, %esi
	xorw	modPXMatrix(%rsi,%rsi), %cx
	cmpw	%cx, %dx
	jne	.LBB13_8
# %bb.9:                                #   in Loop: Header=BB13_5 Depth=1
	incq	%r12
	cmpq	$100000, %r12                   # imm = 0x186A0
	jne	.LBB13_5
	jmp	.LBB13_10
.LBB13_8:
	movl	$_ZSt4cout, %edi
	movl	$.L.str.11, %esi
	movl	$5, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	$_ZSt4cout, %edi
	movl	%r12d, %esi
	callq	_ZNSolsEi
	movq	%rax, %r12
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%r12, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%r12, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
.LBB13_10:                              # %.loopexit
	movq	%rbx, %rdi
	callq	free
	movq	%r14, %rdi
	callq	free
	movq	%r15, %rdi
	callq	free
	movl	$_ZSt4cout, %edi
	movl	$.L.str.12, %esi
	movl	$16, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	24(%rsp), %rax
	subq	40(%rsp), %rax
	imulq	$1000000, %rax, %rsi            # imm = 0xF4240
	addq	32(%rsp), %rsi
	subq	48(%rsp), %rsi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo9_M_insertIlEERSoT_
	movq	%rax, %rbx
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%rbx, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movq	%rax, %rbx
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%rbx, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	addq	$56, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end13:
	.size	_Z11GPUMultiplyv, .Lfunc_end13-_Z11GPUMultiplyv
	.cfi_endproc
                                        # -- End function
	.globl	_Z12GPUInversionv               # -- Begin function _Z12GPUInversionv
	.type	_Z12GPUInversionv,@function
_Z12GPUInversionv:                      # @_Z12GPUInversionv
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r12
	.cfi_def_cfa_offset 32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	subq	$40, %rsp
	.cfi_def_cfa_offset 80
	.cfi_offset %rbx, -40
	.cfi_offset %r12, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movl	$_ZSt4cout, %ebx
	movl	$_ZSt4cout, %edi
	movl	$.L.str.13, %esi
	movl	$37, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	$200512, %edi                   # imm = 0x30F40
	callq	malloc
	movq	%rax, %rbx
	movl	$200512, %edi                   # imm = 0x30F40
	callq	malloc
	movq	%rax, %r14
	xorl	%r15d, %r15d
.LBB14_1:                               # =>This Inner Loop Header: Depth=1
	callq	rand
                                        # kill: def $eax killed $eax def $rax
	leal	255(%rax), %ecx
	testl	%eax, %eax
	cmovnsl	%eax, %ecx
	andl	$65280, %ecx                    # imm = 0xFF00
	subl	%ecx, %eax
	movw	%ax, (%rbx,%r15,2)
	incq	%r15
	cmpq	$100000, %r15                   # imm = 0x186A0
	jne	.LBB14_1
# %bb.2:                                # %.preheader
	movq	%rbx, %rdi
	addq	$200000, %rdi                   # imm = 0x30D40
	movl	$inversionMatrix, %esi
	movl	$512, %edx                      # imm = 0x200
	callq	memcpy@PLT
	xorl	%r12d, %r12d
	leaq	24(%rsp), %rdi
	xorl	%esi, %esi
	callq	gettimeofday
	movq	%rsp, %r15
	movl	$200512, %esi                   # imm = 0x30F40
	movq	%r15, %rdi
	callq	hipMalloc
	movq	(%r15), %rdi
	movl	$200512, %edx                   # imm = 0x30F40
	movq	%rbx, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movabsq	$42949672970, %rdi              # imm = 0xA0000000A
	movabsq	$137438953504, %rdx             # imm = 0x2000000020
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB14_4
# %bb.3:
	movq	(%rsp), %rdi
	callq	_Z26__device_stub__d_inversionPt
.LBB14_4:
	callq	hipDeviceSynchronize
	movq	(%rsp), %rsi
	movl	$200512, %edx                   # imm = 0x30F40
	movq	%r14, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movq	(%rsp), %rdi
	callq	hipFree
	leaq	8(%rsp), %rdi
	xorl	%esi, %esi
	callq	gettimeofday
	movl	$255, %eax
	xorl	%r15d, %r15d
.LBB14_5:                               # =>This Loop Header: Depth=1
                                        #     Child Loop BB14_9 Depth 2
	movzwl	(%rbx,%r15,2), %esi
	movzwl	(%r14,%r15,2), %edi
	testl	%esi, %esi
	je	.LBB14_6
# %bb.8:                                #   in Loop: Header=BB14_5 Depth=1
	xorl	%ecx, %ecx
	xorl	%edx, %edx
.LBB14_9:                               #   Parent Loop BB14_5 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movl	%edi, %r8d
	shll	%cl, %r8d
	btl	%ecx, %esi
	cmovael	%r12d, %r8d
	xorl	%r8d, %edx
	incl	%ecx
	cmpl	$8, %ecx
	jne	.LBB14_9
# %bb.10:                               # %_Z8multiplytt.exit
                                        #   in Loop: Header=BB14_5 Depth=1
	movl	%edx, %ecx
	andl	%eax, %ecx
	shrl	$8, %edx
	andl	%eax, %edx
	xorw	modPXMatrix(%rdx,%rdx), %cx
	cmpw	$1, %cx
	je	.LBB14_7
	jmp	.LBB14_11
.LBB14_6:                               #   in Loop: Header=BB14_5 Depth=1
	testl	%edi, %edi
	jne	.LBB14_11
.LBB14_7:                               # %.thread
                                        #   in Loop: Header=BB14_5 Depth=1
	incq	%r15
	cmpq	$100000, %r15                   # imm = 0x186A0
	jne	.LBB14_5
	jmp	.LBB14_12
.LBB14_11:
	movl	$_ZSt4cout, %edi
	movl	$.L.str.14, %esi
	movl	$6, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	$_ZSt4cout, %edi
	movl	%r15d, %esi
	callq	_ZNSolsEi
	movq	%rax, %r15
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%r15, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%r15, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
.LBB14_12:                              # %.loopexit
	movq	%rbx, %rdi
	callq	free
	movq	%r14, %rdi
	callq	free
	movl	$_ZSt4cout, %edi
	movl	$.L.str.12, %esi
	movl	$16, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	8(%rsp), %rax
	subq	24(%rsp), %rax
	imulq	$1000000, %rax, %rsi            # imm = 0xF4240
	addq	16(%rsp), %rsi
	subq	32(%rsp), %rsi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo9_M_insertIlEERSoT_
	movq	%rax, %rbx
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%rbx, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movq	%rax, %rbx
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%rbx, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	addq	$40, %rsp
	.cfi_def_cfa_offset 40
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end14:
	.size	_Z12GPUInversionv, .Lfunc_end14-_Z12GPUInversionv
	.cfi_endproc
                                        # -- End function
	.globl	main                            # -- Begin function main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	subq	$16, %rsp
	.cfi_def_cfa_offset 32
	.cfi_offset %rbx, -16
	xorl	%edi, %edi
	callq	time
	movl	%eax, %edi
	callq	srand
	leaq	8(%rsp), %rbx
	movl	$200, %esi
	movq	%rbx, %rdi
	callq	hipMalloc
	movq	(%rbx), %rdi
	callq	hipFree
	callq	_Z11GPUMultiplyv
	callq	_Z12GPUInversionv
	callq	_Z12sBoxAnalysisv
	xorl	%eax, %eax
	addq	$16, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end15:
	.size	main, .Lfunc_end15-main
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_ctor
	.type	__hip_module_ctor,@function
__hip_module_ctor:                      # @__hip_module_ctor
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	subq	$32, %rsp
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -16
	movq	__hip_gpubin_handle(%rip), %rbx
	testq	%rbx, %rbx
	jne	.LBB16_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB16_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z5countPtS_, %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	$_Z10d_multiplyPtS_, %esi
	movl	$.L__unnamed_2, %edx
	movl	$.L__unnamed_2, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z11d_inversionPt, %esi
	movl	$.L__unnamed_3, %edx
	movl	$.L__unnamed_3, %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_end16:
	.size	__hip_module_ctor, .Lfunc_end16-__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	.LBB17_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
.LBB17_2:
	retq
.Lfunc_end17:
	.size	__hip_module_dtor, .Lfunc_end17-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	inversionMatrix,@object         # @inversionMatrix
	.data
	.globl	inversionMatrix
	.p2align	4, 0x0
inversionMatrix:
	.short	0                               # 0x0
	.short	1                               # 0x1
	.short	141                             # 0x8d
	.short	246                             # 0xf6
	.short	203                             # 0xcb
	.short	82                              # 0x52
	.short	123                             # 0x7b
	.short	209                             # 0xd1
	.short	232                             # 0xe8
	.short	79                              # 0x4f
	.short	41                              # 0x29
	.short	192                             # 0xc0
	.short	176                             # 0xb0
	.short	225                             # 0xe1
	.short	229                             # 0xe5
	.short	199                             # 0xc7
	.short	116                             # 0x74
	.short	180                             # 0xb4
	.short	170                             # 0xaa
	.short	75                              # 0x4b
	.short	153                             # 0x99
	.short	43                              # 0x2b
	.short	96                              # 0x60
	.short	95                              # 0x5f
	.short	88                              # 0x58
	.short	63                              # 0x3f
	.short	253                             # 0xfd
	.short	204                             # 0xcc
	.short	255                             # 0xff
	.short	64                              # 0x40
	.short	238                             # 0xee
	.short	178                             # 0xb2
	.short	58                              # 0x3a
	.short	110                             # 0x6e
	.short	90                              # 0x5a
	.short	241                             # 0xf1
	.short	85                              # 0x55
	.short	77                              # 0x4d
	.short	168                             # 0xa8
	.short	201                             # 0xc9
	.short	193                             # 0xc1
	.short	10                              # 0xa
	.short	152                             # 0x98
	.short	21                              # 0x15
	.short	48                              # 0x30
	.short	68                              # 0x44
	.short	162                             # 0xa2
	.short	194                             # 0xc2
	.short	44                              # 0x2c
	.short	69                              # 0x45
	.short	146                             # 0x92
	.short	108                             # 0x6c
	.short	243                             # 0xf3
	.short	57                              # 0x39
	.short	102                             # 0x66
	.short	66                              # 0x42
	.short	242                             # 0xf2
	.short	53                              # 0x35
	.short	32                              # 0x20
	.short	111                             # 0x6f
	.short	119                             # 0x77
	.short	187                             # 0xbb
	.short	89                              # 0x59
	.short	25                              # 0x19
	.short	29                              # 0x1d
	.short	254                             # 0xfe
	.short	55                              # 0x37
	.short	103                             # 0x67
	.short	45                              # 0x2d
	.short	49                              # 0x31
	.short	245                             # 0xf5
	.short	105                             # 0x69
	.short	167                             # 0xa7
	.short	100                             # 0x64
	.short	171                             # 0xab
	.short	19                              # 0x13
	.short	84                              # 0x54
	.short	37                              # 0x25
	.short	233                             # 0xe9
	.short	9                               # 0x9
	.short	237                             # 0xed
	.short	92                              # 0x5c
	.short	5                               # 0x5
	.short	202                             # 0xca
	.short	76                              # 0x4c
	.short	36                              # 0x24
	.short	135                             # 0x87
	.short	191                             # 0xbf
	.short	24                              # 0x18
	.short	62                              # 0x3e
	.short	34                              # 0x22
	.short	240                             # 0xf0
	.short	81                              # 0x51
	.short	236                             # 0xec
	.short	97                              # 0x61
	.short	23                              # 0x17
	.short	22                              # 0x16
	.short	94                              # 0x5e
	.short	175                             # 0xaf
	.short	211                             # 0xd3
	.short	73                              # 0x49
	.short	166                             # 0xa6
	.short	54                              # 0x36
	.short	67                              # 0x43
	.short	244                             # 0xf4
	.short	71                              # 0x47
	.short	145                             # 0x91
	.short	223                             # 0xdf
	.short	51                              # 0x33
	.short	147                             # 0x93
	.short	33                              # 0x21
	.short	59                              # 0x3b
	.short	121                             # 0x79
	.short	183                             # 0xb7
	.short	151                             # 0x97
	.short	133                             # 0x85
	.short	16                              # 0x10
	.short	181                             # 0xb5
	.short	186                             # 0xba
	.short	60                              # 0x3c
	.short	182                             # 0xb6
	.short	112                             # 0x70
	.short	208                             # 0xd0
	.short	6                               # 0x6
	.short	161                             # 0xa1
	.short	250                             # 0xfa
	.short	129                             # 0x81
	.short	130                             # 0x82
	.short	131                             # 0x83
	.short	126                             # 0x7e
	.short	127                             # 0x7f
	.short	128                             # 0x80
	.short	150                             # 0x96
	.short	115                             # 0x73
	.short	190                             # 0xbe
	.short	86                              # 0x56
	.short	155                             # 0x9b
	.short	158                             # 0x9e
	.short	149                             # 0x95
	.short	217                             # 0xd9
	.short	247                             # 0xf7
	.short	2                               # 0x2
	.short	185                             # 0xb9
	.short	164                             # 0xa4
	.short	222                             # 0xde
	.short	106                             # 0x6a
	.short	50                              # 0x32
	.short	109                             # 0x6d
	.short	216                             # 0xd8
	.short	138                             # 0x8a
	.short	132                             # 0x84
	.short	114                             # 0x72
	.short	42                              # 0x2a
	.short	20                              # 0x14
	.short	159                             # 0x9f
	.short	136                             # 0x88
	.short	249                             # 0xf9
	.short	220                             # 0xdc
	.short	137                             # 0x89
	.short	154                             # 0x9a
	.short	251                             # 0xfb
	.short	124                             # 0x7c
	.short	46                              # 0x2e
	.short	195                             # 0xc3
	.short	143                             # 0x8f
	.short	184                             # 0xb8
	.short	101                             # 0x65
	.short	72                              # 0x48
	.short	38                              # 0x26
	.short	200                             # 0xc8
	.short	18                              # 0x12
	.short	74                              # 0x4a
	.short	206                             # 0xce
	.short	231                             # 0xe7
	.short	210                             # 0xd2
	.short	98                              # 0x62
	.short	12                              # 0xc
	.short	224                             # 0xe0
	.short	31                              # 0x1f
	.short	239                             # 0xef
	.short	17                              # 0x11
	.short	117                             # 0x75
	.short	120                             # 0x78
	.short	113                             # 0x71
	.short	165                             # 0xa5
	.short	142                             # 0x8e
	.short	118                             # 0x76
	.short	61                              # 0x3d
	.short	189                             # 0xbd
	.short	188                             # 0xbc
	.short	134                             # 0x86
	.short	87                              # 0x57
	.short	11                              # 0xb
	.short	40                              # 0x28
	.short	47                              # 0x2f
	.short	163                             # 0xa3
	.short	218                             # 0xda
	.short	212                             # 0xd4
	.short	228                             # 0xe4
	.short	15                              # 0xf
	.short	169                             # 0xa9
	.short	39                              # 0x27
	.short	83                              # 0x53
	.short	4                               # 0x4
	.short	27                              # 0x1b
	.short	252                             # 0xfc
	.short	172                             # 0xac
	.short	230                             # 0xe6
	.short	122                             # 0x7a
	.short	7                               # 0x7
	.short	174                             # 0xae
	.short	99                              # 0x63
	.short	197                             # 0xc5
	.short	219                             # 0xdb
	.short	226                             # 0xe2
	.short	234                             # 0xea
	.short	148                             # 0x94
	.short	139                             # 0x8b
	.short	196                             # 0xc4
	.short	213                             # 0xd5
	.short	157                             # 0x9d
	.short	248                             # 0xf8
	.short	144                             # 0x90
	.short	107                             # 0x6b
	.short	177                             # 0xb1
	.short	13                              # 0xd
	.short	214                             # 0xd6
	.short	235                             # 0xeb
	.short	198                             # 0xc6
	.short	14                              # 0xe
	.short	207                             # 0xcf
	.short	173                             # 0xad
	.short	8                               # 0x8
	.short	78                              # 0x4e
	.short	215                             # 0xd7
	.short	227                             # 0xe3
	.short	93                              # 0x5d
	.short	80                              # 0x50
	.short	30                              # 0x1e
	.short	179                             # 0xb3
	.short	91                              # 0x5b
	.short	35                              # 0x23
	.short	56                              # 0x38
	.short	52                              # 0x34
	.short	104                             # 0x68
	.short	70                              # 0x46
	.short	3                               # 0x3
	.short	140                             # 0x8c
	.short	221                             # 0xdd
	.short	156                             # 0x9c
	.short	125                             # 0x7d
	.short	160                             # 0xa0
	.short	205                             # 0xcd
	.short	26                              # 0x1a
	.short	65                              # 0x41
	.short	28                              # 0x1c
	.size	inversionMatrix, 512

	.type	modPXMatrix,@object             # @modPXMatrix
	.globl	modPXMatrix
	.p2align	4, 0x0
modPXMatrix:
	.short	0                               # 0x0
	.short	27                              # 0x1b
	.short	54                              # 0x36
	.short	45                              # 0x2d
	.short	108                             # 0x6c
	.short	119                             # 0x77
	.short	90                              # 0x5a
	.short	65                              # 0x41
	.short	216                             # 0xd8
	.short	195                             # 0xc3
	.short	238                             # 0xee
	.short	245                             # 0xf5
	.short	180                             # 0xb4
	.short	175                             # 0xaf
	.short	130                             # 0x82
	.short	153                             # 0x99
	.short	171                             # 0xab
	.short	176                             # 0xb0
	.short	157                             # 0x9d
	.short	134                             # 0x86
	.short	199                             # 0xc7
	.short	220                             # 0xdc
	.short	241                             # 0xf1
	.short	234                             # 0xea
	.short	115                             # 0x73
	.short	104                             # 0x68
	.short	69                              # 0x45
	.short	94                              # 0x5e
	.short	31                              # 0x1f
	.short	4                               # 0x4
	.short	41                              # 0x29
	.short	50                              # 0x32
	.short	77                              # 0x4d
	.short	86                              # 0x56
	.short	123                             # 0x7b
	.short	96                              # 0x60
	.short	33                              # 0x21
	.short	58                              # 0x3a
	.short	23                              # 0x17
	.short	12                              # 0xc
	.short	149                             # 0x95
	.short	142                             # 0x8e
	.short	163                             # 0xa3
	.short	184                             # 0xb8
	.short	249                             # 0xf9
	.short	226                             # 0xe2
	.short	207                             # 0xcf
	.short	212                             # 0xd4
	.short	230                             # 0xe6
	.short	253                             # 0xfd
	.short	208                             # 0xd0
	.short	203                             # 0xcb
	.short	138                             # 0x8a
	.short	145                             # 0x91
	.short	188                             # 0xbc
	.short	167                             # 0xa7
	.short	62                              # 0x3e
	.short	37                              # 0x25
	.short	8                               # 0x8
	.short	19                              # 0x13
	.short	82                              # 0x52
	.short	73                              # 0x49
	.short	100                             # 0x64
	.short	127                             # 0x7f
	.short	154                             # 0x9a
	.short	129                             # 0x81
	.short	172                             # 0xac
	.short	183                             # 0xb7
	.short	246                             # 0xf6
	.short	237                             # 0xed
	.short	192                             # 0xc0
	.short	219                             # 0xdb
	.short	66                              # 0x42
	.short	89                              # 0x59
	.short	116                             # 0x74
	.short	111                             # 0x6f
	.short	46                              # 0x2e
	.short	53                              # 0x35
	.short	24                              # 0x18
	.short	3                               # 0x3
	.short	49                              # 0x31
	.short	42                              # 0x2a
	.short	7                               # 0x7
	.short	28                              # 0x1c
	.short	93                              # 0x5d
	.short	70                              # 0x46
	.short	107                             # 0x6b
	.short	112                             # 0x70
	.short	233                             # 0xe9
	.short	242                             # 0xf2
	.short	223                             # 0xdf
	.short	196                             # 0xc4
	.short	133                             # 0x85
	.short	158                             # 0x9e
	.short	179                             # 0xb3
	.short	168                             # 0xa8
	.short	215                             # 0xd7
	.short	204                             # 0xcc
	.short	225                             # 0xe1
	.short	250                             # 0xfa
	.short	187                             # 0xbb
	.short	160                             # 0xa0
	.short	141                             # 0x8d
	.short	150                             # 0x96
	.short	15                              # 0xf
	.short	20                              # 0x14
	.short	57                              # 0x39
	.short	34                              # 0x22
	.short	99                              # 0x63
	.short	120                             # 0x78
	.short	85                              # 0x55
	.short	78                              # 0x4e
	.short	124                             # 0x7c
	.short	103                             # 0x67
	.short	74                              # 0x4a
	.short	81                              # 0x51
	.short	16                              # 0x10
	.short	11                              # 0xb
	.short	38                              # 0x26
	.short	61                              # 0x3d
	.short	164                             # 0xa4
	.short	191                             # 0xbf
	.short	146                             # 0x92
	.short	137                             # 0x89
	.short	200                             # 0xc8
	.short	211                             # 0xd3
	.short	254                             # 0xfe
	.short	229                             # 0xe5
	.short	47                              # 0x2f
	.short	52                              # 0x34
	.short	25                              # 0x19
	.short	2                               # 0x2
	.short	67                              # 0x43
	.short	88                              # 0x58
	.short	117                             # 0x75
	.short	110                             # 0x6e
	.short	247                             # 0xf7
	.short	236                             # 0xec
	.short	193                             # 0xc1
	.short	218                             # 0xda
	.short	155                             # 0x9b
	.short	128                             # 0x80
	.short	173                             # 0xad
	.short	182                             # 0xb6
	.short	132                             # 0x84
	.short	159                             # 0x9f
	.short	178                             # 0xb2
	.short	169                             # 0xa9
	.short	232                             # 0xe8
	.short	243                             # 0xf3
	.short	222                             # 0xde
	.short	197                             # 0xc5
	.short	92                              # 0x5c
	.short	71                              # 0x47
	.short	106                             # 0x6a
	.short	113                             # 0x71
	.short	48                              # 0x30
	.short	43                              # 0x2b
	.short	6                               # 0x6
	.short	29                              # 0x1d
	.short	98                              # 0x62
	.short	121                             # 0x79
	.short	84                              # 0x54
	.short	79                              # 0x4f
	.short	14                              # 0xe
	.short	21                              # 0x15
	.short	56                              # 0x38
	.short	35                              # 0x23
	.short	186                             # 0xba
	.short	161                             # 0xa1
	.short	140                             # 0x8c
	.short	151                             # 0x97
	.short	214                             # 0xd6
	.short	205                             # 0xcd
	.short	224                             # 0xe0
	.short	251                             # 0xfb
	.short	201                             # 0xc9
	.short	210                             # 0xd2
	.short	255                             # 0xff
	.short	228                             # 0xe4
	.short	165                             # 0xa5
	.short	190                             # 0xbe
	.short	147                             # 0x93
	.short	136                             # 0x88
	.short	17                              # 0x11
	.short	10                              # 0xa
	.short	39                              # 0x27
	.short	60                              # 0x3c
	.short	125                             # 0x7d
	.short	102                             # 0x66
	.short	75                              # 0x4b
	.short	80                              # 0x50
	.short	181                             # 0xb5
	.short	174                             # 0xae
	.short	131                             # 0x83
	.short	152                             # 0x98
	.short	217                             # 0xd9
	.short	194                             # 0xc2
	.short	239                             # 0xef
	.short	244                             # 0xf4
	.short	109                             # 0x6d
	.short	118                             # 0x76
	.short	91                              # 0x5b
	.short	64                              # 0x40
	.short	1                               # 0x1
	.short	26                              # 0x1a
	.short	55                              # 0x37
	.short	44                              # 0x2c
	.short	30                              # 0x1e
	.short	5                               # 0x5
	.short	40                              # 0x28
	.short	51                              # 0x33
	.short	114                             # 0x72
	.short	105                             # 0x69
	.short	68                              # 0x44
	.short	95                              # 0x5f
	.short	198                             # 0xc6
	.short	221                             # 0xdd
	.short	240                             # 0xf0
	.short	235                             # 0xeb
	.short	170                             # 0xaa
	.short	177                             # 0xb1
	.short	156                             # 0x9c
	.short	135                             # 0x87
	.short	248                             # 0xf8
	.short	227                             # 0xe3
	.short	206                             # 0xce
	.short	213                             # 0xd5
	.short	148                             # 0x94
	.short	143                             # 0x8f
	.short	162                             # 0xa2
	.short	185                             # 0xb9
	.short	32                              # 0x20
	.short	59                              # 0x3b
	.short	22                              # 0x16
	.short	13                              # 0xd
	.short	76                              # 0x4c
	.short	87                              # 0x57
	.short	122                             # 0x7a
	.short	97                              # 0x61
	.short	83                              # 0x53
	.short	72                              # 0x48
	.short	101                             # 0x65
	.short	126                             # 0x7e
	.short	63                              # 0x3f
	.short	36                              # 0x24
	.short	9                               # 0x9
	.short	18                              # 0x12
	.short	139                             # 0x8b
	.short	144                             # 0x90
	.short	189                             # 0xbd
	.short	166                             # 0xa6
	.short	231                             # 0xe7
	.short	252                             # 0xfc
	.short	209                             # 0xd1
	.short	202                             # 0xca
	.size	modPXMatrix, 512

	.type	sBox,@object                    # @sBox
	.globl	sBox
	.p2align	4, 0x0
sBox:
	.ascii	"c|w{\362ko\3050\001g+\376\327\253v\312\202\311}\372YG\360\255\324\242\257\234\244r\300\267\375\223&6?\367\3144\245\345\361q\3301\025\004\307#\303\030\226\005\232\007\022\200\342\353'\262u\t\203,\032\033nZ\240R;\326\263)\343/\204S\321\000\355 \374\261[j\313\2769JLX\317\320\357\252\373CM3\205E\371\002\177P<\237\250Q\243@\217\222\2358\365\274\266\332!\020\377\363\322\315\f\023\354_\227D\027\304\247~=d]\031s`\201O\334\"*\220\210F\356\270\024\336^\013\333\3402:\nI\006$\\\302\323\254b\221\225\344y\347\3107m\215\325N\251lV\364\352ez\256\b\272x%.\034\246\264\306\350\335t\037K\275\213\212p>\265fH\003\366\016a5W\271\206\301\035\236\341\370\230\021i\331\216\224\233\036\207\351\316U(\337\214\241\211\r\277\346BhA\231-\017\260T\273\026"
	.size	sBox, 256

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

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

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

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

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

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

	.type	.L.str.6,@object                # @.str.6
.L.str.6:
	.asciz	"Multiply for 100,000 times."
	.size	.L.str.6, 28

	.type	.L.str.7,@object                # @.str.7
.L.str.7:
	.asciz	"uSeconds: "
	.size	.L.str.7, 11

	.type	.L.str.8,@object                # @.str.8
.L.str.8:
	.asciz	"Inversion for 100,000 times."
	.size	.L.str.8, 29

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

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

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

	.type	.L.str.9,@object                # @.str.9
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str.9:
	.asciz	"sBox linear analysis with CUDA"
	.size	.L.str.9, 31

	.type	.L.str.10,@object               # @.str.10
.L.str.10:
	.asciz	"Multiply for 100,000 times with GPU."
	.size	.L.str.10, 37

	.type	.L.str.11,@object               # @.str.11
.L.str.11:
	.asciz	"Error"
	.size	.L.str.11, 6

	.type	.L.str.12,@object               # @.str.12
.L.str.12:
	.asciz	"Time(uSeconds): "
	.size	.L.str.12, 17

	.type	.L.str.13,@object               # @.str.13
.L.str.13:
	.asciz	"Inversion for 100,000 times with GPU."
	.size	.L.str.13, 38

	.type	.L.str.14,@object               # @.str.14
.L.str.14:
	.asciz	"Error "
	.size	.L.str.14, 7

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

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"_Z10d_multiplyPtS_"
	.size	.L__unnamed_2, 19

	.type	.L__unnamed_3,@object           # @2
.L__unnamed_3:
	.asciz	"_Z11d_inversionPt"
	.size	.L__unnamed_3, 18

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

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

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym _Z20__device_stub__countPtS_
	.addrsig_sym _Z25__device_stub__d_multiplyPtS_
	.addrsig_sym _Z26__device_stub__d_inversionPt
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _ZSt4cout
	.addrsig_sym _Z5countPtS_
	.addrsig_sym _Z10d_multiplyPtS_
	.addrsig_sym _Z11d_inversionPt
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z5countPtS_            ; -- Begin function _Z5countPtS_
	.globl	_Z5countPtS_
	.p2align	8
	.type	_Z5countPtS_,@function
_Z5countPtS_:                           ; @_Z5countPtS_
; %bb.0:
	s_clause 0x1
	s_load_b32 s4, s[0:1], 0x1c
	s_load_b128 s[0:3], s[0:1], 0x0
	v_dual_mov_b32 v2, 0 :: v_dual_mov_b32 v1, 0
	s_mov_b32 s5, 0
	s_mov_b32 s7, 0
	s_waitcnt lgkmcnt(0)
	s_and_b32 s4, s4, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_mul_i32 s15, s15, s4
	v_add_nc_u32_e32 v0, s15, v0
	s_and_b32 s6, s15, 0xffff
.LBB0_1:                                ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB0_2 Depth 2
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_add_i32 s4, s7, s6
	s_lshl_b64 s[8:9], s[4:5], 1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_add_u32 s8, s0, s8
	s_addc_u32 s9, s1, s9
	s_lshl_b32 s4, s7, 8
	global_load_u16 v3, v2, s[8:9]
	s_waitcnt vmcnt(0)
	v_xor_b32_e32 v3, s4, v3
	s_mov_b32 s4, 1
	s_delay_alu instid0(VALU_DEP_1)
	v_and_b32_e32 v3, v3, v0
.LBB0_2:                                ;   Parent Loop BB0_1 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_and_b32_e32 v4, 0xffff, v3
	v_lshrrev_b32_e32 v4, s4, v4
	s_add_i32 s4, s4, 1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	s_cmp_eq_u32 s4, 16
	v_and_b32_e32 v4, 1, v4
	s_delay_alu instid0(VALU_DEP_1)
	v_xor_b32_e32 v3, v3, v4
	s_cbranch_scc0 .LBB0_2
; %bb.3:                                ;   in Loop: Header=BB0_1 Depth=1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	v_xor_b32_e32 v3, -1, v3
	s_add_i32 s7, s7, 1
	s_cmpk_eq_i32 s7, 0x100
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_and_b32_e32 v3, 1, v3
	v_add_nc_u16 v1, v3, v1
	s_cbranch_scc0 .LBB0_1
; %bb.4:
	v_and_b32_e32 v0, 0xffff, v0
	s_delay_alu instid0(VALU_DEP_1)
	v_lshlrev_b32_e32 v0, 1, v0
	global_store_b16 v0, v1, s[2:3]
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z5countPtS_
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 272
		.amdhsa_user_sgpr_count 15
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 5
		.amdhsa_next_free_sgpr 16
		.amdhsa_reserve_vcc 0
		.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	_Z5countPtS_, .Lfunc_end0-_Z5countPtS_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 240
; NumSgprs: 16
; NumVgprs: 5
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 1
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 16
; NumVGPRsForWavesPerEU: 5
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.text
	.protected	_Z10d_multiplyPtS_      ; -- Begin function _Z10d_multiplyPtS_
	.globl	_Z10d_multiplyPtS_
	.p2align	8
	.type	_Z10d_multiplyPtS_,@function
_Z10d_multiplyPtS_:                     ; @_Z10d_multiplyPtS_
; %bb.0:
	v_and_b32_e32 v1, 0x3ff, v0
	v_bfe_u32 v0, v0, 10, 10
	s_mulk_i32 s15, 0x2800
	s_lshl_b32 s2, s14, 5
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mad_u32_u24 v0, 0x140, v0, v1
	v_add3_u32 v0, s2, s15, v0
	s_mov_b32 s2, exec_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_gt_i32_e32 0x186a0, v0
	s_cbranch_execz .LBB1_6
; %bb.1:                                ; %.preheader
	s_load_b128 s[0:3], s[0:1], 0x0
	v_ashrrev_i32_e32 v1, 31, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[0:1], 1, v[0:1]
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v2, vcc_lo, s2, v0
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v3, vcc_lo, s3, v1, vcc_lo
	v_add_co_u32 v0, vcc_lo, s0, v0
	v_add_co_ci_u32_e32 v1, vcc_lo, s1, v1, vcc_lo
	global_load_u16 v2, v[2:3], off
	v_mov_b32_e32 v3, 0
	s_mov_b32 s0, 7
.LBB1_2:                                ; =>This Inner Loop Header: Depth=1
	s_delay_alu instid0(VALU_DEP_1)
	v_lshlrev_b16 v4, 1, v3
	v_and_b32_e32 v3, 0x80, v3
	s_waitcnt vmcnt(0)
	v_lshrrev_b32_e32 v5, s0, v2
	s_mov_b32 s1, exec_lo
	v_xor_b32_e32 v6, 0x11b, v4
	v_cmp_eq_u16_e32 vcc_lo, 0, v3
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_and_b32_e32 v5, 1, v5
	v_cndmask_b32_e32 v3, v6, v4, vcc_lo
	s_delay_alu instid0(VALU_DEP_2)
	v_cmpx_eq_u32_e32 1, v5
	s_cbranch_execz .LBB1_4
; %bb.3:                                ;   in Loop: Header=BB1_2 Depth=1
	global_load_u16 v4, v[0:1], off
	s_waitcnt vmcnt(0)
	v_xor_b32_e32 v3, v4, v3
.LBB1_4:                                ;   in Loop: Header=BB1_2 Depth=1
	s_or_b32 exec_lo, exec_lo, s1
	s_add_i32 s0, s0, -1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmp_eq_u32 s0, -1
	s_cbranch_scc0 .LBB1_2
; %bb.5:
	global_store_b16 v[0:1], v3, off
.LBB1_6:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z10d_multiplyPtS_
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 16
		.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 7
		.amdhsa_next_free_sgpr 16
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end1:
	.size	_Z10d_multiplyPtS_, .Lfunc_end1-_Z10d_multiplyPtS_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 264
; NumSgprs: 18
; NumVgprs: 7
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 7
; 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	_Z11d_inversionPt       ; -- Begin function _Z11d_inversionPt
	.globl	_Z11d_inversionPt
	.p2align	8
	.type	_Z11d_inversionPt,@function
_Z11d_inversionPt:                      ; @_Z11d_inversionPt
; %bb.0:
	s_load_b64 s[0:1], s[0:1], 0x0
	v_bfe_u32 v1, v0, 10, 10
	v_and_b32_e32 v0, 0x3ff, v0
	s_mov_b32 s2, exec_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_mad_u32_u24 v0, 0x140, v1, v0
	v_cmpx_gt_u32_e32 8, v1
	s_cbranch_execz .LBB2_2
; %bb.1:
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mad_i32_i24 v1, 0xfffffee0, v1, v0
	v_mov_b32_e32 v2, 0
	v_lshlrev_b64 v[2:3], 1, v[1:2]
	v_lshlrev_b32_e32 v1, 1, v1
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_u32 v2, vcc_lo, s0, v2
	v_add_co_ci_u32_e32 v3, vcc_lo, s1, v3, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v2, vcc_lo, 0x30000, v2
	v_add_co_ci_u32_e32 v3, vcc_lo, 0, v3, vcc_lo
	global_load_u16 v2, v[2:3], off offset:3392
	s_waitcnt vmcnt(0)
	ds_store_b16 v1, v2
.LBB2_2:
	s_or_b32 exec_lo, exec_lo, s2
	s_mulk_i32 s15, 0x2800
	s_lshl_b32 s2, s14, 5
	s_waitcnt lgkmcnt(0)
	v_add3_u32 v0, s2, s15, v0
	s_barrier
	buffer_gl0_inv
	s_mov_b32 s2, exec_lo
	v_cmpx_gt_i32_e32 0x186a0, v0
	s_cbranch_execz .LBB2_4
; %bb.3:
	v_ashrrev_i32_e32 v1, 31, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[0:1], 1, v[0:1]
	v_add_co_u32 v0, vcc_lo, s0, v0
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v1, vcc_lo, s1, v1, vcc_lo
	global_load_u16 v2, v[0:1], off
	s_waitcnt vmcnt(0)
	v_lshlrev_b32_e32 v2, 1, v2
	ds_load_u16 v2, v2
	s_waitcnt lgkmcnt(0)
	global_store_b16 v[0:1], v2, off
.LBB2_4:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z11d_inversionPt
		.amdhsa_group_segment_fixed_size 512
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 8
		.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 4
		.amdhsa_next_free_sgpr 16
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end2:
	.size	_Z11d_inversionPt, .Lfunc_end2-_Z11d_inversionPt
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 276
; NumSgprs: 18
; NumVgprs: 4
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 512 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 4
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 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:     hidden_block_count_x
      - .offset:         20
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         24
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         28
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         30
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         32
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         34
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         36
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         38
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         56
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         64
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         80
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 272
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z5countPtS_
    .private_segment_fixed_size: 0
    .sgpr_count:     16
    .sgpr_spill_count: 0
    .symbol:         _Z5countPtS_.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     5
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 16
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z10d_multiplyPtS_
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z10d_multiplyPtS_.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     7
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
    .group_segment_fixed_size: 512
    .kernarg_segment_align: 8
    .kernarg_segment_size: 8
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z11d_inversionPt
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z11d_inversionPt.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     4
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
