#include <stdio.h>
#include <iostream>
#include <stdlib.h>

#include <math.h>
#include "aes_wrapper_gpu.cuh"

#define KEY_SIZE 16
#define SUBKEYS_SIZE 176
#define SBOX_SIZE 256

unsigned char *d_buffer_key = NULL;
unsigned char *d_buffer_sbox = NULL;
unsigned char *d_buffer_inv_sbox = NULL;

unsigned char SubKeys[SUBKEYS_SIZE];

/* Test Key */
unsigned char Key[KEY_SIZE] = {
  0x0f, 0x15, 0x71, 0xc9, 0x47, 0xd9, 0xe8, 0x59, 0x0c, 0xb7, 0xad, 0xd6,
  0xaf, 0x7f, 0x67, 0x98
};

__device__ unsigned char mulGaloisField2_8 (unsigned char a, unsigned char b) {
  register unsigned char p = 0;
  register unsigned char hi_bit_set;
  register unsigned char counter;

  for(counter = 0; counter < 8; counter++) {
    if ((b & 1) == 1)
      p ^= a;
    hi_bit_set = (a & 0x80);
    a <<= 1;
    if (hi_bit_set == 0x80)
      a ^= 0x1b;
    b >>= 1;
  }

  return p;
}

__device__ void aes_KeyAddition (unsigned char * internBuffer, unsigned char * SubKeys, int round) {
  for(register int k = 0; k < 16; k++) {
    internBuffer[k] = internBuffer[k] ^ SubKeys[(16 * round) + k];
  }
}

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

__device__ void aes_SubstitutionBox (unsigned char * internBuffer, unsigned char * Sbox) {
  for(register int k = 0; k < 16; k++) {
    internBuffer[k] = Sbox[internBuffer[k]];
  }
}

__device__ void aes_ShiftRows (unsigned char * internBuffer) {
  register unsigned char tmpBuffer;
  /*      State
     0    4       8       12      no shift
     1    5       9       13      1 left shift
     2    6       10      14      2 left shift
     3    7       11      15      3 left shift
   */

     tmpBuffer         = internBuffer[1];
     internBuffer[1]   = internBuffer[5];
     internBuffer[5]   = internBuffer[9];
     internBuffer[9]   = internBuffer[13];
     internBuffer[13]  = tmpBuffer;

     tmpBuffer         = internBuffer[2];
     internBuffer[2]   = internBuffer[10];
     internBuffer[10]  = tmpBuffer;
     tmpBuffer         = internBuffer[6];
     internBuffer[6]   = internBuffer[14];
     internBuffer[14]  = tmpBuffer;

     tmpBuffer = internBuffer[15];
     internBuffer[15] = internBuffer[11];
     internBuffer[11] = internBuffer[7];
     internBuffer[7] = internBuffer[3];
     internBuffer[3] = tmpBuffer;
   }

   __device__ void mixColumn (unsigned char * column) {
    register unsigned char i;
    register unsigned char cpy[4];

    for (i = 0; i < 4; i++) {
      cpy[i] = column[i];
    }

    column[0] = mulGaloisField2_8 (cpy[0], 2) ^
    mulGaloisField2_8 (cpy[1], 3) ^
    mulGaloisField2_8 (cpy[2], 1) ^
    mulGaloisField2_8 (cpy[3], 1);

    column[1] = mulGaloisField2_8 (cpy[0], 1) ^
    mulGaloisField2_8 (cpy[1], 2) ^
    mulGaloisField2_8 (cpy[2], 3) ^
    mulGaloisField2_8 (cpy[3], 1);

    column[2] = mulGaloisField2_8 (cpy[0], 1) ^
    mulGaloisField2_8 (cpy[1], 1) ^
    mulGaloisField2_8 (cpy[2], 2) ^
    mulGaloisField2_8 (cpy[3], 3);

    column[3] = mulGaloisField2_8 (cpy[0], 3) ^
    mulGaloisField2_8 (cpy[1], 1) ^
    mulGaloisField2_8 (cpy[2], 1) ^
    mulGaloisField2_8 (cpy[3], 2);
  }

  __device__ void aes_MixColumns (unsigned char * internBuffer) {
    register int i, j;
    register unsigned char column[4];

    for(i = 0; i < 4; i++) {
      for(j = 0; j < 4; j++) {
        column[j] = internBuffer[(i * 4) + j];
      }

      mixColumn (column);

      for(j = 0; j < 4; j++) {
        internBuffer[(i * 4) + j] = column[j];
      }
    }
  }

  __global__ void aes_encryption (unsigned char * SBOX, unsigned char * BufferData, unsigned char * SubKeys) {
  /* each thread progresses 16 bytes */
    register int id = (blockDim.x * blockIdx.x + threadIdx.x) * 16;

  /* IF threadId == 0, copy shared S_BOX & copy shared SubKeys 
   * Wait for Sync.
   */
   __shared__ unsigned char sharedSbox[SBOX_SIZE];
   __shared__ unsigned char sharedSubKeys[SUBKEYS_SIZE];
   if(threadIdx.x == 0) {
    for(int i = 0; i < SBOX_SIZE; i++) {
      sharedSbox[i] = SBOX[i];
    }
    for (int i = 0; i < SUBKEYS_SIZE; i++) {
      sharedSubKeys[i] = SubKeys[i];
    }
  }

  __syncthreads ();
  /* Now we are sync with all threads */
  /* Copy 16 bytes to intern buffer */
  register unsigned char internBuffer[16];
  for(register int i = 0; i < 16; i++) {
    internBuffer[i] = BufferData[id + i];
  }
  /* Now run complete AES */

  /*Initial XOR */
  aes_KeyAddition (internBuffer, sharedSubKeys, 0);

  /* Round loop */
  for(register int i = 1; i < 11; i++) {
    /* SubBytes */
    aes_SubstitutionBox (internBuffer, sharedSbox);

    /* ShiftRows */
    aes_ShiftRows (internBuffer);

    if(i != 10) {
  	  /* MixColumns, skip in last round */
     aes_MixColumns (internBuffer);
   }

    /* Key Addition */
   aes_KeyAddition (internBuffer, sharedSubKeys, i);
 }

  /* Copy everything back to the buffer */
 for(register int i = 0; i < 16; i++) {
  BufferData[id + i] = internBuffer[i];
}
}

