Implementation of AES using CUDA
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.


You said you call AESROUND&lt;&lt;&lt; 8, 16, 16 &gt;&gt;&gt;(1024);, but you present __global__ void AESROUND_AD(const int SIZE). Are AESROUND() and AESROUND_AD() the same kernels? Is that a typo?
Calling AESROUND&lt;&lt;&lt; 8, 16, 16 &gt;&gt;&gt;(1024); means (amongst other) that you dynamically allocate 16 (third parameter of the &lt;&lt;&lt; &gt;&gt;&gt;) 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?
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?
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.