Changes

# [mailto:nmayuranathan@myseneca.ca?subject=gpu610 Natesh Mayuranathan ][mailto:pAziz1@myseneca.ca,smohammad8@myseneca.ca,nmayuranathan@myseneca.ca?subject=dps901-gpu610 Email All]

==[[User:Philip_A_Aziz | Philip Aziz]]:==

 ==[[User:Saad_Mohammad| Saad Mohammad]]:==

==Assignment 3 ==<pre> diff --git a/accessjpeg.cpp b/accessjpeg.cpp index f967465..2bfb7eb 100644 --- a/accessjpeg.cpp +++ b/accessjpeg.cpp @@ -1,7 +1,7 @@  /********************************************************  Winter 2013   GPU610 Assignment - Author: Saad Mohammad, Phillip Aziz, Natesh Mayuranathan + Author: Saad Mohammad  Team: NullPointerException    Original Project: <https://github.com/markwatson/Image-Encrypt> @@ -93,3 +93,13 @@ char * accessJpeg::accessBlock()  cursor += BLOCK_SIZE;  return ret;  } + +int accessJpeg::getNumberOfBlocks(){ + return ((size - cursor)/BLOCK_SIZE) + 1; +} + +void accessJpeg::goToStart() +{ + cursor += 0; + jumpToStart(); +} diff --git a/accessjpeg.h b/accessjpeg.h index 5293bd9..fe87d0d 100644 --- a/accessjpeg.h +++ b/accessjpeg.h @@ -1,7 +1,7 @@  /********************************************************  Winter 2013   GPU610 Assignment - Author: Saad Mohammad, Phillip Aziz, Natesh Mayuranathan + Author: Saad Mohammad  Team: NullPointerException    Original Project: <https://github.com/markwatson/Image-Encrypt> @@ -29,6 +29,8 @@ public:  bool jumpToStart();  char * accessBlock();  bool hasMore(); + int getNumberOfBlocks(); + void goToStart();  };    #endif // ACCESSJPEG_H diff --git a/aesencrypt.cu b/aesencrypt.cu index 4c89090..fc4393b 100644 --- a/aesencrypt.cu +++ b/aesencrypt.cu @@ -1,7 +1,7 @@  /********************************************************  Winter 2013   GPU610 Assignment  -Author: Saad Mohammad, Phillip Aziz, Natesh Mayuranathan +Author: Saad Mohammad  Team: NullPointerException    Original Project: <https://github.com/markwatson/Image-Encrypt> @@ -19,6 +19,8 @@ Original Author: Mark Watson  #include "aesencrypt.cuh"  using namespace std;   +const int NTPB = 1024; //number of threads per block +  aesEncrypt::aesEncrypt()  {  // set values @@ -71,105 +73,135 @@ aesEncrypt::~aesEncrypt()  cleanDeviceMemory();  }   -__global__ void xorRoundKeyGlobal(char * state, char * key) -{ - int idx = threadIdx.x; - state[idx] = (int)((int)state[idx] ^ (int)key[idx]); -} - -char* aesEncrypt::copyChar(char* source, int size){ +__global__ void encryptKernel(char **state, char *key, unsigned char *sbBox, int Nb, int Nr, int size){ + int position = 0; + int idx = blockIdx.x * blockDim.x + threadIdx.x;   + if(idx < size){ + char* block = state[idx]; +  + //CALL for: xorRoundKey(block, expandedkey) + for(int i=0;i<Nb;i++){ + block[i] = (int)((int)block[i] ^ (int)key[i + position]); + } +  + for(int cnt=1; cnt <= (Nr-1); cnt++){ + //CALL for: subBytes(block)) + for(int i=0;i<Nb;i++){ + int n = static_cast<unsigned char>(block[i]); + block[i] = sbBox[n]; + }   - return NULL; -} + //CALL for: shiftRows(block); + for(int i=0;i<4;i++){ + int row = i + 1; + int j = row *4; + unsigned char *tmp = (unsigned char*) block +j; +  + for(int i=1; i < (row+1); i++) + { + unsigned char trimmings; + trimmings = tmp[0]; + tmp[0] = tmp[1]; + tmp[1] = tmp[2]; + tmp[2] = tmp[3]; + tmp[3] = trimmings; + }  + }   -__global__ void subBytesCuda(char * state, unsigned char * sbBox) -{ - int idx = threadIdx.x; - int n = static_cast<unsigned char>(state[idx]); - state[idx] = sbBox[n]; -} + //CALL for: xorRoundKey(block, expandedkey) + position = cnt * Nb; + for(int i=0;i<Nb;i++){ + block[i] = (int)((int)block[i] ^ (int)key[i + position]); + }   -__global__ void shiftRowsCuda(char* state) -{ - int row = threadIdx.x + 1; - int j = row *4; - unsigned char *tmp = (unsigned char*) state +j; + position = 0; + }   - for(int i=1; i < (row+1); i++) - { - unsigned char trimmings; - trimmings = tmp[0]; - tmp[0] = tmp[1]; - tmp[1] = tmp[2]; - tmp[2] = tmp[3]; - tmp[3] = trimmings; - - }  -} + //CALL for: subBytes(block)) + for(int i=0;i<Nb;i++){ + int n = static_cast<unsigned char>(block[i]); + block[i] = sbBox[n]; + }   -bool aesEncrypt::encryptBlock(char* block) -{ - //Initalize device memory - char* deviceBlock; - char* deviceTempChar; - cudaMalloc((void**)&deviceBlock, sizeof(char) * (BLOCK_SIZE)); -  - //Copy memory to device - cudaMemcpy(deviceBlock, block,sizeof( char) * (BLOCK_SIZE), cudaMemcpyHostToDevice); + //CALL for: shiftRows(block); + for(int i=0;i<4;i++){ + int row = i + 1; + int j = row *4; + unsigned char *tmp = (unsigned char*) block +j;   - //CUDA CALL for: xorRoundKey(block, expandedkey) - xorRoundKeyGlobal<<<1, Nb>>>(deviceBlock, deviceExpandedKey);  - cudaDeviceSynchronize(); + for(int i=1; i < (row+1); i++) + { + unsigned char trimmings; + trimmings = tmp[0]; + tmp[0] = tmp[1]; + tmp[1] = tmp[2]; + tmp[2] = tmp[3 ]; + tmp[3] = trimmings; + }  + }   - for(int cnt = 1;cnt <= (Nr-1);cnt++) - { - //CUDA CALL for: subBytes(block)) - subBytesCuda<<<1,Nb>>>(deviceBlock, deviceBoxify); - cudaDeviceSynchronize(); -  - //CUDA CALL for: shiftRows(block); - shiftRowsCuda<<<1,4>>>(deviceBlock); - cudaDeviceSynchronize(); -  - //mixColumns(block); // Disabled -  - //CUDA CALL for: xorRoundKey(block, expandedkey + (cnt * Nb))) - cudaMalloc((void**)&deviceTempChar, sizeof(char)*(Nb)); - cudaMemcpy(deviceTempChar, expandedkey + (cnt * Nb), sizeof(char) * Nb, cudaMemcpyHostToDevice); + //CALL for: xorRoundKey(block, expandedkey) + position = Nr * Nb; + for(int i=0;i<Nb;i++){ + block[i] = (int)((int)block[i] ^ (int)key[i + position]); + } + } +}   - xorRoundKeyGlobal<<<1, Nb>>>(deviceBlock, deviceTempChar); //xorRoundKey(block, expandedkey); - cudaDeviceSynchronize(); -  - cudaFree(deviceTempChar); +/* + * Assigns the location of eachBlock into the 2D pointer (block) + */ +__global__ void copyToBlocks(char** block, char* eachBlock, int index){ + char *b = eachBlock; + block[index] = b; +} + +/* + * Returns the value of block[index] into eachBlock + */ +__global__ void copyFromBlocks(char** block, char* eachBlock, int index){ + for(int i=0; i < BLOCK_SIZE; i++){ + eachBlock[i] = block[index][i];  } - //CUDA CALL for: subBytes(block)) - subBytesCuda<<<1,Nb>>>(deviceBlock, deviceBoxify); - cudaDeviceSynchronize(); +}   - //CUDA CALL for: shiftRows(block); - shiftRowsCuda<<<1,4>>>(deviceBlock); - cudaDeviceSynchronize(); +void aesEncrypt::encryptBlock(char** blocks, int size) +{ + int NBLKS = (size + NTPB - 1) / NTPB;   - //CUDA CALL for: xorRoundKey(block, expandedkey + (Nr * Nb)) - cudaMalloc((void**)&deviceTempChar, sizeof(char)*(Nb+1)); - cudaMemcpy(deviceTempChar, expandedkey + (Nr * Nb), sizeof(char) * (Nb +1), cudaMemcpyHostToDevice); + //Initalize device memory + char** deviceBlock; + cudaMalloc((void**)&deviceBlock, size * sizeof(char*));   - xorRoundKeyGlobal<<<1, Nb>>>(deviceBlock, deviceTempChar); //xorRoundKey(block, expandedkey); ////////////////////////////////////////////////////////IMPORTANT - cudaDeviceSynchronize(); + //Copy host block to device block + for(int i=0; i < size; i++){ + char* eachBlock;   + cudaMalloc((void**)&eachBlock, BLOCK_SIZE * sizeof(char)); + cudaMemcpy(eachBlock, blocks[i] ,sizeof(char) * (BLOCK_SIZE), cudaMemcpyHostToDevice); + copyToBlocks<<<1,1>>>(deviceBlock, eachBlock, i); + }   - //END_FUNCTION - cudaMemcpy(block, deviceBlock, sizeof(char)* (BLOCK_SIZE), cudaMemcpyDeviceToHost); - - //Free memory - cudaFree(deviceBlock); - cudaFree(deviceTempChar); + //Encrypt blocks in parallel + encryptKernel<<<NBLKS, NTPB>>>(deviceBlock, deviceExpandedKey, deviceBoxify, Nb, Nr, size);   + //Copy back to host block + for(int i=0; i < size; i++){ + char *hostBlock = new char [BLOCK_SIZE]; + char* retBlock;   - return true; // could error check in future -} + cudaMalloc((void**)&retBlock, BLOCK_SIZE * sizeof(char)); + copyFromBlocks<<<1,1>>>(deviceBlock, retBlock, i); + cudaMemcpy(hostBlock, retBlock ,sizeof(char) * (BLOCK_SIZE), cudaMemcpyDeviceToHost); +  + blocks[i] = hostBlock; + + cudaFree(retBlock); //Clean from memory + }   + cudaFree(deviceBlock); //Clean from memory +}    bool aesEncrypt::decryptBlock(char* block)  { @@ -181,7 +213,6 @@ bool aesEncrypt::decryptBlock(char* block)  invShiftRows(block);  invSubBytes(block);  xorRoundKey(block, expandedkey + cnt * Nb); - //invMixColumns(block); // Disabled  }    invShiftRows(block); @@ -215,10 +246,7 @@ void aesEncrypt::setTextKey(std::string key)  }    expandKey(); -  - //Save host expanded key to device - cudaMalloc((void**)&deviceExpandedKey, sizeof(char) * (strlen(expandedkey) +1)); - cudaMemcpy(deviceExpandedKey, expandedkey, sizeof(char) * (strlen(expandedkey) +1), cudaMemcpyHostToDevice); +  }    void aesEncrypt::expandKey() @@ -283,6 +311,10 @@ void aesEncrypt::expandKey()  }    delete[] temp2; // cleanup + + //Save host expanded key to device + cudaMalloc((void**)&deviceExpandedKey, sizeof(char) * (4 * limit)); + cudaMemcpy(deviceExpandedKey, expandedkey, sizeof(char) * (4 * limit), cudaMemcpyHostToDevice);  }    void aesEncrypt::subBytes(char * state) @@ -370,13 +402,9 @@ void aesEncrypt::mixColumns(char * state)    void aesEncrypt::xorRoundKey(char * state, char * key)  { - //cout << "EQ: \t\t" ;  for (int cnt = 0;cnt < Nb;cnt++)  { - //cout << (int) state[cnt] << " ^ "  - //cout << (int) key[cnt] << " ";//"= ";  state[cnt] ^= key[cnt]; - //cout << (int) state[cnt] << endl;  }    //cout << endl; diff --git a/aesencrypt.cuh b/aesencrypt.cuh index 7de2276..d340dcb 100644 --- a/aesencrypt.cuh +++ b/aesencrypt.cuh @@ -1,7 +1,7 @@  /********************************************************  Winter 2013   GPU610 Assignment - Author: Saad Mohammad, Phillip Aziz, Natesh Mayuranathan + Author: Saad Mohammad  Team: NullPointerException    Original Project: <https://github.com/markwatson/Image-Encrypt> @@ -11,10 +11,10 @@    #ifndef AESENCRYPT_H  #define AESENCRYPT_H -#define BLOCK_SIZE 16    #include <string>  #include <bitset> +#include "accessjpeg.h"  #include "stdio.h"    class aesEncrypt @@ -37,7 +37,7 @@ private:  public:  aesEncrypt();  ~aesEncrypt(); - bool encryptBlock(char*); + void encryptBlock(char**, int);  bool decryptBlock(char*);  void setTextKey(std::string);  void expandKey(); diff --git a/encryptjpeg.cpp b/encryptjpeg.cpp index a722b0d..2692aa5 100644 --- a/encryptjpeg.cpp +++ b/encryptjpeg.cpp @@ -1,7 +1,7 @@  /********************************************************  Winter 2013   GPU610 Assignment - Author: Saad Mohammad, Phillip Aziz, Natesh Mayuranathan + Author: Saad Mohammad  Team: NullPointerException    Original Project: <https://github.com/markwatson/Image-Encrypt> @@ -14,6 +14,8 @@  #include "accessjpeg.h"  #include "aesencrypt.cuh"  #include "encryptjpeg.h" +#include "cuda_runtime.h" +using namespace std;    bool encryptJpeg::process(char action)  { @@ -31,26 +33,45 @@ bool encryptJpeg::process(char action)    // set the key  encrypter.setTextKey(key); -  - // do the encryption - while (file.hasMore()) + + if (action == 'e') + { + int size = file.getNumberOfBlocks(); + char **state = new char*[size]; + int count = 0; + + for(int i=0; i <size; i++) + { + char *temp = new char[BLOCK_SIZE+8]; + ptr = file.accessBlock(); + memcpy(temp, ptr, BLOCK_SIZE); + state[i] = temp; + } + + encrypter.encryptBlock(state, size); // run the encryption +  + file.goToStart(); + while (file.hasMore()) + { + ptr = file.accessBlock(); + memcpy(ptr, state[count], BLOCK_SIZE); + delete state[count++]; //Delete from memory + } + + delete state; //Delete from memory + + } else if (action == 'd') + { + while (file.hasMore()){ + ptr = file.accessBlock(); + encrypter.decryptBlock(ptr); + } + } else  { - ptr = file.accessBlock(); // get a block of the image - if (action == 'e') - { - encrypter.encryptBlock(ptr); // run the encryption - } - else if (action == 'd') - { - encrypter.decryptBlock(ptr); - } - else - { - throw invalidAction(); - return false; - } - } -  + throw invalidAction(); + return false; + } +  // write out the file  out_success = file.writeOutFile(out_file);  if (!out_success) diff --git a/encryptjpeg.h b/encryptjpeg.h index f412ec0..208b0d5 100644 --- a/encryptjpeg.h +++ b/encryptjpeg.h @@ -1,7 +1,7 @@  /********************************************************  Winter 2013   GPU610 Assignment - Author: Saad Mohammad, Phillip Aziz, Natesh Mayuranathan + Author: Saad Mohammad  Team: NullPointerException    Original Project: <https://github.com/markwatson/Image-Encrypt> diff --git a/main.cu b/main.cu index 37707d2..33a9a54 100644 --- a/main.cu +++ b/main.cu @@ -1,7 +1,7 @@  /********************************************************  Winter 2013   GPU610 Assignment - Author: Saad Mohammad, Phillip Aziz, Natesh Mayuranathan + Author: Saad Mohammad  Team: NullPointerException    Original Project: <https://github.com/markwatson/Image-Encrypt> @@ -15,6 +15,7 @@  #include <ctime>  #include "encryptjpeg.h"  #include "aesencrypt.cuh" +#include "cuda_runtime.h"  using namespace std;    int main(int argc, char *argv[]) @@ -72,6 +73,7 @@ int main(int argc, char *argv[])  cout << "The program encountered an internal error..." << endl;  }   + cudaDeviceReset();    return 0;  } </pre> ==Final Results ==  [[Image:Smohammad_project_outcome.jpg]]