/* END AES ENCRYPTION FUNCTIONS */

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

__device__ void aes_InvSubstitutionBox (unsigned char * internBuffer, unsigned char * InvSbox) {
  for(register int k = 0; k < 16; k++) {
    internBuffer[k] = InvSbox[internBuffer[k]];
  }
}

__device__ void aes_InvShiftRows (unsigned char * internBuffer) {
  register unsigned char tmpBuffer;
  /*      State
     0    4       8       12      no shift
     1    5       9       13      3 left shift
     2    6       10      14      2 left shift
     3    7       11      15      1 left shift
   */

     tmpBuffer         = internBuffer[13];
     internBuffer[13]  = internBuffer[9];
     internBuffer[9]   = internBuffer[5];
     internBuffer[5]   = internBuffer[1];
     internBuffer[1]   = tmpBuffer;

     tmpBuffer         = internBuffer[2];
     internBuffer[2]   = internBuffer[10];
     internBuffer[10]  = tmpBuffer;
     tmpBuffer         = internBuffer[6];
     internBuffer[6]   = internBuffer[14];
     internBuffer[14]  = tmpBuffer;

     tmpBuffer         = internBuffer[3];
     internBuffer[3]   = internBuffer[7];
     internBuffer[7]   = internBuffer[11];
     internBuffer[11]  = internBuffer[15];
     internBuffer[15]  = tmpBuffer;
   }


   __device__ void InvMixColumn (unsigned char * column) {
    register unsigned char i;
    register unsigned char cpy[4];

    for(i = 0; i < 4; i++) {
      cpy[i] = column[i];
    }

    column[0] = mulGaloisField2_8 (cpy[0], 0x0E) ^
    mulGaloisField2_8 (cpy[1], 0x0B) ^
    mulGaloisField2_8 (cpy[2], 0x0D) ^
    mulGaloisField2_8 (cpy[3], 0x09);

    column[1] = mulGaloisField2_8 (cpy[0], 0x09) ^
    mulGaloisField2_8 (cpy[1], 0x0E) ^
    mulGaloisField2_8 (cpy[2], 0x0B) ^
    mulGaloisField2_8 (cpy[3], 0x0D);

    column[2] = mulGaloisField2_8 (cpy[0], 0x0D) ^
    mulGaloisField2_8 (cpy[1], 0x09) ^
    mulGaloisField2_8 (cpy[2], 0x0E) ^
    mulGaloisField2_8 (cpy[3], 0x0B);

    column[3] = mulGaloisField2_8 (cpy[0], 0x0B) ^
    mulGaloisField2_8 (cpy[1], 0x0D) ^
    mulGaloisField2_8 (cpy[2], 0x09) ^
    mulGaloisField2_8 (cpy[3], 0x0E);
  }

  __device__ void aes_InvMixColumns (unsigned char * internBuffer) {
    register int i, j;
    register unsigned char column[4];

    for(i = 0; i < 4; i++) {
      for(j = 0; j < 4; j++) {
        column[j] = internBuffer[(i * 4) + j];
      }

      InvMixColumn (column);

      for(j = 0; j < 4; j++) {
        internBuffer[(i * 4) + j] = column[j];
      }
    }
  }

extern "C" int validate_pad_block(int pad, unsigned char* pad_block) {
  int i;

  for(i=AES_BLOCK_SIZE-1; i>=AES_BLOCK_SIZE-pad; i--) {
    if(pad != (int) pad_block[i]) return 0;
  }

  return 1;
}

extern "C" void init_aes_data_on_device() {

  cudaError_t status;

  /*
   *  allocate d_buffer_key, d_buffer_sbox and d_buffer_inv_sbox on device
   */

  status = cudaMalloc((void **) &d_buffer_key, sizeof(unsigned char) * SUBKEYS_SIZE);
  if(cudaSuccess != status) {
    printf("[CUDA] Error allocating device memory: %s\n", cudaGetErrorString(status));
  }

  status = cudaMalloc((void **) &d_buffer_sbox, sizeof(unsigned char) * SBOX_SIZE);
  if(cudaSuccess != status) {
    printf("[CUDA] Error allocating device memory: %s\n", cudaGetErrorString(status));
  }

  status = cudaMalloc((void **) &d_buffer_inv_sbox, sizeof(unsigned char) * SBOX_SIZE);
  if(cudaSuccess != status) {
    printf("[CUDA] Error allocating device memory: %s\n", cudaGetErrorString(status));
  }

  /*
   *  copy d_buffer_key, d_buffer_sbox and d_buffer_inv_sbox on device
   */
  status = cudaMemcpy(d_buffer_key, SubKeys, sizeof(unsigned char) * SUBKEYS_SIZE, cudaMemcpyHostToDevice);
  if(cudaSuccess != status) {
    printf("[CUDA] Error allocating device memory: %s\n", cudaGetErrorString(status));
  } 

  status = cudaMemcpy(d_buffer_sbox, Sbox, sizeof(unsigned char) * SBOX_SIZE, cudaMemcpyHostToDevice);
  if(cudaSuccess != status) {
    printf("[CUDA] Error allocating device memory: %s\n", cudaGetErrorString(status));
  }

  status = cudaMemcpy(d_buffer_inv_sbox, InvSbox, sizeof(unsigned char) * SBOX_SIZE, cudaMemcpyHostToDevice);
  if(cudaSuccess != status) {
    printf("[CUDA] Error allocating device memory: %s\n", cudaGetErrorString(status));
  }
}

