#include <stdio.h>
#include <stdlib.h>
#include <time.h>
#include <unistd.h>
#include <sys/time.h>
#include <string.h>

#define R_size 129
#define k 1024
#define n_size 128

__global__ void get_square(unsigned char a[], unsigned int accumulator[], unsigned int n);
void square(unsigned char *a, unsigned char *c, unsigned int size);
__global__ void get_products(unsigned char a[], unsigned char b[], unsigned int accumulator[], unsigned int n);
void barrett_reduction(unsigned char *buf_cpu, unsigned char *buf_gpu, unsigned char *gpu_r, unsigned char *gpu_n, unsigned char *reduction, unsigned char *temp_cpu, unsigned char *temp_gpu, unsigned char *shifted_cpu, unsigned char *shifted_gpu, unsigned char *xprime_cpu, unsigned char *xprime_gpu, unsigned char *result, unsigned char *tmp, unsigned int exp_size, unsigned int *kernel_buf, unsigned int *transfer, unsigned char *n);
void multiplication(unsigned char *a, unsigned char *b, unsigned char *c, unsigned int size); 
char checkbit(unsigned char *exponent, unsigned int index_of_bit); 
void subtraction(unsigned char *a, unsigned char *b, unsigned char *c, unsigned int size); 
void bit_shift(unsigned char *a, unsigned char *b, unsigned int shift, unsigned int size_of_a); 
void exponentiation(unsigned char *gpu_message, unsigned char *exponent, unsigned char *gpu_ciphertext, unsigned char *m0_copy, unsigned char *reduction, unsigned char *buf_cpu, unsigned char *buf_gpu, unsigned char *temp_cpu, unsigned char *temp_gpu, unsigned char *shifted_cpu, unsigned char *shifted_gpu, unsigned char *xprime_cpu, unsigned char *xprime_gpu, unsigned char *result, unsigned char *tmp, unsigned int exponent_size, unsigned char *gpu_r, unsigned char *gpu_n, unsigned int *kernel_buf, unsigned int *transfer, unsigned char *n);

int main(int argc, char *argv[]) {

unsigned char *n = (unsigned char *) calloc((4*n_size + n_size), sizeof(char));

    n[0] = 0xcf;
    n[1] = 0x82;
    n[2] = 0x69;
    n[3] = 0x57;
    n[4] = 0x4d;
    n[5] = 0xe7;
    n[6] = 0x82;
    n[7] = 0x1a;
    n[8] = 0xe4;
    n[9] = 0x20;
    n[10] = 0x14;
    n[11] = 0x47;
    n[12] = 0x39;
    n[13] = 0x52;
    n[14] = 0x55;
    n[15] = 0x28;
    n[16] = 0xed;
    n[17] = 0x3f;
    n[18] = 0xa4;
    n[19] = 0x61;
    n[20] = 0xd3;
    n[21] = 0xf4;
    n[22] = 0xf2;
    n[23] = 0x34;
    n[24] = 0x6a;
    n[25] = 0x54;
    n[26] = 0xd1;
    n[27] = 0x15;
    n[28] = 0x7d;
    n[29] = 0x67;
    n[30] = 0xb;
    n[31] = 0xc7;
    n[32] = 0x8c;
    n[33] = 0xfe;
    n[34] = 0x1b;
    n[35] = 0x68;
    n[36] = 0x44;
    n[37] = 0x7;
    n[38] = 0x26;
    n[39] = 0x99;
    n[40] = 0xb;
    n[41] = 0x4d;
    n[42] = 0xc7;
    n[43] = 0x3f;
    n[44] = 0x52;
    n[45] = 0x90;
    n[46] = 0x2;
    n[47] = 0x68;
    n[48] = 0x3d;
    n[49] = 0x83;
    n[50] = 0x1d;
    n[51] = 0x79;
    n[52] = 0x7a;
    n[53] = 0x3f;
    n[54] = 0x36;
    n[55] = 0xf3;
    n[56] = 0x41;
    n[57] = 0x8b;
    n[58] = 0x7c;
    n[59] = 0xdf;
    n[60] = 0x64;
    n[61] = 0xac;
    n[62] = 0x74;
    n[63] = 0x7c;
    n[64] = 0x8;
    n[65] = 0xdb;
    n[66] = 0xa0;
    n[67] = 0x6f;
    n[68] = 0x10;
    n[69] = 0x71;
    n[70] = 0x13;
    n[71] = 0x86;
    n[72] = 0xaf;
    n[73] = 0xb8;
    n[74] = 0x71;
    n[75] = 0xf8;
    n[76] = 0xf0;
    n[77] = 0x45;
    n[78] = 0xa7;
    n[79] = 0x94;
    n[80] = 0xb3;
    n[81] = 0x6b;
    n[82] = 0x1e;
    n[83] = 0xff;
    n[84] = 0x8e;
    n[85] = 0x13;
    n[86] = 0xae;
    n[87] = 0xc2;
    n[88] = 0x59;
    n[89] = 0x56;
    n[90] = 0xd3;
    n[91] = 0xd;
    n[92] = 0x20;
    n[93] = 0x62;
    n[94] = 0x21;
    n[95] = 0x30;
    n[96] = 0x1d;
    n[97] = 0x6b;
    n[98] = 0x5e;
    n[99] = 0xc;
    n[100] = 0x0;
    n[101] = 0x35;
    n[102] = 0xae;
    n[103] = 0xbd;
    n[104] = 0xa5;
    n[105] = 0xc2;
    n[106] = 0x25;
    n[107] = 0x98;
    n[108] = 0xe7;
    n[109] = 0x57;
    n[110] = 0x89;
    n[111] = 0xc;
    n[112] = 0x12;
    n[113] = 0xf9;
    n[114] = 0x33;
    n[115] = 0x3d;
    n[116] = 0xa;
    n[117] = 0xac;
    n[118] = 0x51;
    n[119] = 0xd8;
    n[120] = 0x5c;
    n[121] = 0x40;
    n[122] = 0x9b;
    n[123] = 0xfa;
    n[124] = 0xf9;
    n[125] = 0xbc;
    n[126] = 0x3;
    n[127] = 0xe6;

    
	unsigned char *gpu_n;
	cudaMalloc(&gpu_n, (4*n_size + n_size));
	cudaMemcpy(gpu_n, n, (4*n_size + n_size), cudaMemcpyHostToDevice);

unsigned char *r = (unsigned char *) calloc(2*n_size, sizeof(char));

    r[0] = 0x7f;
    r[1] = 0x9d;
    r[2] = 0xe9;
    r[3] = 0x40;
    r[4] = 0x57;
    r[5] = 0x2;
    r[6] = 0x6e;
    r[7] = 0x93;
    r[8] = 0x2b;
    r[9] = 0xb4;
    r[10] = 0xe3;
    r[11] = 0xfd;
    r[12] = 0xba;
    r[13] = 0xc;
    r[14] = 0xcd;
    r[15] = 0x78;
    r[16] = 0x7d;
    r[17] = 0xae;
    r[18] = 0x8d;
    r[19] = 0x80;
    r[20] = 0xff;
    r[21] = 0x66;
    r[22] = 0x33;
    r[23] = 0xb;
    r[24] = 0x28;
    r[25] = 0x4c;
    r[26] = 0x93;
    r[27] = 0x30;
    r[28] = 0x2;
    r[29] = 0x92;
    r[30] = 0xa0;
    r[31] = 0x7c;
    r[32] = 0xf1;
    r[33] = 0xc;
    r[34] = 0xa;
    r[35] = 0x5e;
    r[36] = 0xf2;
    r[37] = 0x9a;
    r[38] = 0x8f;
    r[39] = 0x17;
    r[40] = 0x4c;
    r[41] = 0x82;
    r[42] = 0x25;
    r[43] = 0xe5;
    r[44] = 0x98;
    r[45] = 0x45;
    r[46] = 0x4d;
    r[47] = 0xc7;
    r[48] = 0xd9;
    r[49] = 0x53;
    r[50] = 0x5e;
    r[51] = 0x5a;
    r[52] = 0x6e;
    r[53] = 0x37;
    r[54] = 0x43;
    r[55] = 0x29;
    r[56] = 0x88;
    r[57] = 0xcb;
    r[58] = 0xe9;
    r[59] = 0x31;
    r[60] = 0x2f;
    r[61] = 0xd7;
    r[62] = 0x6;
    r[63] = 0xfb;
    r[64] = 0xf1;
    r[65] = 0x38;
    r[66] = 0xdf;
    r[67] = 0xc4;
    r[68] = 0xda;
    r[69] = 0x7c;
    r[70] = 0x9;
    r[71] = 0x5c;
    r[72] = 0xf9;
    r[73] = 0x2b;
    r[74] = 0x81;
    r[75] = 0x30;
    r[76] = 0xe9;
    r[77] = 0x29;
    r[78] = 0xcd;
    r[79] = 0x45;
    r[80] = 0xee;
    r[81] = 0xff;
    r[82] = 0x5b;
    r[83] = 0x3c;
    r[84] = 0x23;
    r[85] = 0x6d;
    r[86] = 0xb9;
    r[87] = 0xa1;
    r[88] = 0x89;
    r[89] = 0x3f;
    r[90] = 0xc3;
    r[91] = 0x9e;
    r[92] = 0xa1;
    r[93] = 0x30;
    r[94] = 0x98;
    r[95] = 0xf8;
    r[96] = 0xc8;
    r[97] = 0x4a;
    r[98] = 0xbe;
    r[99] = 0xc6;
    r[100] = 0x49;
    r[101] = 0xf7;
    r[102] = 0xb3;
    r[103] = 0xff;
    r[104] = 0x9;
    r[105] = 0x3b;
    r[106] = 0x94;
    r[107] = 0x9d;
    r[108] = 0x2f;
    r[109] = 0x5c;
    r[110] = 0x68;
    r[111] = 0xe1;
    r[112] = 0x6;
    r[113] = 0xf1;
    r[114] = 0x33;
    r[115] = 0xeb;
    r[116] = 0xc5;
    r[117] = 0x88;
    r[118] = 0xa5;
    r[119] = 0x1c;
    r[120] = 0xde;
    r[121] = 0x2c;
    r[122] = 0x64;
    r[123] = 0xad;
    r[124] = 0x5c;
    r[125] = 0xc9;
    r[126] = 0xeb;
    r[127] = 0x1c;
    r[128] = 0x1;

	unsigned char *gpu_r;
	cudaMalloc(&gpu_r, (2*n_size));
	cudaMemcpy(gpu_r, r, 2*n_size, cudaMemcpyHostToDevice);

	unsigned char *message = (unsigned char *) calloc(n_size, sizeof(char));
	message[0] = 0x68;//h
	message[1] = 0x65;//e
	message[2] = 0x6c;//l
	message[3] = 0x6c;//l
	message[4] = 0x6f;//o


	unsigned char *gpu_message;
	cudaMalloc(&gpu_message, n_size);
	cudaMemcpy(gpu_message, message, n_size, cudaMemcpyHostToDevice);

	unsigned char *exponent = (unsigned char *) malloc(3);
	exponent[0] = 0x01;
	exponent[1] = 0x00;
	exponent[2] = 0x01;
	unsigned int exponent_size = 3;

	//exponentiate m^e mod n
	//parameters: 
	//message(m)
	//exponent(e)
	//precomputation of r = floor((4^k)/n) where k is found by where (2^k) > n
	//modulus (n)

	unsigned char *cpu_ciphertext = (unsigned char *) calloc(n_size, sizeof(char));
	unsigned char *gpu_ciphertext;
	cudaMalloc(&gpu_ciphertext, n_size);
	cudaMemset(gpu_ciphertext, 0x00, n_size);

	unsigned char *m0_copy;
	cudaMalloc(&m0_copy, n_size);

	unsigned char *reduction = (unsigned char *) calloc(n_size, sizeof(char));

	unsigned char *buf_cpu = (unsigned char *) calloc((n_size * 2) + 1, sizeof(char));
	unsigned char *buf_gpu;
	cudaMalloc(&buf_gpu, ((n_size * 2) + 1));
	cudaMemset(buf_gpu, 0x00, (n_size * 2) + 1);

	unsigned char *temp_cpu = (unsigned char *) calloc(3*n_size, sizeof(char));
	unsigned char *temp_gpu;
	cudaMalloc(&temp_gpu, (3*n_size));
	
	unsigned char *shifted_cpu = (unsigned char *) calloc(n_size, sizeof(char));
	unsigned char *shifted_gpu;
	cudaMalloc(&shifted_gpu, n_size);
	cudaMemset(shifted_gpu, 0x00, n_size);

	unsigned char *xprime_cpu = (unsigned char *) calloc(2*n_size, sizeof(char));
	unsigned char *xprime_gpu;
	cudaMalloc(&xprime_gpu, (2*n_size));
	cudaMemset(xprime_gpu, 0x00, 2*n_size);

	unsigned char *result = (unsigned char *) calloc(n_size + 1, sizeof(char));
	
	unsigned char *tmp = (unsigned char *) calloc(n_size + 1, sizeof(char));

	unsigned int *transfer = (unsigned int *) calloc(4*n_size, sizeof(int));
	unsigned int *kernel_buf;
	cudaMalloc(&kernel_buf, 4*n_size*sizeof(int));
	cudaMemset(kernel_buf, 0x00, 4*n_size*sizeof(int));

	cudaError_t error;
	cudaEvent_t start;

	error = cudaEventCreate(&start);
	if(error != cudaSuccess)
	printf("error\n");

	cudaEvent_t stop;
	error = cudaEventCreate(&stop);
	if(error != cudaSuccess)
	printf("error\n");

	error = cudaEventRecord(start, NULL);


	exponentiation(gpu_message, exponent, gpu_ciphertext, m0_copy, reduction, buf_cpu, buf_gpu, temp_cpu, temp_gpu, shifted_cpu, shifted_gpu, xprime_cpu, xprime_gpu, result, tmp, exponent_size, gpu_r, gpu_n, kernel_buf, transfer, n);

	unsigned int d_exponent_size = 128;
	unsigned char *d_exponent = (unsigned char *) malloc(128);
    
    d_exponent[0] = 0x91;
    d_exponent[1] = 0xa;
    d_exponent[2] = 0xb3;
    d_exponent[3] = 0x66;
    d_exponent[4] = 0xbd;
    d_exponent[5] = 0x6f;
    d_exponent[6] = 0x18;
    d_exponent[7] = 0xde;
    d_exponent[8] = 0xd5;
    d_exponent[9] = 0x1;
    d_exponent[10] = 0x61;
    d_exponent[11] = 0x36;
    d_exponent[12] = 0x95;
    d_exponent[13] = 0x6d;
    d_exponent[14] = 0xdd;
    d_exponent[15] = 0x33;
    d_exponent[16] = 0xdb;
    d_exponent[17] = 0x26;
    d_exponent[18] = 0x3;
    d_exponent[19] = 0xe;
    d_exponent[20] = 0x68;
    d_exponent[21] = 0x54;
    d_exponent[22] = 0x73;
    d_exponent[23] = 0xa0;
    d_exponent[24] = 0xe0;
    d_exponent[25] = 0x6e;
    d_exponent[26] = 0x70;
    d_exponent[27] = 0x74;
    d_exponent[28] = 0x25;
    d_exponent[29] = 0x8b;
    d_exponent[30] = 0x2b;
    d_exponent[31] = 0xfb;
    d_exponent[32] = 0x9e;
    d_exponent[33] = 0x3c;
    d_exponent[34] = 0x34;
    d_exponent[35] = 0x2e;
    d_exponent[36] = 0x45;
    d_exponent[37] = 0x10;
    d_exponent[38] = 0x10;
    d_exponent[39] = 0x6c;
    d_exponent[40] = 0xfb;
    d_exponent[41] = 0xb7;
    d_exponent[42] = 0x9b;
    d_exponent[43] = 0xc8;
    d_exponent[44] = 0xcf;
    d_exponent[45] = 0x71;
    d_exponent[46] = 0xd9;
    d_exponent[47] = 0x96;
    d_exponent[48] = 0xb7;
    d_exponent[49] = 0xbb;
    d_exponent[50] = 0x5f;
    d_exponent[51] = 0x19;
    d_exponent[52] = 0x76;
    d_exponent[53] = 0x36;
    d_exponent[54] = 0x49;
    d_exponent[55] = 0x6a;
    d_exponent[56] = 0xb3;
    d_exponent[57] = 0x83;
    d_exponent[58] = 0xc3;
    d_exponent[59] = 0x59;
    d_exponent[60] = 0x2e;
    d_exponent[61] = 0x62;
    d_exponent[62] = 0x87;
    d_exponent[63] = 0xa2;
    d_exponent[64] = 0x5a;
    d_exponent[65] = 0x2f;
    d_exponent[66] = 0x60;
    d_exponent[67] = 0x75;
    d_exponent[68] = 0x1;
    d_exponent[69] = 0xf0;
    d_exponent[70] = 0x3f;
    d_exponent[71] = 0xdb;
    d_exponent[72] = 0x5a;
    d_exponent[73] = 0x70;
    d_exponent[74] = 0x1f;
    d_exponent[75] = 0x44;
    d_exponent[76] = 0x6a;
    d_exponent[77] = 0x9c;
    d_exponent[78] = 0x77;
    d_exponent[79] = 0x63;
    d_exponent[80] = 0xba;
    d_exponent[81] = 0xcb;
    d_exponent[82] = 0xcd;
    d_exponent[83] = 0x1f;
    d_exponent[84] = 0x99;
    d_exponent[85] = 0x70;
    d_exponent[86] = 0x89;
    d_exponent[87] = 0x94;
    d_exponent[88] = 0x31;
    d_exponent[89] = 0x2;
    d_exponent[90] = 0xa;
    d_exponent[91] = 0x32;
    d_exponent[92] = 0x96;
    d_exponent[93] = 0x65;
    d_exponent[94] = 0x21;
    d_exponent[95] = 0x21;
    d_exponent[96] = 0x59;
    d_exponent[97] = 0x55;
    d_exponent[98] = 0x8a;
    d_exponent[99] = 0xd0;
    d_exponent[100] = 0x7a;
    d_exponent[101] = 0x1c;
    d_exponent[102] = 0xd2;
    d_exponent[103] = 0x66;
    d_exponent[104] = 0x48;
    d_exponent[105] = 0x95;
    d_exponent[106] = 0x8;
    d_exponent[107] = 0xd3;
    d_exponent[108] = 0x6b;
    d_exponent[109] = 0xe7;
    d_exponent[110] = 0x9c;
    d_exponent[111] = 0xb9;
    d_exponent[112] = 0x96;
    d_exponent[113] = 0x20;
    d_exponent[114] = 0x20;
    d_exponent[115] = 0x8a;
    d_exponent[116] = 0xe5;
    d_exponent[117] = 0x4d;
    d_exponent[118] = 0x3e;
    d_exponent[119] = 0x53;
    d_exponent[120] = 0x4b;
    d_exponent[121] = 0xd8;
    d_exponent[122] = 0x21;
    d_exponent[123] = 0x4;
    d_exponent[124] = 0x81;
    d_exponent[125] = 0x7d;
    d_exponent[126] = 0x29;
    d_exponent[127] = 0x38;


	memset(message, 0x00, n_size);
	cudaMemset(gpu_message, 0x00, n_size);

	exponentiation(gpu_ciphertext, d_exponent, gpu_message, m0_copy, reduction, buf_cpu, buf_gpu, temp_cpu, temp_gpu, shifted_cpu, shifted_gpu, xprime_cpu, xprime_gpu, result, tmp, d_exponent_size, gpu_r, gpu_n, kernel_buf, transfer, n);


	error = cudaEventRecord(stop, NULL);

	error = cudaEventSynchronize(stop);

	if(error != cudaSuccess)
		printf("error\n");

	float msecTotal = 0.0f;
	error = cudaEventElapsedTime(&msecTotal, start, stop);
	printf("GPU time: %.6f\n", msecTotal / 1000);


	cudaMemcpy(message, gpu_message, n_size, cudaMemcpyDeviceToHost);
	int z = 0;
	while (z < n_size) {
		printf("message[%d] = %x\n", z, message[z]);
		z++;
	}

	return 0;
}

