summaryrefslogtreecommitdiff
path: root/benchmarks/CUDA/AES/aesDecrypt128_kernel.h
diff options
context:
space:
mode:
Diffstat (limited to 'benchmarks/CUDA/AES/aesDecrypt128_kernel.h')
-rw-r--r--benchmarks/CUDA/AES/aesDecrypt128_kernel.h324
1 files changed, 0 insertions, 324 deletions
diff --git a/benchmarks/CUDA/AES/aesDecrypt128_kernel.h b/benchmarks/CUDA/AES/aesDecrypt128_kernel.h
deleted file mode 100644
index 5634ef8..0000000
--- a/benchmarks/CUDA/AES/aesDecrypt128_kernel.h
+++ /dev/null
@@ -1,324 +0,0 @@
-
-/***************************************************************************
- * 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 <[email protected]>
- */
-
-/* aes decryption operation:
- * Device code.
- *
- */
-
-#ifndef _AESDECRYPT128_KERNEL_H_
-#define _AESDECRYPT128_KERNEL_H_
-
-#include <stdio.h>
-
-// 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<unsigned, 1, cudaReadModeElementType> texDKey128;
-
-__global__ void aesDecrypt128( 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(texDKey128, 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(texDKey128, 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(texDKey128, 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(texDKey128, 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(texDKey128, 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(texDKey128, 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(texDKey128, 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(texDKey128, 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(texDKey128, 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(texDKey128, 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];
-
- x = mod4tx+40;
- keyElem = tex1Dfetch(texDKey128, 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 11th stage --------------------------------
-
- result[BSIZE * bx + tx] = STAGEBLOCK2(tx).uival;
- // end of AES
-
-}
-
-#endif // #ifndef _AESDECRYPT_KERNEL_H_
-
generated by cgit v1.3 (git 2.53.0) at 2026-07-10 22:28:51 +0000