extern "C" void free_aes_data_on_device() {

  cudaError_t status;

  status = cudaFree(d_buffer_key);
  if(cudaSuccess != status) {
    printf("[CUDA] Error freeing device memory: %s\n", cudaGetErrorString(status));
  }

  status = cudaFree(d_buffer_key);
  if(cudaSuccess != status) {
    printf("[CUDA] Error freeing device memory: %s\n", cudaGetErrorString(status));
  }

  status = cudaFree(d_buffer_key);
  if(cudaSuccess != status) {
    printf("[CUDA] Error freeing device memory: %s\n", cudaGetErrorString(status));
  }
}

  __global__ void aes_decryption (unsigned char * InvSbox, unsigned char * BufferData, unsigned char * SubKeys) {
  /* each thread progresses 16 bytes */
    register int id = (blockDim.x * blockIdx.x + threadIdx.x) * 16;

  /* IF threadId == 0, copy shared S_BOX & copy shared SubKeys 
   * Wait for Sync.
   */
   __shared__ unsigned char sharedInvSbox[SBOX_SIZE];
   __shared__ unsigned char sharedSubKeys[SUBKEYS_SIZE];
   if(threadIdx.x == 0) {
    for(int i = 0; i < SBOX_SIZE; i++) {
      sharedInvSbox[i] = InvSbox[i];
    }
    for(int i = 0; i < SUBKEYS_SIZE; i++) {
      sharedSubKeys[i] = SubKeys[i];
    }
  }
  __syncthreads ();

  /* Now we are synced with all threads */
  /* Copy 16 bytes to intern buffer */
  register unsigned char internBuffer[16];
  for(register int i = 0; i < 16; i++) {
    internBuffer[i] = BufferData[id + i];
  }
  /* Now run complete AES */

  /*Initial XOR */
  aes_KeyAddition(internBuffer, sharedSubKeys, 10);

  /* Round loop */
  for(register int i = 9; i >= 0; i--) {
    /* Inverted ShiftRows */
    aes_InvShiftRows(internBuffer);

    /* Inverted SubBytes */
    aes_InvSubstitutionBox(internBuffer, sharedInvSbox);

    /* Key Addition */
    aes_KeyAddition(internBuffer, sharedSubKeys, i);

    if(i != 0) {
      /* Inverted MixColumns, skip in last round */
      aes_InvMixColumns (internBuffer);
    }
  }

  /* Copy everything back to the buffer */
  for(register int i = 0; i < 16; i++) {
    BufferData[id + i] = internBuffer[i];
  }
}

extern "C" void aes_encryption_wrapper(unsigned char * BufferData, ssize_t buffer_size) {
  size_t d_buffer_data_size = buffer_size * sizeof(unsigned char);

  int NBlocks           = 1;
  int NThreadsPerBlock  = buffer_size/AES_BLOCK_SIZE;

  //printf("[CUDA] Using %d NBlocks and %d NThreadsPerBlock to process %jd unsigned char* (%jd bytes)\n", NBlocks, NThreadsPerBlock, buffer_size, d_buffer_data_size);

  // cudaEvent_t start, stop;
  // cudaEventCreate(&start);
  // cudaEventCreate(&stop);

  // cudaEventRecord(start);
  // ensure sbox, subkeys and bufferdata is in GPU memory
  // sbox can be in constant memory it won't be changed
  unsigned char* d_buffer_data;
  cudaError_t status;

  status = cudaMalloc(&d_buffer_data, d_buffer_data_size);

  if(cudaSuccess != status) {
    printf("[CUDA] Error allocating device memory: %s\n", cudaGetErrorString(status));
  }

  status = cudaMemcpy(d_buffer_data, BufferData, d_buffer_data_size, cudaMemcpyHostToDevice);
  if(cudaSuccess != status) {
    printf("[CUDA] Error copying data to device memory: %s\n", cudaGetErrorString(status));
  }

  //printf("Launching aes_encryption<<<%d, %d>>>\n", NBlocks, NThreadsPerBlock);

  aes_encryption<<<NBlocks, NThreadsPerBlock>>>(d_buffer_sbox, d_buffer_data, d_buffer_key);

  status = cudaMemcpy(BufferData, d_buffer_data, d_buffer_data_size, cudaMemcpyDeviceToHost);
  if(cudaSuccess != status) {
    printf("[CUDA] Error copying data to host memory: %s\n", cudaGetErrorString(status));
  }

  /*
   *  free allocated memory
   */
  status = cudaFree(d_buffer_data);
  if(cudaSuccess != status) {
    printf("[CUDA] Error freeing device memory: %s\n", cudaGetErrorString(status));
  }

  // cudaEventRecord(stop);
}

extern "C" void aes_decryption_wrapper(unsigned char * BufferData, ssize_t buffer_size) {
  size_t d_buffer_data_size = buffer_size * sizeof(unsigned char);
  unsigned char* d_buffer_data;
  cudaError_t status;

  int NBlocks           = 1;
  int NThreadsPerBlock  = buffer_size/AES_BLOCK_SIZE;

  //printf("[CUDA] Using %d NBlocks and %d NThreadsPerBlock\n", NBlocks, NThreadsPerBlock);

  // cudaEvent_t start, stop;
  // cudaEventCreate(&start);
  // cudaEventCreate(&stop);

  // cudaEventRecord(start);
  // ensure invsbox, subkeys and bufferdata is in GPU memory
  // invsbox can be in constant memory it won't be changed

  status = cudaMalloc(&d_buffer_data, d_buffer_data_size);

  if(cudaSuccess != status) {
    printf("[CUDA] Error allocating device memory: %s\n", cudaGetErrorString(status));
  }

  status = cudaMemcpy(d_buffer_data, BufferData, d_buffer_data_size, cudaMemcpyHostToDevice);
  if(cudaSuccess != status) {
    printf("[CUDA] Error copying data to device memory: %s\n", cudaGetErrorString(status));
  }

  aes_decryption<<<NBlocks, NThreadsPerBlock>>>(d_buffer_inv_sbox, d_buffer_data, d_buffer_key);

  status = cudaMemcpy(BufferData, d_buffer_data, d_buffer_data_size, cudaMemcpyDeviceToHost);
  if(cudaSuccess != status) {
    printf("[CUDA] Error copying data to host memory\n: %s\n", cudaGetErrorString(status));
  }


  /*
   *  free allocated memory
   */
  status = cudaFree(d_buffer_data);
  if(cudaSuccess != status) {
    printf("[CUDA] Error freeing device memory: %s\n", cudaGetErrorString(status));
  }
  // cudaEventRecord(stop);
}