void exponentiation(unsigned char *gpu_message, unsigned char *exponent, unsigned char *gpu_ciphertext, unsigned char *m0_copy, unsigned char *reduction, unsigned char *buf_cpu, unsigned char *buf_gpu, unsigned char *temp_cpu, unsigned char *temp_gpu, unsigned char *shifted_cpu, unsigned char *shifted_gpu, unsigned char *xprime_cpu, unsigned char *xprime_gpu, unsigned char *result, unsigned char *tmp, unsigned int exponent_size, unsigned char *gpu_r, unsigned char *gpu_n, unsigned int *kernel_buf, unsigned int *transfer, unsigned char *n) {


	dim3 blocksPerGrid(2);
	dim3 threadsPerBlock(64);


	//get the total amount of bits in strlen(exponent) zero based
	//not including the final char index msb (byte)
	unsigned int total_bits = exponent_size * 8 - 1;
	
	//find the most signinficant bit in the most significant byte (char index)
	//find most significant bit in exponent[exp_size - 1]
	unsigned char mask = 0x80; //10000000 in binary
	unsigned char msb = 0;
	int i = 0;
	while(i < 8) {
		if((exponent[exponent_size - 1] & (mask >> i)) == (mask >> i)) {
			msb = i;
			break;
		}
		i++;
	}

	//subtract most significant bit from total_bits to know total amount of significant bits
	//for loop of exponent in binary
	unsigned int exp_bits = (total_bits - msb);

	//keep copy of original message m0
	cudaMemcpy(m0_copy, gpu_message, n_size, cudaMemcpyDeviceToDevice);

	//compute m^e where e is in binary 
	//RULES:
		//iterate over the values of msb to 0 bit by bit
		//msb is amount of relevent bits to check for exponentiation
		//total bits is the amount of total bits in exponent lenth
		//square m(current) for each itteration
		//check if current bit is 1
		//current bit is 1: m(current) * m0
		//curent bit is 0: return to loop


	//subtract one from total because to exponentiate in binary
	//start at the second bit after the most significant bit
	//each bit equals m^2 and when the current bit is 1 it is
	//(m^2)*m0 or if it is 0 then m^2
	int index_of_bit = exp_bits - 1; //subtraction of 1 is becuase msb is zero based
	while (index_of_bit >= 0) {

		//allocate space for reduction to hold a value strickly less than n
		//buf holds value at most m^2 which is less than n^2

		//calculate m^2
		get_square<<<blocksPerGrid, threadsPerBlock>>>(gpu_message, kernel_buf, n_size);
		cudaMemcpy(transfer, kernel_buf, 2*n_size*sizeof(unsigned int), cudaMemcpyDeviceToHost);

		unsigned int index = 0;
		while(index < 2*n_size) {
			buf_cpu[index] = (unsigned char) transfer[index];
			transfer[index + 1] += (unsigned int) (transfer[index]>>8);
			index++;
		}
	
		memset(transfer, 0x00, 4*n_size*sizeof(int));	
		cudaMemset(kernel_buf, 0x00, 4*n_size*sizeof(int));
		cudaMemcpy(buf_gpu, buf_cpu, 2*n_size, cudaMemcpyHostToDevice);

		//calculate m^2 mod n
		barrett_reduction(buf_cpu, buf_gpu, gpu_r, gpu_n, reduction, temp_cpu, temp_gpu, shifted_cpu, shifted_gpu, xprime_cpu, xprime_gpu, result, tmp, exponent_size, kernel_buf, transfer, n);

		cudaMemcpy(gpu_message, reduction, n_size, cudaMemcpyHostToDevice);

		cudaMemset(buf_gpu, 0x00, 2*n_size);
		memset(buf_cpu, 0x00, 2*n_size);
		memset(reduction, 0x00, n_size);

		char bit;
		if ((bit = checkbit(exponent, index_of_bit)) == 1) {
			//m * m0
			get_products<<<blocksPerGrid, threadsPerBlock>>>(gpu_message, m0_copy, kernel_buf, n_size);
			cudaMemcpy(transfer, kernel_buf, 2*n_size*sizeof(unsigned int), cudaMemcpyDeviceToHost);

			index = 0;
			while(index < 2*n_size) {
				buf_cpu[index] = (unsigned char) transfer[index];
				transfer[index + 1] += (unsigned int) (transfer[index]>>8);
				index++;
			}
		
			memset(transfer, 0x00, 4*n_size*sizeof(int));
			cudaMemset(kernel_buf, 0x00, 4*n_size*sizeof(int));
			cudaMemcpy(buf_gpu, buf_cpu, 2*n_size, cudaMemcpyHostToDevice);



			barrett_reduction(buf_cpu, buf_gpu, gpu_r, gpu_n, reduction, temp_cpu, temp_gpu, shifted_cpu, shifted_gpu, xprime_cpu, xprime_gpu, result, tmp, exponent_size, kernel_buf, transfer, n);

			cudaMemcpy(gpu_message, reduction, n_size, cudaMemcpyHostToDevice);
			cudaMemset(buf_gpu, 0x00, 2*n_size);
			memset(buf_cpu, 0x00, 2*n_size);
			memset(reduction, 0x00, n_size);
		
		}

		index_of_bit--;
	}
	
	//copy back final value of message to ciphertext for decryption
	cudaMemcpy(gpu_ciphertext, gpu_message, n_size, cudaMemcpyDeviceToHost);

	cudaMemset(m0_copy, 0x00, n_size);

	return;
}
void barrett_reduction(unsigned char *buf_cpu, unsigned char *buf_gpu, unsigned char *gpu_r, unsigned char *gpu_n, unsigned char *reduction, unsigned char *temp_cpu, unsigned char *temp_gpu, unsigned char *shifted_cpu, unsigned char *shifted_gpu, unsigned char *xprime_cpu, unsigned char *xprime_gpu, unsigned char *result, unsigned char *tmp, unsigned int exp_size, unsigned int *kernel_buf, unsigned int *transfer, unsigned char *n) {
	////////////////////////////////////////////////////////////////////////////////////

	////////////////////  calculate: t = x - ((x*r)/(4^k))*n  /////////////////////////

	///////////////////////////////////////////////////////////////////////////////////
	///////////////////////////////////////////////////////////////////////////////////
	//multiply: x * r = temp
	//size of x is assumed to be the largest value which is = largest value of 2*n
	//size of r is precomputed
	dim3 blocksPerGrid_two(4);
	dim3 threadsPerBlock_two(64);

	get_products<<<blocksPerGrid_two, threadsPerBlock_two>>>(gpu_r, buf_gpu, kernel_buf, 2*n_size);
	cudaMemcpy(transfer, kernel_buf, 3*n_size*sizeof(unsigned int), cudaMemcpyDeviceToHost);

	unsigned int index = 0;
	while(index < 3*n_size) {
		temp_cpu[index] = (unsigned char) transfer[index];
		transfer[index + 1] += (unsigned int) (transfer[index]>>8);
		index++;
	}

	cudaMemset(kernel_buf, 0x00, 4*n_size*sizeof(int));

	//shift bits by (4^k) or (2^(2*k))
	//shift temp by 2*k store to shifted
	//size of shifted is 2*n + sizeof(r)
	
	//find the actual amount of bits/bytes left in the value of temp
	//which is equal to x * r so that the correct size of the value
	//can be used in the bit_shift function

	//first find the amount of bytes from most significant byte
	//to least and then when one char does not equal to 0x00

	unsigned int zero_bytes = 0;
	int count = (3*n_size) - 1;
	while((count >= 0) && (temp_cpu[count] == 0x00)) {
		count--;
		zero_bytes++;
	}


	bit_shift(temp_cpu, shifted_cpu, k, (3*n_size) - zero_bytes);

	//multiply: shifted * n = xprime
	//xprime is the size of 2*n + R_size - (k >> 0x07) + n
	cudaMemcpy(shifted_gpu, shifted_cpu, n_size, cudaMemcpyHostToDevice);

	dim3 blocksPerGrid_one(2);
	dim3 threadsPerBlock_one(64);
	get_products<<<blocksPerGrid_one, threadsPerBlock_one>>>(shifted_gpu, gpu_n, kernel_buf, n_size);
	cudaMemcpy(transfer, kernel_buf, 2*n_size*sizeof(unsigned int), cudaMemcpyDeviceToHost);

	index = 0;
	while(index < 2*n_size) {
		xprime_cpu[index] = (unsigned char) transfer[index];
		transfer[index + 1] += (unsigned int) (transfer[index]>>8);
		index++;
	}



	//subtract xprime from x^2
	cudaMemset(kernel_buf, 0x00, 4*n_size*sizeof(int));
	subtraction(buf_cpu, xprime_cpu, result, 2*n_size);

	//the field of n, if the value is not within the field of n then reduce the value by subtracting
	//the value of result = t - n which is guaranteed to be in the field of n
	if ((result[n_size] == 0x00) && (result[n_size - 1] < n[n_size - 1])) {
		memcpy(reduction, result, n_size);
	}

	else {
		unsigned char *tmp = (unsigned char *) calloc(n_size + 1, sizeof(char));
		subtraction(result, n, tmp, n_size + 1);
		memcpy(reduction, tmp, n_size);
		memset(tmp, 0x00, n_size + 1);
	}


	memset(temp_cpu, 0x00, 3*n_size);
	cudaMemset(temp_gpu, 0x00, 3*n_size);
	memset(shifted_cpu, 0x00, n_size);
	cudaMemset(shifted_gpu, 0x00, n_size);
	memset(xprime_cpu, 0x00, 2*n_size);
	cudaMemset(xprime_gpu, 0x00, 2*n_size);
	memset(result, 0x00, n_size + 1);
	memset(transfer, 0x00, 4*n_size*sizeof(int));


	return;
}

void subtraction(unsigned char *a, unsigned char *b, unsigned char *c, unsigned int size) {

	//borrow represents the value 1 or 0 for the current index
	//indecating if the current index has been borrowed from by
	//the previous index, borrow = 1 true, 0 false
	unsigned char borrow = 0x00; //value is 0 or 1

	//loop through array a size and subtract a - b,
	//a is guaranted to be greater than b in 
	//barrett reduction
	unsigned int i = 0;
	while(i < size) {
		//check current value of a to make sure that it is
		//not 0 when the previous index has borrowed
		if (a[i] == 0 && borrow == 1) {
			//borrow from next sequential index with
			//0x100 and subtract 0x01 for the
			//previous borrow which is = 0xff
			c[i] = 0xff - b[i];
			//turn on borrow for next index
			borrow = 0x01;
			i++;
			continue;
		}

		//calculate current value of a along with if the 
		//previous index has borrowed
		a[i] = a[i] - borrow;
	
		//calculate the value of a - b only when a - b >= 0
		//borrow has already been accounted for
		if (a[i] >= b[i]) {
			c[i] = a[i] - b[i];
			borrow = 0x00;
		}

		//a - b !> 0, borrow from next sequential index by 
		//taking the value 0x100 and adding to a[i] and 
		//subtracting b[i] which will give a value between
		//{0x01...0xff} and turn on borrow for next index
		else {
			c[i] = 0x100 + a[i] - b[i];
			borrow = 0x01;
		}
		i++;
	}
	return;
}

char checkbit(unsigned char *exponent, unsigned int index_of_bit) {

	unsigned char bit;
	//get the characters index of which the bit is located in by 
	//taking index_of_bit which is the size of the bits left to 
	//check and divide by 8 giving the location index of the
	//current bit to be checked
	unsigned int quotient = (index_of_bit >> 0x03); // index_of_bit / 8

	//find the bit within the index previously found by finding the 
	//remainder of 8 % index of bit, this will locate the exact
	//bit to be checked
	unsigned int remainder = index_of_bit & (0x07); // index_of_bit % 8
	
	//mask is equivelent to 1 in order to compare a single bit with a
	//the current bit to be checked
	unsigned char mask = 0x01; // use single bit to mask with selected bit
	
	//use the remainder by knowing the index of the character and 
	//the remainder allows the bit to be shifted to the position of
	//the current bit to be checked
	mask = mask << remainder; // shift single bit to bit_in_index position
	
	//bit is now located at index_of_bit character index of array
	//and bit location bit_in_index in group of 8 bits at index
	bit = (exponent[quotient] & mask); // & to see if single bit is on or off

	//shift bit back to the 1 position to represent value 1 or 0
	bit = bit >> remainder; //shift bit back to value of one or zero
	
	return bit;
}

//b is expected to be completely zero before shift
void bit_shift(unsigned char *a, unsigned char *b, unsigned int k_val, unsigned int size_of_a) {

        //expected that k will be equivlent to some power of 2
        //represents the division of (4^k) which is = (2^(2*k))
        unsigned int shift = k_val * 2;

        //quotient represents groups of 8 bits that equal 0 as in >> 8 in single char
        //leaving it to be the value of 0x00
        unsigned int quotient = shift >> 0x03; // k / 8 as integer

	//in case that the shift is greater than the actual value of the
	//number being shifted
	if(quotient > size_of_a) {
		return;
	}

        //printf("quotient = %d\n", quotient);
        //remainder will find final char index shift value = {0...7}
        //the specific bits to be shifted in the last group which is not greater than 7
        unsigned int remainder = shift & 0x07; // k % 8 
        //printf("remainder = %d\n", remainder);
        //move a to b by shifting the characters an index of quotient amount
        //and then use the remainder to shift the final index to correct 
        //position
        unsigned int constant = (size_of_a - quotient);

        unsigned int j = 0;
        while (j < constant) {
                b[j] = a[quotient + j] >> remainder;
                unsigned char cpy_bits = a[quotient + j + 1] << (8 - remainder);
                b[j] = b[j] | cpy_bits;
                j++;
        }

        return;
}

__global__ void get_products(unsigned char a[], unsigned char b[], unsigned int accumulator[], unsigned int n) {

        int multiplier = 0;
        unsigned int index = (blockIdx.x * blockDim.x) + threadIdx.x;
        unsigned int multiplicand = index;
        unsigned int product = 0;

        while(multiplier < n) {
                product = (unsigned int) a[multiplier] * b[multiplicand];
                atomicAdd(&accumulator[multiplier + index], product<<24>>24);
                atomicAdd(&accumulator[multiplier + index + 1], product>>8);
                multiplier++;
        }

	return;
}


__global__ void get_square(unsigned char a[], unsigned int accumulator[], unsigned int n) {

        int multiplier = 0;
        unsigned int index = (blockIdx.x * blockDim.x) + threadIdx.x;
        unsigned int multiplicand = index;
        unsigned int product = 0;

        while(multiplier < n) {
                product = (unsigned int) a[multiplier] * a[multiplicand];
                atomicAdd(&accumulator[multiplier + index], product<<24>>24);
                atomicAdd(&accumulator[multiplier + index + 1], product>>8);
                multiplier++;
        }
	return;
}

	.file	"tmpxft_0033e2df_00000000-6_GPU_RSA.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2048:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2048:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.type	_Z10cudaMallocIhE9cudaErrorPPT_m.isra.0, @function
_Z10cudaMallocIhE9cudaErrorPPT_m.isra.0:
.LFB2193:
	.cfi_startproc
	jmp	cudaMalloc@PLT
	.cfi_endproc
.LFE2193:
	.size	_Z10cudaMallocIhE9cudaErrorPPT_m.isra.0, .-_Z10cudaMallocIhE9cudaErrorPPT_m.isra.0
	.globl	_Z11subtractionPhS_S_j
	.type	_Z11subtractionPhS_S_j, @function
_Z11subtractionPhS_S_j:
.LFB2043:
	.cfi_startproc
	endbr64
	movq	%rdi, %r9
	xorl	%r8d, %r8d
	xorl	%edi, %edi
.L4:
	cmpl	%ecx, %edi
	jnb	.L18
	movb	(%r9,%rdi), %al
	testb	%al, %al
	jne	.L5
	testb	%r8b, %r8b
	je	.L5
	movb	(%rsi,%rdi), %al
	notl	%eax
	jmp	.L16
.L5:
	subl	%r8d, %eax
	movb	%al, (%r9,%rdi)
	movb	(%rsi,%rdi), %r8b
	cmpb	%r8b, %al
	jb	.L7
	subl	%r8d, %eax
	xorl	%r8d, %r8d
	jmp	.L6
.L7:
	subl	%r8d, %eax
.L16:
	movb	$1, %r8b
.L6:
	movb	%al, (%rdx,%rdi)
	incq	%rdi
	jmp	.L4
.L18:
	ret
	.cfi_endproc
.LFE2043:
	.size	_Z11subtractionPhS_S_j, .-_Z11subtractionPhS_S_j
	.globl	_Z8checkbitPhj
	.type	_Z8checkbitPhj, @function
_Z8checkbitPhj:
.LFB2044:
	.cfi_startproc
	endbr64
	movl	%esi, %ecx
	movl	$1, %eax
	shrl	$3, %esi
	andl	$7, %ecx
	sall	%cl, %eax
	andb	(%rdi,%rsi), %al
	movzbl	%al, %eax
	sarl	%cl, %eax
	ret
	.cfi_endproc
.LFE2044:
	.size	_Z8checkbitPhj, .-_Z8checkbitPhj
	.globl	_Z9bit_shiftPhS_jj
	.type	_Z9bit_shiftPhS_jj, @function
_Z9bit_shiftPhS_jj:
.LFB2045:
	.cfi_startproc
	endbr64
	addl	%edx, %edx
	movl	%edx, %r8d
	shrl	$3, %r8d
	cmpl	%r8d, %ecx
	jb	.L26
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	movq	%rdi, %r9
	andl	$7, %edx
	movl	%ecx, %edi
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	movl	$8, %r11d
	movq	%rsi, %r10
	subl	%r8d, %edi
	xorl	%eax, %eax
	leal	1(%r8), %ebp
	subl	%edx, %r11d
.L22:
	cmpl	%edi, %eax
	jnb	.L29
	leal	(%r8,%rax), %ecx
	movzbl	(%r9,%rcx), %esi
	movl	%edx, %ecx
	sarl	%cl, %esi
	leal	0(%rbp,%rax), %ecx
	movb	%sil, (%r10,%rax)
	movzbl	(%r9,%rcx), %ebx
	movl	%r11d, %ecx
	sall	%cl, %ebx
	orl	%ebx, %esi
	movb	%sil, (%r10,%rax)
	incq	%rax
	jmp	.L22
.L29:
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	ret
.L26:
	.cfi_restore 3
	.cfi_restore 6
	ret
	.cfi_endproc
.LFE2045:
	.size	_Z9bit_shiftPhS_jj, .-_Z9bit_shiftPhS_jj
	.globl	_Z37__device_stub__Z12get_productsPhS_PjjPhS_Pjj
	.type	_Z37__device_stub__Z12get_productsPhS_PjjPhS_Pjj, @function
