I am trying to implement AES on GPU using CUDA programming. I use 4 TBoxes in my implementation that requires 4kB of GPU Memory. I have used a 1KB array for 1KB plaintext. first all of plaintext would be copied to Gpu memory,then encryption would be started ,

cudaMemcpyToSymbolAsync(DEV_message, H_Message, 1024, 0, cudaMemcpyHostToDevice);

then, the global kernel will run.

AESROUND << <8, 16, 16 >> >(1024);

AESROUND performs one round of AES algorithm on a 1024bit State, then 32 bytes of The array will XOR with 32Bs of The State.

__global__ void AESROUND_AD(const int SIZE)
{
    __shared__ unsigned char dev_rkey[16];
    __shared__ unsigned char dev_sh_state[16];

int tid = blockIdx.x * 16 + threadIdx.x;
if (tid < 128)
{
    for (long long i = 0; i <SIZE/32; i++)
    {   
        dev_sh_state[threadIdx.x] = dev_state[(tid + 112) % 128];
        dev_rkey[threadIdx.x] = dev_state[tid];
        __syncthreads();

        if (threadIdx.x < 4)
        {
            U32  v1 = ((U32*)dev_TE0)[*(dev_sh_state + threadIdx.x * 4)];
            U32 v2 = ((U32*)dev_TE1)[*(dev_sh_state + ((threadIdx.x * 4 + 5) % 16))];
            U32 v3 = ((U32*)dev_TE2)[*(dev_sh_state + ((threadIdx.x * 4 + 10) % 16))];
            U32 v4 = ((U32*)dev_TE3)[*(dev_sh_state + ((threadIdx.x * 4 + 15) % 16))];

            ((U32*)dev_sh_state)[threadIdx.x] = v1 ^ v2 ^ v3 ^ v4 ^ ((U32*)dev_rkey)[threadIdx.x];
        }
        __syncthreads();

        dev_state[tid] = dev_sh_state[threadIdx.x];
        __syncthreads();

    }
}

}

there is no problem in code and also I have no Error in my program. But The big problem is its slow speed: just 1 Megabyte Per Second. I dont know why this program is too slow? I dont know what is the problem? I use Microsoft visual C++ in windows 8.1. CUDA Toolkit 7 is used.

My Gpu specifications are : GeForce GT 720M, Compute Capability:2.1 , Number of SMs:2 , Graphic clock (MHz) 625, Processor clock (MHz) 1250 Memory Clock (MHz) 800. my processor is Core(TM) i5-3337U CPU @ 1.80GHz (4 CPUs) microprocessor,

source code of the Algorithm can be found here