/* END AES DECRYPTION FUNCTIONS */

/* START AES KEYSCHEDULE */
unsigned char Rcon[11] = {
  0x00, 0x01, 0x02, 0x04, 0x08, 0x10, 0x20, 0x40, 0x80, 0x1b, 0x36
};

void aes_keyschedule_gFunction (unsigned char * input, int round) {
  /* G Function in AES Keyschedule */

  /* ROTATE / 1 SHIFT LEFT */
  unsigned char tmpBuffer;
  tmpBuffer = input[0];
  input[0]  = input[1];
  input[1]  = input[2];
  input[2]  = input[3];
  input[3]  = tmpBuffer;

  /* SBOX */
  input[0]  = Sbox[input[0]];
  input[1]  = Sbox[input[1]];
  input[2]  = Sbox[input[2]];
  input[3]  = Sbox[input[3]];

  /* XOR Rcon */
  input[0] ^= Rcon[round];
}

extern "C" void aes_keyschedule() {
  /* First SubKey = Key */
  unsigned char Buffer[16];
  for(int i = 0; i < 16; i++) {
    SubKeys[i] = Key[i];
    Buffer[i] = Key[i];
  }

  /* Calculate remaining SubKeys */
  for(int round = 1; round < 11; round++) {
    /* Buffer for g function */
    unsigned char g_Buffer[4];
    for(int i = 0; i < 4; i++) {
      g_Buffer[i] = Buffer[12 + i];
    }

    aes_keyschedule_gFunction(g_Buffer, round);
    /* BUFFER[0] XOR Output_G (32 bit) */
    for(int i = 0; i < 4; i++) {
      Buffer[i] ^= g_Buffer[i];
    }

    /* Buffer[0] XOR BUFFER[1] (32-bit) */
    for(int i = 0; i < 4; i++) {
      Buffer[4 + i] ^= Buffer[i];
    }

    /* Buffer[1] XOR BUFFER[2] (32-bit) */
    for(int i = 0; i < 4; i++) {
      Buffer[8 + i] ^= Buffer[4 + i];
    }

    /* Buffer[2] XOR BUFFER[3] (32-bit) */
    for(int i = 0; i < 4; i++) {
      Buffer[12 + i] ^= Buffer[8 + i];
    }

    for(int i = 0; i < 16; i++) {
      SubKeys[(round * 16) + i] = Buffer[i];
    }
  }
}

/* END AES KEYSCHEDULE */

	.file	"tmpxft_0021afc1_00000000-6_aes_wrapper_gpu.cudafe1.cpp"
	.text
#APP
	.globl _ZSt21ios_base_library_initv

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB3671:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE3671:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z17mulGaloisField2_8hh
	.type	_Z17mulGaloisField2_8hh, @function
_Z17mulGaloisField2_8hh:
.LFB3652:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE3652:
	.size	_Z17mulGaloisField2_8hh, .-_Z17mulGaloisField2_8hh
	.globl	_Z15aes_KeyAdditionPhS_i
	.type	_Z15aes_KeyAdditionPhS_i, @function
_Z15aes_KeyAdditionPhS_i:
.LFB3653:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE3653:
	.size	_Z15aes_KeyAdditionPhS_i, .-_Z15aes_KeyAdditionPhS_i
	.globl	_Z19aes_SubstitutionBoxPhS_
	.type	_Z19aes_SubstitutionBoxPhS_, @function
_Z19aes_SubstitutionBoxPhS_:
.LFB3654:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE3654:
	.size	_Z19aes_SubstitutionBoxPhS_, .-_Z19aes_SubstitutionBoxPhS_
	.globl	_Z13aes_ShiftRowsPh
	.type	_Z13aes_ShiftRowsPh, @function
_Z13aes_ShiftRowsPh:
.LFB3655:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE3655:
	.size	_Z13aes_ShiftRowsPh, .-_Z13aes_ShiftRowsPh
	.globl	_Z9mixColumnPh
	.type	_Z9mixColumnPh, @function
_Z9mixColumnPh:
.LFB4321:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	pushq	%rax
	.cfi_def_cfa_offset 16
	call	_Z13aes_ShiftRowsPh
	.cfi_endproc
.LFE4321:
	.size	_Z9mixColumnPh, .-_Z9mixColumnPh
	.globl	_Z14aes_MixColumnsPh
	.type	_Z14aes_MixColumnsPh, @function
_Z14aes_MixColumnsPh:
.LFB4323:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	pushq	%rax
	.cfi_def_cfa_offset 16
	call	_Z13aes_ShiftRowsPh
	.cfi_endproc
.LFE4323:
	.size	_Z14aes_MixColumnsPh, .-_Z14aes_MixColumnsPh
	.globl	_Z22aes_InvSubstitutionBoxPhS_
	.type	_Z22aes_InvSubstitutionBoxPhS_, @function
_Z22aes_InvSubstitutionBoxPhS_:
.LFB4319:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	pushq	%rax
	.cfi_def_cfa_offset 16
	call	_Z19aes_SubstitutionBoxPhS_
	.cfi_endproc
.LFE4319:
	.size	_Z22aes_InvSubstitutionBoxPhS_, .-_Z22aes_InvSubstitutionBoxPhS_
	.globl	_Z16aes_InvShiftRowsPh
	.type	_Z16aes_InvShiftRowsPh, @function
_Z16aes_InvShiftRowsPh:
.LFB4325:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	pushq	%rax
	.cfi_def_cfa_offset 16
	call	_Z13aes_ShiftRowsPh
	.cfi_endproc
.LFE4325:
	.size	_Z16aes_InvShiftRowsPh, .-_Z16aes_InvShiftRowsPh
	.globl	_Z12InvMixColumnPh
	.type	_Z12InvMixColumnPh, @function
_Z12InvMixColumnPh:
.LFB4327:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	pushq	%rax
	.cfi_def_cfa_offset 16
	call	_Z13aes_ShiftRowsPh
	.cfi_endproc
.LFE4327:
	.size	_Z12InvMixColumnPh, .-_Z12InvMixColumnPh
	.globl	_Z17aes_InvMixColumnsPh
	.type	_Z17aes_InvMixColumnsPh, @function
