From 69f2911e04ffb1b19eef1fafb8c040af271f656e Mon Sep 17 00:00:00 2001 From: Tor Aamodt Date: Thu, 15 Jul 2010 18:09:46 -0800 Subject: creating branch for adding support for CUDA 3.x and Fermi [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 6829] --- benchmarks/CUDA/AES/aesDecrypt256_kernel.h | 416 +++++++++++++++++++++++++++++ 1 file changed, 416 insertions(+) create mode 100644 benchmarks/CUDA/AES/aesDecrypt256_kernel.h (limited to 'benchmarks/CUDA/AES/aesDecrypt256_kernel.h') diff --git a/benchmarks/CUDA/AES/aesDecrypt256_kernel.h b/benchmarks/CUDA/AES/aesDecrypt256_kernel.h new file mode 100644 index 0000000..243b549 --- /dev/null +++ b/benchmarks/CUDA/AES/aesDecrypt256_kernel.h @@ -0,0 +1,416 @@ + +/*************************************************************************** + * Copyright (C) 2006 * + * * + * * + * This program is free software; you can redistribute it and/or modify * + * it under the terms of the GNU General Public License as published by * + * the Free Software Foundation; either version 2 of the License, or * + * (at your option) any later version. * + * * + * This program is distributed in the hope that it will be useful, * + * but WITHOUT ANY WARRANTY; without even the implied warranty of * + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the * + * GNU General Public License for more details. * + * * + * You should have received a copy of the GNU General Public License * + * along with this program; if not, write to the * + * Free Software Foundation, Inc., * + * 59 Temple Place - Suite 330, Boston, MA 02111-1307, USA. * + ***************************************************************************/ + + +/** + @author Svetlin Manavski + */ + +/* aes decryption operation: + * Device code. + * + */ + +#ifndef _AESDECRYPT_KERNEL_H_ +#define _AESDECRYPT_KERNEL_H_ + +#include + +// Thread block size +#define BSIZE 256 + +#define STAGEBLOCK1(index) CUT_BANK_CHECKER( stageBlock1, index ) +#define STAGEBLOCK2(index) CUT_BANK_CHECKER( stageBlock2, index ) + +#define TBOXE0(index) CUT_BANK_CHECKER( tBox0Block, index ) +#define TBOXE1(index) CUT_BANK_CHECKER( tBox1Block, index ) +#define TBOXE2(index) CUT_BANK_CHECKER( tBox2Block, index ) +#define TBOXE3(index) CUT_BANK_CHECKER( tBox3Block, index ) + +#define INVSBOX(index) CUT_BANK_CHECKER( invSBoxBlock, index ) + +texture texDKey; + +__global__ void aesDecrypt256( unsigned * result, unsigned * inData, int inputSize) +{ + unsigned bx = blockIdx.x; + unsigned tx = threadIdx.x; + unsigned mod4tx = tx%4; + unsigned int4tx = tx/4; + unsigned idx2 = int4tx*4; + int x; + unsigned keyElem; + + __shared__ UByte4 stageBlock1[BSIZE]; + __shared__ UByte4 stageBlock2[BSIZE]; + __shared__ UByte4 tBox0Block[BSIZE]; + __shared__ UByte4 tBox1Block[BSIZE]; + __shared__ UByte4 tBox2Block[BSIZE]; + __shared__ UByte4 tBox3Block[BSIZE]; + __shared__ UByte4 invSBoxBlock[BSIZE]; + + // input caricati in memoria + STAGEBLOCK1(tx).uival = inData[BSIZE * bx + tx ]; + TBOXE0(tx).uival = TBoxi0[tx]; + TBOXE1(tx).uival = TBoxi1[tx]; + TBOXE2(tx).uival = TBoxi2[tx]; + TBOXE3(tx).uival = TBoxi3[tx]; + INVSBOX(tx).ubval[0] = inv_SBox[tx]; + + __syncthreads(); + + + //----------------------------------- 1st stage ----------------------------------- + + x = mod4tx; + keyElem = tex1Dfetch(texDKey, x); + STAGEBLOCK2(tx).uival = STAGEBLOCK1(tx).uival ^ keyElem; + + __syncthreads(); + + //-------------------------------- end of 1st stage -------------------------------- + + + //----------------------------------- 2nd stage ----------------------------------- + + unsigned op1 = STAGEBLOCK2( posIdx_D[16 + mod4tx*4] + idx2 ).ubval[0]; + unsigned op2 = STAGEBLOCK2( posIdx_D[16 + mod4tx*4+1] + idx2 ).ubval[1]; + unsigned op3 = STAGEBLOCK2( posIdx_D[16 + mod4tx*4+2] + idx2 ).ubval[2]; + unsigned op4 = STAGEBLOCK2( posIdx_D[16 + mod4tx*4+3] + idx2 ).ubval[3]; + + op1 = TBOXE0(op1).uival; + + op2 = TBOXE1(op2).uival; + + op3 = TBOXE2(op3).uival; + + op4 = TBOXE3(op4).uival; + + x = mod4tx+4; + keyElem = tex1Dfetch(texDKey, x); + STAGEBLOCK1(tx).uival = op1^op2^op3^op4^keyElem; + + __syncthreads(); + + //-------------------------------- end of 2nd stage -------------------------------- + + //----------------------------------- 3th stage ----------------------------------- + + op1 = STAGEBLOCK1( posIdx_D[16 + mod4tx*4] + idx2 ).ubval[0]; + op2 = STAGEBLOCK1( posIdx_D[16 + mod4tx*4+1] + idx2 ).ubval[1]; + op3 = STAGEBLOCK1( posIdx_D[16 + mod4tx*4+2] + idx2 ).ubval[2]; + op4 = STAGEBLOCK1( posIdx_D[16 + mod4tx*4+3] + idx2 ).ubval[3]; + + op1 = TBOXE0(op1).uival; + + op2 = TBOXE1(op2).uival; + + op3 = TBOXE2(op3).uival; + + op4 = TBOXE3(op4).uival; + + x = mod4tx+8; + keyElem = tex1Dfetch(texDKey, x); + STAGEBLOCK2(tx).uival = op1^op2^op3^op4^keyElem; + + __syncthreads(); + + //-------------------------------- end of 3th stage -------------------------------- + + //----------------------------------- 4th stage ----------------------------------- + + op1 = STAGEBLOCK2( posIdx_D[16 + mod4tx*4] + idx2 ).ubval[0]; + op2 = STAGEBLOCK2( posIdx_D[16 + mod4tx*4+1] + idx2 ).ubval[1]; + op3 = STAGEBLOCK2( posIdx_D[16 + mod4tx*4+2] + idx2 ).ubval[2]; + op4 = STAGEBLOCK2( posIdx_D[16 + mod4tx*4+3] + idx2 ).ubval[3]; + + op1 = TBOXE0(op1).uival; + + op2 = TBOXE1(op2).uival; + + op3 = TBOXE2(op3).uival; + + op4 = TBOXE3(op4).uival; + + x = mod4tx+12; + keyElem = tex1Dfetch(texDKey, x); + STAGEBLOCK1(tx).uival = op1^op2^op3^op4^keyElem; + + __syncthreads(); + + //-------------------------------- end of 4th stage -------------------------------- + + //----------------------------------- 5th stage ----------------------------------- + + op1 = STAGEBLOCK1( posIdx_D[16 + mod4tx*4] + idx2 ).ubval[0]; + op2 = STAGEBLOCK1( posIdx_D[16 + mod4tx*4+1] + idx2 ).ubval[1]; + op3 = STAGEBLOCK1( posIdx_D[16 + mod4tx*4+2] + idx2 ).ubval[2]; + op4 = STAGEBLOCK1( posIdx_D[16 + mod4tx*4+3] + idx2 ).ubval[3]; + + op1 = TBOXE0(op1).uival; + + op2 = TBOXE1(op2).uival; + + op3 = TBOXE2(op3).uival; + + op4 = TBOXE3(op4).uival; + + x = mod4tx+16; + keyElem = tex1Dfetch(texDKey, x); + STAGEBLOCK2(tx).uival = op1^op2^op3^op4^keyElem; + + __syncthreads(); + + //-------------------------------- end of 5th stage -------------------------------- + + //----------------------------------- 6th stage ----------------------------------- + + op1 = STAGEBLOCK2( posIdx_D[16 + mod4tx*4] + idx2 ).ubval[0]; + op2 = STAGEBLOCK2( posIdx_D[16 + mod4tx*4+1] + idx2 ).ubval[1]; + op3 = STAGEBLOCK2( posIdx_D[16 + mod4tx*4+2] + idx2 ).ubval[2]; + op4 = STAGEBLOCK2( posIdx_D[16 + mod4tx*4+3] + idx2 ).ubval[3]; + + op1 = TBOXE0(op1).uival; + + op2 = TBOXE1(op2).uival; + + op3 = TBOXE2(op3).uival; + + op4 = TBOXE3(op4).uival; + + x = mod4tx+20; + keyElem = tex1Dfetch(texDKey, x); + STAGEBLOCK1(tx).uival = op1^op2^op3^op4^keyElem; + + __syncthreads(); + + //-------------------------------- end of 6th stage -------------------------------- + + //----------------------------------- 7th stage ----------------------------------- + + op1 = STAGEBLOCK1( posIdx_D[16 + mod4tx*4] + idx2 ).ubval[0]; + op2 = STAGEBLOCK1( posIdx_D[16 + mod4tx*4+1] + idx2 ).ubval[1]; + op3 = STAGEBLOCK1( posIdx_D[16 + mod4tx*4+2] + idx2 ).ubval[2]; + op4 = STAGEBLOCK1( posIdx_D[16 + mod4tx*4+3] + idx2 ).ubval[3]; + + op1 = TBOXE0(op1).uival; + + op2 = TBOXE1(op2).uival; + + op3 = TBOXE2(op3).uival; + + op4 = TBOXE3(op4).uival; + + x = mod4tx+24; + keyElem = tex1Dfetch(texDKey, x); + STAGEBLOCK2(tx).uival = op1^op2^op3^op4^keyElem; + + __syncthreads(); + + //-------------------------------- end of 7th stage -------------------------------- + + //----------------------------------- 8th stage ----------------------------------- + + op1 = STAGEBLOCK2( posIdx_D[16 + mod4tx*4] + idx2 ).ubval[0]; + op2 = STAGEBLOCK2( posIdx_D[16 + mod4tx*4+1] + idx2 ).ubval[1]; + op3 = STAGEBLOCK2( posIdx_D[16 + mod4tx*4+2] + idx2 ).ubval[2]; + op4 = STAGEBLOCK2( posIdx_D[16 + mod4tx*4+3] + idx2 ).ubval[3]; + + op1 = TBOXE0(op1).uival; + + op2 = TBOXE1(op2).uival; + + op3 = TBOXE2(op3).uival; + + op4 = TBOXE3(op4).uival; + + x = mod4tx+28; + keyElem = tex1Dfetch(texDKey, x); + STAGEBLOCK1(tx).uival = op1^op2^op3^op4^keyElem; + + __syncthreads(); + + //-------------------------------- end of 8th stage -------------------------------- + + //----------------------------------- 9th stage ----------------------------------- + + op1 = STAGEBLOCK1( posIdx_D[16 + mod4tx*4] + idx2 ).ubval[0]; + op2 = STAGEBLOCK1( posIdx_D[16 + mod4tx*4+1] + idx2 ).ubval[1]; + op3 = STAGEBLOCK1( posIdx_D[16 + mod4tx*4+2] + idx2 ).ubval[2]; + op4 = STAGEBLOCK1( posIdx_D[16 + mod4tx*4+3] + idx2 ).ubval[3]; + + op1 = TBOXE0(op1).uival; + + op2 = TBOXE1(op2).uival; + + op3 = TBOXE2(op3).uival; + + op4 = TBOXE3(op4).uival; + + x = mod4tx+32; + keyElem = tex1Dfetch(texDKey, x); + STAGEBLOCK2(tx).uival = op1^op2^op3^op4^keyElem; + + __syncthreads(); + + //-------------------------------- end of 9th stage -------------------------------- + + //----------------------------------- 10th stage ----------------------------------- + + op1 = STAGEBLOCK2( posIdx_D[16 + mod4tx*4] + idx2 ).ubval[0]; + op2 = STAGEBLOCK2( posIdx_D[16 + mod4tx*4+1] + idx2 ).ubval[1]; + op3 = STAGEBLOCK2( posIdx_D[16 + mod4tx*4+2] + idx2 ).ubval[2]; + op4 = STAGEBLOCK2( posIdx_D[16 + mod4tx*4+3] + idx2 ).ubval[3]; + + op1 = TBOXE0(op1).uival; + + op2 = TBOXE1(op2).uival; + + op3 = TBOXE2(op3).uival; + + op4 = TBOXE3(op4).uival; + + x = mod4tx+36; + keyElem = tex1Dfetch(texDKey, x); + STAGEBLOCK1(tx).uival = op1^op2^op3^op4^keyElem; + + __syncthreads(); + + //-------------------------------- end of 10th stage -------------------------------- + + //----------------------------------- 11th stage ----------------------------------- + + op1 = STAGEBLOCK1( posIdx_D[16 + mod4tx*4] + idx2 ).ubval[0]; + op2 = STAGEBLOCK1( posIdx_D[16 + mod4tx*4+1] + idx2 ).ubval[1]; + op3 = STAGEBLOCK1( posIdx_D[16 + mod4tx*4+2] + idx2 ).ubval[2]; + op4 = STAGEBLOCK1( posIdx_D[16 + mod4tx*4+3] + idx2 ).ubval[3]; + + op1 = TBOXE0(op1).uival; + + op2 = TBOXE1(op2).uival; + + op3 = TBOXE2(op3).uival; + + op4 = TBOXE3(op4).uival; + + x = mod4tx+40; + keyElem = tex1Dfetch(texDKey, x); + STAGEBLOCK2(tx).uival = op1^op2^op3^op4^keyElem; + + __syncthreads(); + + //-------------------------------- end of 11th stage -------------------------------- + + //----------------------------------- 12th stage ----------------------------------- + + op1 = STAGEBLOCK2( posIdx_D[16 + mod4tx*4] + idx2 ).ubval[0]; + op2 = STAGEBLOCK2( posIdx_D[16 + mod4tx*4+1] + idx2 ).ubval[1]; + op3 = STAGEBLOCK2( posIdx_D[16 + mod4tx*4+2] + idx2 ).ubval[2]; + op4 = STAGEBLOCK2( posIdx_D[16 + mod4tx*4+3] + idx2 ).ubval[3]; + + op1 = TBOXE0(op1).uival; + + op2 = TBOXE1(op2).uival; + + op3 = TBOXE2(op3).uival; + + op4 = TBOXE3(op4).uival; + + x = mod4tx+44; + keyElem = tex1Dfetch(texDKey, x); + STAGEBLOCK1(tx).uival = op1^op2^op3^op4^keyElem; + + __syncthreads(); + + //-------------------------------- end of 12th stage -------------------------------- + + //----------------------------------- 13th stage ----------------------------------- + + op1 = STAGEBLOCK1( posIdx_D[16 + mod4tx*4] + idx2 ).ubval[0]; + op2 = STAGEBLOCK1( posIdx_D[16 + mod4tx*4+1] + idx2 ).ubval[1]; + op3 = STAGEBLOCK1( posIdx_D[16 + mod4tx*4+2] + idx2 ).ubval[2]; + op4 = STAGEBLOCK1( posIdx_D[16 + mod4tx*4+3] + idx2 ).ubval[3]; + + op1 = TBOXE0(op1).uival; + + op2 = TBOXE1(op2).uival; + + op3 = TBOXE2(op3).uival; + + op4 = TBOXE3(op4).uival; + + x = mod4tx+48; + keyElem = tex1Dfetch(texDKey, x); + STAGEBLOCK2(tx).uival = op1^op2^op3^op4^keyElem; + + __syncthreads(); + + //-------------------------------- end of 13th stage -------------------------------- + + //----------------------------------- 14th stage ----------------------------------- + + op1 = STAGEBLOCK2( posIdx_D[16 + mod4tx*4] + idx2 ).ubval[0]; + op2 = STAGEBLOCK2( posIdx_D[16 + mod4tx*4+1] + idx2 ).ubval[1]; + op3 = STAGEBLOCK2( posIdx_D[16 + mod4tx*4+2] + idx2 ).ubval[2]; + op4 = STAGEBLOCK2( posIdx_D[16 + mod4tx*4+3] + idx2 ).ubval[3]; + + op1 = TBOXE0(op1).uival; + + op2 = TBOXE1(op2).uival; + + op3 = TBOXE2(op3).uival; + + op4 = TBOXE3(op4).uival; + + x = mod4tx+52; + keyElem = tex1Dfetch(texDKey, x); + STAGEBLOCK1(tx).uival = op1^op2^op3^op4^keyElem; + + __syncthreads(); + + //-------------------------------- end of 14th stage -------------------------------- + + //----------------------------------- 15th stage ----------------------------------- + + op1 = STAGEBLOCK1( posIdx_D[16 + mod4tx*4] + idx2 ).ubval[0]; + op2 = STAGEBLOCK1( posIdx_D[16 + mod4tx*4+1] + idx2 ).ubval[1]; + op3 = STAGEBLOCK1( posIdx_D[16 + mod4tx*4+2] + idx2 ).ubval[2]; + op4 = STAGEBLOCK1( posIdx_D[16 + mod4tx*4+3] + idx2 ).ubval[3]; + + x = mod4tx+56; + keyElem = tex1Dfetch(texDKey, x); + + STAGEBLOCK2(tx).ubval[3] = INVSBOX(op4).ubval[0]^( keyElem>>24); + STAGEBLOCK2(tx).ubval[2] = INVSBOX(op3).ubval[0]^( keyElem>>16 & 0x000000FF); + STAGEBLOCK2(tx).ubval[1] = INVSBOX(op2).ubval[0]^( keyElem>>8 & 0x000000FF); + STAGEBLOCK2(tx).ubval[0] = INVSBOX(op1).ubval[0]^( keyElem & 0x000000FF); + + __syncthreads(); + + //-------------------------------- end of 15th stage -------------------------------- + + result[BSIZE * bx + tx] = STAGEBLOCK2(tx).uival; + // end of AES + +} + +#endif // #ifndef _AESDECRYPT_KERNEL_H_ + -- cgit v1.3