_Z37__device_stub__Z12get_productsPhS_PjjPhS_Pjj:
.LFB2070:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movq	%rdi, 24(%rsp)
	leaq	56(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	68(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	40(%rsp), %rdx
	movl	%ecx, 4(%rsp)
	leaq	48(%rsp), %rcx
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 104(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 128(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 56(%rsp)
	movq	%rax, 68(%rsp)
	movl	$1, 76(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L30
	pushq	48(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z12get_productsPhS_Pjj(%rip), %rdi
	pushq	48(%rsp)
	.cfi_def_cfa_offset 176
	movq	84(%rsp), %rcx
	movl	92(%rsp), %r8d
	movq	72(%rsp), %rsi
	movl	80(%rsp), %edx
	leaq	120(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 168
	popq	%rdx
	.cfi_def_cfa_offset 160
.L30:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L32
	call	__stack_chk_fail@PLT
.L32:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2070:
	.size	_Z37__device_stub__Z12get_productsPhS_PjjPhS_Pjj, .-_Z37__device_stub__Z12get_productsPhS_PjjPhS_Pjj
	.globl	_Z12get_productsPhS_Pjj
	.type	_Z12get_productsPhS_Pjj, @function
_Z12get_productsPhS_Pjj:
.LFB2071:
	.cfi_startproc
	endbr64
	jmp	_Z37__device_stub__Z12get_productsPhS_PjjPhS_Pjj
	.cfi_endproc
.LFE2071:
	.size	_Z12get_productsPhS_Pjj, .-_Z12get_productsPhS_Pjj
	.globl	_Z17barrett_reductionPhS_S_S_S_S_S_S_S_S_S_S_S_jPjS0_S_
	.type	_Z17barrett_reductionPhS_S_S_S_S_S_S_S_S_S_S_S_jPjS0_S_, @function
_Z17barrett_reductionPhS_S_S_S_S_S_S_S_S_S_S_S_jPjS0_S_:
.LFB2042:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movq	%rsi, %r15
	movl	$1, %esi
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	movq	%rdx, %r14
	movl	$67108865, %edx
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	salq	$6, %rdx
	movq	%rcx, %r13
	movl	$1, %ecx
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movq	%rdi, %r12
	movl	$1073741825, %edi
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	salq	$2, %rdi
	movq	%r9, %rbp
	xorl	%r9d, %r9d
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	subq	$72, %rsp
	.cfi_def_cfa_offset 128
	movq	%r8, 8(%rsp)
	xorl	%r8d, %r8d
	movq	200(%rsp), %rbx
	movq	%rdi, 16(%rsp)
	movl	$1, 24(%rsp)
	movq	%rdx, 28(%rsp)
	movl	$1, 36(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L36
	movq	192(%rsp), %rdx
	movl	$256, %ecx
	movq	%r15, %rsi
	movq	%r14, %rdi
	call	_Z37__device_stub__Z12get_productsPhS_PjjPhS_Pjj
.L36:
	movq	192(%rsp), %rsi
	movl	$2, %ecx
	movl	$1536, %edx
	movq	%rbx, %rdi
	call	cudaMemcpy@PLT
	xorl	%eax, %eax
.L37:
	movl	(%rbx,%rax,4), %edx
	movb	%dl, 0(%rbp,%rax)
	movd	(%rbx,%rax,4), %xmm0
	movd	4(%rbx,%rax,4), %xmm1
	psrld	$8, %xmm0
	paddd	%xmm1, %xmm0
	movd	%xmm0, 4(%rbx,%rax,4)
	incq	%rax
	cmpq	$384, %rax
	jne	.L37
	movq	192(%rsp), %rdi
	movl	$2048, %edx
	xorl	%esi, %esi
	call	cudaMemset@PLT
	leaq	383(%rbp), %rdx
	xorl	%eax, %eax
.L38:
	cmpb	$0, (%rdx)
	jne	.L39
	incl	%eax
	decq	%rdx
	cmpl	$384, %eax
	jne	.L38
.L39:
	movq	136(%rsp), %rsi
	movl	$384, %ecx
	movl	$1024, %edx
	movq	%rbp, %rdi
	subl	%eax, %ecx
	call	_Z9bit_shiftPhS_jj
	movl	$1, %ecx
	movl	$128, %edx
	movq	136(%rsp), %rsi
	movq	144(%rsp), %rdi
	call	cudaMemcpy@PLT
	movl	$2147483649, %edi
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$67108865, %edx
	addq	%rdi, %rdi
	movl	$1, %ecx
	movl	$1, %esi
	salq	$6, %rdx
	movq	%rdi, 40(%rsp)
	movl	$1, 48(%rsp)
	movq	%rdx, 52(%rsp)
	movl	$1, 60(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L41
	movq	192(%rsp), %rdx
	movq	144(%rsp), %rdi
	movl	$128, %ecx
	movq	%r13, %rsi
	call	_Z37__device_stub__Z12get_productsPhS_PjjPhS_Pjj
.L41:
	movq	192(%rsp), %rsi
	movl	$2, %ecx
	movl	$1024, %edx
	movq	%rbx, %rdi
	call	cudaMemcpy@PLT
	xorl	%eax, %eax
.L42:
	movl	(%rbx,%rax,4), %edx
	movq	152(%rsp), %rsi
	movb	%dl, (%rsi,%rax)
	movd	(%rbx,%rax,4), %xmm0
	movd	4(%rbx,%rax,4), %xmm1
	psrld	$8, %xmm0
	paddd	%xmm1, %xmm0
	movd	%xmm0, 4(%rbx,%rax,4)
	incq	%rax
	cmpq	$256, %rax
	jne	.L42
	movq	192(%rsp), %rdi
	xorl	%esi, %esi
	movl	$2048, %edx
	call	cudaMemset@PLT
	movq	168(%rsp), %rdx
	movl	$256, %ecx
	movq	%r12, %rdi
	movq	152(%rsp), %rsi
	call	_Z11subtractionPhS_S_j
	movq	168(%rsp), %rax
	cmpb	$0, 128(%rax)
	jne	.L43
	movq	208(%rsp), %rsi
	movb	127(%rsi), %sil
	cmpb	%sil, 127(%rax)
	jnb	.L43
	movq	8(%rsp), %rdi
	movl	$32, %ecx
	movq	%rax, %rsi
	jmp	.L51
.L43:
	movl	$1, %esi
	movl	$129, %edi
	call	calloc@PLT
	movq	208(%rsp), %rsi
	movl	$129, %ecx
	movq	168(%rsp), %rdi
	movq	%rax, %rdx
	call	_Z11subtractionPhS_S_j
	movq	8(%rsp), %rdi
	movl	$32, %ecx
	movq	%rdx, %rsi
.L51:
	rep movsl
	xorl	%r12d, %r12d
	movq	%rbp, %rdi
	movl	$96, %ecx
	movl	%r12d, %eax
	movl	$384, %edx
	xorl	%esi, %esi
	rep stosl
	movq	128(%rsp), %rdi
	call	cudaMemset@PLT
	movl	%r12d, %eax
	movl	$32, %ecx
	xorl	%esi, %esi
	movq	136(%rsp), %rdi
	movl	$128, %edx
	rep stosl
	movq	144(%rsp), %rdi
	call	cudaMemset@PLT
	movl	%r12d, %eax
	movl	$64, %ecx
	xorl	%esi, %esi
	movq	152(%rsp), %rdi
	movl	$256, %edx
	rep stosl
	movq	160(%rsp), %rdi
	call	cudaMemset@PLT
	movq	168(%rsp), %rdi
	xorl	%eax, %eax
	movl	$129, %ecx
	rep stosb
	movq	%rbx, %rdi
	movl	%r12d, %eax
	movl	$512, %ecx
	rep stosl
	addq	$72, %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
.LFE2042:
	.size	_Z17barrett_reductionPhS_S_S_S_S_S_S_S_S_S_S_S_jPjS0_S_, .-_Z17barrett_reductionPhS_S_S_S_S_S_S_S_S_S_S_S_jPjS0_S_
	.globl	_Z33__device_stub__Z10get_squarePhPjjPhPjj
	.type	_Z33__device_stub__Z10get_squarePhPjjPhPjj, @function
_Z33__device_stub__Z10get_squarePhPjjPhPjj:
.LFB2072:
	.cfi_startproc
	endbr64
	subq	$136, %rsp
	.cfi_def_cfa_offset 144
	movq	%rdi, 24(%rsp)
	leaq	40(%rsp), %rcx
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movl	%edx, 12(%rsp)
	leaq	32(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 120(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 96(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 112(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L52
	pushq	40(%rsp)
	.cfi_def_cfa_offset 152
	leaq	_Z10get_squarePhPjj(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 160
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	112(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 152
	popq	%rdx
	.cfi_def_cfa_offset 144
.L52:
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L54
	call	__stack_chk_fail@PLT
.L54:
	addq	$136, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2072:
	.size	_Z33__device_stub__Z10get_squarePhPjjPhPjj, .-_Z33__device_stub__Z10get_squarePhPjjPhPjj
	.globl	_Z10get_squarePhPjj
	.type	_Z10get_squarePhPjj, @function
_Z10get_squarePhPjj:
.LFB2073:
	.cfi_startproc
	endbr64
	jmp	_Z33__device_stub__Z10get_squarePhPjjPhPjj
	.cfi_endproc
.LFE2073:
	.size	_Z10get_squarePhPjj, .-_Z10get_squarePhPjj
	.globl	_Z14exponentiationPhS_S_S_S_S_S_S_S_S_S_S_S_S_S_jS_S_PjS0_S_
	.type	_Z14exponentiationPhS_S_S_S_S_S_S_S_S_S_S_S_S_S_jS_S_PjS0_S_, @function
_Z14exponentiationPhS_S_S_S_S_S_S_S_S_S_S_S_S_S_jS_S_PjS0_S_:
.LFB2041:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	subq	$184, %rsp
	.cfi_def_cfa_offset 240
	movq	240(%rsp), %rax
	movq	%rdx, 136(%rsp)
	movq	%rcx, 40(%rsp)
	xorl	%ecx, %ecx
	movq	%rax, 32(%rsp)
	movq	248(%rsp), %rax
	movq	%rdi, 16(%rsp)
	movq	%rax, 56(%rsp)
	movq	256(%rsp), %rax
	movq	%rsi, 48(%rsp)
	movq	%rax, 64(%rsp)
	movq	264(%rsp), %rax
	movq	%r8, 24(%rsp)
	movq	%rax, 72(%rsp)
	movq	272(%rsp), %rax
	movq	%r9, 8(%rsp)
	movq	%rax, 80(%rsp)
	movq	280(%rsp), %rax
	movq	%rax, 88(%rsp)
	movq	288(%rsp), %rax
	movl	312(%rsp), %r13d
	movq	336(%rsp), %r12
	movl	$1, 160(%rsp)
	movq	%rax, 96(%rsp)
	movq	296(%rsp), %rax
	movq	344(%rsp), %rbx
	movq	352(%rsp), %r15
	movl	$1, 172(%rsp)
	leal	0(,%r13,8), %ebp
	movq	%rax, 104(%rsp)
	movq	304(%rsp), %rax
	movq	%rax, 112(%rsp)
	movq	320(%rsp), %rax
	movq	%rax, 120(%rsp)
	movq	328(%rsp), %rax
	movq	%rax, 128(%rsp)
	movl	$2147483649, %eax
	addq	%rax, %rax
	movq	%rax, 152(%rsp)
	addq	$62, %rax
	movq	%rax, 164(%rsp)
	leal	-1(%r13), %eax
	movzbl	(%rsi,%rax), %edx
	movl	$128, %eax
	testb	%dl, %dl
	js	.L58
.L59:
	incl	%ecx
	cmpl	$8, %ecx
	je	.L74
	movl	%eax, %esi
	movl	%edx, %edi
	sarl	%cl, %esi
	notl	%edi
	testl	%esi, %edi
	jne	.L59
.L58:
	movl	%ecx, %r14d
	jmp	.L60
.L74:
	xorl	%r14d, %r14d
.L60:
	movq	16(%rsp), %rsi
	movq	40(%rsp), %rdi
	movl	$3, %ecx
	subl	$2, %ebp
	movl	$128, %edx
	call	cudaMemcpy@PLT
	movzbl	%r14b, %ecx
	subl	%ecx, %ebp
.L62:
	testl	%ebp, %ebp
	js	.L75
	movl	172(%rsp), %ecx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movq	164(%rsp), %rdx
	movq	152(%rsp), %rdi
	movl	160(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L63
	movq	16(%rsp), %rdi
	movl	$128, %edx
	movq	%r12, %rsi
	call	_Z33__device_stub__Z10get_squarePhPjjPhPjj
.L63:
	movl	$2, %ecx
	movl	$1024, %edx
	movq	%r12, %rsi
	movq	%rbx, %rdi
	call	cudaMemcpy@PLT
	xorl	%eax, %eax
.L64:
	movl	(%rbx,%rax,4), %edx
	movq	8(%rsp), %rsi
	movb	%dl, (%rsi,%rax)
	movd	(%rbx,%rax,4), %xmm0
	movd	4(%rbx,%rax,4), %xmm1
	psrld	$8, %xmm0
	paddd	%xmm1, %xmm0
	movd	%xmm0, 4(%rbx,%rax,4)
	incq	%rax
	cmpq	$256, %rax
	jne	.L64
	xorl	%eax, %eax
	movl	$512, %ecx
	movq	%rbx, %rdi
	xorl	%esi, %esi
	rep stosl
	movl	$2048, %edx
	movq	%r12, %rdi
	call	cudaMemset@PLT
	movq	8(%rsp), %rsi
	movq	32(%rsp), %rdi
	movl	$256, %edx
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	pushq	%rdx
	.cfi_def_cfa_offset 248
	pushq	%r15
	.cfi_def_cfa_offset 256
	pushq	%rbx
	.cfi_def_cfa_offset 264
	pushq	%r12
	.cfi_def_cfa_offset 272
	pushq	%r13
	.cfi_def_cfa_offset 280
	pushq	152(%rsp)
	.cfi_def_cfa_offset 288
	pushq	152(%rsp)
	.cfi_def_cfa_offset 296
	pushq	152(%rsp)
	.cfi_def_cfa_offset 304
	pushq	152(%rsp)
	.cfi_def_cfa_offset 312
	pushq	152(%rsp)
	.cfi_def_cfa_offset 320
	pushq	152(%rsp)
	.cfi_def_cfa_offset 328
	pushq	152(%rsp)
	.cfi_def_cfa_offset 336
	movq	152(%rsp), %r9
	movq	120(%rsp), %r8
	movq	224(%rsp), %rcx
	movq	216(%rsp), %rdx
	movq	128(%rsp), %rsi
	movq	104(%rsp), %rdi
	call	_Z17barrett_reductionPhS_S_S_S_S_S_S_S_S_S_S_S_jPjS0_S_
	movq	120(%rsp), %rsi
	movq	112(%rsp), %rdi
	addq	$96, %rsp
	.cfi_def_cfa_offset 240
	movl	$1, %ecx
	movl	$128, %edx
	call	cudaMemcpy@PLT
	movq	32(%rsp), %rdi
	xorl	%esi, %esi
	movl	$256, %edx
	call	cudaMemset@PLT
	movq	8(%rsp), %rdi
	xorl	%eax, %eax
	movl	%ebp, %esi
	movl	$64, %ecx
	rep stosl
	movq	24(%rsp), %rdi
	movl	$32, %ecx
	rep stosl
	movq	48(%rsp), %rdi
	call	_Z8checkbitPhj
	decb	%al
	jne	.L65
	movl	172(%rsp), %ecx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movq	164(%rsp), %rdx
	movq	152(%rsp), %rdi
	movl	160(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L66
	movq	40(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$128, %ecx
	movq	%r12, %rdx
	call	_Z37__device_stub__Z12get_productsPhS_PjjPhS_Pjj
.L66:
	movl	$2, %ecx
	movl	$1024, %edx
	movq	%r12, %rsi
	movq	%rbx, %rdi
	call	cudaMemcpy@PLT
	xorl	%eax, %eax
.L67:
	movl	(%rbx,%rax,4), %edx
	movq	8(%rsp), %rsi
	movb	%dl, (%rsi,%rax)
	movd	(%rbx,%rax,4), %xmm0
	movd	4(%rbx,%rax,4), %xmm1
	psrld	$8, %xmm0
	paddd	%xmm1, %xmm0
	movd	%xmm0, 4(%rbx,%rax,4)
	incq	%rax
	cmpq	$256, %rax
	jne	.L67
	xorl	%r14d, %r14d
	movl	$512, %ecx
	movq	%rbx, %rdi
	movl	$2048, %edx
	movl	%r14d, %eax
	xorl	%esi, %esi
	rep stosl
	movq	%r12, %rdi
	call	cudaMemset@PLT
	movq	8(%rsp), %rsi
	movq	32(%rsp), %rdi
	movl	$1, %ecx
	movl	$256, %edx
	call	cudaMemcpy@PLT
	pushq	%rax
	.cfi_def_cfa_offset 248
	pushq	%r15
	.cfi_def_cfa_offset 256
	pushq	%rbx
	.cfi_def_cfa_offset 264
	pushq	%r12
	.cfi_def_cfa_offset 272
	pushq	%r13
	.cfi_def_cfa_offset 280
	pushq	152(%rsp)
	.cfi_def_cfa_offset 288
	pushq	152(%rsp)
	.cfi_def_cfa_offset 296
	pushq	152(%rsp)
	.cfi_def_cfa_offset 304
	pushq	152(%rsp)
	.cfi_def_cfa_offset 312
	pushq	152(%rsp)
	.cfi_def_cfa_offset 320
	pushq	152(%rsp)
	.cfi_def_cfa_offset 328
	pushq	152(%rsp)
	.cfi_def_cfa_offset 336
	movq	152(%rsp), %r9
	movq	120(%rsp), %r8
	movq	224(%rsp), %rcx
	movq	216(%rsp), %rdx
	movq	128(%rsp), %rsi
	movq	104(%rsp), %rdi
	call	_Z17barrett_reductionPhS_S_S_S_S_S_S_S_S_S_S_S_jPjS0_S_
	movq	120(%rsp), %rsi
	movq	112(%rsp), %rdi
	addq	$96, %rsp
	.cfi_def_cfa_offset 240
	movl	$1, %ecx
	movl	$128, %edx
	call	cudaMemcpy@PLT
	movq	32(%rsp), %rdi
	movl	$256, %edx
	xorl	%esi, %esi
	call	cudaMemset@PLT
	movq	8(%rsp), %rdi
	movl	%r14d, %eax
	movl	$64, %ecx
	rep stosl
	movq	24(%rsp), %rdi
	movl	$32, %ecx
	rep stosl
.L65:
	decl	%ebp
	jmp	.L62
.L75:
	movq	16(%rsp), %rsi
	movl	$128, %edx
	movl	$2, %ecx
	movq	136(%rsp), %rdi
	call	cudaMemcpy@PLT
	movq	40(%rsp), %rdi
	addq	$184, %rsp
	.cfi_def_cfa_offset 56
	xorl	%esi, %esi
	popq	%rbx
	.cfi_def_cfa_offset 48
	movl	$128, %edx
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	jmp	cudaMemset@PLT
	.cfi_endproc
.LFE2041:
	.size	_Z14exponentiationPhS_S_S_S_S_S_S_S_S_S_S_S_S_S_jS_S_PjS0_S_, .-_Z14exponentiationPhS_S_S_S_S_S_S_S_S_S_S_S_S_S_jS_S_PjS0_S_
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"error\n"
.LC3:
	.string	"GPU time: %.6f\n"
.LC4:
	.string	"message[%d] = %x\n"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2040:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movl	$1, %esi
	movl	$640, %edi
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	subq	$168, %rsp
	.cfi_def_cfa_offset 224
	movq	%fs:40, %rax
	movq	%rax, 152(%rsp)
	xorl	%eax, %eax
	call	calloc@PLT
	leaq	56(%rsp), %rdi
	movl	$640, %esi
	movq	%rax, %r12
	movabsq	$1910343511323542223, %rax
	movq	%rax, (%r12)
	movabsq	$2906319540496638180, %rax
	movq	%rax, 8(%r12)
	movabsq	$3815381023049662445, %rax
	movq	%rax, 16(%r12)
	movabsq	$-4104072848483462038, %rax
	movq	%rax, 24(%r12)
	movabsq	$-7411228146405736820, %rax
	movq	%rax, 32(%r12)
	movabsq	$7494711412829670667, %rax
	movq	%rax, 40(%r12)
	movabsq	$-921479278500150467, %rax
	movq	%rax, 48(%r12)
	movabsq	$8967982307247688513, %rax
	movq	%rax, 56(%r12)
	movabsq	$-8785554132663477496, %rax
	movq	%rax, 64(%r12)
	movabsq	$-7735136933722867537, %rax
	movq	%rax, 72(%r12)
	movabsq	$-4418572679517410381, %rax
	movq	%rax, 80(%r12)
	movabsq	$3468161077862422105, %rax
	movq	%rax, 88(%r12)
	movabsq	$-4778823880269731043, %rax
	movq	%rax, 96(%r12)
	movabsq	$903349852466168485, %rax
	movq	%rax, 104(%r12)
	movabsq	$-2859315128427087598, %rax
	movq	%rax, 112(%r12)
	movabsq	$-1872445238218637220, %rax
	movq	%rax, 120(%r12)
	call	_Z10cudaMallocIhE9cudaErrorPPT_m.isra.0
	movq	56(%rsp), %rdi
	movl	$1, %ecx
	movq	%r12, %rsi
	movl	$640, %edx
	call	cudaMemcpy@PLT
	movl	$1, %esi
	movl	$256, %edi
	call	calloc@PLT
	leaq	64(%rsp), %rdi
	movl	$256, %esi
	movq	%rax, %rbx
	movabsq	$-7823312928921510529, %rax
	movq	%rax, (%rbx)
	movabsq	$8704627652040045611, %rax
	movq	%rax, 8(%rbx)
	movabsq	$807102005788913277, %rax
	movq	%rax, 16(%rbx)
	movabsq	$8980338195079318568, %rax
	movq	%rax, 24(%rbx)
	movabsq	$1697745950292446449, %rax
	movq	%rax, 32(%rbx)
	movabsq	$-4085532763973385652, %rax
	movq	%rax, 40(%rbx)
	movabsq	$2973281126096720857, %rax
	movq	%rax, 48(%rbx)
	movabsq	$-358362522628535416, %rax
	movq	%rax, 56(%rbx)
	movabsq	$6631969205327444209, %rax
	movq	%rax, 64(%rbx)
	movabsq	$5029722440360602617, %rax
	movq	%rax, 72(%rbx)
	movabsq	$-6793278564807737362, %rax
	movq	%rax, 80(%rbx)
	movabsq	$-533623085131939959, %rax
	movq	%rax, 88(%rbx)
	movabsq	$-21401676967687480, %rax
	movq	%rax, 96(%rbx)
	movabsq	$-2204410658020902135, %rax
	movq	%rax, 104(%rbx)
	movabsq	$2064206387855225094, %rax
	movq	%rax, 112(%rbx)
	movabsq	$2083980652472184030, %rax
	movq	%rax, 120(%rbx)
	movb	$1, 128(%rbx)
	call	_Z10cudaMallocIhE9cudaErrorPPT_m.isra.0
	movq	64(%rsp), %rdi
	movl	$1, %ecx
	movq	%rbx, %rsi
	movl	$256, %edx
	call	cudaMemcpy@PLT
	movl	$1, %esi
	movl	$128, %edi
	call	calloc@PLT
	leaq	72(%rsp), %rdi
	movl	$128, %esi
	movl	$1819043176, (%rax)
	movq	%rax, %rbx
	movb	$111, 4(%rax)
	call	_Z10cudaMallocIhE9cudaErrorPPT_m.isra.0
	movq	72(%rsp), %rdi
	movl	$1, %ecx
	movq	%rbx, %rsi
	movl	$128, %edx
	call	cudaMemcpy@PLT
	movl	$3, %edi
	call	malloc@PLT
	leaq	80(%rsp), %rdi
	movl	$128, %esi
	movw	$1, (%rax)
	movq	%rax, %rbp
	movb	$1, 2(%rax)
	call	_Z10cudaMallocIhE9cudaErrorPPT_m.isra.0
	movq	80(%rsp), %rdi
	movl	$128, %edx
	xorl	%esi, %esi
	call	cudaMemset@PLT
	leaq	88(%rsp), %rdi
	movl	$128, %esi
	call	_Z10cudaMallocIhE9cudaErrorPPT_m.isra.0
	movl	$1, %esi
	movl	$128, %edi
	call	calloc@PLT
	movl	$1, %esi
	movl	$257, %edi
	movq	%rax, 8(%rsp)
	call	calloc@PLT
	leaq	96(%rsp), %rdi
	movl	$257, %esi
	movq	%rax, 16(%rsp)
	call	_Z10cudaMallocIhE9cudaErrorPPT_m.isra.0
	movq	96(%rsp), %rdi
	movl	$257, %edx
	xorl	%esi, %esi
	call	cudaMemset@PLT
	movl	$1, %esi
	movl	$384, %edi
	call	calloc@PLT
	leaq	104(%rsp), %rdi
	movl	$384, %esi
	movq	%rax, 24(%rsp)
	call	_Z10cudaMallocIhE9cudaErrorPPT_m.isra.0
	movl	$1, %esi
	movl	$128, %edi
	call	calloc@PLT
	leaq	112(%rsp), %rdi
	movl	$128, %esi
	movq	%rax, 32(%rsp)
	call	_Z10cudaMallocIhE9cudaErrorPPT_m.isra.0
	movq	112(%rsp), %rdi
	movl	$128, %edx
	xorl	%esi, %esi
	call	cudaMemset@PLT
	movl	$1, %esi
	movl	$256, %edi
	call	calloc@PLT
	leaq	120(%rsp), %rdi
	movl	$256, %esi
	movq	%rax, 40(%rsp)
	call	_Z10cudaMallocIhE9cudaErrorPPT_m.isra.0
	movq	120(%rsp), %rdi
	movl	$256, %edx
	xorl	%esi, %esi
	call	cudaMemset@PLT
	movl	$1, %esi
	movl	$129, %edi
	call	calloc@PLT
	movl	$1, %esi
	movl	$129, %edi
	movq	%rax, %r13
	call	calloc@PLT
	movl	$4, %esi
	movl	$512, %edi
	movq	%rax, %r14
	call	calloc@PLT
	leaq	128(%rsp), %rdi
	movl	$2048, %esi
	movq	%rax, %r15
	call	_Z10cudaMallocIhE9cudaErrorPPT_m.isra.0
	movq	128(%rsp), %rdi
	xorl	%esi, %esi
	movl	$2048, %edx
	call	cudaMemset@PLT
	leaq	136(%rsp), %rdi
	call	cudaEventCreate@PLT
	testl	%eax, %eax
	je	.L77
	leaq	.LC0(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
.L77:
	leaq	144(%rsp), %rdi
	call	cudaEventCreate@PLT
	testl	%eax, %eax
	je	.L78
	leaq	.LC0(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
.L78:
	movq	136(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movq	%rbp, %rsi
	pushq	%rax
	.cfi_def_cfa_offset 232
	pushq	%r12
	.cfi_def_cfa_offset 240
	pushq	%r15
	.cfi_def_cfa_offset 248
	pushq	152(%rsp)
	.cfi_def_cfa_offset 256
	pushq	88(%rsp)
	.cfi_def_cfa_offset 264
	pushq	104(%rsp)
	.cfi_def_cfa_offset 272
	pushq	$3
	.cfi_def_cfa_offset 280
	pushq	%r14
	.cfi_def_cfa_offset 288
	pushq	%r13
	.cfi_def_cfa_offset 296
	pushq	192(%rsp)
	.cfi_def_cfa_offset 304
	pushq	120(%rsp)
	.cfi_def_cfa_offset 312
	pushq	200(%rsp)
	.cfi_def_cfa_offset 320
	pushq	128(%rsp)
	.cfi_def_cfa_offset 328
	pushq	208(%rsp)
	.cfi_def_cfa_offset 336
	pushq	136(%rsp)
	.cfi_def_cfa_offset 344
	pushq	216(%rsp)
	.cfi_def_cfa_offset 352
	movq	144(%rsp), %r9
	movq	136(%rsp), %r8
	movq	216(%rsp), %rcx
	movq	208(%rsp), %rdx
	movq	200(%rsp), %rdi
	call	_Z14exponentiationPhS_S_S_S_S_S_S_S_S_S_S_S_S_S_jS_S_PjS0_S_
	subq	$-128, %rsp
	.cfi_def_cfa_offset 224
	movl	$128, %edi
	call	malloc@PLT
	movl	$32, %ecx
	movq	%rbx, %rdi
	xorl	%esi, %esi
	movq	%rax, %rbp
	movl	$128, %edx
	movabsq	$-2443079938585982319, %rax
	movq	%rax, 0(%rbp)
	movabsq	$3737263753417261525, %rax
	movq	%rax, 8(%rbp)
	movabsq	$-6885066599430936869, %rax
	movq	%rax, 16(%rbp)
	movabsq	$-348031553207505184, %rax
	movq	%rax, 24(%rbp)
	movabsq	$7786741645037550750, %rax
	movq	%rax, 32(%rbp)
	movabsq	$-7576899760836331525, %rax
	movq	%rax, 40(%rbp)
	movabsq	$7658712522179984311, %rax
	movq	%rax, 48(%rbp)
	movabsq	$-6735306766495284301, %rax
	movq	%rax, 56(%rbp)
	movabsq	$-2648134166815690918, %rax
	movq	%rax, 64(%rbp)
	movabsq	$7167369312206811226, %rax
	movq	%rax, 72(%rbp)
	movabsq	$-7743534281320969286, %rax
	movq	%rax, 80(%rbp)
	movabsq	$2387300973242090033, %rax
	movq	%rax, 88(%rbp)
	movabsq	$7409015650788201817, %rax
	movq	%rax, 96(%rbp)
	movabsq	$-5071924630037949112, %rax
	movq	%rax, 104(%rbp)
	movabsq	$5998317401964290198, %rax
	movq	%rax, 112(%rbp)
	movabsq	$4046903733242681419, %rax
	movq	%rax, 120(%rbp)
	xorl	%eax, %eax
	rep stosl
	movq	72(%rsp), %rdi
	call	cudaMemset@PLT
	movq	%rbp, %rsi
	pushq	%rdx
	.cfi_def_cfa_offset 232
	pushq	%r12
	.cfi_def_cfa_offset 240
	pushq	%r15
	.cfi_def_cfa_offset 248
	pushq	152(%rsp)
	.cfi_def_cfa_offset 256
	pushq	88(%rsp)
	.cfi_def_cfa_offset 264
	pushq	104(%rsp)
	.cfi_def_cfa_offset 272
	pushq	$128
	.cfi_def_cfa_offset 280
	pushq	%r14
	.cfi_def_cfa_offset 288
	pushq	%r13
	.cfi_def_cfa_offset 296
	pushq	192(%rsp)
	.cfi_def_cfa_offset 304
	pushq	120(%rsp)
	.cfi_def_cfa_offset 312
	pushq	200(%rsp)
	.cfi_def_cfa_offset 320
	pushq	128(%rsp)
	.cfi_def_cfa_offset 328
	pushq	208(%rsp)
	.cfi_def_cfa_offset 336
	pushq	136(%rsp)
	.cfi_def_cfa_offset 344
	pushq	216(%rsp)
	.cfi_def_cfa_offset 352
	movq	144(%rsp), %r9
	movq	136(%rsp), %r8
	movq	216(%rsp), %rcx
	movq	200(%rsp), %rdx
	movq	208(%rsp), %rdi
	call	_Z14exponentiationPhS_S_S_S_S_S_S_S_S_S_S_S_S_S_jS_S_PjS0_S_
	movq	272(%rsp), %rdi
	subq	$-128, %rsp
	.cfi_def_cfa_offset 224
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movq	144(%rsp), %rdi
	call	cudaEventSynchronize@PLT
	testl	%eax, %eax
	je	.L79
	leaq	.LC0(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
.L79:
	movq	144(%rsp), %rdx
	leaq	52(%rsp), %rdi
	xorl	%ebp, %ebp
	movq	136(%rsp), %rsi
	movl	$0x00000000, 52(%rsp)
	leaq	.LC4(%rip), %r12
	call	cudaEventElapsedTime@PLT
	leaq	.LC3(%rip), %rsi
	movl	$2, %edi
	movb	$1, %al
	movss	52(%rsp), %xmm0
	divss	.LC2(%rip), %xmm0
	cvtss2sd	%xmm0, %xmm0
	call	__printf_chk@PLT
	movq	72(%rsp), %rsi
	movl	$2, %ecx
	movq	%rbx, %rdi
	movl	$128, %edx
	call	cudaMemcpy@PLT
.L80:
	movzbl	(%rbx,%rbp), %ecx
	movl	%ebp, %edx
	movq	%r12, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	incq	%rbp
	call	__printf_chk@PLT
	cmpq	$128, %rbp
	jne	.L80
	movq	152(%rsp), %rax
	subq	%fs:40, %rax
	je	.L81
	call	__stack_chk_fail@PLT
.L81:
	addq	$168, %rsp
	.cfi_def_cfa_offset 56
	xorl	%eax, %eax
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2040:
	.size	main, .-main
	.section	.rodata.str1.1
.LC5:
	.string	"_Z10get_squarePhPjj"
.LC6:
	.string	"_Z12get_productsPhS_Pjj"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2075:
	.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	.LC5(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z10get_squarePhPjj(%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	.LC6(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z12get_productsPhS_Pjj(%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
.LFE2075:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.nvFatBinSegment,"aw"
	.align 8
	.type	_ZL15__fatDeviceText, @object
	.size	_ZL15__fatDeviceText, 24
_ZL15__fatDeviceText:
	.long	1180844977
	.long	1
	.quad	fatbinData
	.quad	0
	.local	_ZL20__cudaFatCubinHandle
	.comm	_ZL20__cudaFatCubinHandle,8,8
	.section	.rodata.cst4,"aM",@progbits,4
	.align 4
.LC2:
	.long	1148846080
	.ident	"GCC: (Ubuntu 13.3.0-6ubuntu2~24.04) 13.3.0"
	.section	.note.GNU-stack,"",@progbits
	.section	.note.gnu.property,"a"
	.align 8
	.long	1f - 0f
	.long	4f - 1f
	.long	5
0:
	.string	"GNU"
1:
	.align 8
	.long	0xc0000002
	.long	3f - 2f
2:
	.long	0x3
3:
	.align 8
4:


#include <hip/hip_runtime.h>
#include <stdio.h>
#include <stdlib.h>
#include <time.h>
#include <unistd.h>
#include <sys/time.h>
#include <string.h>

#define R_size 129
#define k 1024
#define n_size 128

__global__ void get_square(unsigned char a[], unsigned int accumulator[], unsigned int n);
void square(unsigned char *a, unsigned char *c, unsigned int size);
__global__ void get_products(unsigned char a[], unsigned char b[], unsigned int accumulator[], unsigned int n);
void barrett_reduction(unsigned char *buf_cpu, unsigned char *buf_gpu, unsigned char *gpu_r, unsigned char *gpu_n, unsigned char *reduction, unsigned char *temp_cpu, unsigned char *temp_gpu, unsigned char *shifted_cpu, unsigned char *shifted_gpu, unsigned char *xprime_cpu, unsigned char *xprime_gpu, unsigned char *result, unsigned char *tmp, unsigned int exp_size, unsigned int *kernel_buf, unsigned int *transfer, unsigned char *n);
void multiplication(unsigned char *a, unsigned char *b, unsigned char *c, unsigned int size); 
char checkbit(unsigned char *exponent, unsigned int index_of_bit); 
void subtraction(unsigned char *a, unsigned char *b, unsigned char *c, unsigned int size); 
void bit_shift(unsigned char *a, unsigned char *b, unsigned int shift, unsigned int size_of_a); 
void exponentiation(unsigned char *gpu_message, unsigned char *exponent, unsigned char *gpu_ciphertext, unsigned char *m0_copy, unsigned char *reduction, unsigned char *buf_cpu, unsigned char *buf_gpu, unsigned char *temp_cpu, unsigned char *temp_gpu, unsigned char *shifted_cpu, unsigned char *shifted_gpu, unsigned char *xprime_cpu, unsigned char *xprime_gpu, unsigned char *result, unsigned char *tmp, unsigned int exponent_size, unsigned char *gpu_r, unsigned char *gpu_n, unsigned int *kernel_buf, unsigned int *transfer, unsigned char *n);

int main(int argc, char *argv[]) {

unsigned char *n = (unsigned char *) calloc((4*n_size + n_size), sizeof(char));

    n[0] = 0xcf;
    n[1] = 0x82;
    n[2] = 0x69;
    n[3] = 0x57;
    n[4] = 0x4d;
    n[5] = 0xe7;
    n[6] = 0x82;
    n[7] = 0x1a;
    n[8] = 0xe4;
    n[9] = 0x20;
    n[10] = 0x14;
    n[11] = 0x47;
    n[12] = 0x39;
    n[13] = 0x52;
    n[14] = 0x55;
    n[15] = 0x28;
    n[16] = 0xed;
    n[17] = 0x3f;
    n[18] = 0xa4;
    n[19] = 0x61;
    n[20] = 0xd3;
    n[21] = 0xf4;
    n[22] = 0xf2;
    n[23] = 0x34;
    n[24] = 0x6a;
    n[25] = 0x54;
    n[26] = 0xd1;
    n[27] = 0x15;
    n[28] = 0x7d;
    n[29] = 0x67;
    n[30] = 0xb;
    n[31] = 0xc7;
    n[32] = 0x8c;
    n[33] = 0xfe;
    n[34] = 0x1b;
    n[35] = 0x68;
    n[36] = 0x44;
    n[37] = 0x7;
    n[38] = 0x26;
    n[39] = 0x99;
    n[40] = 0xb;
    n[41] = 0x4d;
    n[42] = 0xc7;
    n[43] = 0x3f;
    n[44] = 0x52;
    n[45] = 0x90;
    n[46] = 0x2;
    n[47] = 0x68;
    n[48] = 0x3d;
    n[49] = 0x83;
    n[50] = 0x1d;
    n[51] = 0x79;
    n[52] = 0x7a;
    n[53] = 0x3f;
    n[54] = 0x36;
    n[55] = 0xf3;
    n[56] = 0x41;
    n[57] = 0x8b;
    n[58] = 0x7c;
    n[59] = 0xdf;
    n[60] = 0x64;
    n[61] = 0xac;
    n[62] = 0x74;
    n[63] = 0x7c;
    n[64] = 0x8;
    n[65] = 0xdb;
    n[66] = 0xa0;
    n[67] = 0x6f;
    n[68] = 0x10;
    n[69] = 0x71;
    n[70] = 0x13;
    n[71] = 0x86;
    n[72] = 0xaf;
    n[73] = 0xb8;
    n[74] = 0x71;
    n[75] = 0xf8;
    n[76] = 0xf0;
    n[77] = 0x45;
    n[78] = 0xa7;
    n[79] = 0x94;
    n[80] = 0xb3;
    n[81] = 0x6b;
    n[82] = 0x1e;
    n[83] = 0xff;
    n[84] = 0x8e;
    n[85] = 0x13;
    n[86] = 0xae;
    n[87] = 0xc2;
    n[88] = 0x59;
    n[89] = 0x56;
    n[90] = 0xd3;
    n[91] = 0xd;
    n[92] = 0x20;
    n[93] = 0x62;
    n[94] = 0x21;
    n[95] = 0x30;
    n[96] = 0x1d;
    n[97] = 0x6b;
    n[98] = 0x5e;
    n[99] = 0xc;
    n[100] = 0x0;
    n[101] = 0x35;
    n[102] = 0xae;
    n[103] = 0xbd;
    n[104] = 0xa5;
    n[105] = 0xc2;
    n[106] = 0x25;
    n[107] = 0x98;
    n[108] = 0xe7;
    n[109] = 0x57;
    n[110] = 0x89;
    n[111] = 0xc;
    n[112] = 0x12;
    n[113] = 0xf9;
    n[114] = 0x33;
    n[115] = 0x3d;
    n[116] = 0xa;
    n[117] = 0xac;
    n[118] = 0x51;
    n[119] = 0xd8;
    n[120] = 0x5c;
    n[121] = 0x40;
    n[122] = 0x9b;
    n[123] = 0xfa;
    n[124] = 0xf9;
    n[125] = 0xbc;
    n[126] = 0x3;
    n[127] = 0xe6;

    
	unsigned char *gpu_n;
	hipMalloc(&gpu_n, (4*n_size + n_size));
	hipMemcpy(gpu_n, n, (4*n_size + n_size), hipMemcpyHostToDevice);

unsigned char *r = (unsigned char *) calloc(2*n_size, sizeof(char));

    r[0] = 0x7f;
    r[1] = 0x9d;
    r[2] = 0xe9;
    r[3] = 0x40;
    r[4] = 0x57;
    r[5] = 0x2;
    r[6] = 0x6e;
    r[7] = 0x93;
    r[8] = 0x2b;
    r[9] = 0xb4;
    r[10] = 0xe3;
    r[11] = 0xfd;
    r[12] = 0xba;
    r[13] = 0xc;
    r[14] = 0xcd;
    r[15] = 0x78;
    r[16] = 0x7d;
    r[17] = 0xae;
    r[18] = 0x8d;
    r[19] = 0x80;
    r[20] = 0xff;
    r[21] = 0x66;
    r[22] = 0x33;
    r[23] = 0xb;
    r[24] = 0x28;
    r[25] = 0x4c;
    r[26] = 0x93;
    r[27] = 0x30;
    r[28] = 0x2;
    r[29] = 0x92;
    r[30] = 0xa0;
    r[31] = 0x7c;
    r[32] = 0xf1;
    r[33] = 0xc;
    r[34] = 0xa;
    r[35] = 0x5e;
    r[36] = 0xf2;
    r[37] = 0x9a;
    r[38] = 0x8f;
    r[39] = 0x17;
    r[40] = 0x4c;
    r[41] = 0x82;
    r[42] = 0x25;
    r[43] = 0xe5;
    r[44] = 0x98;
    r[45] = 0x45;
    r[46] = 0x4d;
    r[47] = 0xc7;
    r[48] = 0xd9;
    r[49] = 0x53;
    r[50] = 0x5e;
    r[51] = 0x5a;
    r[52] = 0x6e;
    r[53] = 0x37;
    r[54] = 0x43;
    r[55] = 0x29;
    r[56] = 0x88;
    r[57] = 0xcb;
    r[58] = 0xe9;
    r[59] = 0x31;
    r[60] = 0x2f;
    r[61] = 0xd7;
    r[62] = 0x6;
    r[63] = 0xfb;
    r[64] = 0xf1;
    r[65] = 0x38;
    r[66] = 0xdf;
    r[67] = 0xc4;
    r[68] = 0xda;
    r[69] = 0x7c;
    r[70] = 0x9;
    r[71] = 0x5c;
    r[72] = 0xf9;
    r[73] = 0x2b;
    r[74] = 0x81;
    r[75] = 0x30;
    r[76] = 0xe9;
    r[77] = 0x29;
    r[78] = 0xcd;
    r[79] = 0x45;
    r[80] = 0xee;
    r[81] = 0xff;
    r[82] = 0x5b;
    r[83] = 0x3c;
    r[84] = 0x23;
    r[85] = 0x6d;
    r[86] = 0xb9;
    r[87] = 0xa1;
    r[88] = 0x89;
    r[89] = 0x3f;
    r[90] = 0xc3;
    r[91] = 0x9e;
    r[92] = 0xa1;
    r[93] = 0x30;
    r[94] = 0x98;
    r[95] = 0xf8;
    r[96] = 0xc8;
    r[97] = 0x4a;
    r[98] = 0xbe;
    r[99] = 0xc6;
    r[100] = 0x49;
    r[101] = 0xf7;
    r[102] = 0xb3;
    r[103] = 0xff;
    r[104] = 0x9;
    r[105] = 0x3b;
    r[106] = 0x94;
    r[107] = 0x9d;
    r[108] = 0x2f;
    r[109] = 0x5c;
    r[110] = 0x68;
    r[111] = 0xe1;
    r[112] = 0x6;
    r[113] = 0xf1;
    r[114] = 0x33;
    r[115] = 0xeb;
    r[116] = 0xc5;
    r[117] = 0x88;
    r[118] = 0xa5;
    r[119] = 0x1c;
    r[120] = 0xde;
    r[121] = 0x2c;
    r[122] = 0x64;
    r[123] = 0xad;
    r[124] = 0x5c;
    r[125] = 0xc9;
    r[126] = 0xeb;
    r[127] = 0x1c;
    r[128] = 0x1;

	unsigned char *gpu_r;
	hipMalloc(&gpu_r, (2*n_size));
	hipMemcpy(gpu_r, r, 2*n_size, hipMemcpyHostToDevice);

	unsigned char *message = (unsigned char *) calloc(n_size, sizeof(char));
	message[0] = 0x68;//h
	message[1] = 0x65;//e
	message[2] = 0x6c;//l
	message[3] = 0x6c;//l
	message[4] = 0x6f;//o


	unsigned char *gpu_message;
	hipMalloc(&gpu_message, n_size);
	hipMemcpy(gpu_message, message, n_size, hipMemcpyHostToDevice);

	unsigned char *exponent = (unsigned char *) malloc(3);
	exponent[0] = 0x01;
	exponent[1] = 0x00;
	exponent[2] = 0x01;
	unsigned int exponent_size = 3;

	//exponentiate m^e mod n
	//parameters: 
	//message(m)
	//exponent(e)
	//precomputation of r = floor((4^k)/n) where k is found by where (2^k) > n
	//modulus (n)

	unsigned char *cpu_ciphertext = (unsigned char *) calloc(n_size, sizeof(char));
	unsigned char *gpu_ciphertext;
	hipMalloc(&gpu_ciphertext, n_size);
	hipMemset(gpu_ciphertext, 0x00, n_size);

	unsigned char *m0_copy;
	hipMalloc(&m0_copy, n_size);

	unsigned char *reduction = (unsigned char *) calloc(n_size, sizeof(char));

	unsigned char *buf_cpu = (unsigned char *) calloc((n_size * 2) + 1, sizeof(char));
	unsigned char *buf_gpu;
	hipMalloc(&buf_gpu, ((n_size * 2) + 1));
	hipMemset(buf_gpu, 0x00, (n_size * 2) + 1);

	unsigned char *temp_cpu = (unsigned char *) calloc(3*n_size, sizeof(char));
	unsigned char *temp_gpu;
	hipMalloc(&temp_gpu, (3*n_size));
	
	unsigned char *shifted_cpu = (unsigned char *) calloc(n_size, sizeof(char));
	unsigned char *shifted_gpu;
	hipMalloc(&shifted_gpu, n_size);
	hipMemset(shifted_gpu, 0x00, n_size);

	unsigned char *xprime_cpu = (unsigned char *) calloc(2*n_size, sizeof(char));
	unsigned char *xprime_gpu;
	hipMalloc(&xprime_gpu, (2*n_size));
	hipMemset(xprime_gpu, 0x00, 2*n_size);

	unsigned char *result = (unsigned char *) calloc(n_size + 1, sizeof(char));
	
	unsigned char *tmp = (unsigned char *) calloc(n_size + 1, sizeof(char));

	unsigned int *transfer = (unsigned int *) calloc(4*n_size, sizeof(int));
	unsigned int *kernel_buf;
	hipMalloc(&kernel_buf, 4*n_size*sizeof(int));
	hipMemset(kernel_buf, 0x00, 4*n_size*sizeof(int));

	hipError_t error;
	hipEvent_t start;

	error = hipEventCreate(&start);
	if(error != hipSuccess)
	printf("error\n");

	hipEvent_t stop;
	error = hipEventCreate(&stop);
	if(error != hipSuccess)
	printf("error\n");

	error = hipEventRecord(start, NULL);


	exponentiation(gpu_message, exponent, gpu_ciphertext, m0_copy, reduction, buf_cpu, buf_gpu, temp_cpu, temp_gpu, shifted_cpu, shifted_gpu, xprime_cpu, xprime_gpu, result, tmp, exponent_size, gpu_r, gpu_n, kernel_buf, transfer, n);

	unsigned int d_exponent_size = 128;
	unsigned char *d_exponent = (unsigned char *) malloc(128);
    
    d_exponent[0] = 0x91;
    d_exponent[1] = 0xa;
    d_exponent[2] = 0xb3;
    d_exponent[3] = 0x66;
    d_exponent[4] = 0xbd;
    d_exponent[5] = 0x6f;
    d_exponent[6] = 0x18;
    d_exponent[7] = 0xde;
    d_exponent[8] = 0xd5;
    d_exponent[9] = 0x1;
    d_exponent[10] = 0x61;
    d_exponent[11] = 0x36;
    d_exponent[12] = 0x95;
    d_exponent[13] = 0x6d;
    d_exponent[14] = 0xdd;
    d_exponent[15] = 0x33;
    d_exponent[16] = 0xdb;
    d_exponent[17] = 0x26;
    d_exponent[18] = 0x3;
    d_exponent[19] = 0xe;
    d_exponent[20] = 0x68;
    d_exponent[21] = 0x54;
    d_exponent[22] = 0x73;
    d_exponent[23] = 0xa0;
    d_exponent[24] = 0xe0;
    d_exponent[25] = 0x6e;
    d_exponent[26] = 0x70;
    d_exponent[27] = 0x74;
    d_exponent[28] = 0x25;
    d_exponent[29] = 0x8b;
    d_exponent[30] = 0x2b;
    d_exponent[31] = 0xfb;
    d_exponent[32] = 0x9e;
    d_exponent[33] = 0x3c;
    d_exponent[34] = 0x34;
    d_exponent[35] = 0x2e;
    d_exponent[36] = 0x45;
    d_exponent[37] = 0x10;
    d_exponent[38] = 0x10;
    d_exponent[39] = 0x6c;
    d_exponent[40] = 0xfb;
    d_exponent[41] = 0xb7;
    d_exponent[42] = 0x9b;
    d_exponent[43] = 0xc8;
    d_exponent[44] = 0xcf;
    d_exponent[45] = 0x71;
    d_exponent[46] = 0xd9;
    d_exponent[47] = 0x96;
    d_exponent[48] = 0xb7;
    d_exponent[49] = 0xbb;
    d_exponent[50] = 0x5f;
    d_exponent[51] = 0x19;
    d_exponent[52] = 0x76;
    d_exponent[53] = 0x36;
    d_exponent[54] = 0x49;
    d_exponent[55] = 0x6a;
    d_exponent[56] = 0xb3;
    d_exponent[57] = 0x83;
    d_exponent[58] = 0xc3;
    d_exponent[59] = 0x59;
    d_exponent[60] = 0x2e;
    d_exponent[61] = 0x62;
    d_exponent[62] = 0x87;
    d_exponent[63] = 0xa2;
    d_exponent[64] = 0x5a;
    d_exponent[65] = 0x2f;
    d_exponent[66] = 0x60;
    d_exponent[67] = 0x75;
    d_exponent[68] = 0x1;
    d_exponent[69] = 0xf0;
    d_exponent[70] = 0x3f;
    d_exponent[71] = 0xdb;
    d_exponent[72] = 0x5a;
    d_exponent[73] = 0x70;
    d_exponent[74] = 0x1f;
    d_exponent[75] = 0x44;
    d_exponent[76] = 0x6a;
    d_exponent[77] = 0x9c;
    d_exponent[78] = 0x77;
    d_exponent[79] = 0x63;
    d_exponent[80] = 0xba;
    d_exponent[81] = 0xcb;
    d_exponent[82] = 0xcd;
    d_exponent[83] = 0x1f;
    d_exponent[84] = 0x99;
    d_exponent[85] = 0x70;
    d_exponent[86] = 0x89;
    d_exponent[87] = 0x94;
    d_exponent[88] = 0x31;
    d_exponent[89] = 0x2;
    d_exponent[90] = 0xa;
    d_exponent[91] = 0x32;
    d_exponent[92] = 0x96;
    d_exponent[93] = 0x65;
    d_exponent[94] = 0x21;
    d_exponent[95] = 0x21;
    d_exponent[96] = 0x59;
    d_exponent[97] = 0x55;
    d_exponent[98] = 0x8a;
    d_exponent[99] = 0xd0;
    d_exponent[100] = 0x7a;
    d_exponent[101] = 0x1c;
    d_exponent[102] = 0xd2;
    d_exponent[103] = 0x66;
    d_exponent[104] = 0x48;
    d_exponent[105] = 0x95;
    d_exponent[106] = 0x8;
    d_exponent[107] = 0xd3;
    d_exponent[108] = 0x6b;
    d_exponent[109] = 0xe7;
    d_exponent[110] = 0x9c;
    d_exponent[111] = 0xb9;
    d_exponent[112] = 0x96;
    d_exponent[113] = 0x20;
    d_exponent[114] = 0x20;
    d_exponent[115] = 0x8a;
    d_exponent[116] = 0xe5;
    d_exponent[117] = 0x4d;
    d_exponent[118] = 0x3e;
    d_exponent[119] = 0x53;
    d_exponent[120] = 0x4b;
    d_exponent[121] = 0xd8;
    d_exponent[122] = 0x21;
    d_exponent[123] = 0x4;
    d_exponent[124] = 0x81;
    d_exponent[125] = 0x7d;
    d_exponent[126] = 0x29;
    d_exponent[127] = 0x38;


	memset(message, 0x00, n_size);
	hipMemset(gpu_message, 0x00, n_size);

	exponentiation(gpu_ciphertext, d_exponent, gpu_message, m0_copy, reduction, buf_cpu, buf_gpu, temp_cpu, temp_gpu, shifted_cpu, shifted_gpu, xprime_cpu, xprime_gpu, result, tmp, d_exponent_size, gpu_r, gpu_n, kernel_buf, transfer, n);


	error = hipEventRecord(stop, NULL);

	error = hipEventSynchronize(stop);

	if(error != hipSuccess)
		printf("error\n");

	float msecTotal = 0.0f;
	error = hipEventElapsedTime(&msecTotal, start, stop);
	printf("GPU time: %.6f\n", msecTotal / 1000);


	hipMemcpy(message, gpu_message, n_size, hipMemcpyDeviceToHost);
	int z = 0;
	while (z < n_size) {
		printf("message[%d] = %x\n", z, message[z]);
		z++;
	}

	return 0;
}

void exponentiation(unsigned char *gpu_message, unsigned char *exponent, unsigned char *gpu_ciphertext, unsigned char *m0_copy, unsigned char *reduction, unsigned char *buf_cpu, unsigned char *buf_gpu, unsigned char *temp_cpu, unsigned char *temp_gpu, unsigned char *shifted_cpu, unsigned char *shifted_gpu, unsigned char *xprime_cpu, unsigned char *xprime_gpu, unsigned char *result, unsigned char *tmp, unsigned int exponent_size, unsigned char *gpu_r, unsigned char *gpu_n, unsigned int *kernel_buf, unsigned int *transfer, unsigned char *n) {


	dim3 blocksPerGrid(2);
	dim3 threadsPerBlock(64);


	//get the total amount of bits in strlen(exponent) zero based
	//not including the final char index msb (byte)
	unsigned int total_bits = exponent_size * 8 - 1;
	
	//find the most signinficant bit in the most significant byte (char index)
	//find most significant bit in exponent[exp_size - 1]
	unsigned char mask = 0x80; //10000000 in binary
	unsigned char msb = 0;
	int i = 0;
	while(i < 8) {
		if((exponent[exponent_size - 1] & (mask >> i)) == (mask >> i)) {
			msb = i;
			break;
		}
		i++;
	}

	//subtract most significant bit from total_bits to know total amount of significant bits
	//for loop of exponent in binary
	unsigned int exp_bits = (total_bits - msb);

	//keep copy of original message m0
	hipMemcpy(m0_copy, gpu_message, n_size, hipMemcpyDeviceToDevice);

	//compute m^e where e is in binary 
	//RULES:
		//iterate over the values of msb to 0 bit by bit
		//msb is amount of relevent bits to check for exponentiation
		//total bits is the amount of total bits in exponent lenth
		//square m(current) for each itteration
		//check if current bit is 1
		//current bit is 1: m(current) * m0
		//curent bit is 0: return to loop


	//subtract one from total because to exponentiate in binary
	//start at the second bit after the most significant bit
	//each bit equals m^2 and when the current bit is 1 it is
	//(m^2)*m0 or if it is 0 then m^2
	int index_of_bit = exp_bits - 1; //subtraction of 1 is becuase msb is zero based
	while (index_of_bit >= 0) {

		//allocate space for reduction to hold a value strickly less than n
		//buf holds value at most m^2 which is less than n^2

		//calculate m^2
		get_square<<<blocksPerGrid, threadsPerBlock>>>(gpu_message, kernel_buf, n_size);
		hipMemcpy(transfer, kernel_buf, 2*n_size*sizeof(unsigned int), hipMemcpyDeviceToHost);

		unsigned int index = 0;
		while(index < 2*n_size) {
			buf_cpu[index] = (unsigned char) transfer[index];
			transfer[index + 1] += (unsigned int) (transfer[index]>>8);
			index++;
		}
	
		memset(transfer, 0x00, 4*n_size*sizeof(int));	
		hipMemset(kernel_buf, 0x00, 4*n_size*sizeof(int));
		hipMemcpy(buf_gpu, buf_cpu, 2*n_size, hipMemcpyHostToDevice);

		//calculate m^2 mod n
		barrett_reduction(buf_cpu, buf_gpu, gpu_r, gpu_n, reduction, temp_cpu, temp_gpu, shifted_cpu, shifted_gpu, xprime_cpu, xprime_gpu, result, tmp, exponent_size, kernel_buf, transfer, n);

		hipMemcpy(gpu_message, reduction, n_size, hipMemcpyHostToDevice);

		hipMemset(buf_gpu, 0x00, 2*n_size);
		memset(buf_cpu, 0x00, 2*n_size);
		memset(reduction, 0x00, n_size);

		char bit;
		if ((bit = checkbit(exponent, index_of_bit)) == 1) {
			//m * m0
			get_products<<<blocksPerGrid, threadsPerBlock>>>(gpu_message, m0_copy, kernel_buf, n_size);
			hipMemcpy(transfer, kernel_buf, 2*n_size*sizeof(unsigned int), hipMemcpyDeviceToHost);

			index = 0;
			while(index < 2*n_size) {
				buf_cpu[index] = (unsigned char) transfer[index];
				transfer[index + 1] += (unsigned int) (transfer[index]>>8);
				index++;
			}
		
			memset(transfer, 0x00, 4*n_size*sizeof(int));
			hipMemset(kernel_buf, 0x00, 4*n_size*sizeof(int));
			hipMemcpy(buf_gpu, buf_cpu, 2*n_size, hipMemcpyHostToDevice);



			barrett_reduction(buf_cpu, buf_gpu, gpu_r, gpu_n, reduction, temp_cpu, temp_gpu, shifted_cpu, shifted_gpu, xprime_cpu, xprime_gpu, result, tmp, exponent_size, kernel_buf, transfer, n);

			hipMemcpy(gpu_message, reduction, n_size, hipMemcpyHostToDevice);
			hipMemset(buf_gpu, 0x00, 2*n_size);
			memset(buf_cpu, 0x00, 2*n_size);
			memset(reduction, 0x00, n_size);
		
		}

		index_of_bit--;
	}
	
	//copy back final value of message to ciphertext for decryption
	hipMemcpy(gpu_ciphertext, gpu_message, n_size, hipMemcpyDeviceToHost);

	hipMemset(m0_copy, 0x00, n_size);

	return;
}
void barrett_reduction(unsigned char *buf_cpu, unsigned char *buf_gpu, unsigned char *gpu_r, unsigned char *gpu_n, unsigned char *reduction, unsigned char *temp_cpu, unsigned char *temp_gpu, unsigned char *shifted_cpu, unsigned char *shifted_gpu, unsigned char *xprime_cpu, unsigned char *xprime_gpu, unsigned char *result, unsigned char *tmp, unsigned int exp_size, unsigned int *kernel_buf, unsigned int *transfer, unsigned char *n) {
	////////////////////////////////////////////////////////////////////////////////////

	////////////////////  calculate: t = x - ((x*r)/(4^k))*n  /////////////////////////

	///////////////////////////////////////////////////////////////////////////////////
	///////////////////////////////////////////////////////////////////////////////////
	//multiply: x * r = temp
	//size of x is assumed to be the largest value which is = largest value of 2*n
	//size of r is precomputed
	dim3 blocksPerGrid_two(4);
	dim3 threadsPerBlock_two(64);

	get_products<<<blocksPerGrid_two, threadsPerBlock_two>>>(gpu_r, buf_gpu, kernel_buf, 2*n_size);
	hipMemcpy(transfer, kernel_buf, 3*n_size*sizeof(unsigned int), hipMemcpyDeviceToHost);

	unsigned int index = 0;
	while(index < 3*n_size) {
		temp_cpu[index] = (unsigned char) transfer[index];
		transfer[index + 1] += (unsigned int) (transfer[index]>>8);
		index++;
	}

	hipMemset(kernel_buf, 0x00, 4*n_size*sizeof(int));

	//shift bits by (4^k) or (2^(2*k))
	//shift temp by 2*k store to shifted
	//size of shifted is 2*n + sizeof(r)
	
	//find the actual amount of bits/bytes left in the value of temp
	//which is equal to x * r so that the correct size of the value
	//can be used in the bit_shift function

	//first find the amount of bytes from most significant byte
	//to least and then when one char does not equal to 0x00

	unsigned int zero_bytes = 0;
	int count = (3*n_size) - 1;
	while((count >= 0) && (temp_cpu[count] == 0x00)) {
		count--;
		zero_bytes++;
	}


	bit_shift(temp_cpu, shifted_cpu, k, (3*n_size) - zero_bytes);

	//multiply: shifted * n = xprime
	//xprime is the size of 2*n + R_size - (k >> 0x07) + n
	hipMemcpy(shifted_gpu, shifted_cpu, n_size, hipMemcpyHostToDevice);

	dim3 blocksPerGrid_one(2);
	dim3 threadsPerBlock_one(64);
	get_products<<<blocksPerGrid_one, threadsPerBlock_one>>>(shifted_gpu, gpu_n, kernel_buf, n_size);
	hipMemcpy(transfer, kernel_buf, 2*n_size*sizeof(unsigned int), hipMemcpyDeviceToHost);

	index = 0;
	while(index < 2*n_size) {
		xprime_cpu[index] = (unsigned char) transfer[index];
		transfer[index + 1] += (unsigned int) (transfer[index]>>8);
		index++;
	}



	//subtract xprime from x^2
	hipMemset(kernel_buf, 0x00, 4*n_size*sizeof(int));
	subtraction(buf_cpu, xprime_cpu, result, 2*n_size);

	//the field of n, if the value is not within the field of n then reduce the value by subtracting
	//the value of result = t - n which is guaranteed to be in the field of n
	if ((result[n_size] == 0x00) && (result[n_size - 1] < n[n_size - 1])) {
		memcpy(reduction, result, n_size);
	}

	else {
		unsigned char *tmp = (unsigned char *) calloc(n_size + 1, sizeof(char));
		subtraction(result, n, tmp, n_size + 1);
		memcpy(reduction, tmp, n_size);
		memset(tmp, 0x00, n_size + 1);
	}


	memset(temp_cpu, 0x00, 3*n_size);
	hipMemset(temp_gpu, 0x00, 3*n_size);
	memset(shifted_cpu, 0x00, n_size);
	hipMemset(shifted_gpu, 0x00, n_size);
	memset(xprime_cpu, 0x00, 2*n_size);
	hipMemset(xprime_gpu, 0x00, 2*n_size);
	memset(result, 0x00, n_size + 1);
	memset(transfer, 0x00, 4*n_size*sizeof(int));


	return;
}

void subtraction(unsigned char *a, unsigned char *b, unsigned char *c, unsigned int size) {

	//borrow represents the value 1 or 0 for the current index
	//indecating if the current index has been borrowed from by
	//the previous index, borrow = 1 true, 0 false
	unsigned char borrow = 0x00; //value is 0 or 1

	//loop through array a size and subtract a - b,
	//a is guaranted to be greater than b in 
	//barrett reduction
	unsigned int i = 0;
	while(i < size) {
		//check current value of a to make sure that it is
		//not 0 when the previous index has borrowed
		if (a[i] == 0 && borrow == 1) {
			//borrow from next sequential index with
			//0x100 and subtract 0x01 for the
			//previous borrow which is = 0xff
			c[i] = 0xff - b[i];
			//turn on borrow for next index
			borrow = 0x01;
			i++;
			continue;
		}

		//calculate current value of a along with if the 
		//previous index has borrowed
		a[i] = a[i] - borrow;
	
		//calculate the value of a - b only when a - b >= 0
		//borrow has already been accounted for
		if (a[i] >= b[i]) {
			c[i] = a[i] - b[i];
			borrow = 0x00;
		}

		//a - b !> 0, borrow from next sequential index by 
		//taking the value 0x100 and adding to a[i] and 
		//subtracting b[i] which will give a value between
		//{0x01...0xff} and turn on borrow for next index
		else {
			c[i] = 0x100 + a[i] - b[i];
			borrow = 0x01;
		}
		i++;
	}
	return;
}

char checkbit(unsigned char *exponent, unsigned int index_of_bit) {

	unsigned char bit;
	//get the characters index of which the bit is located in by 
	//taking index_of_bit which is the size of the bits left to 
	//check and divide by 8 giving the location index of the
	//current bit to be checked
	unsigned int quotient = (index_of_bit >> 0x03); // index_of_bit / 8

	//find the bit within the index previously found by finding the 
	//remainder of 8 % index of bit, this will locate the exact
	//bit to be checked
	unsigned int remainder = index_of_bit & (0x07); // index_of_bit % 8
	
	//mask is equivelent to 1 in order to compare a single bit with a
	//the current bit to be checked
	unsigned char mask = 0x01; // use single bit to mask with selected bit
	
	//use the remainder by knowing the index of the character and 
	//the remainder allows the bit to be shifted to the position of
	//the current bit to be checked
	mask = mask << remainder; // shift single bit to bit_in_index position
	
	//bit is now located at index_of_bit character index of array
	//and bit location bit_in_index in group of 8 bits at index
	bit = (exponent[quotient] & mask); // & to see if single bit is on or off

	//shift bit back to the 1 position to represent value 1 or 0
	bit = bit >> remainder; //shift bit back to value of one or zero
	
	return bit;
}

//b is expected to be completely zero before shift
void bit_shift(unsigned char *a, unsigned char *b, unsigned int k_val, unsigned int size_of_a) {

        //expected that k will be equivlent to some power of 2
        //represents the division of (4^k) which is = (2^(2*k))
        unsigned int shift = k_val * 2;

        //quotient represents groups of 8 bits that equal 0 as in >> 8 in single char
        //leaving it to be the value of 0x00
        unsigned int quotient = shift >> 0x03; // k / 8 as integer

	//in case that the shift is greater than the actual value of the
	//number being shifted
	if(quotient > size_of_a) {
		return;
	}

        //printf("quotient = %d\n", quotient);
        //remainder will find final char index shift value = {0...7}
        //the specific bits to be shifted in the last group which is not greater than 7
        unsigned int remainder = shift & 0x07; // k % 8 
        //printf("remainder = %d\n", remainder);
        //move a to b by shifting the characters an index of quotient amount
        //and then use the remainder to shift the final index to correct 
        //position
        unsigned int constant = (size_of_a - quotient);

        unsigned int j = 0;
        while (j < constant) {
                b[j] = a[quotient + j] >> remainder;
                unsigned char cpy_bits = a[quotient + j + 1] << (8 - remainder);
                b[j] = b[j] | cpy_bits;
                j++;
        }

        return;
}

__global__ void get_products(unsigned char a[], unsigned char b[], unsigned int accumulator[], unsigned int n) {

        int multiplier = 0;
        unsigned int index = (blockIdx.x * blockDim.x) + threadIdx.x;
        unsigned int multiplicand = index;
        unsigned int product = 0;

        while(multiplier < n) {
                product = (unsigned int) a[multiplier] * b[multiplicand];
                atomicAdd(&accumulator[multiplier + index], product<<24>>24);
                atomicAdd(&accumulator[multiplier + index + 1], product>>8);
                multiplier++;
        }

	return;
}


__global__ void get_square(unsigned char a[], unsigned int accumulator[], unsigned int n) {

        int multiplier = 0;
        unsigned int index = (blockIdx.x * blockDim.x) + threadIdx.x;
        unsigned int multiplicand = index;
        unsigned int product = 0;

        while(multiplier < n) {
                product = (unsigned int) a[multiplier] * a[multiplicand];
                atomicAdd(&accumulator[multiplier + index], product<<24>>24);
                atomicAdd(&accumulator[multiplier + index + 1], product>>8);
                multiplier++;
        }
	return;
}

	.text
	.file	"GPU_RSA.hip"
	.section	.rodata.cst16,"aM",@progbits,16
	.p2align	4, 0x0                          # -- Begin function main
.LCPI0_0:
	.byte	207                             # 0xcf
	.byte	130                             # 0x82
	.byte	105                             # 0x69
	.byte	87                              # 0x57
	.byte	77                              # 0x4d
	.byte	231                             # 0xe7
	.byte	130                             # 0x82
	.byte	26                              # 0x1a
	.byte	228                             # 0xe4
	.byte	32                              # 0x20
	.byte	20                              # 0x14
	.byte	71                              # 0x47
	.byte	57                              # 0x39
	.byte	82                              # 0x52
	.byte	85                              # 0x55
	.byte	40                              # 0x28
.LCPI0_1:
	.byte	237                             # 0xed
	.byte	63                              # 0x3f
	.byte	164                             # 0xa4
	.byte	97                              # 0x61
	.byte	211                             # 0xd3
	.byte	244                             # 0xf4
	.byte	242                             # 0xf2
	.byte	52                              # 0x34
	.byte	106                             # 0x6a
	.byte	84                              # 0x54
	.byte	209                             # 0xd1
	.byte	21                              # 0x15
	.byte	125                             # 0x7d
	.byte	103                             # 0x67
	.byte	11                              # 0xb
	.byte	199                             # 0xc7
.LCPI0_2:
	.byte	140                             # 0x8c
	.byte	254                             # 0xfe
	.byte	27                              # 0x1b
	.byte	104                             # 0x68
	.byte	68                              # 0x44
	.byte	7                               # 0x7
	.byte	38                              # 0x26
	.byte	153                             # 0x99
	.byte	11                              # 0xb
	.byte	77                              # 0x4d
	.byte	199                             # 0xc7
	.byte	63                              # 0x3f
	.byte	82                              # 0x52
	.byte	144                             # 0x90
	.byte	2                               # 0x2
	.byte	104                             # 0x68
.LCPI0_3:
	.byte	61                              # 0x3d
	.byte	131                             # 0x83
	.byte	29                              # 0x1d
	.byte	121                             # 0x79
	.byte	122                             # 0x7a
	.byte	63                              # 0x3f
	.byte	54                              # 0x36
	.byte	243                             # 0xf3
	.byte	65                              # 0x41
	.byte	139                             # 0x8b
	.byte	124                             # 0x7c
	.byte	223                             # 0xdf
	.byte	100                             # 0x64
	.byte	172                             # 0xac
	.byte	116                             # 0x74
	.byte	124                             # 0x7c
.LCPI0_4:
	.byte	8                               # 0x8
	.byte	219                             # 0xdb
	.byte	160                             # 0xa0
	.byte	111                             # 0x6f
	.byte	16                              # 0x10
	.byte	113                             # 0x71
	.byte	19                              # 0x13
	.byte	134                             # 0x86
	.byte	175                             # 0xaf
	.byte	184                             # 0xb8
	.byte	113                             # 0x71
	.byte	248                             # 0xf8
	.byte	240                             # 0xf0
	.byte	69                              # 0x45
	.byte	167                             # 0xa7
	.byte	148                             # 0x94
.LCPI0_5:
	.byte	179                             # 0xb3
	.byte	107                             # 0x6b
	.byte	30                              # 0x1e
	.byte	255                             # 0xff
	.byte	142                             # 0x8e
	.byte	19                              # 0x13
	.byte	174                             # 0xae
	.byte	194                             # 0xc2
	.byte	89                              # 0x59
	.byte	86                              # 0x56
	.byte	211                             # 0xd3
	.byte	13                              # 0xd
	.byte	32                              # 0x20
	.byte	98                              # 0x62
	.byte	33                              # 0x21
	.byte	48                              # 0x30
.LCPI0_6:
	.byte	29                              # 0x1d
	.byte	107                             # 0x6b
	.byte	94                              # 0x5e
	.byte	12                              # 0xc
	.byte	0                               # 0x0
	.byte	53                              # 0x35
	.byte	174                             # 0xae
	.byte	189                             # 0xbd
	.byte	165                             # 0xa5
	.byte	194                             # 0xc2
	.byte	37                              # 0x25
	.byte	152                             # 0x98
	.byte	231                             # 0xe7
	.byte	87                              # 0x57
	.byte	137                             # 0x89
	.byte	12                              # 0xc
.LCPI0_7:
	.byte	18                              # 0x12
	.byte	249                             # 0xf9
	.byte	51                              # 0x33
	.byte	61                              # 0x3d
	.byte	10                              # 0xa
	.byte	172                             # 0xac
	.byte	81                              # 0x51
	.byte	216                             # 0xd8
	.byte	92                              # 0x5c
	.byte	64                              # 0x40
	.byte	155                             # 0x9b
	.byte	250                             # 0xfa
	.byte	249                             # 0xf9
	.byte	188                             # 0xbc
	.byte	3                               # 0x3
	.byte	230                             # 0xe6
.LCPI0_8:
	.byte	127                             # 0x7f
	.byte	157                             # 0x9d
	.byte	233                             # 0xe9
	.byte	64                              # 0x40
	.byte	87                              # 0x57
	.byte	2                               # 0x2
	.byte	110                             # 0x6e
	.byte	147                             # 0x93
	.byte	43                              # 0x2b
	.byte	180                             # 0xb4
	.byte	227                             # 0xe3
	.byte	253                             # 0xfd
	.byte	186                             # 0xba
	.byte	12                              # 0xc
	.byte	205                             # 0xcd
	.byte	120                             # 0x78
.LCPI0_9:
	.byte	125                             # 0x7d
	.byte	174                             # 0xae
	.byte	141                             # 0x8d
	.byte	128                             # 0x80
	.byte	255                             # 0xff
	.byte	102                             # 0x66
	.byte	51                              # 0x33
	.byte	11                              # 0xb
	.byte	40                              # 0x28
	.byte	76                              # 0x4c
	.byte	147                             # 0x93
	.byte	48                              # 0x30
	.byte	2                               # 0x2
	.byte	146                             # 0x92
	.byte	160                             # 0xa0
	.byte	124                             # 0x7c
.LCPI0_10:
	.byte	241                             # 0xf1
	.byte	12                              # 0xc
	.byte	10                              # 0xa
	.byte	94                              # 0x5e
	.byte	242                             # 0xf2
	.byte	154                             # 0x9a
	.byte	143                             # 0x8f
	.byte	23                              # 0x17
	.byte	76                              # 0x4c
	.byte	130                             # 0x82
	.byte	37                              # 0x25
	.byte	229                             # 0xe5
	.byte	152                             # 0x98
	.byte	69                              # 0x45
	.byte	77                              # 0x4d
	.byte	199                             # 0xc7
.LCPI0_11:
	.byte	217                             # 0xd9
	.byte	83                              # 0x53
	.byte	94                              # 0x5e
	.byte	90                              # 0x5a
	.byte	110                             # 0x6e
	.byte	55                              # 0x37
	.byte	67                              # 0x43
	.byte	41                              # 0x29
	.byte	136                             # 0x88
	.byte	203                             # 0xcb
	.byte	233                             # 0xe9
	.byte	49                              # 0x31
	.byte	47                              # 0x2f
	.byte	215                             # 0xd7
	.byte	6                               # 0x6
	.byte	251                             # 0xfb
.LCPI0_12:
	.byte	241                             # 0xf1
	.byte	56                              # 0x38
	.byte	223                             # 0xdf
	.byte	196                             # 0xc4
	.byte	218                             # 0xda
	.byte	124                             # 0x7c
	.byte	9                               # 0x9
	.byte	92                              # 0x5c
	.byte	249                             # 0xf9
	.byte	43                              # 0x2b
	.byte	129                             # 0x81
	.byte	48                              # 0x30
	.byte	233                             # 0xe9
	.byte	41                              # 0x29
	.byte	205                             # 0xcd
	.byte	69                              # 0x45
.LCPI0_13:
	.byte	238                             # 0xee
	.byte	255                             # 0xff
	.byte	91                              # 0x5b
	.byte	60                              # 0x3c
	.byte	35                              # 0x23
	.byte	109                             # 0x6d
	.byte	185                             # 0xb9
	.byte	161                             # 0xa1
	.byte	137                             # 0x89
	.byte	63                              # 0x3f
	.byte	195                             # 0xc3
	.byte	158                             # 0x9e
	.byte	161                             # 0xa1
	.byte	48                              # 0x30
	.byte	152                             # 0x98
	.byte	248                             # 0xf8
.LCPI0_14:
	.byte	200                             # 0xc8
	.byte	74                              # 0x4a
	.byte	190                             # 0xbe
	.byte	198                             # 0xc6
	.byte	73                              # 0x49
	.byte	247                             # 0xf7
	.byte	179                             # 0xb3
	.byte	255                             # 0xff
	.byte	9                               # 0x9
	.byte	59                              # 0x3b
	.byte	148                             # 0x94
	.byte	157                             # 0x9d
	.byte	47                              # 0x2f
	.byte	92                              # 0x5c
	.byte	104                             # 0x68
	.byte	225                             # 0xe1
.LCPI0_15:
	.byte	6                               # 0x6
	.byte	241                             # 0xf1
	.byte	51                              # 0x33
	.byte	235                             # 0xeb
	.byte	197                             # 0xc5
	.byte	136                             # 0x88
	.byte	165                             # 0xa5
	.byte	28                              # 0x1c
	.byte	222                             # 0xde
	.byte	44                              # 0x2c
	.byte	100                             # 0x64
	.byte	173                             # 0xad
	.byte	92                              # 0x5c
	.byte	201                             # 0xc9
	.byte	235                             # 0xeb
	.byte	28                              # 0x1c
.LCPI0_16:
	.byte	145                             # 0x91
	.byte	10                              # 0xa
	.byte	179                             # 0xb3
	.byte	102                             # 0x66
	.byte	189                             # 0xbd
	.byte	111                             # 0x6f
	.byte	24                              # 0x18
	.byte	222                             # 0xde
	.byte	213                             # 0xd5
	.byte	1                               # 0x1
	.byte	97                              # 0x61
	.byte	54                              # 0x36
	.byte	149                             # 0x95
	.byte	109                             # 0x6d
	.byte	221                             # 0xdd
	.byte	51                              # 0x33
.LCPI0_17:
	.byte	219                             # 0xdb
	.byte	38                              # 0x26
	.byte	3                               # 0x3
	.byte	14                              # 0xe
	.byte	104                             # 0x68
	.byte	84                              # 0x54
	.byte	115                             # 0x73
	.byte	160                             # 0xa0
	.byte	224                             # 0xe0
	.byte	110                             # 0x6e
	.byte	112                             # 0x70
	.byte	116                             # 0x74
	.byte	37                              # 0x25
	.byte	139                             # 0x8b
	.byte	43                              # 0x2b
	.byte	251                             # 0xfb
.LCPI0_18:
	.byte	158                             # 0x9e
	.byte	60                              # 0x3c
	.byte	52                              # 0x34
	.byte	46                              # 0x2e
	.byte	69                              # 0x45
	.byte	16                              # 0x10
	.byte	16                              # 0x10
	.byte	108                             # 0x6c
	.byte	251                             # 0xfb
	.byte	183                             # 0xb7
	.byte	155                             # 0x9b
	.byte	200                             # 0xc8
	.byte	207                             # 0xcf
	.byte	113                             # 0x71
	.byte	217                             # 0xd9
	.byte	150                             # 0x96
.LCPI0_19:
	.byte	183                             # 0xb7
	.byte	187                             # 0xbb
	.byte	95                              # 0x5f
	.byte	25                              # 0x19
	.byte	118                             # 0x76
	.byte	54                              # 0x36
	.byte	73                              # 0x49
	.byte	106                             # 0x6a
	.byte	179                             # 0xb3
	.byte	131                             # 0x83
	.byte	195                             # 0xc3
	.byte	89                              # 0x59
	.byte	46                              # 0x2e
	.byte	98                              # 0x62
	.byte	135                             # 0x87
	.byte	162                             # 0xa2
.LCPI0_20:
	.byte	90                              # 0x5a
	.byte	47                              # 0x2f
	.byte	96                              # 0x60
	.byte	117                             # 0x75
	.byte	1                               # 0x1
	.byte	240                             # 0xf0
	.byte	63                              # 0x3f
	.byte	219                             # 0xdb
	.byte	90                              # 0x5a
	.byte	112                             # 0x70
	.byte	31                              # 0x1f
	.byte	68                              # 0x44
	.byte	106                             # 0x6a
	.byte	156                             # 0x9c
	.byte	119                             # 0x77
	.byte	99                              # 0x63
.LCPI0_21:
	.byte	186                             # 0xba
	.byte	203                             # 0xcb
	.byte	205                             # 0xcd
	.byte	31                              # 0x1f
	.byte	153                             # 0x99
	.byte	112                             # 0x70
	.byte	137                             # 0x89
	.byte	148                             # 0x94
	.byte	49                              # 0x31
	.byte	2                               # 0x2
	.byte	10                              # 0xa
	.byte	50                              # 0x32
	.byte	150                             # 0x96
	.byte	101                             # 0x65
	.byte	33                              # 0x21
	.byte	33                              # 0x21
.LCPI0_22:
	.byte	89                              # 0x59
	.byte	85                              # 0x55
	.byte	138                             # 0x8a
	.byte	208                             # 0xd0
	.byte	122                             # 0x7a
	.byte	28                              # 0x1c
	.byte	210                             # 0xd2
	.byte	102                             # 0x66
	.byte	72                              # 0x48
	.byte	149                             # 0x95
	.byte	8                               # 0x8
	.byte	211                             # 0xd3
	.byte	107                             # 0x6b
	.byte	231                             # 0xe7
	.byte	156                             # 0x9c
	.byte	185                             # 0xb9
.LCPI0_23:
	.byte	150                             # 0x96
	.byte	32                              # 0x20
	.byte	32                              # 0x20
	.byte	138                             # 0x8a
	.byte	229                             # 0xe5
	.byte	77                              # 0x4d
	.byte	62                              # 0x3e
	.byte	83                              # 0x53
	.byte	75                              # 0x4b
	.byte	216                             # 0xd8
	.byte	33                              # 0x21
	.byte	4                               # 0x4
	.byte	129                             # 0x81
	.byte	125                             # 0x7d
	.byte	41                              # 0x29
	.byte	56                              # 0x38
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0
.LCPI0_24:
	.long	0x447a0000                      # float 1000
	.text
	.globl	main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	subq	$168, %rsp
	.cfi_def_cfa_offset 224
	.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	$640, %edi                      # imm = 0x280
	movl	$1, %esi
	callq	calloc
	movq	%rax, %r14
	movaps	.LCPI0_0(%rip), %xmm0           # xmm0 = [207,130,105,87,77,231,130,26,228,32,20,71,57,82,85,40]
	movups	%xmm0, (%rax)
	movaps	.LCPI0_1(%rip), %xmm0           # xmm0 = [237,63,164,97,211,244,242,52,106,84,209,21,125,103,11,199]
	movups	%xmm0, 16(%rax)
	movaps	.LCPI0_2(%rip), %xmm0           # xmm0 = [140,254,27,104,68,7,38,153,11,77,199,63,82,144,2,104]
	movups	%xmm0, 32(%rax)
	movaps	.LCPI0_3(%rip), %xmm0           # xmm0 = [61,131,29,121,122,63,54,243,65,139,124,223,100,172,116,124]
	movups	%xmm0, 48(%rax)
	movaps	.LCPI0_4(%rip), %xmm0           # xmm0 = [8,219,160,111,16,113,19,134,175,184,113,248,240,69,167,148]
	movups	%xmm0, 64(%rax)
	movaps	.LCPI0_5(%rip), %xmm0           # xmm0 = [179,107,30,255,142,19,174,194,89,86,211,13,32,98,33,48]
	movups	%xmm0, 80(%rax)
	movaps	.LCPI0_6(%rip), %xmm0           # xmm0 = [29,107,94,12,0,53,174,189,165,194,37,152,231,87,137,12]
	movups	%xmm0, 96(%rax)
	movaps	.LCPI0_7(%rip), %xmm0           # xmm0 = [18,249,51,61,10,172,81,216,92,64,155,250,249,188,3,230]
	movups	%xmm0, 112(%rax)
	leaq	112(%rsp), %rbx
	movl	$640, %esi                      # imm = 0x280
	movq	%rbx, %rdi
	callq	hipMalloc
	movq	(%rbx), %rdi
	movl	$640, %edx                      # imm = 0x280
	movq	%r14, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movl	$256, %edi                      # imm = 0x100
	movl	$1, %esi
	callq	calloc
	movq	%rax, %rbx
	movaps	.LCPI0_8(%rip), %xmm0           # xmm0 = [127,157,233,64,87,2,110,147,43,180,227,253,186,12,205,120]
	movups	%xmm0, (%rax)
	movaps	.LCPI0_9(%rip), %xmm0           # xmm0 = [125,174,141,128,255,102,51,11,40,76,147,48,2,146,160,124]
	movups	%xmm0, 16(%rax)
	movaps	.LCPI0_10(%rip), %xmm0          # xmm0 = [241,12,10,94,242,154,143,23,76,130,37,229,152,69,77,199]
	movups	%xmm0, 32(%rax)
	movaps	.LCPI0_11(%rip), %xmm0          # xmm0 = [217,83,94,90,110,55,67,41,136,203,233,49,47,215,6,251]
	movups	%xmm0, 48(%rax)
	movaps	.LCPI0_12(%rip), %xmm0          # xmm0 = [241,56,223,196,218,124,9,92,249,43,129,48,233,41,205,69]
	movups	%xmm0, 64(%rax)
	movaps	.LCPI0_13(%rip), %xmm0          # xmm0 = [238,255,91,60,35,109,185,161,137,63,195,158,161,48,152,248]
	movups	%xmm0, 80(%rax)
	movaps	.LCPI0_14(%rip), %xmm0          # xmm0 = [200,74,190,198,73,247,179,255,9,59,148,157,47,92,104,225]
	movups	%xmm0, 96(%rax)
	movaps	.LCPI0_15(%rip), %xmm0          # xmm0 = [6,241,51,235,197,136,165,28,222,44,100,173,92,201,235,28]
	movups	%xmm0, 112(%rax)
	movb	$1, %bpl
	movb	%bpl, 128(%rax)
	leaq	104(%rsp), %r15
	movl	$256, %esi                      # imm = 0x100
	movq	%r15, %rdi
	callq	hipMalloc
	movq	(%r15), %rdi
	movl	$256, %edx                      # imm = 0x100
	movq	%rbx, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movl	$128, %edi
	movl	$1, %esi
	callq	calloc
	movq	%rax, %rbx
	movl	$1819043176, (%rax)             # imm = 0x6C6C6568
	movb	$111, 4(%rax)
	movq	%rsp, %r15
	movl	$128, %esi
	movq	%r15, %rdi
	callq	hipMalloc
	movq	(%r15), %rdi
	movl	$128, %edx
	movq	%rbx, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movl	$3, %edi
	callq	malloc
	movq	%rax, %r12
	movw	$1, (%rax)
	movb	%bpl, 2(%rax)
	leaq	96(%rsp), %r15
	movl	$128, %esi
	movq	%r15, %rdi
	callq	hipMalloc
	movq	(%r15), %rdi
	movl	$128, %edx
	xorl	%esi, %esi
	callq	hipMemset
	leaq	88(%rsp), %rdi
	movl	$128, %esi
	callq	hipMalloc
	movl	$128, %edi
	movl	$1, %esi
	callq	calloc
	movq	%rax, 152(%rsp)                 # 8-byte Spill
	movl	$257, %edi                      # imm = 0x101
	movl	$1, %esi
	callq	calloc
	movq	%rax, 24(%rsp)                  # 8-byte Spill
	leaq	80(%rsp), %r15
	movl	$257, %esi                      # imm = 0x101
	movq	%r15, %rdi
	callq	hipMalloc
	movq	(%r15), %rdi
	movl	$257, %edx                      # imm = 0x101
	xorl	%esi, %esi
	callq	hipMemset
	movl	$384, %edi                      # imm = 0x180
	movl	$1, %esi
	callq	calloc
	movq	%rax, 144(%rsp)                 # 8-byte Spill
	leaq	72(%rsp), %rdi
	movl	$384, %esi                      # imm = 0x180
	callq	hipMalloc
	movl	$128, %edi
	movl	$1, %esi
	callq	calloc
	movq	%rax, 136(%rsp)                 # 8-byte Spill
	leaq	64(%rsp), %r15
	movl	$128, %esi
	movq	%r15, %rdi
	callq	hipMalloc
	movq	(%r15), %rdi
	movl	$128, %edx
	xorl	%esi, %esi
	callq	hipMemset
	movl	$256, %edi                      # imm = 0x100
	movl	$1, %esi
	callq	calloc
	movq	%rax, 32(%rsp)                  # 8-byte Spill
	leaq	56(%rsp), %r15
	movl	$256, %esi                      # imm = 0x100
	movq	%r15, %rdi
	callq	hipMalloc
	movq	(%r15), %rdi
	movl	$256, %edx                      # imm = 0x100
	xorl	%esi, %esi
	callq	hipMemset
	movl	$129, %edi
	movl	$1, %esi
	callq	calloc
	movq	%rax, %r15
	movl	$129, %edi
	movl	$1, %esi
	callq	calloc
	movq	%rax, 16(%rsp)                  # 8-byte Spill
	movl	$512, %edi                      # imm = 0x200
	movl	$4, %esi
	callq	calloc
	movq	%rax, %rbp
	leaq	48(%rsp), %r13
	movl	$2048, %esi                     # imm = 0x800
	movq	%r13, %rdi
	callq	hipMalloc
	movq	(%r13), %rdi
	movl	$2048, %edx                     # imm = 0x800
	xorl	%esi, %esi
	callq	hipMemset
	leaq	40(%rsp), %rdi
	callq	hipEventCreate
	testl	%eax, %eax
	je	.LBB0_2
# %bb.1:
	movl	$.Lstr.2, %edi
	callq	puts@PLT
.LBB0_2:
	leaq	8(%rsp), %rdi
	callq	hipEventCreate
	testl	%eax, %eax
	je	.LBB0_4
# %bb.3:
	movl	$.Lstr.2, %edi
	callq	puts@PLT
.LBB0_4:
	movq	40(%rsp), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	(%rsp), %rdi
	movq	96(%rsp), %rdx
	movq	88(%rsp), %rcx
	subq	$8, %rsp
	.cfi_adjust_cfa_offset 8
	movq	%r12, %rsi
	movq	160(%rsp), %r13                 # 8-byte Reload
	movq	%r13, %r8
	movq	32(%rsp), %r9                   # 8-byte Reload
	pushq	%r14
	.cfi_adjust_cfa_offset 8
	movq	%rbp, 144(%rsp)                 # 8-byte Spill
	pushq	%rbp
	.cfi_adjust_cfa_offset 8
	pushq	72(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	144(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	144(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	$3
	.cfi_adjust_cfa_offset 8
	pushq	72(%rsp)                        # 8-byte Folded Reload
	.cfi_adjust_cfa_offset 8
	pushq	%r15
	.cfi_adjust_cfa_offset 8
	pushq	128(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	112(%rsp)                       # 8-byte Folded Reload
	.cfi_adjust_cfa_offset 8
	pushq	152(%rsp)
	.cfi_adjust_cfa_offset 8
	movq	%r15, 216(%rsp)                 # 8-byte Spill
	movq	232(%rsp), %r15                 # 8-byte Reload
	pushq	%r15
	.cfi_adjust_cfa_offset 8
	pushq	176(%rsp)
	.cfi_adjust_cfa_offset 8
	movq	256(%rsp), %rbp                 # 8-byte Reload
	pushq	%rbp
	.cfi_adjust_cfa_offset 8
	pushq	200(%rsp)
	.cfi_adjust_cfa_offset 8
	callq	_Z14exponentiationPhS_S_S_S_S_S_S_S_S_S_S_S_S_S_jS_S_PjS0_S_
	addq	$128, %rsp
	.cfi_adjust_cfa_offset -128
	movl	$128, %edi
	callq	malloc
	movq	%rax, %r12
	movaps	.LCPI0_16(%rip), %xmm0          # xmm0 = [145,10,179,102,189,111,24,222,213,1,97,54,149,109,221,51]
	movups	%xmm0, (%rax)
	movaps	.LCPI0_17(%rip), %xmm0          # xmm0 = [219,38,3,14,104,84,115,160,224,110,112,116,37,139,43,251]
	movups	%xmm0, 16(%rax)
	movaps	.LCPI0_18(%rip), %xmm0          # xmm0 = [158,60,52,46,69,16,16,108,251,183,155,200,207,113,217,150]
	movups	%xmm0, 32(%rax)
	movaps	.LCPI0_19(%rip), %xmm0          # xmm0 = [183,187,95,25,118,54,73,106,179,131,195,89,46,98,135,162]
	movups	%xmm0, 48(%rax)
	movaps	.LCPI0_20(%rip), %xmm0          # xmm0 = [90,47,96,117,1,240,63,219,90,112,31,68,106,156,119,99]
	movups	%xmm0, 64(%rax)
	movaps	.LCPI0_21(%rip), %xmm0          # xmm0 = [186,203,205,31,153,112,137,148,49,2,10,50,150,101,33,33]
	movups	%xmm0, 80(%rax)
	movaps	.LCPI0_22(%rip), %xmm0          # xmm0 = [89,85,138,208,122,28,210,102,72,149,8,211,107,231,156,185]
	movups	%xmm0, 96(%rax)
	movaps	.LCPI0_23(%rip), %xmm0          # xmm0 = [150,32,32,138,229,77,62,83,75,216,33,4,129,125,41,56]
	movups	%xmm0, 112(%rax)
	xorps	%xmm0, %xmm0
	movups	%xmm0, (%rbx)
	movups	%xmm0, 16(%rbx)
	movups	%xmm0, 32(%rbx)
	movups	%xmm0, 48(%rbx)
	movups	%xmm0, 64(%rbx)
	movups	%xmm0, 80(%rbx)
	movups	%xmm0, 96(%rbx)
	movups	%xmm0, 112(%rbx)
	movq	(%rsp), %rdi
	movl	$128, %edx
	xorl	%esi, %esi
	callq	hipMemset
	movq	96(%rsp), %rdi
	movq	(%rsp), %rdx
	movq	88(%rsp), %rcx
	subq	$8, %rsp
	.cfi_adjust_cfa_offset 8
	movq	%r12, %rsi
	movq	%r13, %r8
	movq	32(%rsp), %r9                   # 8-byte Reload
	pushq	%r14
	.cfi_adjust_cfa_offset 8
	pushq	144(%rsp)                       # 8-byte Folded Reload
	.cfi_adjust_cfa_offset 8
	pushq	72(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	144(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	144(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	$128
	.cfi_adjust_cfa_offset 8
	pushq	72(%rsp)                        # 8-byte Folded Reload
	.cfi_adjust_cfa_offset 8
	pushq	184(%rsp)                       # 8-byte Folded Reload
	.cfi_adjust_cfa_offset 8
	pushq	128(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	112(%rsp)                       # 8-byte Folded Reload
	.cfi_adjust_cfa_offset 8
	pushq	152(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	%r15
	.cfi_adjust_cfa_offset 8
	pushq	176(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	%rbp
	.cfi_adjust_cfa_offset 8
	pushq	200(%rsp)
	.cfi_adjust_cfa_offset 8
	callq	_Z14exponentiationPhS_S_S_S_S_S_S_S_S_S_S_S_S_S_jS_S_PjS0_S_
	addq	$128, %rsp
	.cfi_adjust_cfa_offset -128
	movq	8(%rsp), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	8(%rsp), %rdi
	callq	hipEventSynchronize
	testl	%eax, %eax
	je	.LBB0_6
# %bb.5:
	movl	$.Lstr.2, %edi
	callq	puts@PLT
.LBB0_6:
	leaq	164(%rsp), %r14
	movl	$0, (%r14)
	movq	40(%rsp), %rsi
	movq	8(%rsp), %rdx
	movq	%r14, %rdi
	callq	hipEventElapsedTime
	movss	(%r14), %xmm0                   # xmm0 = mem[0],zero,zero,zero
	divss	.LCPI0_24(%rip), %xmm0
	cvtss2sd	%xmm0, %xmm0
	movl	$.L.str.1, %edi
	movb	$1, %al
	callq	printf
	movq	(%rsp), %rsi
	movl	$128, %edx
	movq	%rbx, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	xorl	%r14d, %r14d
.LBB0_7:                                # =>This Inner Loop Header: Depth=1
	movzbl	(%rbx,%r14), %edx
	movl	$.L.str.2, %edi
	movl	%r14d, %esi
	xorl	%eax, %eax
	callq	printf
	incq	%r14
	cmpq	$128, %r14
	jne	.LBB0_7
# %bb.8:
	xorl	%eax, %eax
	addq	$168, %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_end0:
	.size	main, .Lfunc_end0-main
	.cfi_endproc
                                        # -- End function
	.globl	_Z14exponentiationPhS_S_S_S_S_S_S_S_S_S_S_S_S_S_jS_S_PjS0_S_ # -- Begin function _Z14exponentiationPhS_S_S_S_S_S_S_S_S_S_S_S_S_S_jS_S_PjS0_S_
	.type	_Z14exponentiationPhS_S_S_S_S_S_S_S_S_S_S_S_S_S_jS_S_PjS0_S_,@function
_Z14exponentiationPhS_S_S_S_S_S_S_S_S_S_S_S_S_S_jS_S_PjS0_S_: # @_Z14exponentiationPhS_S_S_S_S_S_S_S_S_S_S_S_S_S_jS_S_PjS0_S_
	.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	$216, %rsp
	.cfi_def_cfa_offset 272
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movq	%r9, %r12
	movq	%r8, %r13
	movq	%rcx, 104(%rsp)                 # 8-byte Spill
	movq	%rdx, 200(%rsp)                 # 8-byte Spill
	movq	%rdi, 96(%rsp)                  # 8-byte Spill
	movq	384(%rsp), %rax
	movq	%rax, 192(%rsp)                 # 8-byte Spill
	movq	376(%rsp), %r14
	movq	368(%rsp), %rbx
	movq	360(%rsp), %rax
	movq	%rax, 184(%rsp)                 # 8-byte Spill
	movq	352(%rsp), %rax
	movq	%rax, 176(%rsp)                 # 8-byte Spill
	movq	328(%rsp), %rax
	movq	%rax, 168(%rsp)                 # 8-byte Spill
	movq	320(%rsp), %rax
	movq	%rax, 160(%rsp)                 # 8-byte Spill
	movl	344(%rsp), %ecx
	leal	(,%rcx,8), %eax
	decl	%ecx
	movq	%rsi, 208(%rsp)                 # 8-byte Spill
	movzbl	(%rsi,%rcx), %edx
	movq	312(%rsp), %rcx
	movq	%rcx, 152(%rsp)                 # 8-byte Spill
	movq	304(%rsp), %rcx
	movq	%rcx, 144(%rsp)                 # 8-byte Spill
	movq	296(%rsp), %rcx
	movq	%rcx, 136(%rsp)                 # 8-byte Spill
	movq	288(%rsp), %rcx
	movq	%rcx, 128(%rsp)                 # 8-byte Spill
	movq	280(%rsp), %rcx
	movq	%rcx, 120(%rsp)                 # 8-byte Spill
	movq	272(%rsp), %rbp
	xorl	%r15d, %r15d
.LBB1_1:                                # =>This Inner Loop Header: Depth=1
	movl	%edx, %esi
	movl	%r15d, %ecx
	shll	%cl, %esi
	testb	%sil, %sil
	js	.LBB1_2
# %bb.3:                                #   in Loop: Header=BB1_1 Depth=1
	incl	%r15d
	cmpl	$8, %r15d
	jne	.LBB1_1
# %bb.4:
	movl	$-1, %r15d
	jmp	.LBB1_5
.LBB1_2:
	notl	%r15d
.LBB1_5:                                # %.loopexit
	addl	%eax, %r15d
	movl	$128, %edx
	movq	104(%rsp), %rdi                 # 8-byte Reload
	movq	96(%rsp), %rsi                  # 8-byte Reload
	movl	$3, %ecx
	callq	hipMemcpy
	decl	%r15d
	js	.LBB1_18
# %bb.6:                                # %.lr.ph.preheader
	movabsq	$4294967298, %rax               # imm = 0x100000002
	addq	$62, %rax
	movq	%rax, 112(%rsp)                 # 8-byte Spill
.LBB1_7:                                # %.lr.ph
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB1_10 Depth 2
                                        #     Child Loop BB1_15 Depth 2
	movabsq	$4294967298, %rdi               # imm = 0x100000002
	movl	$1, %esi
	movq	112(%rsp), %rdx                 # 8-byte Reload
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB1_9
# %bb.8:                                #   in Loop: Header=BB1_7 Depth=1
	movq	96(%rsp), %rdi                  # 8-byte Reload
	movq	%rbx, %rsi
	movl	$128, %edx
	callq	_Z25__device_stub__get_squarePhPjj
.LBB1_9:                                #   in Loop: Header=BB1_7 Depth=1
	movl	$1024, %edx                     # imm = 0x400
	movq	%r14, %rdi
	movq	%rbx, %rsi
	movl	$2, %ecx
	callq	hipMemcpy
	movl	(%r14), %ecx
	movl	$1, %eax
.LBB1_10:                               #   Parent Loop BB1_7 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movb	%cl, -1(%r12,%rax)
	movl	-4(%r14,%rax,4), %ecx
	shrl	$8, %ecx
	addl	(%r14,%rax,4), %ecx
	movl	%ecx, (%r14,%rax,4)
	incq	%rax
	cmpq	$257, %rax                      # imm = 0x101
	jne	.LBB1_10
# %bb.11:                               #   in Loop: Header=BB1_7 Depth=1
	movl	$2048, %edx                     # imm = 0x800
	movq	%r14, %rdi
	xorl	%esi, %esi
	callq	memset@PLT
	movl	$2048, %edx                     # imm = 0x800
	movq	%rbx, %rdi
	xorl	%esi, %esi
	callq	hipMemset
	movl	$256, %edx                      # imm = 0x100
	movq	%rbp, %rdi
	movq	%r12, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movq	192(%rsp), %rax                 # 8-byte Reload
	movq	%rax, 80(%rsp)
	movq	%r14, 72(%rsp)
	movq	%rbx, 64(%rsp)
	movq	168(%rsp), %rax                 # 8-byte Reload
	movq	%rax, 40(%rsp)
	movq	160(%rsp), %rax                 # 8-byte Reload
	movq	%rax, 32(%rsp)
	movq	152(%rsp), %rax                 # 8-byte Reload
	movq	%rax, 24(%rsp)
	movq	144(%rsp), %rax                 # 8-byte Reload
	movq	%rax, 16(%rsp)
	movq	136(%rsp), %rax                 # 8-byte Reload
	movq	%rax, 8(%rsp)
	movq	128(%rsp), %rax                 # 8-byte Reload
	movq	%rax, (%rsp)
	movq	%r12, %rdi
	movq	%rbp, %rsi
	movq	176(%rsp), %rdx                 # 8-byte Reload
	movq	184(%rsp), %rcx                 # 8-byte Reload
	movq	%r13, %r8
	movq	120(%rsp), %r9                  # 8-byte Reload
	callq	_Z17barrett_reductionPhS_S_S_S_S_S_S_S_S_S_S_S_jPjS0_S_
	movl	$128, %edx
	movq	96(%rsp), %rdi                  # 8-byte Reload
	movq	%r13, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movl	$256, %edx                      # imm = 0x100
	movq	%rbp, %rdi
	xorl	%esi, %esi
	callq	hipMemset
	movl	$256, %edx                      # imm = 0x100
	movq	%r12, %rdi
	xorl	%esi, %esi
	callq	memset@PLT
	xorps	%xmm0, %xmm0
	movups	%xmm0, (%r13)
	movups	%xmm0, 16(%r13)
	movups	%xmm0, 32(%r13)
	movups	%xmm0, 48(%r13)
	movups	%xmm0, 64(%r13)
	movups	%xmm0, 80(%r13)
	movups	%xmm0, 96(%r13)
	movups	%xmm0, 112(%r13)
	movl	%r15d, %eax
	movl	%r15d, %ecx
	andb	$7, %cl
	movl	$1, %edx
	shll	%cl, %edx
	shrl	$3, %eax
	movq	208(%rsp), %rsi                 # 8-byte Reload
	movzbl	(%rsi,%rax), %eax
	andl	%edx, %eax
	shrl	%cl, %eax
	cmpl	$1, %eax
	jne	.LBB1_17
# %bb.12:                               #   in Loop: Header=BB1_7 Depth=1
	movabsq	$4294967298, %rdi               # imm = 0x100000002
	movl	$1, %esi
	movq	112(%rsp), %rdx                 # 8-byte Reload
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB1_14
# %bb.13:                               #   in Loop: Header=BB1_7 Depth=1
	movq	96(%rsp), %rdi                  # 8-byte Reload
	movq	104(%rsp), %rsi                 # 8-byte Reload
	movq	%rbx, %rdx
	movl	$128, %ecx
	callq	_Z27__device_stub__get_productsPhS_Pjj
.LBB1_14:                               #   in Loop: Header=BB1_7 Depth=1
	movl	$1024, %edx                     # imm = 0x400
	movq	%r14, %rdi
	movq	%rbx, %rsi
	movl	$2, %ecx
	callq	hipMemcpy
	movl	(%r14), %ecx
	movl	$1, %eax
.LBB1_15:                               #   Parent Loop BB1_7 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movb	%cl, -1(%r12,%rax)
	movl	-4(%r14,%rax,4), %ecx
	shrl	$8, %ecx
	addl	(%r14,%rax,4), %ecx
	movl	%ecx, (%r14,%rax,4)
	incq	%rax
	cmpq	$257, %rax                      # imm = 0x101
	jne	.LBB1_15
# %bb.16:                               #   in Loop: Header=BB1_7 Depth=1
	movl	$2048, %edx                     # imm = 0x800
	movq	%r14, %rdi
	xorl	%esi, %esi
	callq	memset@PLT
	movl	$2048, %edx                     # imm = 0x800
	movq	%rbx, %rdi
	xorl	%esi, %esi
	callq	hipMemset
	movl	$256, %edx                      # imm = 0x100
	movq	%rbp, %rdi
	movq	%r12, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movq	192(%rsp), %rax                 # 8-byte Reload
	movq	%rax, 80(%rsp)
	movq	%r14, 72(%rsp)
	movq	%rbx, 64(%rsp)
	movq	168(%rsp), %rax                 # 8-byte Reload
	movq	%rax, 40(%rsp)
	movq	160(%rsp), %rax                 # 8-byte Reload
	movq	%rax, 32(%rsp)
	movq	152(%rsp), %rax                 # 8-byte Reload
	movq	%rax, 24(%rsp)
	movq	144(%rsp), %rax                 # 8-byte Reload
	movq	%rax, 16(%rsp)
	movq	136(%rsp), %rax                 # 8-byte Reload
	movq	%rax, 8(%rsp)
	movq	128(%rsp), %rax                 # 8-byte Reload
	movq	%rax, (%rsp)
	movq	%r12, %rdi
	movq	%rbp, %rsi
	movq	176(%rsp), %rdx                 # 8-byte Reload
	movq	184(%rsp), %rcx                 # 8-byte Reload
	movq	%r13, %r8
	movq	120(%rsp), %r9                  # 8-byte Reload
	callq	_Z17barrett_reductionPhS_S_S_S_S_S_S_S_S_S_S_S_jPjS0_S_
	movl	$128, %edx
	movq	96(%rsp), %rdi                  # 8-byte Reload
	movq	%r13, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movl	$256, %edx                      # imm = 0x100
	movq	%rbp, %rdi
	xorl	%esi, %esi
	callq	hipMemset
	movl	$256, %edx                      # imm = 0x100
	movq	%r12, %rdi
	xorl	%esi, %esi
	callq	memset@PLT
	xorps	%xmm0, %xmm0
	movups	%xmm0, (%r13)
	movups	%xmm0, 16(%r13)
	movups	%xmm0, 32(%r13)
	movups	%xmm0, 48(%r13)
	movups	%xmm0, 64(%r13)
	movups	%xmm0, 80(%r13)
	movups	%xmm0, 96(%r13)
	movups	%xmm0, 112(%r13)
.LBB1_17:                               #   in Loop: Header=BB1_7 Depth=1
	leal	-1(%r15), %eax
	testl	%r15d, %r15d
	movl	%eax, %r15d
	jg	.LBB1_7
.LBB1_18:                               # %._crit_edge
	movl	$128, %edx
	movq	200(%rsp), %rdi                 # 8-byte Reload
	movq	96(%rsp), %rsi                  # 8-byte Reload
	movl	$2, %ecx
	callq	hipMemcpy
	movl	$128, %edx
	movq	104(%rsp), %rdi                 # 8-byte Reload
	xorl	%esi, %esi
	addq	$216, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	jmp	hipMemset                       # TAILCALL
.Lfunc_end1:
	.size	_Z14exponentiationPhS_S_S_S_S_S_S_S_S_S_S_S_S_S_jS_S_PjS0_S_, .Lfunc_end1-_Z14exponentiationPhS_S_S_S_S_S_S_S_S_S_S_S_S_S_jS_S_PjS0_S_
	.cfi_endproc
                                        # -- End function
	.globl	_Z25__device_stub__get_squarePhPjj # -- Begin function _Z25__device_stub__get_squarePhPjj
	.type	_Z25__device_stub__get_squarePhPjj,@function
_Z25__device_stub__get_squarePhPjj:     # @_Z25__device_stub__get_squarePhPjj
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$112, %rsp
	.cfi_def_cfa_offset 160
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	40(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	32(%rsp), %rcx
	movq	%rsi, (%rcx)
	leaq	12(%rsp), %rsi
	movl	%edx, (%rsi)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%rbx)
	movq	%rsi, 16(%rbx)
	leaq	64(%rsp), %r14
	leaq	48(%rsp), %r15
	leaq	24(%rsp), %r12
	leaq	16(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z10get_squarePhPjj, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$128, %rsp
	.cfi_adjust_cfa_offset -128
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end2:
	.size	_Z25__device_stub__get_squarePhPjj, .Lfunc_end2-_Z25__device_stub__get_squarePhPjj
	.cfi_endproc
                                        # -- End function
	.globl	_Z17barrett_reductionPhS_S_S_S_S_S_S_S_S_S_S_S_jPjS0_S_ # -- Begin function _Z17barrett_reductionPhS_S_S_S_S_S_S_S_S_S_S_S_jPjS0_S_
	.type	_Z17barrett_reductionPhS_S_S_S_S_S_S_S_S_S_S_S_jPjS0_S_,@function
_Z17barrett_reductionPhS_S_S_S_S_S_S_S_S_S_S_S_jPjS0_S_: # @_Z17barrett_reductionPhS_S_S_S_S_S_S_S_S_S_S_S_jPjS0_S_
	.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	$72, %rsp
	.cfi_def_cfa_offset 128
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movq	%r9, %r15
	movq	%r8, 48(%rsp)                   # 8-byte Spill
	movq	%rcx, 32(%rsp)                  # 8-byte Spill
	movq	%rdx, 24(%rsp)                  # 8-byte Spill
	movq	%rsi, (%rsp)                    # 8-byte Spill
	movq	%rdi, %r13
	movq	208(%rsp), %rax
	movq	%rax, 40(%rsp)                  # 8-byte Spill
	movq	200(%rsp), %rbx
	movq	192(%rsp), %rbp
	movq	168(%rsp), %r14
	movq	160(%rsp), %rax
	movq	%rax, 64(%rsp)                  # 8-byte Spill
	movq	152(%rsp), %r12
	movq	144(%rsp), %rax
	movq	%rax, 16(%rsp)                  # 8-byte Spill
	movq	136(%rsp), %rax
	movq	%rax, 8(%rsp)                   # 8-byte Spill
	movq	128(%rsp), %rax
	movq	%rax, 56(%rsp)                  # 8-byte Spill
	movabsq	$4294967360, %rdx               # imm = 0x100000040
	leaq	-60(%rdx), %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB3_2
# %bb.1:
	movq	24(%rsp), %rdi                  # 8-byte Reload
	movq	(%rsp), %rsi                    # 8-byte Reload
	movq	%rbp, %rdx
	movl	$256, %ecx                      # imm = 0x100
	callq	_Z27__device_stub__get_productsPhS_Pjj
.LBB3_2:
	movl	$1536, %edx                     # imm = 0x600
	movq	%rbx, %rdi
	movq	%rbp, (%rsp)                    # 8-byte Spill
	movq	%rbp, %rsi
	movl	$2, %ecx
	callq	hipMemcpy
	movl	(%rbx), %ecx
	movl	$1, %eax
.LBB3_3:                                # =>This Inner Loop Header: Depth=1
	movb	%cl, -1(%r15,%rax)
	movl	-4(%rbx,%rax,4), %ecx
	shrl	$8, %ecx
	addl	(%rbx,%rax,4), %ecx
	movl	%ecx, (%rbx,%rax,4)
	incq	%rax
	cmpq	$385, %rax                      # imm = 0x181
	jne	.LBB3_3
# %bb.4:
	movl	$2048, %edx                     # imm = 0x800
	movq	(%rsp), %rdi                    # 8-byte Reload
	xorl	%esi, %esi
	callq	hipMemset
	movl	$383, %ecx                      # imm = 0x17F
	movl	$128, %eax
	movq	8(%rsp), %rsi                   # 8-byte Reload
	movq	16(%rsp), %rbp                  # 8-byte Reload
.LBB3_5:                                # =>This Inner Loop Header: Depth=1
	movl	%ecx, %edx
	cmpb	$0, (%r15,%rdx)
	jne	.LBB3_7
# %bb.6:                                #   in Loop: Header=BB3_5 Depth=1
	decl	%eax
	addl	$-1, %ecx
	jb	.LBB3_5
	jmp	.LBB3_10
.LBB3_7:                                # %.critedge
	movl	%eax, %ecx
	negl	%ecx
	cmpl	$257, %ecx                      # imm = 0x101
	jb	.LBB3_10
# %bb.8:                                # %.lr.ph.i
	movl	%eax, %eax
	negq	%rax
	movl	$256, %ecx                      # imm = 0x100
.LBB3_9:                                # =>This Inner Loop Header: Depth=1
	movl	%ecx, %edx
	movb	(%r15,%rdx), %dl
	movb	%dl, -256(%rsi,%rcx)
	leaq	(%rax,%rcx), %rdx
	incq	%rdx
	incq	%rcx
	cmpq	$256, %rdx                      # imm = 0x100
	jne	.LBB3_9
.LBB3_10:                               # %_Z9bit_shiftPhS_jj.exit
	movl	$128, %edx
	movq	%rbp, %rdi
	movl	$1, %ecx
	callq	hipMemcpy
	movabsq	$4294967360, %rdx               # imm = 0x100000040
	leaq	-62(%rdx), %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB3_12
# %bb.11:
	movq	%rbp, %rdi
	movq	32(%rsp), %rsi                  # 8-byte Reload
	movq	(%rsp), %rdx                    # 8-byte Reload
	movl	$128, %ecx
	callq	_Z27__device_stub__get_productsPhS_Pjj
.LBB3_12:
	movl	$1024, %edx                     # imm = 0x400
	movq	%rbx, %rdi
	movq	(%rsp), %rsi                    # 8-byte Reload
	movl	$2, %ecx
	callq	hipMemcpy
	movl	(%rbx), %ecx
	movl	$1, %eax
.LBB3_13:                               # =>This Inner Loop Header: Depth=1
	movb	%cl, -1(%r12,%rax)
	movl	-4(%rbx,%rax,4), %ecx
	shrl	$8, %ecx
	addl	(%rbx,%rax,4), %ecx
	movl	%ecx, (%rbx,%rax,4)
	incq	%rax
	cmpq	$257, %rax                      # imm = 0x101
	jne	.LBB3_13
# %bb.14:
	xorl	%ebp, %ebp
	movl	$2048, %edx                     # imm = 0x800
	movq	(%rsp), %rdi                    # 8-byte Reload
	xorl	%esi, %esi
	callq	hipMemset
	xorl	%eax, %eax
.LBB3_15:                               # %.lr.ph.i80
                                        # =>This Inner Loop Header: Depth=1
	movb	(%r13,%rbp), %cl
	cmpl	$1, %eax
	jne	.LBB3_27
# %bb.16:                               # %.lr.ph.i80
                                        #   in Loop: Header=BB3_15 Depth=1
	testb	%cl, %cl
	jne	.LBB3_27
# %bb.17:                               #   in Loop: Header=BB3_15 Depth=1
	movb	(%r12,%rbp), %cl
	notb	%cl
	movl	$1, %eax
	jmp	.LBB3_18
.LBB3_27:                               #   in Loop: Header=BB3_15 Depth=1
	subb	%al, %cl
	movb	%cl, (%r13,%rbp)
	xorl	%eax, %eax
	subb	(%r12,%rbp), %cl
	setb	%al
.LBB3_18:                               # %.backedge.i
                                        #   in Loop: Header=BB3_15 Depth=1
	movb	%cl, (%r14,%rbp)
	incq	%rbp
	cmpq	$256, %rbp                      # imm = 0x100
	jne	.LBB3_15
# %bb.19:                               # %_Z11subtractionPhS_S_j.exit
	cmpb	$0, 128(%r14)
	movq	16(%rsp), %r13                  # 8-byte Reload
	movq	40(%rsp), %rbp                  # 8-byte Reload
	jne	.LBB3_21
# %bb.20:
	movb	127(%r14), %al
	movq	%r14, %rsi
	cmpb	127(%rbp), %al
	jb	.LBB3_26
.LBB3_21:
	movl	$129, %edi
	movl	$1, %esi
	callq	calloc
	movq	%rax, %rsi
	xorl	%eax, %eax
	xorl	%ecx, %ecx
.LBB3_22:                               # %.lr.ph.i85
                                        # =>This Inner Loop Header: Depth=1
	movb	(%r14,%rax), %dl
	cmpl	$1, %ecx
	jne	.LBB3_28
# %bb.23:                               # %.lr.ph.i85
                                        #   in Loop: Header=BB3_22 Depth=1
	testb	%dl, %dl
	jne	.LBB3_28
# %bb.24:                               #   in Loop: Header=BB3_22 Depth=1
	movb	(%rbp,%rax), %dl
	notb	%dl
	movl	$1, %ecx
	jmp	.LBB3_25
.LBB3_28:                               #   in Loop: Header=BB3_22 Depth=1
	subb	%cl, %dl
	movb	%dl, (%r14,%rax)
	xorl	%ecx, %ecx
	subb	(%rbp,%rax), %dl
	setb	%cl
.LBB3_25:                               # %.backedge.i92
                                        #   in Loop: Header=BB3_22 Depth=1
	movb	%dl, (%rsi,%rax)
	incq	%rax
	cmpq	$129, %rax
	jne	.LBB3_22
.LBB3_26:                               # %_Z11subtractionPhS_S_j.exit97
	movl	$128, %edx
	movq	48(%rsp), %rdi                  # 8-byte Reload
	callq	memcpy@PLT
	movl	$384, %edx                      # imm = 0x180
	movq	%r15, %rdi
	xorl	%esi, %esi
	callq	memset@PLT
	movl	$384, %edx                      # imm = 0x180
	movq	56(%rsp), %rdi                  # 8-byte Reload
	xorl	%esi, %esi
	callq	hipMemset
	xorps	%xmm0, %xmm0
	movq	8(%rsp), %rax                   # 8-byte Reload
	movups	%xmm0, (%rax)
	movups	%xmm0, 16(%rax)
	movups	%xmm0, 32(%rax)
	movups	%xmm0, 48(%rax)
	movups	%xmm0, 64(%rax)
	movups	%xmm0, 80(%rax)
	movups	%xmm0, 96(%rax)
	movups	%xmm0, 112(%rax)
	movl	$128, %edx
	movq	%r13, %rdi
	xorl	%esi, %esi
	callq	hipMemset
	movl	$256, %edx                      # imm = 0x100
	movq	%r12, %rdi
	xorl	%esi, %esi
	callq	memset@PLT
	movl	$256, %edx                      # imm = 0x100
	movq	64(%rsp), %rdi                  # 8-byte Reload
	xorl	%esi, %esi
	callq	hipMemset
	movl	$129, %edx
	movq	%r14, %rdi
	xorl	%esi, %esi
	callq	memset@PLT
	movl	$2048, %edx                     # imm = 0x800
	movq	%rbx, %rdi
	xorl	%esi, %esi
	addq	$72, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	jmp	memset@PLT                      # TAILCALL
.Lfunc_end3:
	.size	_Z17barrett_reductionPhS_S_S_S_S_S_S_S_S_S_S_S_jPjS0_S_, .Lfunc_end3-_Z17barrett_reductionPhS_S_S_S_S_S_S_S_S_S_S_S_jPjS0_S_
	.cfi_endproc
                                        # -- End function
	.globl	_Z8checkbitPhj                  # -- Begin function _Z8checkbitPhj
	.type	_Z8checkbitPhj,@function
_Z8checkbitPhj:                         # @_Z8checkbitPhj
	.cfi_startproc
# %bb.0:
	movl	%esi, %ecx
	movl	%esi, %eax
	andb	$7, %cl
	movl	$1, %edx
	shll	%cl, %edx
	shrl	$3, %eax
	movzbl	(%rdi,%rax), %eax
	andl	%edx, %eax
                                        # kill: def $cl killed $cl killed $ecx
	shrl	%cl, %eax
                                        # kill: def $al killed $al killed $eax
	retq
.Lfunc_end4:
	.size	_Z8checkbitPhj, .Lfunc_end4-_Z8checkbitPhj
	.cfi_endproc
                                        # -- End function
	.globl	_Z27__device_stub__get_productsPhS_Pjj # -- Begin function _Z27__device_stub__get_productsPhS_Pjj
	.type	_Z27__device_stub__get_productsPhS_Pjj,@function
_Z27__device_stub__get_productsPhS_Pjj: # @_Z27__device_stub__get_productsPhS_Pjj
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$112, %rsp
	.cfi_def_cfa_offset 160
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	40(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	32(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	24(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	4(%rsp), %rdx
	movl	%ecx, (%rdx)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	leaq	64(%rsp), %r14
	leaq	48(%rsp), %r15
	leaq	16(%rsp), %r12
	leaq	8(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z12get_productsPhS_Pjj, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$128, %rsp
	.cfi_adjust_cfa_offset -128
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end5:
	.size	_Z27__device_stub__get_productsPhS_Pjj, .Lfunc_end5-_Z27__device_stub__get_productsPhS_Pjj
	.cfi_endproc
                                        # -- End function
	.globl	_Z9bit_shiftPhS_jj              # -- Begin function _Z9bit_shiftPhS_jj
	.type	_Z9bit_shiftPhS_jj,@function
_Z9bit_shiftPhS_jj:                     # @_Z9bit_shiftPhS_jj
	.cfi_startproc
# %bb.0:
	addl	%edx, %edx
	movl	%edx, %r9d
	shrl	$3, %r9d
	cmpl	%r9d, %ecx
	jb	.LBB6_5
# %bb.1:
	movl	%ecx, %eax
	subl	%r9d, %eax
	je	.LBB6_5
# %bb.2:                                # %.lr.ph
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset %rbx, -24
	.cfi_offset %rbp, -16
	andl	$6, %edx
	movl	$8, %r8d
	subl	%edx, %r8d
	leaq	1(%r9), %r10
	xorl	%r11d, %r11d
.LBB6_3:                                # =>This Inner Loop Header: Depth=1
	leal	(%r9,%r11), %ecx
	movzbl	(%rdi,%rcx), %ebx
	movl	%edx, %ecx
	shrl	%cl, %ebx
	movb	%bl, (%rsi,%r11)
	leal	(%r10,%r11), %ecx
	movzbl	(%rdi,%rcx), %ebp
	movl	%r8d, %ecx
	shll	%cl, %ebp
	orl	%ebx, %ebp
	movb	%bpl, (%rsi,%r11)
	incq	%r11
	cmpq	%r11, %rax
	jne	.LBB6_3
# %bb.4:
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	.cfi_restore %rbx
	.cfi_restore %rbp
.LBB6_5:                                # %.loopexit
	retq
.Lfunc_end6:
	.size	_Z9bit_shiftPhS_jj, .Lfunc_end6-_Z9bit_shiftPhS_jj
	.cfi_endproc
                                        # -- End function
	.globl	_Z11subtractionPhS_S_j          # -- Begin function _Z11subtractionPhS_S_j
	.type	_Z11subtractionPhS_S_j,@function
_Z11subtractionPhS_S_j:                 # @_Z11subtractionPhS_S_j
	.cfi_startproc
# %bb.0:
	testl	%ecx, %ecx
	je	.LBB7_6
# %bb.1:                                # %.lr.ph.preheader
	movl	%ecx, %eax
	xorl	%ecx, %ecx
	xorl	%r8d, %r8d
.LBB7_2:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movb	(%rdi,%rcx), %r9b
	cmpl	$1, %r8d
	jne	.LBB7_7
# %bb.3:                                # %.lr.ph
                                        #   in Loop: Header=BB7_2 Depth=1
	testb	%r9b, %r9b
	jne	.LBB7_7
# %bb.4:                                #   in Loop: Header=BB7_2 Depth=1
	movb	(%rsi,%rcx), %r9b
	notb	%r9b
	movl	$1, %r8d
	jmp	.LBB7_5
.LBB7_7:                                #   in Loop: Header=BB7_2 Depth=1
	subb	%r8b, %r9b
	movb	%r9b, (%rdi,%rcx)
	xorl	%r8d, %r8d
	subb	(%rsi,%rcx), %r9b
	setb	%r8b
.LBB7_5:                                # %.backedge
                                        #   in Loop: Header=BB7_2 Depth=1
	movb	%r9b, (%rdx,%rcx)
	incq	%rcx
	cmpq	%rcx, %rax
	jne	.LBB7_2
.LBB7_6:                                # %._crit_edge
	retq
.Lfunc_end7:
	.size	_Z11subtractionPhS_S_j, .Lfunc_end7-_Z11subtractionPhS_S_j
	.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	.LBB8_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB8_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z12get_productsPhS_Pjj, %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	$_Z10get_squarePhPjj, %esi
	movl	$.L__unnamed_2, %edx
	movl	$.L__unnamed_2, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end8:
	.size	__hip_module_ctor, .Lfunc_end8-__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	.LBB9_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
.LBB9_2:
	retq
.Lfunc_end9:
	.size	__hip_module_dtor, .Lfunc_end9-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	.L.str.1,@object                # @.str.1
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str.1:
	.asciz	"GPU time: %.6f\n"
	.size	.L.str.1, 16

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	"message[%d] = %x\n"
	.size	.L.str.2, 18

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

	.type	_Z12get_productsPhS_Pjj,@object # @_Z12get_productsPhS_Pjj
	.globl	_Z12get_productsPhS_Pjj
	.p2align	3, 0x0
_Z12get_productsPhS_Pjj:
	.quad	_Z27__device_stub__get_productsPhS_Pjj
	.size	_Z12get_productsPhS_Pjj, 8

	.type	.L__unnamed_1,@object           # @0
	.section	.rodata.str1.1,"aMS",@progbits,1
.L__unnamed_1:
	.asciz	"_Z12get_productsPhS_Pjj"
	.size	.L__unnamed_1, 24

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"_Z10get_squarePhPjj"
	.size	.L__unnamed_2, 20

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

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

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

	.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 _Z25__device_stub__get_squarePhPjj
	.addrsig_sym _Z27__device_stub__get_productsPhS_Pjj
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z10get_squarePhPjj
	.addrsig_sym _Z12get_productsPhS_Pjj
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z12get_productsPhS_Pjj ; -- Begin function _Z12get_productsPhS_Pjj
	.globl	_Z12get_productsPhS_Pjj
	.p2align	8
	.type	_Z12get_productsPhS_Pjj,@function
_Z12get_productsPhS_Pjj:                ; @_Z12get_productsPhS_Pjj
; %bb.0:
	s_load_b32 s2, s[0:1], 0x18
	s_waitcnt lgkmcnt(0)
	s_cmp_eq_u32 s2, 0
	s_cbranch_scc1 .LBB0_3
; %bb.1:                                ; %.lr.ph
	s_clause 0x2
	s_load_b32 s3, s[0:1], 0x2c
	s_load_b128 s[4:7], s[0:1], 0x0
	s_load_b64 s[0:1], s[0:1], 0x10
	s_waitcnt lgkmcnt(0)
	s_and_b32 s3, s3, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_mad_u64_u32 v[1:2], null, s15, s3, v[0:1]
	v_mov_b32_e32 v2, 0
	v_add_co_u32 v3, s3, s6, v1
	s_delay_alu instid0(VALU_DEP_1)
	v_add_co_ci_u32_e64 v4, null, s7, 0, s3
.LBB0_2:                                ; =>This Inner Loop Header: Depth=1
	global_load_u8 v0, v2, s[4:5]
	global_load_u8 v9, v[3:4], off
	v_lshlrev_b64 v[5:6], 2, v[1:2]
	v_add_nc_u32_e32 v1, 1, v1
	s_add_i32 s2, s2, -1
	s_add_u32 s4, s4, 1
	s_addc_u32 s5, s5, 0
	s_cmp_lg_u32 s2, 0
	v_lshlrev_b64 v[7:8], 2, v[1:2]
	v_add_co_u32 v5, vcc_lo, s0, v5
	v_add_co_ci_u32_e32 v6, vcc_lo, s1, v6, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_co_u32 v7, vcc_lo, s0, v7
	v_add_co_ci_u32_e32 v8, vcc_lo, s1, v8, vcc_lo
	s_waitcnt vmcnt(0)
	v_mul_u32_u24_e32 v0, v9, v0
	s_delay_alu instid0(VALU_DEP_1)
	v_and_b32_e32 v9, 0xff, v0
	v_lshrrev_b32_e32 v0, 8, v0
	s_clause 0x1
	global_atomic_add_u32 v[5:6], v9, off
	global_atomic_add_u32 v[7:8], v0, off
	s_cbranch_scc1 .LBB0_2
.LBB0_3:                                ; %._crit_edge
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z12get_productsPhS_Pjj
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 288
		.amdhsa_user_sgpr_count 15
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 10
		.amdhsa_next_free_sgpr 16
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end0:
	.size	_Z12get_productsPhS_Pjj, .Lfunc_end0-_Z12get_productsPhS_Pjj
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 236
; NumSgprs: 18
; NumVgprs: 10
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 10
; 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	_Z10get_squarePhPjj     ; -- Begin function _Z10get_squarePhPjj
	.globl	_Z10get_squarePhPjj
	.p2align	8
	.type	_Z10get_squarePhPjj,@function
_Z10get_squarePhPjj:                    ; @_Z10get_squarePhPjj
; %bb.0:
	s_load_b32 s4, s[0:1], 0x10
	s_waitcnt lgkmcnt(0)
	s_cmp_eq_u32 s4, 0
	s_cbranch_scc1 .LBB1_3
; %bb.1:                                ; %.lr.ph
	s_clause 0x1
	s_load_b32 s5, s[0:1], 0x24
	s_load_b128 s[0:3], s[0:1], 0x0
	s_waitcnt lgkmcnt(0)
	s_and_b32 s5, s5, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_mad_u64_u32 v[1:2], null, s15, s5, v[0:1]
	v_mov_b32_e32 v2, 0
	v_add_co_u32 v3, s5, s0, v1
	s_delay_alu instid0(VALU_DEP_1)
	v_add_co_ci_u32_e64 v4, null, s1, 0, s5
.LBB1_2:                                ; =>This Inner Loop Header: Depth=1
	global_load_u8 v0, v2, s[0:1]
	global_load_u8 v9, v[3:4], off
	v_lshlrev_b64 v[5:6], 2, v[1:2]
	v_add_nc_u32_e32 v1, 1, v1
	s_add_i32 s4, s4, -1
	s_add_u32 s0, s0, 1
	s_addc_u32 s1, s1, 0
	s_cmp_lg_u32 s4, 0
	v_lshlrev_b64 v[7:8], 2, v[1:2]
	v_add_co_u32 v5, vcc_lo, s2, v5
	v_add_co_ci_u32_e32 v6, vcc_lo, s3, v6, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_co_u32 v7, vcc_lo, s2, v7
	v_add_co_ci_u32_e32 v8, vcc_lo, s3, v8, vcc_lo
	s_waitcnt vmcnt(0)
	v_mul_u32_u24_e32 v0, v9, v0
	s_delay_alu instid0(VALU_DEP_1)
	v_and_b32_e32 v9, 0xff, v0
	v_lshrrev_b32_e32 v0, 8, v0
	s_clause 0x1
	global_atomic_add_u32 v[5:6], v9, off
	global_atomic_add_u32 v[7:8], v0, off
	s_cbranch_scc1 .LBB1_2
.LBB1_3:                                ; %._crit_edge
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z10get_squarePhPjj
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 280
		.amdhsa_user_sgpr_count 15
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 10
		.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	_Z10get_squarePhPjj, .Lfunc_end1-_Z10get_squarePhPjj
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 228
; NumSgprs: 18
; NumVgprs: 10
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 10
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

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

	.end_amdgpu_metadata