_Z17aes_InvMixColumnsPh:
.LFB4329:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	pushq	%rax
	.cfi_def_cfa_offset 16
	call	_Z13aes_ShiftRowsPh
	.cfi_endproc
.LFE4329:
	.size	_Z17aes_InvMixColumnsPh, .-_Z17aes_InvMixColumnsPh
	.globl	validate_pad_block
	.type	validate_pad_block, @function
validate_pad_block:
.LFB3662:
	.cfi_startproc
	endbr64
	movl	$16, %edx
	movl	$15, %eax
	subl	%edi, %edx
.L23:
	cmpl	%eax, %edx
	jg	.L28
	movzbl	(%rsi,%rax), %ecx
	decq	%rax
	cmpl	%edi, %ecx
	je	.L23
	xorl	%eax, %eax
	ret
.L28:
	movl	$1, %eax
	ret
	.cfi_endproc
.LFE3662:
	.size	validate_pad_block, .-validate_pad_block
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"[CUDA] Error allocating device memory: %s\n"
	.text
	.globl	init_aes_data_on_device
	.type	init_aes_data_on_device, @function
init_aes_data_on_device:
.LFB3663:
	.cfi_startproc
	endbr64
	pushq	%rsi
	.cfi_def_cfa_offset 16
	leaq	d_buffer_key(%rip), %rdi
	movl	$176, %esi
	call	cudaMalloc@PLT
	testl	%eax, %eax
	je	.L30
	movl	%eax, %edi
	call	cudaGetErrorString@PLT
	leaq	.LC0(%rip), %rsi
	movl	$2, %edi
	movq	%rax, %rdx
	xorl	%eax, %eax
	call	__printf_chk@PLT
.L30:
	leaq	d_buffer_sbox(%rip), %rdi
	movl	$256, %esi
	call	cudaMalloc@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L31
	call	cudaGetErrorString@PLT
	leaq	.LC0(%rip), %rsi
	movl	$2, %edi
	movq	%rax, %rdx
	xorl	%eax, %eax
	call	__printf_chk@PLT
.L31:
	leaq	d_buffer_inv_sbox(%rip), %rdi
	movl	$256, %esi
	call	cudaMalloc@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L32
	call	cudaGetErrorString@PLT
	leaq	.LC0(%rip), %rsi
	movl	$2, %edi
	movq	%rax, %rdx
	xorl	%eax, %eax
	call	__printf_chk@PLT
.L32:
	movq	d_buffer_key(%rip), %rdi
	movl	$1, %ecx
	movl	$176, %edx
	leaq	SubKeys(%rip), %rsi
	call	cudaMemcpy@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L33
	call	cudaGetErrorString@PLT
	leaq	.LC0(%rip), %rsi
	movl	$2, %edi
	movq	%rax, %rdx
	xorl	%eax, %eax
	call	__printf_chk@PLT
.L33:
	movq	d_buffer_sbox(%rip), %rdi
	movl	$1, %ecx
	movl	$256, %edx
	leaq	Sbox(%rip), %rsi
	call	cudaMemcpy@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L34
	call	cudaGetErrorString@PLT
	leaq	.LC0(%rip), %rsi
	movl	$2, %edi
	movq	%rax, %rdx
	xorl	%eax, %eax
	call	__printf_chk@PLT
.L34:
	movq	d_buffer_inv_sbox(%rip), %rdi
	movl	$1, %ecx
	movl	$256, %edx
	leaq	InvSbox(%rip), %rsi
	call	cudaMemcpy@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L29
	call	cudaGetErrorString@PLT
	leaq	.LC0(%rip), %rsi
	movl	$2, %edi
	popq	%rcx
	.cfi_remember_state
	.cfi_def_cfa_offset 8
	movq	%rax, %rdx
	xorl	%eax, %eax
	jmp	__printf_chk@PLT
.L29:
	.cfi_restore_state
	popq	%rax
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3663:
	.size	init_aes_data_on_device, .-init_aes_data_on_device
	.section	.rodata.str1.1
.LC1:
	.string	"[CUDA] Error freeing device memory: %s\n"
	.text
	.globl	free_aes_data_on_device
	.type	free_aes_data_on_device, @function
free_aes_data_on_device:
.LFB3664:
	.cfi_startproc
	endbr64
	pushq	%rsi
	.cfi_def_cfa_offset 16
	movq	d_buffer_key(%rip), %rdi
	call	cudaFree@PLT
	testl	%eax, %eax
	je	.L53
	movl	%eax, %edi
	call	cudaGetErrorString@PLT
	leaq	.LC1(%rip), %rsi
	movl	$2, %edi
	movq	%rax, %rdx
	xorl	%eax, %eax
	call	__printf_chk@PLT
.L53:
	movq	d_buffer_key(%rip), %rdi
	call	cudaFree@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L54
	call	cudaGetErrorString@PLT
	leaq	.LC1(%rip), %rsi
	movl	$2, %edi
	movq	%rax, %rdx
	xorl	%eax, %eax
	call	__printf_chk@PLT
.L54:
	movq	d_buffer_key(%rip), %rdi
	call	cudaFree@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L52
	call	cudaGetErrorString@PLT
	leaq	.LC1(%rip), %rsi
	movl	$2, %edi
	popq	%rcx
	.cfi_remember_state
	.cfi_def_cfa_offset 8
	movq	%rax, %rdx
	xorl	%eax, %eax
	jmp	__printf_chk@PLT
.L52:
	.cfi_restore_state
	popq	%rax
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3664:
	.size	free_aes_data_on_device, .-free_aes_data_on_device
	.globl	_Z25aes_keyschedule_gFunctionPhi
	.type	_Z25aes_keyschedule_gFunctionPhi, @function
