From 11b308e7363e937966b035b4891db32b4eece3bf Mon Sep 17 00:00:00 2001 From: Tor Aamodt Date: Fri, 1 Oct 2010 08:55:28 -0800 Subject: integrating recent changes from fermi-test into fermi (i'll use "fermi" for more disruptive changes to the pipeline model such as updating the MSHRs and getting rid of the warp tracker, ripping out DWF, etc...) [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7805] --- benchmarks/CUDA/AES/aesEncrypt256_kernel.h | 416 ----------------------------- 1 file changed, 416 deletions(-) delete mode 100644 benchmarks/CUDA/AES/aesEncrypt256_kernel.h (limited to 'benchmarks/CUDA/AES/aesEncrypt256_kernel.h') diff --git a/benchmarks/CUDA/AES/aesEncrypt256_kernel.h b/benchmarks/CUDA/AES/aesEncrypt256_kernel.h deleted file mode 100644 index 47f82a4..0000000 --- a/benchmarks/CUDA/AES/aesEncrypt256_kernel.h +++ /dev/null @@ -1,416 +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 - */ - -/* aes encryption operation: - * Device code. - * - */ - -#ifndef _AESENCRYPT_KERNEL_H_ -#define _AESENCRYPT_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 ) - -texture texEKey; - -__global__ void aesEncrypt256( 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[256]; - __shared__ UByte4 tBox1Block[256]; - __shared__ UByte4 tBox2Block[256]; - __shared__ UByte4 tBox3Block[256]; - - // input caricati in memoria - STAGEBLOCK1(tx).uival = inData[BSIZE * bx + tx ]; - - unsigned elemPerThread = 256/BSIZE; - for (unsigned cnt=0; cnt>24); - STAGEBLOCK2(tx).ubval[2] = TBOXE1(op3).ubval[3]^( (keyElem>>16) & 0x000000FF); - STAGEBLOCK2(tx).ubval[1] = TBOXE1(op2).ubval[3]^( (keyElem>>8) & 0x000000FF); - STAGEBLOCK2(tx).ubval[0] = TBOXE1(op1).ubval[3]^( keyElem & 0x000000FF); - - __syncthreads(); - - //-------------------------------- end of 15th stage -------------------------------- - - result[BSIZE * bx + tx] = STAGEBLOCK2(tx).uival; - // end of AES - -} - -#endif // #ifndef _AESENCRYPT_KERNEL_H_ -- cgit v1.3