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();

        }
    }
}

I use Microsoft visual C++ in Windows 8.1 with CUDA Toolkit 7.

My GPU is a GeForce GT 720M with 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,

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 1024-bit state, then 32 bytes of the array will XOR with 32 bits 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();

        }
    }
}

• I use Microsoft Visual C++ in Windows 8.1 with CUDA Toolkit 7.
• My GPU is a 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 Intel Core(TM) i5-3337U CPU @ 1.80GHz (4 CPUs) microprocessor.

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

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();

    }
}

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

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();

        }
    }
}

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 ,

}

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

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 using

There is no problem in code or in my program, but the big problem is its speed: just 1 megabyte per second. I don't know why this program is so slow.

I use Microsoft visual C++ in Windows 8.1 with CUDA Toolkit 7.

My GPU is a GeForce GT 720M with 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.