_Z25aes_keyschedule_gFunctionPhi:
.LFB3667:
	.cfi_startproc
	endbr64
	movzbl	1(%rdi), %edx
	movzbl	2(%rdi), %r8d
	movq	%rdi, %rax
	movl	%esi, %r9d
	movzbl	(%rdi), %ecx
	leaq	Sbox(%rip), %rsi
	movb	%dl, (%rdi)
	movb	%r8b, 1(%rdi)
	movzbl	3(%rdi), %edi
	movb	%cl, 3(%rax)
	movb	%dil, 2(%rax)
	movb	(%rsi,%rdx), %dl
	movb	%dl, (%rax)
	movb	(%rsi,%r8), %r8b
	movb	%r8b, 1(%rax)
	movb	(%rsi,%rdi), %dil
	movb	%dil, 2(%rax)
	movb	(%rsi,%rcx), %cl
	movslq	%r9d, %rsi
	movb	%cl, 3(%rax)
	leaq	Rcon(%rip), %rcx
	xorb	(%rcx,%rsi), %dl
	movb	%dl, (%rax)
	ret
	.cfi_endproc
.LFE3667:
	.size	_Z25aes_keyschedule_gFunctionPhi, .-_Z25aes_keyschedule_gFunctionPhi
	.globl	aes_keyschedule
	.type	aes_keyschedule, @function
