1
edit
Changes
→Assignment 3
=== 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>