Implementation of AES using CUDA

Question

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

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

        }
    }
}

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
  • 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.

Source code of the algorithm can be found here.

Answer 1

There are a few issues that I see in your code. I'm not sure if all matter, but I'll try to list them anyway.

1. You said you call AESROUND<<< 8, 16, 16 >>>(1024);, but you present __global__ void AESROUND_AD(const int SIZE). Are AESROUND() and AESROUND_AD() the same kernels? Is that a typo?

2. Calling AESROUND<<< 8, 16, 16 >>>(1024); means (amongst other) that you dynamically allocate 16 (third parameter of the <<< >>>) elements of a shared memory segment of type to-be-defined inside the kernel (declared extern __shared__ whatever mySegment[];). So you allocate that segment in your kernel call, but do not actually declare in the kernel body. Yet another typo?

3. In the kernel itself, you reference dev_state which isn't declared anywhere. Is that a global variable? Shouldn't that be better put a input parameter of the kernel?

4. The computation of tid explicitly assumes that the block size is 16, which is true here. However, if this changes when calling the kernel, then the formula would become false. So if this isn't supposed to change, then use a macro for the block size (such as #define BLOCKSIZE 16) and use it in your definition of tid, and the call to AESROUND(). But if this size might change in the call of the kernel, then use tid = blockIdx.x * blockDim.x + threadIdx.x;. And as a more general comment, try to avoid "magic numbers" in your code, or at least put them into macros or const variables, so that you give them a sensible name and therefore a meaning.

Now, the initial question was about performance:
Well, I guess that this all boils down to one simple issue: you do not maximise neither your core occupancy, nor your memory bandwidth. Indeed, you only request 16 threads per block, which isn't even enough for a warp. Therefore, irrespective of your actual GPU architecture, at any moment, at most only half of your cores are utilised. Moreover, once the initial data loaded in shared memory, you only compute using 4 threads, so only utilising 1/4 of the available cores at most.

Considering what I can see (and sort of understand) of your kernel, I have the feeling that you'd better be using one single block of 128 threads per message to compute, and to load all of it at once in your shared memory. Then to compute as you do in your kernel, but with now 4 times more threads involved. This would have the double benefit of permitting you to avoid having to store you partial results in global memory at each round, and to increase your occupancy.

And to fully utilise your GPU, you would need to compute several messages in parallel, one per block.