aes_keyschedule:
.LFB3668:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	leaq	Key(%rip), %rsi
	leaq	SubKeys(%rip), %rcx
	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	$40, %rsp
	.cfi_def_cfa_offset 96
	movq	%fs:40, %rax
	movq	%rax, 24(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %r13
.L65:
	movb	(%rsi,%rax), %dl
	movb	%dl, (%rcx,%rax)
	incq	%rax
	movb	%dl, -1(%rax,%r13)
	cmpq	$16, %rax
	jne	.L65
	leaq	SubKeys(%rip), %r14
	movl	$1, %r12d
	leaq	4(%rsp), %r15
.L66:
	movl	20(%rsp), %eax
	movl	%r12d, %esi
	movq	%r15, %rdi
	movb	20(%rsp), %bpl
	movb	21(%rsp), %bl
	movb	22(%rsp), %r11b
	movl	%eax, 4(%rsp)
	movb	23(%rsp), %r10b
	call	_Z25aes_keyschedule_gFunctionPhi
	xorl	%eax, %eax
.L67:
	movb	0(%r13,%rax), %dl
	xorb	(%r15,%rax), %dl
	movb	%dl, 0(%r13,%rax)
	incq	%rax
	cmpq	$4, %rax
	jne	.L67
	movq	%r13, %rax
	leaq	12(%rsp), %rdx
.L68:
	movb	(%rax), %cl
	xorb	%cl, 4(%rax)
	incq	%rax
	cmpq	%rax, %rdx
	jne	.L68
	movb	16(%rsp), %sil
	movb	17(%rsp), %cl
	incl	%r12d
	addq	$16, %r14
	movb	18(%rsp), %dl
	movb	19(%rsp), %al
	xorl	12(%rsp), %esi
	xorb	13(%rsp), %cl
	xorb	14(%rsp), %dl
	xorb	15(%rsp), %al
	movb	%sil, 16(%rsp)
	xorl	%ebp, %esi
	movb	%cl, 17(%rsp)
	xorl	%ebx, %ecx
	movb	%dl, 18(%rsp)
	xorl	%r11d, %edx
	movb	%al, 19(%rsp)
	xorl	%r10d, %eax
	movb	%sil, 20(%rsp)
	movb	%cl, 21(%rsp)
	movb	%dl, 22(%rsp)
	movb	%al, 23(%rsp)
	movups	8(%rsp), %xmm0
	movaps	%xmm0, (%r14)
	cmpl	$11, %r12d
	jne	.L66
	movq	24(%rsp), %rax
	subq	%fs:40, %rax
	je	.L70
	call	__stack_chk_fail@PLT
.L70:
	addq	$40, %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
.LFE3668:
	.size	aes_keyschedule, .-aes_keyschedule
	.globl	_Z38__device_stub__Z14aes_encryptionPhS_S_PhS_S_
	.type	_Z38__device_stub__Z14aes_encryptionPhS_S_PhS_S_, @function
_Z38__device_stub__Z14aes_encryptionPhS_S_PhS_S_:
.LFB3693:
	.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
	movq	%rdx, 8(%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	8(%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	.L75
	pushq	40(%rsp)
	.cfi_def_cfa_offset 152
	leaq	_Z14aes_encryptionPhS_S_(%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
.L75:
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L77
	call	__stack_chk_fail@PLT
.L77:
	addq	$136, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3693:
	.size	_Z38__device_stub__Z14aes_encryptionPhS_S_PhS_S_, .-_Z38__device_stub__Z14aes_encryptionPhS_S_PhS_S_
	.globl	_Z14aes_encryptionPhS_S_
	.type	_Z14aes_encryptionPhS_S_, @function
_Z14aes_encryptionPhS_S_:
.LFB3694:
	.cfi_startproc
	endbr64
	jmp	_Z38__device_stub__Z14aes_encryptionPhS_S_PhS_S_
	.cfi_endproc
.LFE3694:
	.size	_Z14aes_encryptionPhS_S_, .-_Z14aes_encryptionPhS_S_
	.section	.rodata.str1.1
.LC2:
	.string	"[CUDA] Error copying data to device memory: %s\n"
.LC3:
	.string	"[CUDA] Error copying data to host memory: %s\n"
	.text
	.globl	aes_encryption_wrapper
	.type	aes_encryption_wrapper, @function
aes_encryption_wrapper:
.LFB3665:
	.cfi_startproc
	endbr64
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	movl	$16, %ecx
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	movq	%rdi, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	movq	%rsi, %rbx
	subq	$48, %rsp
	.cfi_def_cfa_offset 80
	movq	%fs:40, %rax
	movq	%rax, 40(%rsp)
	xorl	%eax, %eax
	movq	%rsi, %rax
	leaq	8(%rsp), %rdi
	cqto
	idivq	%rcx
	movq	%rax, %r12
	call	cudaMalloc@PLT
	testl	%eax, %eax
	je	.L81
	movl	%eax, %edi
	call	cudaGetErrorString@PLT
	leaq	.LC0(%rip), %rsi
	movl	$2, %edi
	movq	%rax, %rdx
	xorl	%eax, %eax
	call	__printf_chk@PLT
.L81:
	movq	8(%rsp), %rdi
	movl	$1, %ecx
	movq	%rbx, %rdx
	movq	%rbp, %rsi
	call	cudaMemcpy@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L82
	call	cudaGetErrorString@PLT
	leaq	.LC2(%rip), %rsi
	movl	$2, %edi
	movq	%rax, %rdx
	xorl	%eax, %eax
	call	__printf_chk@PLT
.L82:
	movl	%r12d, 28(%rsp)
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$1, %esi
	movabsq	$4294967297, %rdi
	movq	%rdi, 32(%rsp)
	movl	36(%rsp), %ecx
	movq	28(%rsp), %rdx
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L83
	movq	d_buffer_key(%rip), %rdx
	movq	8(%rsp), %rsi
	movq	d_buffer_sbox(%rip), %rdi
	call	_Z38__device_stub__Z14aes_encryptionPhS_S_PhS_S_
.L83:
	movq	8(%rsp), %rsi
	movq	%rbp, %rdi
	movl	$2, %ecx
	movq	%rbx, %rdx
	call	cudaMemcpy@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L84
	call	cudaGetErrorString@PLT
	leaq	.LC3(%rip), %rsi
	movl	$2, %edi
	movq	%rax, %rdx
	xorl	%eax, %eax
	call	__printf_chk@PLT
.L84:
	movq	8(%rsp), %rdi
	call	cudaFree@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L80
	call	cudaGetErrorString@PLT
	leaq	.LC1(%rip), %rsi
	movl	$2, %edi
	movq	%rax, %rdx
	xorl	%eax, %eax
	call	__printf_chk@PLT
.L80:
	movq	40(%rsp), %rax
	subq	%fs:40, %rax
	je	.L86
	call	__stack_chk_fail@PLT
.L86:
	addq	$48, %rsp
	.cfi_def_cfa_offset 32
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3665:
	.size	aes_encryption_wrapper, .-aes_encryption_wrapper
	.globl	_Z38__device_stub__Z14aes_decryptionPhS_S_PhS_S_
	.type	_Z38__device_stub__Z14aes_decryptionPhS_S_PhS_S_, @function
_Z38__device_stub__Z14aes_decryptionPhS_S_PhS_S_:
.LFB3695:
	.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
	movq	%rdx, 8(%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	8(%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	.L100
	pushq	40(%rsp)
	.cfi_def_cfa_offset 152
	leaq	_Z14aes_decryptionPhS_S_(%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
.L100:
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L102
	call	__stack_chk_fail@PLT
.L102:
	addq	$136, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3695:
	.size	_Z38__device_stub__Z14aes_decryptionPhS_S_PhS_S_, .-_Z38__device_stub__Z14aes_decryptionPhS_S_PhS_S_
	.globl	_Z14aes_decryptionPhS_S_
	.type	_Z14aes_decryptionPhS_S_, @function
_Z14aes_decryptionPhS_S_:
.LFB3696:
	.cfi_startproc
	endbr64
	jmp	_Z38__device_stub__Z14aes_decryptionPhS_S_PhS_S_
	.cfi_endproc
.LFE3696:
	.size	_Z14aes_decryptionPhS_S_, .-_Z14aes_decryptionPhS_S_
	.section	.rodata.str1.1
.LC4:
	.string	"[CUDA] Error copying data to host memory\n: %s\n"
	.text
	.globl	aes_decryption_wrapper
	.type	aes_decryption_wrapper, @function
aes_decryption_wrapper:
.LFB3666:
	.cfi_startproc
	endbr64
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	movl	$16, %ecx
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	movq	%rdi, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	movq	%rsi, %rbx
	subq	$48, %rsp
	.cfi_def_cfa_offset 80
	movq	%fs:40, %rax
	movq	%rax, 40(%rsp)
	xorl	%eax, %eax
	movq	%rsi, %rax
	leaq	8(%rsp), %rdi
	cqto
	idivq	%rcx
	movq	%rax, %r12
	call	cudaMalloc@PLT
	testl	%eax, %eax
	je	.L106
	movl	%eax, %edi
	call	cudaGetErrorString@PLT
	leaq	.LC0(%rip), %rsi
	movl	$2, %edi
	movq	%rax, %rdx
	xorl	%eax, %eax
	call	__printf_chk@PLT
.L106:
	movq	8(%rsp), %rdi
	movl	$1, %ecx
	movq	%rbx, %rdx
	movq	%rbp, %rsi
	call	cudaMemcpy@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L107
	call	cudaGetErrorString@PLT
	leaq	.LC2(%rip), %rsi
	movl	$2, %edi
	movq	%rax, %rdx
	xorl	%eax, %eax
	call	__printf_chk@PLT
.L107:
	movl	%r12d, 28(%rsp)
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$1, %esi
	movabsq	$4294967297, %rdi
	movq	%rdi, 32(%rsp)
	movl	36(%rsp), %ecx
	movq	28(%rsp), %rdx
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L108
	movq	d_buffer_key(%rip), %rdx
	movq	8(%rsp), %rsi
	movq	d_buffer_inv_sbox(%rip), %rdi
	call	_Z38__device_stub__Z14aes_decryptionPhS_S_PhS_S_
.L108:
	movq	8(%rsp), %rsi
	movq	%rbp, %rdi
	movl	$2, %ecx
	movq	%rbx, %rdx
	call	cudaMemcpy@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L109
	call	cudaGetErrorString@PLT
	leaq	.LC4(%rip), %rsi
	movl	$2, %edi
	movq	%rax, %rdx
	xorl	%eax, %eax
	call	__printf_chk@PLT
.L109:
	movq	8(%rsp), %rdi
	call	cudaFree@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L105
	call	cudaGetErrorString@PLT
	leaq	.LC1(%rip), %rsi
	movl	$2, %edi
	movq	%rax, %rdx
	xorl	%eax, %eax
	call	__printf_chk@PLT
.L105:
	movq	40(%rsp), %rax
	subq	%fs:40, %rax
	je	.L111
	call	__stack_chk_fail@PLT
.L111:
	addq	$48, %rsp
	.cfi_def_cfa_offset 32
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3666:
	.size	aes_decryption_wrapper, .-aes_decryption_wrapper
	.section	.rodata.str1.1
.LC5:
	.string	"_Z14aes_decryptionPhS_S_"
.LC6:
	.string	"_Z14aes_encryptionPhS_S_"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB3698:
	.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	_Z14aes_decryptionPhS_S_(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC6(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z14aes_encryptionPhS_S_(%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
.LFE3698:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.nvFatBinSegment,"aw"
	.align 8
	.type	_ZL15__fatDeviceText, @object
	.size	_ZL15__fatDeviceText, 24
_ZL15__fatDeviceText:
	.long	1180844977
	.long	1
	.quad	fatbinData
	.quad	0
	.local	_ZL20__cudaFatCubinHandle
	.comm	_ZL20__cudaFatCubinHandle,8,8
	.globl	Rcon
	.data
	.align 8
	.type	Rcon, @object
	.size	Rcon, 11
Rcon:
	.string	""
	.ascii	"\001\002\004\b\020 @\200\0336"
	.globl	InvSbox
	.align 32
	.type	InvSbox, @object
	.size	InvSbox, 256
InvSbox:
	.string	"R\tj\32506\2458\277@\243\236\201\363\327\373|\3439\202\233/\377\2074\216CD\304\336\351\313T{\2242\246\302#=\356L\225\013B\372\303N\b.\241f(\331$\262v[\242Im\213\321%r\370\366d\206h\230\026\324\244\\\314]e\266\222lpHP\375\355\271\332^\025FW\247\215\235\204\220\330\253"
	.ascii	"\214\274\323\n\367\344X\005\270\263E\006\320,\036\217\312?\017"
	.ascii	"\002\301\257\275\003\001\023\212k:\221\021AOg\334\352\227\362"
	.ascii	"\317\316\360\264\346s\226\254t\"\347\2555\205\342\3717\350\034"
	.ascii	"u\337nG\361\032q\035)\305\211o\267b\016\252\030\276\033\374V"
	.ascii	">K\306\322y \232\333\300\376x\315Z\364\037\335\2503\210\007\307"
	.ascii	"1\261\022\020Y'\200\354_`Q\177\251\031\265J\r-\345z\237\223\311"
	.ascii	"\234\357\240\340;M\256*\365\260\310\353\273<\203S\231a\027+\004"
	.ascii	"~\272w\326&\341i\024cU!\f}"
	.globl	Sbox
	.align 32
	.type	Sbox, @object
	.size	Sbox, 256
Sbox:
	.string	"c|w{\362ko\3050\001g+\376\327\253v\312\202\311}\372YG\360\255\324\242\257\234\244r\300\267\375\223&6?\367\3144\245\345\361q\3301\025\004\307#\303\030\226\005\232\007\022\200\342\353'\262u\t\203,\032\033nZ\240R;\326\263)\343/\204S\321"
	.ascii	"\355 \374\261[j\313\2769JLX\317\320\357\252\373CM3\205E\371\002"
	.ascii	"\177P<\237\250Q\243@\217\222\2358\365\274\266\332!\020\377\363"
	.ascii	"\322\315\f\023\354_\227D\027\304\247~=d]\031s`\201O\334\"*\220"
	.ascii	"\210F\356\270\024\336^\013\333\3402:\nI\006$\\\302\323\254b\221"
	.ascii	"\225\344y\347\3107m\215\325N\251lV\364\352ez\256\b\272x%.\034"
	.ascii	"\246\264\306\350\335t\037K\275\213\212p>\265fH\003\366\016a5"
	.ascii	"W\271\206\301\035\236\341\370\230\021i\331\216\224\233\036\207"
	.ascii	"\351\316U(\337\214\241\211\r\277\346BhA\231-\017\260T\273\026"
	.globl	Key
	.align 16
	.type	Key, @object
	.size	Key, 16
Key:
	.ascii	"\017\025q\311G\331\350Y\f\267\255\326\257\177g\230"
	.globl	SubKeys
	.bss
	.align 32
	.type	SubKeys, @object
	.size	SubKeys, 176
SubKeys:
	.zero	176
	.globl	d_buffer_inv_sbox
	.align 8
	.type	d_buffer_inv_sbox, @object
	.size	d_buffer_inv_sbox, 8
d_buffer_inv_sbox:
	.zero	8
	.globl	d_buffer_sbox
	.align 8
	.type	d_buffer_sbox, @object
	.size	d_buffer_sbox, 8
d_buffer_sbox:
	.zero	8
	.globl	d_buffer_key
	.align 8
	.type	d_buffer_key, @object
	.size	d_buffer_key, 8
d_buffer_key:
	.zero	8
	.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:

#ifndef WORKING_AES_H_
#define WORKING_AES_H_

#include <sys/types.h>
#include <sys/stat.h>
#include <fcntl.h>

#include <stdio.h>
//#include <iostream>
#include <stdlib.h>

#include <math.h>
#include <string.h>
#include <unistd.h>

//#include <helper_timer.h>
//#include <helper_cuda.h>

#define AES_BLOCK_SIZE 16

#ifdef __cplusplus
extern "C" {
#endif


void aes_decryption_wrapper(unsigned char * BufferData, ssize_t buffer_size);
void aes_encryption_wrapper(unsigned char * BufferData, ssize_t buffer_size);
int validate_pad_block(int pad, unsigned char* pad_block);
void aes_keyschedule();
void init_aes_data_on_device();
void free_aes_data_on_device();

#ifdef __cplusplus
}
#endif

int compute_pad_value(int block_length);
int get_pad_value(unsigned char* block_data);
int pad_block(unsigned char* aes_block, char* plain_block, int pad_value);


int file_exists(const char *filename);
void dump_to_file(unsigned char* data, const char* path, size_t size);
void printAsState(unsigned char byteArray[], int length);
#endif

	.text
	.file	"aes_wrapper_gpu.hip"
	.type	__hip_cuid_,@object             # @__hip_cuid_
	.bss
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               # 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_

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

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

	.end_amdgpu_metadata
