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/STO/Makefile | 52 -- benchmarks/CUDA/STO/README.GPGPU-Sim | 2 - benchmarks/CUDA/STO/cust.h | 153 ----- benchmarks/CUDA/STO/main.cu | 867 ------------------------ benchmarks/CUDA/STO/md5_cpu.c | 397 ----------- benchmarks/CUDA/STO/md5_cpu.h | 98 --- benchmarks/CUDA/STO/md5_kernel.cu | 995 ---------------------------- benchmarks/CUDA/STO/sha1_cpu.c | 429 ------------ benchmarks/CUDA/STO/sha1_cpu.h | 98 --- benchmarks/CUDA/STO/sha1_kernel.cu | 1140 -------------------------------- benchmarks/CUDA/STO/storeCPU.c | 1114 ------------------------------- benchmarks/CUDA/STO/storeCPU.h | 195 ------ benchmarks/CUDA/STO/storeGPU.cu | 1211 ---------------------------------- benchmarks/CUDA/STO/storeGPU.h | 224 ------- 14 files changed, 6975 deletions(-) delete mode 100644 benchmarks/CUDA/STO/Makefile delete mode 100644 benchmarks/CUDA/STO/README.GPGPU-Sim delete mode 100644 benchmarks/CUDA/STO/cust.h delete mode 100644 benchmarks/CUDA/STO/main.cu delete mode 100644 benchmarks/CUDA/STO/md5_cpu.c delete mode 100644 benchmarks/CUDA/STO/md5_cpu.h delete mode 100644 benchmarks/CUDA/STO/md5_kernel.cu delete mode 100644 benchmarks/CUDA/STO/sha1_cpu.c delete mode 100644 benchmarks/CUDA/STO/sha1_cpu.h delete mode 100644 benchmarks/CUDA/STO/sha1_kernel.cu delete mode 100644 benchmarks/CUDA/STO/storeCPU.c delete mode 100644 benchmarks/CUDA/STO/storeCPU.h delete mode 100644 benchmarks/CUDA/STO/storeGPU.cu delete mode 100644 benchmarks/CUDA/STO/storeGPU.h (limited to 'benchmarks/CUDA/STO') diff --git a/benchmarks/CUDA/STO/Makefile b/benchmarks/CUDA/STO/Makefile deleted file mode 100644 index 44e53f3..0000000 --- a/benchmarks/CUDA/STO/Makefile +++ /dev/null @@ -1,52 +0,0 @@ -################################################################################ -# -# Copyright 1993-2006 NVIDIA Corporation. All rights reserved. -# -# NOTICE TO USER: -# -# This source code is subject to NVIDIA ownership rights under U.S. and -# international Copyright laws. -# -# NVIDIA MAKES NO REPRESENTATION ABOUT THE SUITABILITY OF THIS SOURCE -# CODE FOR ANY PURPOSE. IT IS PROVIDED "AS IS" WITHOUT EXPRESS OR -# IMPLIED WARRANTY OF ANY KIND. NVIDIA DISCLAIMS ALL WARRANTIES WITH -# REGARD TO THIS SOURCE CODE, INCLUDING ALL IMPLIED WARRANTIES OF -# MERCHANTABILITY, NONINFRINGEMENT, AND FITNESS FOR A PARTICULAR PURPOSE. -# IN NO EVENT SHALL NVIDIA BE LIABLE FOR ANY SPECIAL, INDIRECT, INCIDENTAL, -# OR CONSEQUENTIAL DAMAGES, OR ANY DAMAGES WHATSOEVER RESULTING FROM LOSS -# OF USE, DATA OR PROFITS, WHETHER IN AN ACTION OF CONTRACT, NEGLIGENCE -# OR OTHER TORTIOUS ACTION, ARISING OUT OF OR IN CONNECTION WITH THE USE -# OR PERFORMANCE OF THIS SOURCE CODE. -# -# U.S. Government End Users. This source code is a "commercial item" as -# that term is defined at 48 C.F.R. 2.101 (OCT 1995), consisting of -# "commercial computer software" and "commercial computer software -# documentation" as such terms are used in 48 C.F.R. 12.212 (SEPT 1995) -# and is provided to the U.S. Government only as a commercial end item. -# Consistent with 48 C.F.R.12.212 and 48 C.F.R. 227.7202-1 through -# 227.7202-4 (JUNE 1995), all U.S. Government End Users acquire the -# source code with only those rights set forth herein. -# -################################################################################ -# -# Build script for project -# -################################################################################ - -# Add source files here -EXECUTABLE := storegpu -# CUDA source files (compiled with cudacc) -CUFILES := main.cu storeGPU.cu -# CUDA dependency files -CU_DEPS := \ -md5_kernel.cu sha1_kernel.cu -# C/C++ source files (compiled with gcc / c++) -CCFILES := \ - storeCPU.c md5_cpu.c sha1_cpu.c -#CUDACCFLAGS := -po maxrregcount=16 - -################################################################################ -# Rules and targets - -GPGPUSIM_ROOT := ../../.. -include ../../../common/common.mk diff --git a/benchmarks/CUDA/STO/README.GPGPU-Sim b/benchmarks/CUDA/STO/README.GPGPU-Sim deleted file mode 100644 index 6056878..0000000 --- a/benchmarks/CUDA/STO/README.GPGPU-Sim +++ /dev/null @@ -1,2 +0,0 @@ -make -./gpgpu_ptx_sim__storegpu diff --git a/benchmarks/CUDA/STO/cust.h b/benchmarks/CUDA/STO/cust.h deleted file mode 100644 index 92219de..0000000 --- a/benchmarks/CUDA/STO/cust.h +++ /dev/null @@ -1,153 +0,0 @@ -#ifndef CUST_H -#define CUST_H -/*=========================================================================== - - CUSTOMIZATION - -* Copyright (c) 2008, NetSysLab at the University of British Columbia -* All rights reserved. -* -* Redistribution and use in source and binary forms, with or without -* modification, are permitted provided that the following conditions are met: -* * Redistributions of source code must retain the above copyright -* notice, this list of conditions and the following disclaimer. -* * Redistributions in binary form must reproduce the above copyright -* notice, this list of conditions and the following disclaimer in the -* documentation and/or other materials provided with the distribution. -* * Neither the name of the University nor the -* names of its contributors may be used to endorse or promote products -* derived from this software without specific prior written permission. -* -* THIS SOFTWARE IS PROVIDED BY NetSysLab ``AS IS'' AND ANY -* EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED -* WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE -* DISCLAIMED. IN NO EVENT SHALL NetSysLab BE LIABLE FOR ANY -* DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES -* (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; -* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND -* ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT -* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS -* SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. - -DESCRIPTION - Configuration file. - -===========================================================================*/ - - -/*=========================================================================== - - PUBLIC DATA DECLARATIONS - -===========================================================================*/ - -/* NOTE: Whenever you change anything in this file, you have to rebuild - the project (Build->RebuildSolution) to take the effects. running directly - using Debug->RunWithoutDebugging will not recompile the solution with the - new parameters. */ - - -/* ------------------------------------------------------------------ -** GENERAL FEATURES -** ------------------------------------------------------------------ */ - -// Use shared memory implementation. Note that the total size of the test -// must not exceed 96MB. This will be debugged later on. -#define FEATURE_SHARED_MEMORY - -// Use pinned memory pages. -#define FEATURE_PINNED_MODE - -// Hash size returned is trimmed to save bandwidth in the copy back. -// Check md5.h for the actual hash size returned. -#define FEATURE_REDUCED_HASH_SIZE - -// This feature guesses at run time the best execution context the GPU will -// run within -//#define FEATURE_DYNAMIC_EXEC_CONTEXT - -#ifdef FEATURE_DYNAMIC_EXEC_CONTEXT -// when in dynamic execution context setting, this feature will enable -// the algorithm that maximizes the number of threads rather than chunk size. -//#define FEATURE_MAXIMIZE_NUM_OF_THREADS -#endif - -// Tests to run -// To run the overlap test instead of the original one. -#define FEATURE_RUN_OVERLAP_TEST - -// Turn on to run SHA1, otherwise it will run MD5 -#define FEATURE_RUN_SHA1 - -// Enable the multi-thread version for the CPU functions on windows -//#define FEATURE_WIN32_THREADS - -// Enable the multi-thread version for the CPU functions on Linux (with pthreads) -//#define FEATURE_PTHREADS - -/* ------------------------------------------------------------------ -** SANITY CHECKS -** ------------------------------------------------------------------ */ - - - -/* ------------------------------------------------------------------ -** CONSTANTS -** ------------------------------------------------------------------ */ -#ifdef FEATURE_DYNAMIC_EXEC_CONTEXT - #define MAX_THREADS_PER_BLOCK 192 - #define MAX_BLOCKS_PER_GRID (32 * 1024) - #define MAX_NUM_OF_THREADS (MAX_THREADS_PER_BLOCK*MAX_BLOCKS_PER_GRID) - #define BASIC_CHUNK_SIZE 64 - #define MAX_CHUNK_SIZE 2048 - #define NUM_OF_MULTIPROCESSORS 4 -#endif /* FEATURE_DYNAMIC_EXEC_CONTEXT */ - -#ifdef FEATURE_REDUCED_HASH_SIZE - #define MD5_HASH_SIZE 4 - #define SHA1_HASH_SIZE 4 -#else - #define MD5_HASH_SIZE 16 - #define SHA1_HASH_SIZE 20 -#endif // FEATURE_REDUCED_HASH_SIZE - - -#ifdef FEATURE_RUN_OVERLAP_TEST - -// THREADS_PER_BLOCK x CHUNK_SIZE should be < 16K if FEATURE_SHARED_MEMORY - #define THREADS_PER_BLOCK 128 - #define BLOCKS_PER_GRID 384//32 //1024//16//384 - #define TOTAL_NUM_OF_THREADS ( THREADS_PER_BLOCK * BLOCKS_PER_GRID ) - -// Chunk size must be multiple of 4. - #define CHUNK_SIZE (52) - //#define CHUNK_SIZE (1024-12) - -// Offset must be multiple of 4. - #define OFFSET (4) - -// Size to be allocated for the original interface tests - #define TEST_MEM_SIZE_OVERLAP (TOTAL_NUM_OF_THREADS * OFFSET + \ - CHUNK_SIZE - OFFSET) - -#else - - #define THREADS_PER_BLOCK 192 - #define BLOCKS_PER_GRID 512 - #define TOTAL_NUM_OF_THREADS ( THREADS_PER_BLOCK * BLOCKS_PER_GRID ) - -// Currently chunk size must be of the form (64*i - 12). 12 is reserved for -// padding issues. - //#define CHUNK_SIZE (64-12) - //#define CHUNK_SIZE (128-12) - //#define CHUNK_SIZE (256-12) - //#define CHUNK_SIZE (512-12) - #define CHUNK_SIZE (1024-12) - //#define CHUNK_SIZE (2048-12) - -// Size to be allocated for the tests - #define TEST_MEM_SIZE (CHUNK_SIZE * TOTAL_NUM_OF_THREADS) - -#endif /* FEATURE_RUN_OVERLAP_TEST */ - -#endif /* CUST_H */ diff --git a/benchmarks/CUDA/STO/main.cu b/benchmarks/CUDA/STO/main.cu deleted file mode 100644 index 31a234d..0000000 --- a/benchmarks/CUDA/STO/main.cu +++ /dev/null @@ -1,867 +0,0 @@ -/*========================================================================== - M A I N - -* Copyright (c) 2008, NetSysLab at the University of British Columbia -* All rights reserved. -* -* Redistribution and use in source and binary forms, with or without -* modification, are permitted provided that the following conditions are met: -* * Redistributions of source code must retain the above copyright -* notice, this list of conditions and the following disclaimer. -* * Redistributions in binary form must reproduce the above copyright -* notice, this list of conditions and the following disclaimer in the -* documentation and/or other materials provided with the distribution. -* * Neither the name of the University nor the -* names of its contributors may be used to endorse or promote products -* derived from this software without specific prior written permission. -* -* THIS SOFTWARE IS PROVIDED BY NetSysLab ``AS IS'' AND ANY -* EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED -* WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE -* DISCLAIMED. IN NO EVENT SHALL NetSysLab BE LIABLE FOR ANY -* DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES -* (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; -* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND -* ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT -* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS -* SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. - -DESCRIPTION - Main entry. - - -==========================================================================*/ - -/*========================================================================== - - INCLUDES - -==========================================================================*/ -// system -#include -#include -#include - -// project -#include -#include -#include - - -/*========================================================================== - - DATA DECLARATIONS - -==========================================================================*/ - -/*-------------------------------------------------------------------------- - TYPE DEFINITIONS ---------------------------------------------------------------------------*/ - -/*-------------------------------------------------------------------------- - FUNCTION PROTOTYPES ---------------------------------------------------------------------------*/ - - -/*-------------------------------------------------------------------------- - CONSTANTS ---------------------------------------------------------------------------*/ - -/*-------------------------------------------------------------------------- - GLOBAL VARIABLES ---------------------------------------------------------------------------*/ - - -/*-------------------------------------------------------------------------- - MACROS ---------------------------------------------------------------------------*/ - -/*========================================================================== - - FUNCTIONS - -==========================================================================*/ - -/*-------------------------------------------------------------------------- - GLOBAL FUNCTIONS ---------------------------------------------------------------------------*/ - -/*=========================================================================== - -FUNCTION SG_PRINT_TIME_BREAKDOWN - -DESCRIPTION - Prints out the given time breakdown parameter - -DEPENDENCIES - None - -RETURN VALUE - None - -===========================================================================*/ -static void print_gpu_time_breakdown( sg_time_breakdown_type* time_breakdown, - float input_buffer_alloc_time, - float gpu_init_time ) { - - printf("\n== GPU Timing ==\n"); - printf("GPU init : %f\n", gpu_init_time); - printf("Host input buffer alloc : %f\n", input_buffer_alloc_time); - printf("-----\n"); - printf("Host output buffer alloc : %f\n", - time_breakdown->host_output_buffer_alloc_time); - printf("GPU memory alloc : %f\n", time_breakdown->device_mem_alloc_time); - printf("Data copy in : %f\n", time_breakdown->copy_in_time); - printf("Kernel execution : %f\n", time_breakdown->exec_time); - printf("Data copy out : %f\n", time_breakdown->copy_out_time); - printf("Last hasing stage : %f\n", time_breakdown->last_stage_time); - -} - -#ifdef FEATURE_RUN_OVERLAP_TEST -/*=========================================================================== - -FUNCTION run_md5_overlap_test - -DESCRIPTION - run the test - -DEPENDENCIES - None - -RETURN VALUE - None - -===========================================================================*/ -void run_md5_overlap_test( ) { - - //**** Variables ****// - float host_input_buffer_alloc_time, gpu_init_time; - sg_time_breakdown_type gpu_time_breakdown; - unsigned char* sc_output; - unsigned char* sg_output; - unsigned char* buffer; - unsigned int timer; - int sg_output_size; - int sc_output_size; - - - printf( "MD5 Overlap Test\n\n" ); - - //**** create the timer ****// - timer = 0; - CUT_SAFE_CALL( cutCreateTimer( &timer)); - - - - //** GPU initialization timing **// - CUT_SAFE_CALL(cutResetTimer(timer)); - CUT_SAFE_CALL(cutStartTimer(timer)); - - /* GPU device initialization */ - sg_init(); - - /* stop the timer */ - CUT_SAFE_CALL(cutStopTimer(timer)); - gpu_init_time = cutGetTimerValue(timer); - - - - //**** Host input buffer allocation timing ****// - CUT_SAFE_CALL(cutResetTimer(timer)); - CUT_SAFE_CALL(cutStartTimer(timer)); - - /* allocate test buffer */ - buffer = (unsigned char*) sg_malloc(TEST_MEM_SIZE_OVERLAP); - - /* stop the timer */ - CUT_SAFE_CALL(cutStopTimer(timer)); - host_input_buffer_alloc_time = cutGetTimerValue(timer); - - - - //**** initialize test buffer with random data ****// - for( unsigned int i = 0; i < TEST_MEM_SIZE_OVERLAP; ++i) { - buffer[i] = i; - } - - - - /***************/ - /***** GPU *****/ - /***************/ - - //** MD5 timing **// - CUT_SAFE_CALL(cutResetTimer(timer)); - CUT_SAFE_CALL(cutStartTimer(timer)); - - /* run GPU version */ - sg_md5_overlap(buffer, TEST_MEM_SIZE_OVERLAP, CHUNK_SIZE, OFFSET, - &sg_output, &sg_output_size, &gpu_time_breakdown); - - /* stop the timer */ - CUT_SAFE_CALL(cutStopTimer(timer)); - print_gpu_time_breakdown( &gpu_time_breakdown, - host_input_buffer_alloc_time, - gpu_init_time ); - printf( "GPU Proc. Time (gpu init and input alloc are not included): %f \n", - cutGetTimerValue(timer)); - - - - - /***************/ - /***** CPU *****/ - /***************/ - //**** start timer for cpu ****// - CUT_SAFE_CALL( cutResetTimer( timer ) ); - CUT_SAFE_CALL( cutStartTimer( timer ) ); - - - //**** run CPU version ****// - sc_md5_overlap(buffer, TEST_MEM_SIZE_OVERLAP, CHUNK_SIZE, OFFSET, - &sc_output, &sc_output_size); - - - //**** stop the timer ****// - CUT_SAFE_CALL( cutStopTimer( timer)); - printf( "CPU Processing time(ms): %f \n", cutGetTimerValue( timer)); - - - if(sc_output_size != sg_output_size){ - printf( "\nGPU and CPU didn't converse to the same output size:\n"); - printf( "\nGPU output size: %d\n", sg_output_size); - printf( "\nCPU output size: %d\n", sc_output_size); - } else { - printf( "\nOutput size: %d\n", sc_output_size); - } - - - //**** check if the results are equivalent ****// - CUTBoolean res = cutCompareub( sg_output, - sc_output, - sg_output_size); - - - //**** print the results ****// - printf( "Test %s\n", (1 == res) ? "PASSED" : "FAILED"); - printf("CPU GPU\n"); - for ( int i = sc_output_size - 4; i < sc_output_size; i++) { - printf("%X %X\n",sc_output[i], sg_output[i]); - } - - - sg_free(buffer); - //sg_free(sg_output); - cudaFreeHost(sg_output ); - free(sc_output); - -} - -/*=========================================================================== - -FUNCTION run_sha1_overlap_test - -DESCRIPTION - run the test - -DEPENDENCIES - None - -RETURN VALUE - None - -===========================================================================*/ -void run_sha1_overlap_test( ) { - - //**** Variables ****// - float host_input_buffer_alloc_time, gpu_init_time; - sg_time_breakdown_type gpu_time_breakdown; - unsigned char* sc_output; - unsigned char* sg_output; - unsigned char* buffer; - unsigned int timer; - int sg_output_size; - int sc_output_size; - - - printf( "SHA1 Overlap Test\n\n" ); - - //**** create the timer ****// - timer = 0; - CUT_SAFE_CALL( cutCreateTimer( &timer)); - - - - //** GPU initialization timing **// - CUT_SAFE_CALL(cutResetTimer(timer)); - CUT_SAFE_CALL(cutStartTimer(timer)); - - /* GPU device initialization */ - sg_init(); - - /* stop the timer */ - CUT_SAFE_CALL(cutStopTimer(timer)); - gpu_init_time = cutGetTimerValue(timer); - - - - //**** Host input buffer allocation timing ****// - CUT_SAFE_CALL(cutResetTimer(timer)); - CUT_SAFE_CALL(cutStartTimer(timer)); - - /* allocate test buffer */ - buffer = (unsigned char*) sg_malloc(TEST_MEM_SIZE_OVERLAP); - - /* stop the timer */ - CUT_SAFE_CALL(cutStopTimer(timer)); - host_input_buffer_alloc_time = cutGetTimerValue(timer); - - - - //**** initialize test buffer with random data ****// - for( unsigned int i = 0; i < TEST_MEM_SIZE_OVERLAP; ++i) { - buffer[i] = i; - } - - - - /***************/ - /***** GPU *****/ - /***************/ - - //** SHA1 timing **// - CUT_SAFE_CALL(cutResetTimer(timer)); - CUT_SAFE_CALL(cutStartTimer(timer)); - - /* run GPU version */ - sg_sha1_overlap(buffer, TEST_MEM_SIZE_OVERLAP, CHUNK_SIZE, OFFSET, - &sg_output, &sg_output_size, &gpu_time_breakdown); - - /* stop the timer */ - CUT_SAFE_CALL(cutStopTimer(timer)); - print_gpu_time_breakdown( &gpu_time_breakdown, - host_input_buffer_alloc_time, - gpu_init_time ); - printf( "GPU Proc. Time (gpu init and input alloc are not included): %f \n", - cutGetTimerValue(timer)); - - - - - /***************/ - /***** CPU *****/ - /***************/ - //**** start timer for cpu ****// - CUT_SAFE_CALL( cutResetTimer( timer ) ); - CUT_SAFE_CALL( cutStartTimer( timer ) ); - - - //**** run CPU version ****// - sc_sha1_overlap(buffer, TEST_MEM_SIZE_OVERLAP, CHUNK_SIZE, OFFSET, - &sc_output, &sc_output_size); - - - //**** stop the timer ****// - CUT_SAFE_CALL( cutStopTimer( timer)); - printf( "CPU Processing time(ms): %f \n", cutGetTimerValue( timer)); - - - if(sc_output_size != sg_output_size){ - printf( "\nGPU and CPU didn't converse to the same output size:\n"); - printf( "\nGPU output size: %d\n", sg_output_size); - printf( "\nCPU output size: %d\n", sc_output_size); - } else { - printf( "\nOutput size: %d\n", sc_output_size); - } - - - //**** check if the results are equivalent ****// - CUTBoolean res = cutCompareub( sg_output, - sc_output, - sg_output_size); - - - //**** print the results ****// - printf( "Test %s\n", (1 == res) ? "PASSED" : "FAILED"); - printf("CPU GPU\n"); - for ( int i = sc_output_size - 4; i < sc_output_size; i++) { - printf("%X %X\n",sc_output[i], sg_output[i]); - } - - - sg_free(buffer); - //sg_free(sg_output); - cudaFreeHost(sg_output ); - free(sc_output); - -} - -#else /* FEATURE_RUN_OVERLAP_TEST */ - - -/*=========================================================================== - -FUNCTION run_md5_test - -DESCRIPTION - run the test - -DEPENDENCIES - None - -RETURN VALUE - None - -===========================================================================*/ -void run_md5_test( ) { - - //**** Variables ****// - unsigned char *sc_output, *sc_single_output; - unsigned char *sg_output; - unsigned char *buffer; - unsigned int timer; - int sg_output_size; - int sc_output_size; - float host_input_buffer_alloc_time, gpu_init_time; - sg_time_breakdown_type gpu_time_breakdown; - - printf( "MD5 Test\n\n" ); - - //**** create the timer ****// - timer = 0; - CUT_SAFE_CALL( cutCreateTimer( &timer)); - - - //** GPU initialization timing **// - CUT_SAFE_CALL(cutResetTimer(timer)); - CUT_SAFE_CALL(cutStartTimer(timer)); - - /* GPU device initialization */ - sg_init(); - - /* stop the timer */ - CUT_SAFE_CALL(cutStopTimer(timer)); - gpu_init_time = cutGetTimerValue(timer); - - - //**** Host input buffer allocation timing ****// - CUT_SAFE_CALL(cutResetTimer(timer)); - CUT_SAFE_CALL(cutStartTimer(timer)); - - /* allocate test buffer */ - buffer = (unsigned char*) sg_malloc(TEST_MEM_SIZE); - - /* stop the timer */ - CUT_SAFE_CALL(cutStopTimer(timer)); - host_input_buffer_alloc_time = cutGetTimerValue(timer); - - - - //**** initialize test buffer with random data ****// - for( unsigned int i = 0; i < TEST_MEM_SIZE; ++i) { - buffer[i] = i; - } - - - - /***************/ - /***** GPU *****/ - /***************/ - - - //** MD5 timing **// - CUT_SAFE_CALL(cutResetTimer(timer)); - CUT_SAFE_CALL(cutStartTimer(timer)); - - /* run GPU version */ - sg_md5(buffer, TEST_MEM_SIZE, &sg_output, &sg_output_size, - &gpu_time_breakdown); - - /* stop the timer */ - CUT_SAFE_CALL(cutStopTimer(timer)); - print_gpu_time_breakdown( &gpu_time_breakdown, - host_input_buffer_alloc_time, - gpu_init_time ); - printf( "GPU Proc. Time (gpu init and input alloc are not included): %f \n", - cutGetTimerValue(timer)); - - - /***************/ - /***** CPU *****/ - /***************/ - //**** start timer for cpu ****// - CUT_SAFE_CALL( cutResetTimer( timer ) ); - CUT_SAFE_CALL( cutStartTimer( timer ) ); - - //**** run CPU version ****// - sc_md5(buffer, TEST_MEM_SIZE, &sc_output, &sc_output_size); - - //**** stop the timer ****// - CUT_SAFE_CALL( cutStopTimer( timer)); - printf( "CPU Processing time(ms): %f \n", cutGetTimerValue( timer)); - - /*****************************/ - /***** CPU Single Thread *****/ - /*****************************/ - //**** start timer for single thread cpu ****// - CUT_SAFE_CALL( cutResetTimer( timer ) ); - CUT_SAFE_CALL( cutStartTimer( timer ) ); - - //**** run Single Thread CPU version ****// - sc_md5_standard(buffer, TEST_MEM_SIZE, &sc_single_output); - - //**** stop the timer ****// - CUT_SAFE_CALL( cutStopTimer( timer)); - printf( "CPU Single Thread Processing time(ms): %f \n", - cutGetTimerValue( timer)); - - - if(sc_output_size != sg_output_size){ - printf( "\nGPU and CPU didn't converse to the same output size:\n"); - printf( "\nGPU output size: %d\n", sg_output_size); - printf( "\nCPU output size: %d\n", sc_output_size); - } else { - printf( "\nOutput size: %d\n", sc_output_size); - } - - - //**** check if the results are equivalent ****// - CUTBoolean res = cutCompareub( sg_output, - sc_output, - sg_output_size); - - - //**** print the results ****// - printf( "Test %s\n", (1 == res) ? "PASSED" : "FAILED"); - printf("CPU GPU\n"); - for ( int i = sc_output_size - 4; i < sc_output_size; i++) { - printf("%X %X\n",sc_output[i], sg_output[i]); - } - - - - sg_free(buffer); - free(sg_output); /* We dont need to free this using sg_free, it will always be - allocated using malloc. will try to come up with a - cleaner way to make things more clear. */ - free(sc_output); - -} - -/*=========================================================================== - -FUNCTION run_sha1_test - -DESCRIPTION - run the sha1 test - -DEPENDENCIES - None - -RETURN VALUE - None - -===========================================================================*/ -void run_sha1_test( ) { - - //**** Variables ****// - unsigned char *sc_output, *sc_single_output; - unsigned char *sg_output; - unsigned char *buffer; - unsigned int timer; - int sg_output_size; - int sc_output_size; - float host_input_buffer_alloc_time, gpu_init_time; - sg_time_breakdown_type gpu_time_breakdown; - - printf( "SHA1 Test\n\n" ); - - //**** create the timer ****// - timer = 0; - CUT_SAFE_CALL( cutCreateTimer( &timer)); - - //** GPU initialization timing **// - CUT_SAFE_CALL(cutResetTimer(timer)); - CUT_SAFE_CALL(cutStartTimer(timer)); - - /* GPU device initialization */ - sg_init(); - - /* stop the timer */ - CUT_SAFE_CALL(cutStopTimer(timer)); - gpu_init_time = cutGetTimerValue(timer); - - - //**** Host input buffer allocation timing ****// - CUT_SAFE_CALL(cutResetTimer(timer)); - CUT_SAFE_CALL(cutStartTimer(timer)); - - /* allocate test buffer */ - buffer = (unsigned char*) sg_malloc(TEST_MEM_SIZE); - - /* stop the timer */ - CUT_SAFE_CALL(cutStopTimer(timer)); - host_input_buffer_alloc_time = cutGetTimerValue(timer); - - - - //**** initialize test buffer with random data ****// - for( unsigned int i = 0; i < TEST_MEM_SIZE; ++i) { - buffer[i] = i; - } - - - - /***************/ - /***** GPU *****/ - /***************/ - - - //** SHA1 timing **// - CUT_SAFE_CALL(cutResetTimer(timer)); - CUT_SAFE_CALL(cutStartTimer(timer)); - - /* run GPU version */ - sg_sha1(buffer, TEST_MEM_SIZE, &sg_output, &sg_output_size, - &gpu_time_breakdown); - - /* stop the timer */ - CUT_SAFE_CALL(cutStopTimer(timer)); - print_gpu_time_breakdown( &gpu_time_breakdown, - host_input_buffer_alloc_time, - gpu_init_time ); - printf( "GPU Proc. Time (gpu init and input alloc are not included): %f \n", - cutGetTimerValue(timer)); - - - /***************/ - /***** CPU *****/ - /***************/ - //**** start timer for cpu ****// - CUT_SAFE_CALL( cutResetTimer( timer ) ); - CUT_SAFE_CALL( cutStartTimer( timer ) ); - - //**** run CPU version ****// - sc_sha1(buffer, TEST_MEM_SIZE, &sc_output, &sc_output_size); - - //**** stop the timer ****// - CUT_SAFE_CALL( cutStopTimer( timer)); - printf( "CPU Processing time(ms): %f \n", cutGetTimerValue( timer)); - - /*****************************/ - /***** CPU Single Thread *****/ - /*****************************/ - //**** start timer for single thread cpu ****// - CUT_SAFE_CALL( cutResetTimer( timer ) ); - CUT_SAFE_CALL( cutStartTimer( timer ) ); - - //**** run Single Thread CPU version ****// - sc_sha1_standard(buffer, TEST_MEM_SIZE, &sc_single_output); - - //**** stop the timer ****// - CUT_SAFE_CALL( cutStopTimer( timer)); - printf( "CPU Single Thread Processing time(ms): %f \n", - cutGetTimerValue( timer)); - - - if(sc_output_size != sg_output_size){ - printf( "\nGPU and CPU didn't converse to the same output size:\n"); - printf( "\nGPU output size: %d\n", sg_output_size); - printf( "\nCPU output size: %d\n", sc_output_size); - } else { - printf( "\nOutput size: %d\n", sc_output_size); - } - - - //**** check if the results are equivalent ****// - CUTBoolean res = cutCompareub( sg_output, - sc_output, - sg_output_size); - - - //**** print the results ****// - printf( "Test %s\n", (1 == res) ? "PASSED" : "FAILED"); - printf("CPU GPU\n"); - for ( int i = sc_output_size - 4; i < sc_output_size; i++) { - printf("%X %X\n",sc_output[i], sg_output[i]); - } - - - - sg_free(buffer); - free(sg_output); /* We dont need to free this using sg_free, it will always be - allocated using malloc. will try to come up with a - cleaner way to make things more clear. */ - free(sc_output); - -} -/* void run_sha1_test( ) { */ - -/* //\**** Variables ****\// */ -/* unsigned int timer = 0; */ -/* unsigned char* buffer; */ -/* unsigned char* sg_output; */ -/* int sg_output_size; */ -/* unsigned char *sc_output, *sc_single_output; */ -/* int sc_output_size; */ - - -/* printf( "SHA1 Test\n\n" ); */ - - -/* //\**** host memory management ****\// */ -/* // allocate test buffer */ -/* buffer = (unsigned char*) sg_malloc(TEST_MEM_SIZE); */ - - -/* //\**** initialize test buffer with random data ****\// */ -/* for( unsigned int i = 0; i < TEST_MEM_SIZE; ++i) { */ -/* buffer[i] = i; */ -/* } */ - -/* //\**** create the timer ****\// */ -/* timer = 0; */ -/* CUT_SAFE_CALL( cutCreateTimer( &timer)); */ - - -/* /\***************\/ */ -/* /\***** GPU *****\/ */ -/* /\***************\/ */ -/* //\**** start timer for GPU timing ****\// */ -/* CUT_SAFE_CALL( cutResetTimer(timer) ); */ -/* CUT_SAFE_CALL( cutStartTimer( timer)); */ - - -/* //\**** run GPU version ****\// */ -/* sg_sha1(buffer, TEST_MEM_SIZE, &sg_output, &sg_output_size); */ - - -/* //\**** stop the timer ****\// */ -/* CUT_SAFE_CALL( cutStopTimer( timer)); */ - - -/* //\**** print results ****\// */ -/* printf( "GPU Processing time(ms): %f \n", cutGetTimerValue( timer)); */ - - -/* /\***************\/ */ -/* /\***** CPU *****\/ */ -/* /\***************\/ */ -/* //\**** start timer for cpu ****\// */ -/* CUT_SAFE_CALL( cutResetTimer( timer ) ); */ -/* CUT_SAFE_CALL( cutStartTimer( timer ) ); */ - - -/* //\**** run CPU version ****\// */ -/* sc_sha1(buffer, TEST_MEM_SIZE, &sc_output, &sc_output_size); */ - - -/* //\**** stop the timer ****\// */ -/* CUT_SAFE_CALL( cutStopTimer( timer)); */ -/* printf( "CPU Processing time(ms): %f \n", cutGetTimerValue( timer)); */ - - -/* /\*****************************\/ */ -/* /\***** CPU Single Thread *****\/ */ -/* /\*****************************\/ */ -/* //\**** start timer for single thread cpu ****\// */ -/* CUT_SAFE_CALL( cutResetTimer( timer ) ); */ -/* CUT_SAFE_CALL( cutStartTimer( timer ) ); */ - -/* //\**** run CPU version ****\// */ -/* sc_sha1_standard(buffer, TEST_MEM_SIZE, &sc_single_output); */ - -/* //\**** stop the timer ****\// */ -/* CUT_SAFE_CALL( cutStopTimer( timer)); */ -/* printf( "CPU Single Thread Processing time(ms): %f \n", */ -/* cutGetTimerValue( timer)); */ - - -/* if(sc_output_size != sg_output_size){ */ -/* printf( "\nGPU and CPU didn't converse to the same output size:\n"); */ -/* printf( "\nGPU output size: %d\n", sg_output_size); */ -/* printf( "\nCPU output size: %d\n", sc_output_size); */ -/* } else { */ -/* printf( "\nOutput size: %d\n", sc_output_size); */ -/* } */ - - -/* //\**** check if the results are equivalent ****\// */ -/* CUTBoolean res = cutCompareub( sg_output, */ -/* sc_output, */ -/* sg_output_size); */ - - -/* //\**** print the results ****\// */ -/* printf( "Test %s\n", (1 == res) ? "PASSED" : "FAILED"); */ -/* printf("CPU GPU\n"); */ -/* for ( int i = sc_output_size - 4; i < sc_output_size; i++) { */ -/* printf("%X %X\n",sc_output[i], sg_output[i]); */ -/* } */ - - -/* sg_free(buffer); */ -/* free(sg_output); /\* We dont need to free this using sg_free, it will always be */ -/* allocated using malloc. I will try to come up with a */ -/* cleaner way to make things more clear. *\/ */ -/* free(sc_output); */ - -/* } */ -#endif /* FEATURE_RUN_OVERLAP_TEST */ - - - -/*-------------------------------------------------------------------------- - GLOBAL FUNCTIONS ---------------------------------------------------------------------------*/ -/*=========================================================================== - -FUNCTION main - -DESCRIPTION - main entry of the program - -DEPENDENCIES - None - -RETURN VALUE - None - -===========================================================================*/ -int main( int argc, char** argv) { - -#ifdef FEATURE_SHARED_MEMORY - printf("Shared Memory Enabled\n"); -#endif - - -#ifdef FEATURE_PINNED_MODE - printf("Pinned Memory Enabled\n"); -#endif - - -#ifdef FEATURE_REDUCED_HASH_SIZE - printf("Reduced Hash Size Enabled\n"); -#endif - -#ifdef FEATURE_RUN_OVERLAP_TEST -#ifdef FEATURE_RUN_SHA1 - printf("Running SHA1 Overlap Test..\n"); - run_sha1_overlap_test( ); -#else - printf("Running MD5 Overlap Test..\n"); - run_md5_overlap_test( ); -#endif // FEATURE_RUN_SHA1 -#else -#ifdef FEATURE_RUN_SHA1 - printf("Running SHA1 Non-Overlap Test..\n"); - run_sha1_test( ); -#else - printf("Running MD5 Non-Overlap Test..\n"); - run_md5_test( ); -#endif // FEATURE_RUN_SHA1 -#endif // FEATURE_RUN_OVERLAP_TEST - -} - diff --git a/benchmarks/CUDA/STO/md5_cpu.c b/benchmarks/CUDA/STO/md5_cpu.c deleted file mode 100644 index 44fbef2..0000000 --- a/benchmarks/CUDA/STO/md5_cpu.c +++ /dev/null @@ -1,397 +0,0 @@ -/*========================================================================== - M D 5 C P U - - CPU implementation of the md5 algorithm. - - * - * FIPS-180-1 compliant SHA-1 implementation - * - * Copyright (C) 2006-2007 Christophe Devine - * - * This library is free software; you can redistribute it and/or - * modify it under the terms of the GNU Lesser General Public - * License, version 2.1 as published by the Free Software Foundation. - * - * This library 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 - * Lesser General Public License for more details. - * - * You should have received a copy of the GNU Lesser General Public - * License along with this library; if not, write to the Free Software - * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, - * MA 02110-1301 USA - * - * The SHA-1 standard was published by NIST in 1993. - * - * http://www.itl.nist.gov/fipspubs/fip180-1.htm - -DESCRIPTION - CPU implementation of the md5 algorithm. - - -==========================================================================*/ - -/*========================================================================== - - INCLUDES - -==========================================================================*/ -#include -#include - -#include "md5_cpu.h" -#include "cust.h" -/*========================================================================== - - DATA DECLARATIONS - -==========================================================================*/ - -/*-------------------------------------------------------------------------- - TYPE DEFINITIONS ---------------------------------------------------------------------------*/ -typedef struct { - unsigned long total[2]; /*!< number of bytes processed */ - unsigned long state[4]; /*!< intermediate digest state */ - unsigned char buffer[64]; /*!< data block being processed */ -} md5_cpu_context; - - -/*-------------------------------------------------------------------------- - FUNCTION PROTOTYPES ---------------------------------------------------------------------------*/ - -/*-------------------------------------------------------------------------- - CONSTANTS ---------------------------------------------------------------------------*/ - -/*-------------------------------------------------------------------------- - GLOBAL VARIABLES ---------------------------------------------------------------------------*/ -static const unsigned char md5_cpu_padding[64] = -{ - 0x80, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, - 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, - 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, - 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0 -}; - -/*-------------------------------------------------------------------------- - MACROS ---------------------------------------------------------------------------*/ -/* - * 32-bit integer manipulation macros (little endian) - */ -#ifndef GET_UINT32_LE -#define GET_UINT32_LE(n,b,i) \ -{ \ - (n) = ( (unsigned long) (b)[(i) ] ) \ - | ( (unsigned long) (b)[(i) + 1] << 8 ) \ - | ( (unsigned long) (b)[(i) + 2] << 16 ) \ - | ( (unsigned long) (b)[(i) + 3] << 24 ); \ -} -#endif - -#ifndef PUT_UINT32_LE -#define PUT_UINT32_LE(n,b,i) \ -{ \ - (b)[(i) ] = (unsigned char) ( (n) ); \ - (b)[(i) + 1] = (unsigned char) ( (n) >> 8 ); \ - (b)[(i) + 2] = (unsigned char) ( (n) >> 16 ); \ - (b)[(i) + 3] = (unsigned char) ( (n) >> 24 ); \ -} -#endif - -/*========================================================================== - - FUNCTIONS - -==========================================================================*/ -/*-------------------------------------------------------------------------- - LOCAL FUNCTIONS ---------------------------------------------------------------------------*/ -/*=========================================================================== - -FUNCTION MD5_CPU_STARTS - -DESCRIPTION - MD5 context setup - -DEPENDENCIES - None - -RETURN VALUE - None - -===========================================================================*/ -static void md5_cpu_starts( md5_cpu_context *ctx ) { - ctx->total[0] = 0; - ctx->total[1] = 0; - - ctx->state[0] = 0x67452301; - ctx->state[1] = 0xEFCDAB89; - ctx->state[2] = 0x98BADCFE; - ctx->state[3] = 0x10325476; -} - -/*=========================================================================== - -FUNCTION MD5_CPU_PROCESS - -DESCRIPTION - MD5 process buffer - -DEPENDENCIES - None - -RETURN VALUE - None - -===========================================================================*/ -static void md5_cpu_process( md5_cpu_context *ctx, unsigned char data[64] ) { - unsigned long X[16], A, B, C, D; - - GET_UINT32_LE( X[ 0], data, 0 ); - GET_UINT32_LE( X[ 1], data, 4 ); - GET_UINT32_LE( X[ 2], data, 8 ); - GET_UINT32_LE( X[ 3], data, 12 ); - GET_UINT32_LE( X[ 4], data, 16 ); - GET_UINT32_LE( X[ 5], data, 20 ); - GET_UINT32_LE( X[ 6], data, 24 ); - GET_UINT32_LE( X[ 7], data, 28 ); - GET_UINT32_LE( X[ 8], data, 32 ); - GET_UINT32_LE( X[ 9], data, 36 ); - GET_UINT32_LE( X[10], data, 40 ); - GET_UINT32_LE( X[11], data, 44 ); - GET_UINT32_LE( X[12], data, 48 ); - GET_UINT32_LE( X[13], data, 52 ); - GET_UINT32_LE( X[14], data, 56 ); - GET_UINT32_LE( X[15], data, 60 ); - -#define S(x,n) ((x << n) | ((x & 0xFFFFFFFF) >> (32 - n))) - -#define P(a,b,c,d,k,s,t) \ - { \ - a += F(b,c,d) + X[k] + t; a = S(a,s) + b; \ - } - - A = ctx->state[0]; - B = ctx->state[1]; - C = ctx->state[2]; - D = ctx->state[3]; - -#define F(x,y,z) (z ^ (x & (y ^ z))) - - P( A, B, C, D, 0, 7, 0xD76AA478 ); - P( D, A, B, C, 1, 12, 0xE8C7B756 ); - P( C, D, A, B, 2, 17, 0x242070DB ); - P( B, C, D, A, 3, 22, 0xC1BDCEEE ); - P( A, B, C, D, 4, 7, 0xF57C0FAF ); - P( D, A, B, C, 5, 12, 0x4787C62A ); - P( C, D, A, B, 6, 17, 0xA8304613 ); - P( B, C, D, A, 7, 22, 0xFD469501 ); - P( A, B, C, D, 8, 7, 0x698098D8 ); - P( D, A, B, C, 9, 12, 0x8B44F7AF ); - P( C, D, A, B, 10, 17, 0xFFFF5BB1 ); - P( B, C, D, A, 11, 22, 0x895CD7BE ); - P( A, B, C, D, 12, 7, 0x6B901122 ); - P( D, A, B, C, 13, 12, 0xFD987193 ); - P( C, D, A, B, 14, 17, 0xA679438E ); - P( B, C, D, A, 15, 22, 0x49B40821 ); - -#undef F - -#define F(x,y,z) (y ^ (z & (x ^ y))) - - P( A, B, C, D, 1, 5, 0xF61E2562 ); - P( D, A, B, C, 6, 9, 0xC040B340 ); - P( C, D, A, B, 11, 14, 0x265E5A51 ); - P( B, C, D, A, 0, 20, 0xE9B6C7AA ); - P( A, B, C, D, 5, 5, 0xD62F105D ); - P( D, A, B, C, 10, 9, 0x02441453 ); - P( C, D, A, B, 15, 14, 0xD8A1E681 ); - P( B, C, D, A, 4, 20, 0xE7D3FBC8 ); - P( A, B, C, D, 9, 5, 0x21E1CDE6 ); - P( D, A, B, C, 14, 9, 0xC33707D6 ); - P( C, D, A, B, 3, 14, 0xF4D50D87 ); - P( B, C, D, A, 8, 20, 0x455A14ED ); - P( A, B, C, D, 13, 5, 0xA9E3E905 ); - P( D, A, B, C, 2, 9, 0xFCEFA3F8 ); - P( C, D, A, B, 7, 14, 0x676F02D9 ); - P( B, C, D, A, 12, 20, 0x8D2A4C8A ); - -#undef F - -#define F(x,y,z) (x ^ y ^ z) - - P( A, B, C, D, 5, 4, 0xFFFA3942 ); - P( D, A, B, C, 8, 11, 0x8771F681 ); - P( C, D, A, B, 11, 16, 0x6D9D6122 ); - P( B, C, D, A, 14, 23, 0xFDE5380C ); - P( A, B, C, D, 1, 4, 0xA4BEEA44 ); - P( D, A, B, C, 4, 11, 0x4BDECFA9 ); - P( C, D, A, B, 7, 16, 0xF6BB4B60 ); - P( B, C, D, A, 10, 23, 0xBEBFBC70 ); - P( A, B, C, D, 13, 4, 0x289B7EC6 ); - P( D, A, B, C, 0, 11, 0xEAA127FA ); - P( C, D, A, B, 3, 16, 0xD4EF3085 ); - P( B, C, D, A, 6, 23, 0x04881D05 ); - P( A, B, C, D, 9, 4, 0xD9D4D039 ); - P( D, A, B, C, 12, 11, 0xE6DB99E5 ); - P( C, D, A, B, 15, 16, 0x1FA27CF8 ); - P( B, C, D, A, 2, 23, 0xC4AC5665 ); - -#undef F - -#define F(x,y,z) (y ^ (x | ~z)) - - P( A, B, C, D, 0, 6, 0xF4292244 ); - P( D, A, B, C, 7, 10, 0x432AFF97 ); - P( C, D, A, B, 14, 15, 0xAB9423A7 ); - P( B, C, D, A, 5, 21, 0xFC93A039 ); - P( A, B, C, D, 12, 6, 0x655B59C3 ); - P( D, A, B, C, 3, 10, 0x8F0CCC92 ); - P( C, D, A, B, 10, 15, 0xFFEFF47D ); - P( B, C, D, A, 1, 21, 0x85845DD1 ); - P( A, B, C, D, 8, 6, 0x6FA87E4F ); - P( D, A, B, C, 15, 10, 0xFE2CE6E0 ); - P( C, D, A, B, 6, 15, 0xA3014314 ); - P( B, C, D, A, 13, 21, 0x4E0811A1 ); - P( A, B, C, D, 4, 6, 0xF7537E82 ); - P( D, A, B, C, 11, 10, 0xBD3AF235 ); - P( C, D, A, B, 2, 15, 0x2AD7D2BB ); - P( B, C, D, A, 9, 21, 0xEB86D391 ); - -#undef F - - ctx->state[0] += A; - ctx->state[1] += B; - ctx->state[2] += C; - ctx->state[3] += D; -} - -/*=========================================================================== - -FUNCTION MD5_CPU_UPDATE - -DESCRIPTION - MD5 update - -DEPENDENCIES - None - -RETURN VALUE - None - -===========================================================================*/ -void md5_cpu_update( md5_cpu_context *ctx, unsigned char *input, int ilen ) { - int fill; - unsigned long left; - - if( ilen <= 0 ) - return; - - left = ctx->total[0] & 0x3F; - fill = 64 - left; - - ctx->total[0] += ilen; - ctx->total[0] &= 0xFFFFFFFF; - - if( ctx->total[0] < (unsigned long) ilen ) - ctx->total[1]++; - - if( left && ilen >= fill ) { - memcpy( (void *) (ctx->buffer + left), - (void *) input, fill ); - md5_cpu_process( ctx, ctx->buffer ); - input += fill; - ilen -= fill; - left = 0; - } - - while( ilen >= 64 ) { - md5_cpu_process( ctx, input ); - input += 64; - ilen -= 64; - } - - if( ilen > 0 ) { - memcpy( (void *) (ctx->buffer + left), - (void *) input, ilen ); - } -} - - -/*=========================================================================== - -FUNCTION MD5_CPU_FINISH - -DESCRIPTION - MD5 final digest - -DEPENDENCIES - None - -RETURN VALUE - None - -===========================================================================*/ -void md5_cpu_finish( md5_cpu_context *ctx, unsigned char *output ) { - - unsigned long last, padn; - unsigned long high, low; - unsigned char msglen[8]; - - high = ( ctx->total[0] >> 29 ) - | ( ctx->total[1] << 3 ); - low = ( ctx->total[0] << 3 ); - - PUT_UINT32_LE( low, msglen, 0 ); - PUT_UINT32_LE( high, msglen, 4 ); - - last = ctx->total[0] & 0x3F; - padn = ( last < 56 ) ? ( 56 - last ) : ( 120 - last ); - - md5_cpu_update( ctx, (unsigned char *) md5_cpu_padding, padn ); - md5_cpu_update( ctx, msglen, 8 ); - - PUT_UINT32_LE( ctx->state[0], output, 0 ); -#ifndef FEATURE_REDUCED_HASH_SIZE - PUT_UINT32_LE( ctx->state[1], output, 4 ); - PUT_UINT32_LE( ctx->state[2], output, 8 ); - PUT_UINT32_LE( ctx->state[3], output, 12 ); -#endif -} - -/*-------------------------------------------------------------------------- - GLOBAL FUNCTIONS ---------------------------------------------------------------------------*/ -/*=========================================================================== - -FUNCTION CPU_MD5_INTERNAL - -DESCRIPTION - CPU implementation of the MD5 algorithm - -DEPENDENCIES - None - -RETURN VALUE - Hash - -===========================================================================*/ -void md5_cpu_internal( unsigned char *input, int ilen, - unsigned char *output ){ - - md5_cpu_context ctx; - - md5_cpu_starts( &ctx ); - md5_cpu_update( &ctx, input, ilen ); - md5_cpu_finish( &ctx, output ); - - memset( &ctx, 0, sizeof( md5_cpu_context ) ); - -} - - diff --git a/benchmarks/CUDA/STO/md5_cpu.h b/benchmarks/CUDA/STO/md5_cpu.h deleted file mode 100644 index eabdf5e..0000000 --- a/benchmarks/CUDA/STO/md5_cpu.h +++ /dev/null @@ -1,98 +0,0 @@ -#ifndef MD5_CPU_H -#define MD5_CPU_H - -/*========================================================================== - M D 5 C P U - -* Copyright (c) 2008, NetSysLab at the University of British Columbia -* All rights reserved. -* -* Redistribution and use in source and binary forms, with or without -* modification, are permitted provided that the following conditions are met: -* * Redistributions of source code must retain the above copyright -* notice, this list of conditions and the following disclaimer. -* * Redistributions in binary form must reproduce the above copyright -* notice, this list of conditions and the following disclaimer in the -* documentation and/or other materials provided with the distribution. -* * Neither the name of the University nor the -* names of its contributors may be used to endorse or promote products -* derived from this software without specific prior written permission. -* -* THIS SOFTWARE IS PROVIDED BY NetSysLab ``AS IS'' AND ANY -* EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED -* WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE -* DISCLAIMED. IN NO EVENT SHALL NetSysLab BE LIABLE FOR ANY -* DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES -* (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; -* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND -* ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT -* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS -* SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. - -DESCRIPTION - CPU implementation of the md5 algorithm. - - -==========================================================================*/ - -/*========================================================================== - - INCLUDES - -==========================================================================*/ - - -/*========================================================================== - - DATA DECLARATIONS - -==========================================================================*/ - -/*-------------------------------------------------------------------------- - TYPE DEFINITIONS ---------------------------------------------------------------------------*/ - - -/*-------------------------------------------------------------------------- - FUNCTION PROTOTYPES ---------------------------------------------------------------------------*/ - -/*-------------------------------------------------------------------------- - CONSTANTS ---------------------------------------------------------------------------*/ - -/*-------------------------------------------------------------------------- - GLOBAL VARIABLES ---------------------------------------------------------------------------*/ - - -/*-------------------------------------------------------------------------- - MACROS ---------------------------------------------------------------------------*/ - -/*========================================================================== - - FUNCTIONS - -==========================================================================*/ - -/*=========================================================================== - -FUNCTION cpu_md5_internal - -DESCRIPTION - CPU implementation of the MD5 algorithm - -DEPENDENCIES - None - -RETURN VALUE - Hash - -===========================================================================*/ -void md5_cpu_internal( unsigned char *input, int ilen, - unsigned char *output ); - - - -#endif /* MD5_CPU_H */ diff --git a/benchmarks/CUDA/STO/md5_kernel.cu b/benchmarks/CUDA/STO/md5_kernel.cu deleted file mode 100644 index 419daa5..0000000 --- a/benchmarks/CUDA/STO/md5_kernel.cu +++ /dev/null @@ -1,995 +0,0 @@ -/*========================================================================== - MD5 KERNEL - -* Copyright (c) 2008, NetSysLab at the University of British Columbia -* All rights reserved. -* -* Redistribution and use in source and binary forms, with or without -* modification, are permitted provided that the following conditions are met: -* * Redistributions of source code must retain the above copyright -* notice, this list of conditions and the following disclaimer. -* * Redistributions in binary form must reproduce the above copyright -* notice, this list of conditions and the following disclaimer in the -* documentation and/or other materials provided with the distribution. -* * Neither the name of the University nor the -* names of its contributors may be used to endorse or promote products -* derived from this software without specific prior written permission. -* -* THIS SOFTWARE IS PROVIDED BY NetSysLab ``AS IS'' AND ANY -* EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED -* WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE -* DISCLAIMED. IN NO EVENT SHALL NetSysLab BE LIABLE FOR ANY -* DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES -* (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; -* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND -* ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT -* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS -* SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. - -DESCRIPTION - CPU version of the storeGPU library. - - -==========================================================================*/ - -/*========================================================================== - - INCLUDES - -==========================================================================*/ -#include -#include -#include "cust.h" - -/*========================================================================== - - DATA DECLARATIONS - -==========================================================================*/ - -/*-------------------------------------------------------------------------- - TYPE DEFINITIONS ---------------------------------------------------------------------------*/ -typedef struct { - unsigned long total[2]; /*!< number of bytes processed */ - unsigned long state[4]; /*!< intermediate digest state */ - unsigned char buffer[64]; /*!< data block being processed */ -} md5_context; - -/*-------------------------------------------------------------------------- - FUNCTION PROTOTYPES ---------------------------------------------------------------------------*/ - -/*-------------------------------------------------------------------------- - CONSTANTS ---------------------------------------------------------------------------*/ - -/*-------------------------------------------------------------------------- - GLOBAL VARIABLES ---------------------------------------------------------------------------*/ - -__device__ -const unsigned char md5_padding[64] = -{ - 0x80, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, - 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, - 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, - 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0 -}; - -/*-------------------------------------------------------------------------- - MACROS ---------------------------------------------------------------------------*/ -// 32-bit integer manipulation macros (little endian) -#ifndef GET_UINT32_LE -#define GET_UINT32_LE(n,b,i) \ -{ \ - (n) = ( (unsigned long) (b)[(i) ] ) \ - | ( (unsigned long) (b)[(i) + 1] << 8 ) \ - | ( (unsigned long) (b)[(i) + 2] << 16 ) \ - | ( (unsigned long) (b)[(i) + 3] << 24 ); \ -} -#endif - -#ifndef PUT_UINT32_LE -#define PUT_UINT32_LE(n,b,i) \ -{ \ - (b)[(i) ] = (unsigned char) ( (n) ); \ - (b)[(i) + 1] = (unsigned char) ( (n) >> 8 ); \ - (b)[(i) + 2] = (unsigned char) ( (n) >> 16 ); \ - (b)[(i) + 3] = (unsigned char) ( (n) >> 24 ); \ -} -#endif - -#ifdef FEATURE_SHARED_MEMORY -// current thread stride. -#define SHARED_MEMORY_INDEX(index) (32 * (index) + (threadIdx.x & 0x1F)) -#endif /* FEATURE_SHARED_MEMORY */ - - - -/*========================================================================== - - FUNCTIONS - -==========================================================================*/ - -/*-------------------------------------------------------------------------- - LOCAL FUNCTIONS ---------------------------------------------------------------------------*/ - - -#ifndef FEATURE_SHARED_MEMORY -/*=========================================================================== - -FUNCTION - -DESCRIPTION - MD5 context setup - -DEPENDENCIES - - -RETURN VALUE - - -===========================================================================*/ -__device__ -static void md5_starts( md5_context *ctx ) { - ctx->total[0] = 0; - ctx->total[1] = 0; - - ctx->state[0] = 0x67452301; - ctx->state[1] = 0xEFCDAB89; - ctx->state[2] = 0x98BADCFE; - ctx->state[3] = 0x10325476; -} - -/*=========================================================================== - -FUNCTION MD5_PROCESS - -DESCRIPTION - - -DEPENDENCIES - - -RETURN VALUE - - -===========================================================================*/ -__device__ -static void md5_process( md5_context *ctx, unsigned char data[64] ) { - - unsigned long A, B, C, D; - unsigned long *X = (unsigned long *)data; - - - GET_UINT32_LE( X[ 0], data, 0 ); - GET_UINT32_LE( X[ 1], data, 4 ); - GET_UINT32_LE( X[ 2], data, 8 ); - GET_UINT32_LE( X[ 3], data, 12 ); - GET_UINT32_LE( X[ 4], data, 16 ); - GET_UINT32_LE( X[ 5], data, 20 ); - GET_UINT32_LE( X[ 6], data, 24 ); - GET_UINT32_LE( X[ 7], data, 28 ); - GET_UINT32_LE( X[ 8], data, 32 ); - GET_UINT32_LE( X[ 9], data, 36 ); - GET_UINT32_LE( X[10], data, 40 ); - GET_UINT32_LE( X[11], data, 44 ); - GET_UINT32_LE( X[12], data, 48 ); - GET_UINT32_LE( X[13], data, 52 ); - GET_UINT32_LE( X[14], data, 56 ); - GET_UINT32_LE( X[15], data, 60 ); - -#undef S -#define S(x,n) ((x << n) | ((x & 0xFFFFFFFF) >> (32 - n))) - -#undef P -#define P(a,b,c,d,k,s,t) { \ - a += F(b,c,d) + X[k] + t; a = S(a,s) + b; \ - } \ - - A = ctx->state[0]; - B = ctx->state[1]; - C = ctx->state[2]; - D = ctx->state[3]; - -#define F(x,y,z) (z ^ (x & (y ^ z))) - - P( A, B, C, D, 0, 7, 0xD76AA478 ); - P( D, A, B, C, 1, 12, 0xE8C7B756 ); - P( C, D, A, B, 2, 17, 0x242070DB ); - P( B, C, D, A, 3, 22, 0xC1BDCEEE ); - P( A, B, C, D, 4, 7, 0xF57C0FAF ); - P( D, A, B, C, 5, 12, 0x4787C62A ); - P( C, D, A, B, 6, 17, 0xA8304613 ); - P( B, C, D, A, 7, 22, 0xFD469501 ); - P( A, B, C, D, 8, 7, 0x698098D8 ); - P( D, A, B, C, 9, 12, 0x8B44F7AF ); - P( C, D, A, B, 10, 17, 0xFFFF5BB1 ); - P( B, C, D, A, 11, 22, 0x895CD7BE ); - P( A, B, C, D, 12, 7, 0x6B901122 ); - P( D, A, B, C, 13, 12, 0xFD987193 ); - P( C, D, A, B, 14, 17, 0xA679438E ); - P( B, C, D, A, 15, 22, 0x49B40821 ); - -#undef F - -#define F(x,y,z) (y ^ (z & (x ^ y))) - - P( A, B, C, D, 1, 5, 0xF61E2562 ); - P( D, A, B, C, 6, 9, 0xC040B340 ); - P( C, D, A, B, 11, 14, 0x265E5A51 ); - P( B, C, D, A, 0, 20, 0xE9B6C7AA ); - P( A, B, C, D, 5, 5, 0xD62F105D ); - P( D, A, B, C, 10, 9, 0x02441453 ); - P( C, D, A, B, 15, 14, 0xD8A1E681 ); - P( B, C, D, A, 4, 20, 0xE7D3FBC8 ); - P( A, B, C, D, 9, 5, 0x21E1CDE6 ); - P( D, A, B, C, 14, 9, 0xC33707D6 ); - P( C, D, A, B, 3, 14, 0xF4D50D87 ); - P( B, C, D, A, 8, 20, 0x455A14ED ); - P( A, B, C, D, 13, 5, 0xA9E3E905 ); - P( D, A, B, C, 2, 9, 0xFCEFA3F8 ); - P( C, D, A, B, 7, 14, 0x676F02D9 ); - P( B, C, D, A, 12, 20, 0x8D2A4C8A ); - -#undef F - -#define F(x,y,z) (x ^ y ^ z) - - P( A, B, C, D, 5, 4, 0xFFFA3942 ); - P( D, A, B, C, 8, 11, 0x8771F681 ); - P( C, D, A, B, 11, 16, 0x6D9D6122 ); - P( B, C, D, A, 14, 23, 0xFDE5380C ); - P( A, B, C, D, 1, 4, 0xA4BEEA44 ); - P( D, A, B, C, 4, 11, 0x4BDECFA9 ); - P( C, D, A, B, 7, 16, 0xF6BB4B60 ); - P( B, C, D, A, 10, 23, 0xBEBFBC70 ); - P( A, B, C, D, 13, 4, 0x289B7EC6 ); - P( D, A, B, C, 0, 11, 0xEAA127FA ); - P( C, D, A, B, 3, 16, 0xD4EF3085 ); - P( B, C, D, A, 6, 23, 0x04881D05 ); - P( A, B, C, D, 9, 4, 0xD9D4D039 ); - P( D, A, B, C, 12, 11, 0xE6DB99E5 ); - P( C, D, A, B, 15, 16, 0x1FA27CF8 ); - P( B, C, D, A, 2, 23, 0xC4AC5665 ); - -#undef F - -#define F(x,y,z) (y ^ (x | ~z)) - - P( A, B, C, D, 0, 6, 0xF4292244 ); - P( D, A, B, C, 7, 10, 0x432AFF97 ); - P( C, D, A, B, 14, 15, 0xAB9423A7 ); - P( B, C, D, A, 5, 21, 0xFC93A039 ); - P( A, B, C, D, 12, 6, 0x655B59C3 ); - P( D, A, B, C, 3, 10, 0x8F0CCC92 ); - P( C, D, A, B, 10, 15, 0xFFEFF47D ); - P( B, C, D, A, 1, 21, 0x85845DD1 ); - P( A, B, C, D, 8, 6, 0x6FA87E4F ); - P( D, A, B, C, 15, 10, 0xFE2CE6E0 ); - P( C, D, A, B, 6, 15, 0xA3014314 ); - P( B, C, D, A, 13, 21, 0x4E0811A1 ); - P( A, B, C, D, 4, 6, 0xF7537E82 ); - P( D, A, B, C, 11, 10, 0xBD3AF235 ); - P( C, D, A, B, 2, 15, 0x2AD7D2BB ); - P( B, C, D, A, 9, 21, 0xEB86D391 ); - -#undef F - - ctx->state[0] += A; - ctx->state[1] += B; - ctx->state[2] += C; - ctx->state[3] += D; -} - -/*=========================================================================== - -FUNCTION MD5_UPDATE - -DESCRIPTION - MD5 process buffer - -DEPENDENCIES - - -RETURN VALUE - - -===========================================================================*/ -__device__ -static void md5_update( md5_context *ctx, unsigned char *input, int ilen ) { - int fill; - unsigned long left; - - if( ilen <= 0 ) - return; - - left = ctx->total[0] & 0x3F; - fill = 64 - left; - - ctx->total[0] += ilen; - ctx->total[0] &= 0xFFFFFFFF; - - if( ctx->total[0] < (unsigned long) ilen ) - ctx->total[1]++; - - if( left && ilen >= fill ) { - - // - /*memcpy( (void *) (ctx->buffer + left), - (void *) input, fill );*/ - for (int i = 0; i < fill; i++) { - ctx->buffer[i+left] = input[i]; - } - // - - md5_process( ctx, ctx->buffer ); - input += fill; - ilen -= fill; - left = 0; - } - - while( ilen >= 64 ) { - md5_process( ctx, input ); - input += 64; - ilen -= 64; - } - - if( ilen > 0 ) { - - // - /* memcpy( (void *) (ctx->buffer + left), - (void *) input, ilen );*/ - for (int i = 0; i < ilen; i++) { - ctx->buffer[i+left] = input[i]; - } - // - - } -} - -/*=========================================================================== - -FUNCTION MD5_FINISH - -DESCRIPTION - MD5 final digest - -DEPENDENCIES - None. - -RETURN VALUE - - -===========================================================================*/ -__device__ -void md5_finish( md5_context *ctx, unsigned char *output ) { - - unsigned long last, padn; - unsigned long high, low; - unsigned char msglen[8]; - - high = ( ctx->total[0] >> 29 ) | ( ctx->total[1] << 3 ); - low = ( ctx->total[0] << 3 ); - - PUT_UINT32_LE( low, msglen, 0 ); - PUT_UINT32_LE( high, msglen, 4 ); - - last = ctx->total[0] & 0x3F; - padn = ( last < 56 ) ? ( 56 - last ) : ( 120 - last ); - - md5_update( ctx, (unsigned char *) md5_padding, padn ); - md5_update( ctx, msglen, 8 ); - - - PUT_UINT32_LE( ctx->state[0], output, 0 ); -#ifndef FEATURE_REDUCED_HASH_SIZE - PUT_UINT32_LE( ctx->state[1], output, 4 ); - PUT_UINT32_LE( ctx->state[2], output, 8 ); - PUT_UINT32_LE( ctx->state[3], output, 12 ); -#endif -} - -/*=========================================================================== - -FUNCTION MD5_INTERNAL - -DESCRIPTION - Does the real md5 algorithm - -DEPENDENCIES - None - -RETURN VALUE - output is the hash result - -===========================================================================*/ -__device__ -static void md5_internal( unsigned char *input, int ilen, - unsigned char *output ) { - md5_context ctx; - - md5_starts( &ctx ); - md5_update( &ctx, input, ilen ); - md5_finish( &ctx, output ); - -} -#endif /* #ifndef FEATURE_SHARED_MEMORY */ - -#ifdef FEATURE_SHARED_MEMORY -/*=========================================================================== - -FUNCTION MD5_INTERNAL - -DESCRIPTION - Does the real md5 algorithm. - -DEPENDENCIES - None - -RETURN VALUE - output is the hash result - -===========================================================================*/ - -__device__ -static void md5_internal( unsigned int *input, unsigned int *sharedMemory, - int chunkSize, unsigned char *output ) { - - /* Number of passes (512 bit blocks) we have to do */ - int numberOfPasses = chunkSize / 64 + 1; - /* Used during the hashing process */ - unsigned long A, B, C, D; - /* Needed to do the little endian stuff */ - unsigned char *data = (unsigned char *)sharedMemory; - - /* Will hold the hash value through the - intermediate stages of MD5 algorithm */ - unsigned int state0 = 0x67452301; - unsigned int state1 = 0xEFCDAB89; - unsigned int state2 = 0x98BADCFE; - unsigned int state3 = 0x10325476; - - - /* Used to cache the shared memory index calculations, but testing showed - that it has no performance effect. */ - int x0 = SHARED_MEMORY_INDEX(0); - int x1 = SHARED_MEMORY_INDEX(1); - int x2 = SHARED_MEMORY_INDEX(2); - int x3 = SHARED_MEMORY_INDEX(3); - int x4 = SHARED_MEMORY_INDEX(4); - int x5 = SHARED_MEMORY_INDEX(5); - int x6 = SHARED_MEMORY_INDEX(6); - int x7 = SHARED_MEMORY_INDEX(7); - int x8 = SHARED_MEMORY_INDEX(8); - int x9 = SHARED_MEMORY_INDEX(9); - int x10 = SHARED_MEMORY_INDEX(10); - int x11 = SHARED_MEMORY_INDEX(11); - int x12 = SHARED_MEMORY_INDEX(12); - int x13 = SHARED_MEMORY_INDEX(13); - int x14 = SHARED_MEMORY_INDEX(14); - int x15 = SHARED_MEMORY_INDEX(15); - -#undef GET_CACHED_INDEX -#define GET_CACHED_INDEX(index) (x##index) - - - for( int index = 0 ; index < (numberOfPasses) ; index++ ) { - - /* Move data to the thread's shared memory space */ - sharedMemory[GET_CACHED_INDEX(0)] = input[0 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(1)] = input[1 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(2)] = input[2 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(3)] = input[3 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(4)] = input[4 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(5)] = input[5 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(6)] = input[6 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(7)] = input[7 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(8)] = input[8 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(9)] = input[9 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(10)] = input[10 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(11)] = input[11 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(12)] = input[12 + 16 * index]; - - /* Testing the code with and without this if statement shows that - it has no effect on performance. */ - if(index == numberOfPasses -1 ) { - /* The last pass will contain the size of the chunk size (according to - official MD5 algorithm). */ - sharedMemory[GET_CACHED_INDEX(13)] = 0x00000080; - sharedMemory[GET_CACHED_INDEX(14)] = chunkSize << 3; - sharedMemory[GET_CACHED_INDEX(15)] = chunkSize >> 29; - } else { - sharedMemory[GET_CACHED_INDEX(13)] = input[13 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(14)] = input[14 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(15)] = input[15 + 16 * index]; - } - - /* Get the little endian stuff done. */ - GET_UINT32_LE( sharedMemory[ GET_CACHED_INDEX(0)], - data, GET_CACHED_INDEX(0) * 4 ); - GET_UINT32_LE( sharedMemory[ GET_CACHED_INDEX(1)], - data, GET_CACHED_INDEX(1) * 4 ); - GET_UINT32_LE( sharedMemory[ GET_CACHED_INDEX(2)], - data, GET_CACHED_INDEX(2) * 4 ); - GET_UINT32_LE( sharedMemory[ GET_CACHED_INDEX(3)], - data, GET_CACHED_INDEX(3) * 4 ); - GET_UINT32_LE( sharedMemory[ GET_CACHED_INDEX(4)], - data, GET_CACHED_INDEX(4) * 4 ); - GET_UINT32_LE( sharedMemory[ GET_CACHED_INDEX(5)], - data, GET_CACHED_INDEX(5) * 4 ); - GET_UINT32_LE( sharedMemory[ GET_CACHED_INDEX(6)], - data, GET_CACHED_INDEX(6) * 4 ); - GET_UINT32_LE( sharedMemory[ GET_CACHED_INDEX(7)], - data, GET_CACHED_INDEX(7) * 4 ); - GET_UINT32_LE( sharedMemory[ GET_CACHED_INDEX(8)], - data, GET_CACHED_INDEX(8) * 4 ); - GET_UINT32_LE( sharedMemory[ GET_CACHED_INDEX(9)], - data, GET_CACHED_INDEX(9) * 4 ); - GET_UINT32_LE( sharedMemory[GET_CACHED_INDEX(10)], - data, GET_CACHED_INDEX(10) * 4 ); - GET_UINT32_LE( sharedMemory[GET_CACHED_INDEX(11)], - data, GET_CACHED_INDEX(11) * 4 ); - GET_UINT32_LE( sharedMemory[GET_CACHED_INDEX(12)], - data, GET_CACHED_INDEX(12) * 4 ); - GET_UINT32_LE( sharedMemory[GET_CACHED_INDEX(13)], - data, GET_CACHED_INDEX(13) * 4 ); - GET_UINT32_LE( sharedMemory[GET_CACHED_INDEX(14)], - data, GET_CACHED_INDEX(14) * 4 ); - GET_UINT32_LE( sharedMemory[GET_CACHED_INDEX(15)], - data, GET_CACHED_INDEX(15) * 4 ); - - - /* Start the MD5 permutations */ -#undef S -#define S(x,n) ((x << n) | ((x & 0xFFFFFFFF) >> (32 - n))) -#undef P -#define P(a,b,c,d,k,s,t) { \ - a += F(b,c,d) + sharedMemory[GET_CACHED_INDEX(k)] + t; a = S(a,s) + b; \ - } \ - - A = state0; - B = state1; - C = state2; - D = state3; - -#undef F - -#define F(x,y,z) (z ^ (x & (y ^ z))) - - P( A, B, C, D, 0, 7, 0xD76AA478 ); - P( D, A, B, C, 1, 12, 0xE8C7B756 ); - P( C, D, A, B, 2, 17, 0x242070DB ); - P( B, C, D, A, 3, 22, 0xC1BDCEEE ); - P( A, B, C, D, 4, 7, 0xF57C0FAF ); - P( D, A, B, C, 5, 12, 0x4787C62A ); - P( C, D, A, B, 6, 17, 0xA8304613 ); - P( B, C, D, A, 7, 22, 0xFD469501 ); - P( A, B, C, D, 8, 7, 0x698098D8 ); - P( D, A, B, C, 9, 12, 0x8B44F7AF ); - P( C, D, A, B, 10, 17, 0xFFFF5BB1 ); - P( B, C, D, A, 11, 22, 0x895CD7BE ); - P( A, B, C, D, 12, 7, 0x6B901122 ); - P( D, A, B, C, 13, 12, 0xFD987193 ); - P( C, D, A, B, 14, 17, 0xA679438E ); - P( B, C, D, A, 15, 22, 0x49B40821 ); - -#undef F - -#define F(x,y,z) (y ^ (z & (x ^ y))) - - P( A, B, C, D, 1, 5, 0xF61E2562 ); - P( D, A, B, C, 6, 9, 0xC040B340 ); - P( C, D, A, B, 11, 14, 0x265E5A51 ); - P( B, C, D, A, 0, 20, 0xE9B6C7AA ); - P( A, B, C, D, 5, 5, 0xD62F105D ); - P( D, A, B, C, 10, 9, 0x02441453 ); - P( C, D, A, B, 15, 14, 0xD8A1E681 ); - P( B, C, D, A, 4, 20, 0xE7D3FBC8 ); - P( A, B, C, D, 9, 5, 0x21E1CDE6 ); - P( D, A, B, C, 14, 9, 0xC33707D6 ); - P( C, D, A, B, 3, 14, 0xF4D50D87 ); - P( B, C, D, A, 8, 20, 0x455A14ED ); - P( A, B, C, D, 13, 5, 0xA9E3E905 ); - P( D, A, B, C, 2, 9, 0xFCEFA3F8 ); - P( C, D, A, B, 7, 14, 0x676F02D9 ); - P( B, C, D, A, 12, 20, 0x8D2A4C8A ); - -#undef F - -#define F(x,y,z) (x ^ y ^ z) - - P( A, B, C, D, 5, 4, 0xFFFA3942 ); - P( D, A, B, C, 8, 11, 0x8771F681 ); - P( C, D, A, B, 11, 16, 0x6D9D6122 ); - P( B, C, D, A, 14, 23, 0xFDE5380C ); - P( A, B, C, D, 1, 4, 0xA4BEEA44 ); - P( D, A, B, C, 4, 11, 0x4BDECFA9 ); - P( C, D, A, B, 7, 16, 0xF6BB4B60 ); - P( B, C, D, A, 10, 23, 0xBEBFBC70 ); - P( A, B, C, D, 13, 4, 0x289B7EC6 ); - P( D, A, B, C, 0, 11, 0xEAA127FA ); - P( C, D, A, B, 3, 16, 0xD4EF3085 ); - P( B, C, D, A, 6, 23, 0x04881D05 ); - P( A, B, C, D, 9, 4, 0xD9D4D039 ); - P( D, A, B, C, 12, 11, 0xE6DB99E5 ); - P( C, D, A, B, 15, 16, 0x1FA27CF8 ); - P( B, C, D, A, 2, 23, 0xC4AC5665 ); - -#undef F - -#define F(x,y,z) (y ^ (x | ~z)) - - P( A, B, C, D, 0, 6, 0xF4292244 ); - P( D, A, B, C, 7, 10, 0x432AFF97 ); - P( C, D, A, B, 14, 15, 0xAB9423A7 ); - P( B, C, D, A, 5, 21, 0xFC93A039 ); - P( A, B, C, D, 12, 6, 0x655B59C3 ); - P( D, A, B, C, 3, 10, 0x8F0CCC92 ); - P( C, D, A, B, 10, 15, 0xFFEFF47D ); - P( B, C, D, A, 1, 21, 0x85845DD1 ); - P( A, B, C, D, 8, 6, 0x6FA87E4F ); - P( D, A, B, C, 15, 10, 0xFE2CE6E0 ); - P( C, D, A, B, 6, 15, 0xA3014314 ); - P( B, C, D, A, 13, 21, 0x4E0811A1 ); - P( A, B, C, D, 4, 6, 0xF7537E82 ); - P( D, A, B, C, 11, 10, 0xBD3AF235 ); - P( C, D, A, B, 2, 15, 0x2AD7D2BB ); - P( B, C, D, A, 9, 21, 0xEB86D391 ); - -#undef F - - state0 += A; - state1 += B; - state2 += C; - state3 += D; - } - - /* Got the hash, store it in the output buffer. */ - PUT_UINT32_LE( state0, output, 0 ); -#ifndef FEATURE_REDUCED_HASH_SIZE - PUT_UINT32_LE( state1, output, 4 ); - PUT_UINT32_LE( state2, output, 8 ); - PUT_UINT32_LE( state3, output, 12 ); -#endif - -} - -__device__ -static void md5_internal_overlap( unsigned int *input, unsigned int *sharedMemory, - int chunkSize, unsigned char *output ) { - - /* Number of passes (512 bit blocks) we have to do */ - int numberOfPasses = chunkSize / 64 + 1; - /* Used during the hashing process */ - unsigned long A, B, C, D; - /* Needed to do the little endian stuff */ - unsigned char *data = (unsigned char *)sharedMemory; - // number of padding bytes. - int numPadBytes = 0; - int numPadInt = 0; - //int numPadRemain = 0; - - /* Will hold the hash value through the - intermediate stages of MD5 algorithm */ - unsigned int state0 = 0x67452301; - unsigned int state1 = 0xEFCDAB89; - unsigned int state2 = 0x98BADCFE; - unsigned int state3 = 0x10325476; - - - /* Used to cache the shared memory index calculations, but testing showed - that it has no performance effect. */ - int x0 = SHARED_MEMORY_INDEX(0); - int x1 = SHARED_MEMORY_INDEX(1); - int x2 = SHARED_MEMORY_INDEX(2); - int x3 = SHARED_MEMORY_INDEX(3); - int x4 = SHARED_MEMORY_INDEX(4); - int x5 = SHARED_MEMORY_INDEX(5); - int x6 = SHARED_MEMORY_INDEX(6); - int x7 = SHARED_MEMORY_INDEX(7); - int x8 = SHARED_MEMORY_INDEX(8); - int x9 = SHARED_MEMORY_INDEX(9); - int x10 = SHARED_MEMORY_INDEX(10); - int x11 = SHARED_MEMORY_INDEX(11); - int x12 = SHARED_MEMORY_INDEX(12); - int x13 = SHARED_MEMORY_INDEX(13); - int x14 = SHARED_MEMORY_INDEX(14); - int x15 = SHARED_MEMORY_INDEX(15); - -#undef GET_CACHED_INDEX -#define GET_CACHED_INDEX(index) (x##index) - - - for( int index = 0 ; index < (numberOfPasses) ; index++ ) { - - if(index == numberOfPasses - 1 ) { - - numPadBytes = (64-12) - (chunkSize - (numberOfPasses-1)*64); - numPadInt = numPadBytes/sizeof(int); - /*numPadRemain = numPadBytes-numPadInt*sizeof(int); - printf("\nLast loop chunkSize = %d, numberOfPasses= %d and \nnumPadBytes = %d, numPadInt =%d, numPadRemain = %d\n", - chunkSize,numberOfPasses,numPadBytes,numPadInt,numPadRemain);*/ - - int i=0; - for(i = 0 ; i < numPadInt ; i++){ - sharedMemory[SHARED_MEMORY_INDEX(13-i)] = 0; - } - - int j=0; - for(j=0;j<(16-3-numPadInt);j++){ - //printf("j= %d\n",j); - sharedMemory[SHARED_MEMORY_INDEX(j)] = input[j + 16 * index]; - } - - - /* The last pass will contain the size of the chunk size (according to - official MD5 algorithm). */ - sharedMemory[SHARED_MEMORY_INDEX(13-i)] = 0x00000080; - //printf("the last one at %d\n",13-i); - - sharedMemory[GET_CACHED_INDEX(14)] = chunkSize << 3; - sharedMemory[GET_CACHED_INDEX(15)] = chunkSize >> 29; - } else { - /* Move data to the thread's shared memory space */ - //printf("Not last loop\n"); - sharedMemory[GET_CACHED_INDEX(0)] = input[0 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(1)] = input[1 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(2)] = input[2 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(3)] = input[3 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(4)] = input[4 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(5)] = input[5 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(6)] = input[6 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(7)] = input[7 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(8)] = input[8 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(9)] = input[9 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(10)] = input[10 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(11)] = input[11 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(12)] = input[12 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(13)] = input[13 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(14)] = input[14 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(15)] = input[15 + 16 * index]; - } - - /* Get the little endian stuff done. */ - GET_UINT32_LE( sharedMemory[ GET_CACHED_INDEX(0)], - data, GET_CACHED_INDEX(0) * 4 ); - GET_UINT32_LE( sharedMemory[ GET_CACHED_INDEX(1)], - data, GET_CACHED_INDEX(1) * 4 ); - GET_UINT32_LE( sharedMemory[ GET_CACHED_INDEX(2)], - data, GET_CACHED_INDEX(2) * 4 ); - GET_UINT32_LE( sharedMemory[ GET_CACHED_INDEX(3)], - data, GET_CACHED_INDEX(3) * 4 ); - GET_UINT32_LE( sharedMemory[ GET_CACHED_INDEX(4)], - data, GET_CACHED_INDEX(4) * 4 ); - GET_UINT32_LE( sharedMemory[ GET_CACHED_INDEX(5)], - data, GET_CACHED_INDEX(5) * 4 ); - GET_UINT32_LE( sharedMemory[ GET_CACHED_INDEX(6)], - data, GET_CACHED_INDEX(6) * 4 ); - GET_UINT32_LE( sharedMemory[ GET_CACHED_INDEX(7)], - data, GET_CACHED_INDEX(7) * 4 ); - GET_UINT32_LE( sharedMemory[ GET_CACHED_INDEX(8)], - data, GET_CACHED_INDEX(8) * 4 ); - GET_UINT32_LE( sharedMemory[ GET_CACHED_INDEX(9)], - data, GET_CACHED_INDEX(9) * 4 ); - GET_UINT32_LE( sharedMemory[GET_CACHED_INDEX(10)], - data, GET_CACHED_INDEX(10) * 4 ); - GET_UINT32_LE( sharedMemory[GET_CACHED_INDEX(11)], - data, GET_CACHED_INDEX(11) * 4 ); - GET_UINT32_LE( sharedMemory[GET_CACHED_INDEX(12)], - data, GET_CACHED_INDEX(12) * 4 ); - GET_UINT32_LE( sharedMemory[GET_CACHED_INDEX(13)], - data, GET_CACHED_INDEX(13) * 4 ); - GET_UINT32_LE( sharedMemory[GET_CACHED_INDEX(14)], - data, GET_CACHED_INDEX(14) * 4 ); - GET_UINT32_LE( sharedMemory[GET_CACHED_INDEX(15)], - data, GET_CACHED_INDEX(15) * 4 ); - - - /* Start the MD5 permutations */ -#undef S -#define S(x,n) ((x << n) | ((x & 0xFFFFFFFF) >> (32 - n))) -#undef P -#define P(a,b,c,d,k,s,t) { \ - a += F(b,c,d) + sharedMemory[GET_CACHED_INDEX(k)] + t; a = S(a,s) + b; \ - } \ - - A = state0; - B = state1; - C = state2; - D = state3; - -#undef F - -#define F(x,y,z) (z ^ (x & (y ^ z))) - - P( A, B, C, D, 0, 7, 0xD76AA478 ); - P( D, A, B, C, 1, 12, 0xE8C7B756 ); - P( C, D, A, B, 2, 17, 0x242070DB ); - P( B, C, D, A, 3, 22, 0xC1BDCEEE ); - P( A, B, C, D, 4, 7, 0xF57C0FAF ); - P( D, A, B, C, 5, 12, 0x4787C62A ); - P( C, D, A, B, 6, 17, 0xA8304613 ); - P( B, C, D, A, 7, 22, 0xFD469501 ); - P( A, B, C, D, 8, 7, 0x698098D8 ); - P( D, A, B, C, 9, 12, 0x8B44F7AF ); - P( C, D, A, B, 10, 17, 0xFFFF5BB1 ); - P( B, C, D, A, 11, 22, 0x895CD7BE ); - P( A, B, C, D, 12, 7, 0x6B901122 ); - P( D, A, B, C, 13, 12, 0xFD987193 ); - P( C, D, A, B, 14, 17, 0xA679438E ); - P( B, C, D, A, 15, 22, 0x49B40821 ); - -#undef F - -#define F(x,y,z) (y ^ (z & (x ^ y))) - - P( A, B, C, D, 1, 5, 0xF61E2562 ); - P( D, A, B, C, 6, 9, 0xC040B340 ); - P( C, D, A, B, 11, 14, 0x265E5A51 ); - P( B, C, D, A, 0, 20, 0xE9B6C7AA ); - P( A, B, C, D, 5, 5, 0xD62F105D ); - P( D, A, B, C, 10, 9, 0x02441453 ); - P( C, D, A, B, 15, 14, 0xD8A1E681 ); - P( B, C, D, A, 4, 20, 0xE7D3FBC8 ); - P( A, B, C, D, 9, 5, 0x21E1CDE6 ); - P( D, A, B, C, 14, 9, 0xC33707D6 ); - P( C, D, A, B, 3, 14, 0xF4D50D87 ); - P( B, C, D, A, 8, 20, 0x455A14ED ); - P( A, B, C, D, 13, 5, 0xA9E3E905 ); - P( D, A, B, C, 2, 9, 0xFCEFA3F8 ); - P( C, D, A, B, 7, 14, 0x676F02D9 ); - P( B, C, D, A, 12, 20, 0x8D2A4C8A ); - -#undef F - -#define F(x,y,z) (x ^ y ^ z) - - P( A, B, C, D, 5, 4, 0xFFFA3942 ); - P( D, A, B, C, 8, 11, 0x8771F681 ); - P( C, D, A, B, 11, 16, 0x6D9D6122 ); - P( B, C, D, A, 14, 23, 0xFDE5380C ); - P( A, B, C, D, 1, 4, 0xA4BEEA44 ); - P( D, A, B, C, 4, 11, 0x4BDECFA9 ); - P( C, D, A, B, 7, 16, 0xF6BB4B60 ); - P( B, C, D, A, 10, 23, 0xBEBFBC70 ); - P( A, B, C, D, 13, 4, 0x289B7EC6 ); - P( D, A, B, C, 0, 11, 0xEAA127FA ); - P( C, D, A, B, 3, 16, 0xD4EF3085 ); - P( B, C, D, A, 6, 23, 0x04881D05 ); - P( A, B, C, D, 9, 4, 0xD9D4D039 ); - P( D, A, B, C, 12, 11, 0xE6DB99E5 ); - P( C, D, A, B, 15, 16, 0x1FA27CF8 ); - P( B, C, D, A, 2, 23, 0xC4AC5665 ); - -#undef F - -#define F(x,y,z) (y ^ (x | ~z)) - - P( A, B, C, D, 0, 6, 0xF4292244 ); - P( D, A, B, C, 7, 10, 0x432AFF97 ); - P( C, D, A, B, 14, 15, 0xAB9423A7 ); - P( B, C, D, A, 5, 21, 0xFC93A039 ); - P( A, B, C, D, 12, 6, 0x655B59C3 ); - P( D, A, B, C, 3, 10, 0x8F0CCC92 ); - P( C, D, A, B, 10, 15, 0xFFEFF47D ); - P( B, C, D, A, 1, 21, 0x85845DD1 ); - P( A, B, C, D, 8, 6, 0x6FA87E4F ); - P( D, A, B, C, 15, 10, 0xFE2CE6E0 ); - P( C, D, A, B, 6, 15, 0xA3014314 ); - P( B, C, D, A, 13, 21, 0x4E0811A1 ); - P( A, B, C, D, 4, 6, 0xF7537E82 ); - P( D, A, B, C, 11, 10, 0xBD3AF235 ); - P( C, D, A, B, 2, 15, 0x2AD7D2BB ); - P( B, C, D, A, 9, 21, 0xEB86D391 ); - -#undef F - - state0 += A; - state1 += B; - state2 += C; - state3 += D; - } - - /* Got the hash, store it in the output buffer. */ - PUT_UINT32_LE( state0, output, 0 ); -#ifndef FEATURE_REDUCED_HASH_SIZE - PUT_UINT32_LE( state1, output, 4 ); - PUT_UINT32_LE( state2, output, 8 ); - PUT_UINT32_LE( state3, output, 12 ); -#endif - -} -#endif - -/*-------------------------------------------------------------------------- - GLOBAL FUNCTIONS ---------------------------------------------------------------------------*/ - -/*=========================================================================== - -FUNCTION MD5 - -DESCRIPTION - Main md5 hash function - -DEPENDENCIES - GPU must be initialized - -RETURN VALUE - output: the hash result - -===========================================================================*/ -__global__ -void md5( unsigned char *input, int chunkSize, int totalThreads, - int padSize, unsigned char *scratch) { - - int threadIndex = threadIdx.x + blockDim.x * blockIdx.x; - int chunkIndex = threadIndex * chunkSize; - int hashIndex = threadIndex * MD5_HASH_SIZE; - - if(threadIndex >= totalThreads) - return; - - if ((threadIndex == (totalThreads - 1)) && (padSize > 0)) { - for(int i = 0 ; i < padSize ; i++) - input[chunkIndex + chunkSize - padSize + i] = 0; - } - - -#ifdef FEATURE_SHARED_MEMORY - - __shared__ unsigned int sharedMemory[4 * 1024 - 32]; - - // 512 words are allocated for every warp of 32 threads - unsigned int *sharedMemoryIndex = sharedMemory + ((threadIdx.x >> 5) * 512); - unsigned int *inputIndex = (unsigned int *)(input + chunkIndex); - - md5_internal(inputIndex, sharedMemoryIndex, chunkSize, - scratch + hashIndex ); - -#else - md5_internal(input + chunkIndex, chunkSize, scratch + hashIndex ); -#endif /* FEATURE_SHARED_MEMORY */ - -} - - -__global__ -void md5_overlap( unsigned char *input, int chunkSize, int offset, - int totalThreads, int padSize, unsigned char *output ) { - - int threadIndex = threadIdx.x + blockDim.x * blockIdx.x; - int chunkIndex = threadIndex * offset; - int hashIndex = threadIndex * MD5_HASH_SIZE; - - - if(threadIndex >= totalThreads) - return; - - if ((threadIndex == (totalThreads - 1))) { - chunkSize-= padSize; - } - - -#ifdef FEATURE_SHARED_MEMORY - - __shared__ unsigned int sharedMemory[4 * 1024 - 32]; - - unsigned int *sharedMemoryIndex = sharedMemory + ((threadIdx.x >> 5) * 512); - unsigned int *inputIndex = (unsigned int *)(input + chunkIndex); - - md5_internal_overlap(inputIndex, sharedMemoryIndex, chunkSize, - output + hashIndex ); - -#else - md5_internal(input + chunkIndex, chunkSize, output + hashIndex ); -#endif /* FEATURE_SHARED_MEMORY */ - - -} - diff --git a/benchmarks/CUDA/STO/sha1_cpu.c b/benchmarks/CUDA/STO/sha1_cpu.c deleted file mode 100644 index f4a6be5..0000000 --- a/benchmarks/CUDA/STO/sha1_cpu.c +++ /dev/null @@ -1,429 +0,0 @@ -/*========================================================================== - S H A 1 C P U - -DESCRIPTION - CPU implementation of the sha1 algorithm. - - * - * FIPS-180-1 compliant SHA-1 implementation - * - * Copyright (C) 2006-2007 Christophe Devine - * - * This library is free software; you can redistribute it and/or - * modify it under the terms of the GNU Lesser General Public - * License, version 2.1 as published by the Free Software Foundation. - * - * This library 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 - * Lesser General Public License for more details. - * - * You should have received a copy of the GNU Lesser General Public - * License along with this library; if not, write to the Free Software - * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, - * MA 02110-1301 USA - * - * The SHA-1 standard was published by NIST in 1993. - * - * http://www.itl.nist.gov/fipspubs/fip180-1.htm - * -==========================================================================*/ - -/*========================================================================== - - INCLUDES - -==========================================================================*/ -#include -#include - -#include "sha1_cpu.h" -#include "cust.h" -/*========================================================================== - - DATA DECLARATIONS - -==========================================================================*/ - -/*-------------------------------------------------------------------------- - TYPE DEFINITIONS ---------------------------------------------------------------------------*/ -typedef struct { - unsigned long total[2]; /*!< number of bytes processed */ - unsigned long state[5]; /*!< intermediate digest state */ - unsigned char buffer[64]; /*!< data block being processed */ -} sha1_cpu_context; - - -/*-------------------------------------------------------------------------- - FUNCTION PROTOTYPES ---------------------------------------------------------------------------*/ - -/*-------------------------------------------------------------------------- - CONSTANTS ---------------------------------------------------------------------------*/ -static const unsigned char sha1_padding[64] = -{ - 0x80, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, - 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, - 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, - 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0 -}; - -/*-------------------------------------------------------------------------- - GLOBAL VARIABLES ---------------------------------------------------------------------------*/ - -/*-------------------------------------------------------------------------- - MACROS ---------------------------------------------------------------------------*/ -/* - * 32-bit integer manipulation macros (big endian) - */ -#ifndef GET_UINT32_BE -#define GET_UINT32_BE(n,b,i) \ -{ \ - (n) = ( (unsigned long) (b)[(i) ] << 24 ) \ - | ( (unsigned long) (b)[(i) + 1] << 16 ) \ - | ( (unsigned long) (b)[(i) + 2] << 8 ) \ - | ( (unsigned long) (b)[(i) + 3] ); \ -} -#endif - -#ifndef PUT_UINT32_BE -#define PUT_UINT32_BE(n,b,i) \ -{ \ - (b)[(i) ] = (unsigned char) ( (n) >> 24 ); \ - (b)[(i) + 1] = (unsigned char) ( (n) >> 16 ); \ - (b)[(i) + 2] = (unsigned char) ( (n) >> 8 ); \ - (b)[(i) + 3] = (unsigned char) ( (n) ); \ -} -#endif - -#ifndef _CRT_SECURE_NO_DEPRECATE -#define _CRT_SECURE_NO_DEPRECATE 1 -#endif - - -/*=========================================================================== - -FUNCTION SHA1_CPU_STARTS - -DESCRIPTION - SHA-1 context setup - -DEPENDENCIES - None - -RETURN VALUE - None - -===========================================================================*/ -void sha1_cpu_starts( sha1_cpu_context* ctx ) { - ctx->total[0] = 0; - ctx->total[1] = 0; - - ctx->state[0] = 0x67452301; - ctx->state[1] = 0xEFCDAB89; - ctx->state[2] = 0x98BADCFE; - ctx->state[3] = 0x10325476; - ctx->state[4] = 0xC3D2E1F0; -} - -/*=========================================================================== - -FUNCTION SHA1_CPU_PROCESS - -DESCRIPTION - SHA1 process buffer - -DEPENDENCIES - None - -RETURN VALUE - None - -===========================================================================*/ -static void sha1_cpu_process( sha1_cpu_context *ctx, unsigned char data[64]) { - unsigned long temp, W[16], A, B, C, D, E; - - GET_UINT32_BE( W[ 0], data, 0 ); - GET_UINT32_BE( W[ 1], data, 4 ); - GET_UINT32_BE( W[ 2], data, 8 ); - GET_UINT32_BE( W[ 3], data, 12 ); - GET_UINT32_BE( W[ 4], data, 16 ); - GET_UINT32_BE( W[ 5], data, 20 ); - GET_UINT32_BE( W[ 6], data, 24 ); - GET_UINT32_BE( W[ 7], data, 28 ); - GET_UINT32_BE( W[ 8], data, 32 ); - GET_UINT32_BE( W[ 9], data, 36 ); - GET_UINT32_BE( W[10], data, 40 ); - GET_UINT32_BE( W[11], data, 44 ); - GET_UINT32_BE( W[12], data, 48 ); - GET_UINT32_BE( W[13], data, 52 ); - GET_UINT32_BE( W[14], data, 56 ); - GET_UINT32_BE( W[15], data, 60 ); - -#define S(x,n) ((x << n) | ((x & 0xFFFFFFFF) >> (32 - n))) - -/* #define R(t) \ */ -/* ( \ */ -/* temp = W[(t - 3) & 0x000F], \ */ -/* ( W[t & 0x0F] = S(temp,1) ) \ */ -/* ) */ - -#define R(t) \ -( \ - temp = W[(t - 3) & 0x0F] ^ W[(t - 8) & 0x0F] ^ \ - W[(t - 14) & 0x0F] ^ W[ t & 0x0F], \ - ( W[t & 0x0F] = S(temp,1) ) \ -) - -#define P(a,b,c,d,e,x) \ -{ \ - e += S(a,5) + F(b,c,d) + K + x; b = S(b,30); \ -} - - A = ctx->state[0]; - B = ctx->state[1]; - C = ctx->state[2]; - D = ctx->state[3]; - E = ctx->state[4]; - - -#define F(x,y,z) (z ^ (x & (y ^ z))) -#define K 0x5A827999 - - P( A, B, C, D, E, W[0] ); - P( E, A, B, C, D, W[1] ); - P( D, E, A, B, C, W[2] ); - P( C, D, E, A, B, W[3] ); - P( B, C, D, E, A, W[4] ); - P( A, B, C, D, E, W[5] ); - P( E, A, B, C, D, W[6] ); - P( D, E, A, B, C, W[7] ); - P( C, D, E, A, B, W[8] ); - P( B, C, D, E, A, W[9] ); - P( A, B, C, D, E, W[10] ); - P( E, A, B, C, D, W[11] ); - P( D, E, A, B, C, W[12] ); - P( C, D, E, A, B, W[13] ); - P( B, C, D, E, A, W[14] ); - P( A, B, C, D, E, W[15] ); - P( E, A, B, C, D, R(16) ); - P( D, E, A, B, C, R(17) ); - P( C, D, E, A, B, R(18) ); - P( B, C, D, E, A, R(19) ); - - -#undef K -#undef F - -#define F(x,y,z) (x ^ y ^ z) -#define K 0x6ED9EBA1 - - P( A, B, C, D, E, R(20) ); - P( E, A, B, C, D, R(21) ); - P( D, E, A, B, C, R(22) ); - P( C, D, E, A, B, R(23) ); - P( B, C, D, E, A, R(24) ); - P( A, B, C, D, E, R(25) ); - P( E, A, B, C, D, R(26) ); - P( D, E, A, B, C, R(27) ); - P( C, D, E, A, B, R(28) ); - P( B, C, D, E, A, R(29) ); - P( A, B, C, D, E, R(30) ); - P( E, A, B, C, D, R(31) ); - P( D, E, A, B, C, R(32) ); - P( C, D, E, A, B, R(33) ); - P( B, C, D, E, A, R(34) ); - P( A, B, C, D, E, R(35) ); - P( E, A, B, C, D, R(36) ); - P( D, E, A, B, C, R(37) ); - P( C, D, E, A, B, R(38) ); - P( B, C, D, E, A, R(39) ); - -#undef K -#undef F - -#define F(x,y,z) ((x & y) | (z & (x | y))) -#define K 0x8F1BBCDC - - P( A, B, C, D, E, R(40) ); - P( E, A, B, C, D, R(41) ); - P( D, E, A, B, C, R(42) ); - P( C, D, E, A, B, R(43) ); - P( B, C, D, E, A, R(44) ); - P( A, B, C, D, E, R(45) ); - P( E, A, B, C, D, R(46) ); - P( D, E, A, B, C, R(47) ); - P( C, D, E, A, B, R(48) ); - P( B, C, D, E, A, R(49) ); - P( A, B, C, D, E, R(50) ); - P( E, A, B, C, D, R(51) ); - P( D, E, A, B, C, R(52) ); - P( C, D, E, A, B, R(53) ); - P( B, C, D, E, A, R(54) ); - P( A, B, C, D, E, R(55) ); - P( E, A, B, C, D, R(56) ); - P( D, E, A, B, C, R(57) ); - P( C, D, E, A, B, R(58) ); - P( B, C, D, E, A, R(59) ); - -#undef K -#undef F - -#define F(x,y,z) (x ^ y ^ z) -#define K 0xCA62C1D6 - - P( A, B, C, D, E, R(60) ); - P( E, A, B, C, D, R(61) ); - P( D, E, A, B, C, R(62) ); - P( C, D, E, A, B, R(63) ); - P( B, C, D, E, A, R(64) ); - P( A, B, C, D, E, R(65) ); - P( E, A, B, C, D, R(66) ); - P( D, E, A, B, C, R(67) ); - P( C, D, E, A, B, R(68) ); - P( B, C, D, E, A, R(69) ); - P( A, B, C, D, E, R(70) ); - P( E, A, B, C, D, R(71) ); - P( D, E, A, B, C, R(72) ); - P( C, D, E, A, B, R(73) ); - P( B, C, D, E, A, R(74) ); - P( A, B, C, D, E, R(75) ); - P( E, A, B, C, D, R(76) ); - P( D, E, A, B, C, R(77) ); - P( C, D, E, A, B, R(78) ); - P( B, C, D, E, A, R(79) ); - -#undef K -#undef F - - ctx->state[0] += A; - ctx->state[1] += B; - ctx->state[2] += C; - ctx->state[3] += D; - ctx->state[4] += E; -} -/*=========================================================================== - -FUNCTION SHA1_CPU_UPDATE - -DESCRIPTION - SHA-1 process buffer - -DEPENDENCIES - None - -RETURN VALUE - None - -===========================================================================*/ -void sha1_cpu_update( sha1_cpu_context *ctx, unsigned char *input, int ilen ) { - int fill; - unsigned long left; - - if ( ilen <= 0 ) - return; - - left = ctx->total[0] & 0x3F; - fill = 64 - left; - - ctx->total[0] += ilen; - ctx->total[0] &= 0xFFFFFFFF; - - if ( ctx->total[0] < (unsigned long) ilen ) - ctx->total[1]++; - - if ( left && ilen >= fill ) { - memcpy( (void *) (ctx->buffer + left), - (void *) input, fill ); - - sha1_cpu_process( ctx, ctx->buffer ); - input += fill; - ilen -= fill; - left = 0; - } - - while ( ilen >= 64 ) { - sha1_cpu_process( ctx, input ); - input += 64; - ilen -= 64; - } - - if ( ilen > 0 ) { - memcpy( (void *) (ctx->buffer + left), - (void *) input, ilen ); - } -} - -/*=========================================================================== - -FUNCTION SHA1_CPU_FINISH - -DESCRIPTION - SHA1 final digest - -DEPENDENCIES - None - -RETURN VALUE - None - -===========================================================================*/ -void sha1_cpu_finish( sha1_cpu_context *ctx, unsigned char *output ) { - unsigned long last, padn; - unsigned long high, low; - unsigned char msglen[8]; - - - high = ( ctx->total[0] >> 29 ) - | ( ctx->total[1] << 3 ); - low = ( ctx->total[0] << 3 ); - - PUT_UINT32_BE( high, msglen, 0 ); - PUT_UINT32_BE( low, msglen, 4 ); - - last = ctx->total[0] & 0x3F; - padn = ( last < 56 ) ? ( 56 - last ) : ( 120 - last ); - - sha1_cpu_update( ctx, (unsigned char *) sha1_padding, padn ); - sha1_cpu_update( ctx, msglen, 8 ); - - PUT_UINT32_BE( ctx->state[0], output, 0 ); -#ifndef FEATURE_REDUCED_HASH_SIZE - PUT_UINT32_BE( ctx->state[1], output, 4 ); - PUT_UINT32_BE( ctx->state[2], output, 8 ); - PUT_UINT32_BE( ctx->state[3], output, 12 ); - PUT_UINT32_BE( ctx->state[4], output, 16 ); -#endif -} - -/*-------------------------------------------------------------------------- - GLOBAL FUNCTIONS ---------------------------------------------------------------------------*/ -/*=========================================================================== - -FUNCTION CPU_SHA1_INTERNAL - -DESCRIPTION - CPU implementation of the SHA1 algorithm - -DEPENDENCIES - None - -RETURN VALUE - Hash - -===========================================================================*/ -void sha1_cpu_internal( unsigned char *input, int ilen, - unsigned char *output ) { - sha1_cpu_context ctx; - - sha1_cpu_starts( &ctx ); - sha1_cpu_update( &ctx, input, ilen ); - sha1_cpu_finish( &ctx, output ); - - memset( &ctx, 0, sizeof( sha1_cpu_context ) ); -} diff --git a/benchmarks/CUDA/STO/sha1_cpu.h b/benchmarks/CUDA/STO/sha1_cpu.h deleted file mode 100644 index 637376a..0000000 --- a/benchmarks/CUDA/STO/sha1_cpu.h +++ /dev/null @@ -1,98 +0,0 @@ -#ifndef SHA1_CPU_H -#define SHA1_CPU_H - -/*========================================================================== - S H A 1 C P U - -* Copyright (c) 2008, NetSysLab at the University of British Columbia -* All rights reserved. -* -* Redistribution and use in source and binary forms, with or without -* modification, are permitted provided that the following conditions are met: -* * Redistributions of source code must retain the above copyright -* notice, this list of conditions and the following disclaimer. -* * Redistributions in binary form must reproduce the above copyright -* notice, this list of conditions and the following disclaimer in the -* documentation and/or other materials provided with the distribution. -* * Neither the name of the University nor the -* names of its contributors may be used to endorse or promote products -* derived from this software without specific prior written permission. -* -* THIS SOFTWARE IS PROVIDED BY NetSysLab ``AS IS'' AND ANY -* EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED -* WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE -* DISCLAIMED. IN NO EVENT SHALL NetSysLab BE LIABLE FOR ANY -* DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES -* (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; -* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND -* ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT -* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS -* SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. - -DESCRIPTION - CPU implementation of the sha1 algorithm. - - -==========================================================================*/ - -/*========================================================================== - - INCLUDES - -==========================================================================*/ - - -/*========================================================================== - - DATA DECLARATIONS - -==========================================================================*/ - -/*-------------------------------------------------------------------------- - TYPE DEFINITIONS ---------------------------------------------------------------------------*/ - - -/*-------------------------------------------------------------------------- - FUNCTION PROTOTYPES ---------------------------------------------------------------------------*/ - -/*-------------------------------------------------------------------------- - CONSTANTS ---------------------------------------------------------------------------*/ - -/*-------------------------------------------------------------------------- - GLOBAL VARIABLES ---------------------------------------------------------------------------*/ - - -/*-------------------------------------------------------------------------- - MACROS ---------------------------------------------------------------------------*/ - -/*========================================================================== - - FUNCTIONS - -==========================================================================*/ - -/*=========================================================================== - -FUNCTION cpu_sha1_internal - -DESCRIPTION - CPU implementation of the SHA1 algorithm - -DEPENDENCIES - None - -RETURN VALUE - Hash - -===========================================================================*/ -void sha1_cpu_internal( unsigned char *input, int ilen, - unsigned char *output ); - - - -#endif /* SHA1_CPU_H */ diff --git a/benchmarks/CUDA/STO/sha1_kernel.cu b/benchmarks/CUDA/STO/sha1_kernel.cu deleted file mode 100644 index 2b70fd2..0000000 --- a/benchmarks/CUDA/STO/sha1_kernel.cu +++ /dev/null @@ -1,1140 +0,0 @@ -/*========================================================================== - SHA1 KERNEL - -* Copyright (c) 2008, NetSysLab at the University of British Columbia -* All rights reserved. -* -* Redistribution and use in source and binary forms, with or without -* modification, are permitted provided that the following conditions are met: -* * Redistributions of source code must retain the above copyright -* notice, this list of conditions and the following disclaimer. -* * Redistributions in binary form must reproduce the above copyright -* notice, this list of conditions and the following disclaimer in the -* documentation and/or other materials provided with the distribution. -* * Neither the name of the University nor the -* names of its contributors may be used to endorse or promote products -* derived from this software without specific prior written permission. -* -* THIS SOFTWARE IS PROVIDED BY NetSysLab ``AS IS'' AND ANY -* EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED -* WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE -* DISCLAIMED. IN NO EVENT SHALL NetSysLab BE LIABLE FOR ANY -* DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES -* (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; -* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND -* ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT -* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS -* SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. - -DESCRIPTION - CPU version of the storeGPU library. - - -==========================================================================*/ - -/*========================================================================== - - INCLUDES - -==========================================================================*/ -#include -#include -#include "cust.h" - -/*========================================================================== - - DATA DECLARATIONS - -==========================================================================*/ - -/*-------------------------------------------------------------------------- - TYPE DEFINITIONS ---------------------------------------------------------------------------*/ -typedef struct { - unsigned long total[2]; /*!< number of bytes processed */ - unsigned long state[5]; /*!< intermediate digest state */ - unsigned char buffer[64]; /*!< data block being processed */ -} sha1_context; - -/*-------------------------------------------------------------------------- - FUNCTION PROTOTYPES ---------------------------------------------------------------------------*/ - -/*-------------------------------------------------------------------------- - CONSTANTS ---------------------------------------------------------------------------*/ - -/*-------------------------------------------------------------------------- - GLOBAL VARIABLES ---------------------------------------------------------------------------*/ -__device__ -static const unsigned char sha1_padding[64] = -{ - 0x80, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, - 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, - 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, - 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0 -}; - -/*-------------------------------------------------------------------------- - MACROS ---------------------------------------------------------------------------*/ - -#ifndef _CRT_SECURE_NO_DEPRECATE -#define _CRT_SECURE_NO_DEPRECATE 1 -#endif - - -/* - * 32-bit integer manipulation macros (big endian) - */ -#ifndef GET_UINT32_BE -#define GET_UINT32_BE(n,b,i) \ -{ \ - (n) = ( (unsigned long) (b)[(i) ] << 24 ) \ - | ( (unsigned long) (b)[(i) + 1] << 16 ) \ - | ( (unsigned long) (b)[(i) + 2] << 8 ) \ - | ( (unsigned long) (b)[(i) + 3] ); \ -} -#endif - -#ifndef PUT_UINT32_BE -#define PUT_UINT32_BE(n,b,i) \ -{ \ - (b)[(i) ] = (unsigned char) ( (n) >> 24 ); \ - (b)[(i) + 1] = (unsigned char) ( (n) >> 16 ); \ - (b)[(i) + 2] = (unsigned char) ( (n) >> 8 ); \ - (b)[(i) + 3] = (unsigned char) ( (n) ); \ -} -#endif - -#ifdef FEATURE_SHARED_MEMORY -// current thread stride. -#undef SHARED_MEMORY_INDEX -#define SHARED_MEMORY_INDEX(index) (32 * (index) + (threadIdx.x & 0x1F)) - -#endif /* FEATURE_SHARED_MEMORY */ - - - - -/*-------------------------------------------------------------------------- - LOCAL FUNCTIONS ---------------------------------------------------------------------------*/ -#ifndef FEATURE_SHARED_MEMORY -/* - * SHA-1 context setup - */ - -/*=========================================================================== - -FUNCTION SHA1_GPU_STARTS - -DESCRIPTION - SHA-1 context setup - -DEPENDENCIES - None - -RETURN VALUE - None - -===========================================================================*/ -__device__ -void sha1_starts( sha1_context *ctx ) { - ctx->total[0] = 0; - ctx->total[1] = 0; - - ctx->state[0] = 0x67452301; - ctx->state[1] = 0xEFCDAB89; - ctx->state[2] = 0x98BADCFE; - ctx->state[3] = 0x10325476; - ctx->state[4] = 0xC3D2E1F0; -} - -/*=========================================================================== - -FUNCTION SHA1_GPU_PROCESS - -DESCRIPTION - SHA1 process buffer - -DEPENDENCIES - None - -RETURN VALUE - None - -===========================================================================*/ -__device__ -void sha1_process( sha1_context *ctx, unsigned char data[64] ) { - - unsigned long temp, W[16], A, B, C, D, E; - - GET_UINT32_BE( W[ 0], data, 0 ); - GET_UINT32_BE( W[ 1], data, 4 ); - GET_UINT32_BE( W[ 2], data, 8 ); - GET_UINT32_BE( W[ 3], data, 12 ); - GET_UINT32_BE( W[ 4], data, 16 ); - GET_UINT32_BE( W[ 5], data, 20 ); - GET_UINT32_BE( W[ 6], data, 24 ); - GET_UINT32_BE( W[ 7], data, 28 ); - GET_UINT32_BE( W[ 8], data, 32 ); - GET_UINT32_BE( W[ 9], data, 36 ); - GET_UINT32_BE( W[10], data, 40 ); - GET_UINT32_BE( W[11], data, 44 ); - GET_UINT32_BE( W[12], data, 48 ); - GET_UINT32_BE( W[13], data, 52 ); - GET_UINT32_BE( W[14], data, 56 ); - GET_UINT32_BE( W[15], data, 60 ); - -#undef S -#define S(x,n) ((x << n) | ((x & 0xFFFFFFFF) >> (32 - n))) - -#undef R -#define R(t) \ -( \ - temp = W[(t - 3) & 0x0F] ^ W[(t - 8) & 0x0F] ^ \ - W[(t - 14) & 0x0F] ^ W[ t & 0x0F], \ - ( W[t & 0x0F] = S(temp,1) ) \ -) - -#undef P -#define P(a,b,c,d,e,x) \ -{ \ - e += S(a,5) + F(b,c,d) + K + x; b = S(b,30); \ -} - - A = ctx->state[0]; - B = ctx->state[1]; - C = ctx->state[2]; - D = ctx->state[3]; - E = ctx->state[4]; - -#define F(x,y,z) (z ^ (x & (y ^ z))) -#define K 0x5A827999 - - P( A, B, C, D, E, W[0] ); - P( E, A, B, C, D, W[1] ); - P( D, E, A, B, C, W[2] ); - P( C, D, E, A, B, W[3] ); - P( B, C, D, E, A, W[4] ); - P( A, B, C, D, E, W[5] ); - P( E, A, B, C, D, W[6] ); - P( D, E, A, B, C, W[7] ); - P( C, D, E, A, B, W[8] ); - P( B, C, D, E, A, W[9] ); - P( A, B, C, D, E, W[10] ); - P( E, A, B, C, D, W[11] ); - P( D, E, A, B, C, W[12] ); - P( C, D, E, A, B, W[13] ); - P( B, C, D, E, A, W[14] ); - P( A, B, C, D, E, W[15] ); - P( E, A, B, C, D, R(16) ); - P( D, E, A, B, C, R(17) ); - P( C, D, E, A, B, R(18) ); - P( B, C, D, E, A, R(19) ); - -#undef K -#undef F - -#define F(x,y,z) (x ^ y ^ z) -#define K 0x6ED9EBA1 - - P( A, B, C, D, E, R(20) ); - P( E, A, B, C, D, R(21) ); - P( D, E, A, B, C, R(22) ); - P( C, D, E, A, B, R(23) ); - P( B, C, D, E, A, R(24) ); - P( A, B, C, D, E, R(25) ); - P( E, A, B, C, D, R(26) ); - P( D, E, A, B, C, R(27) ); - P( C, D, E, A, B, R(28) ); - P( B, C, D, E, A, R(29) ); - P( A, B, C, D, E, R(30) ); - P( E, A, B, C, D, R(31) ); - P( D, E, A, B, C, R(32) ); - P( C, D, E, A, B, R(33) ); - P( B, C, D, E, A, R(34) ); - P( A, B, C, D, E, R(35) ); - P( E, A, B, C, D, R(36) ); - P( D, E, A, B, C, R(37) ); - P( C, D, E, A, B, R(38) ); - P( B, C, D, E, A, R(39) ); - -#undef K -#undef F - -#define F(x,y,z) ((x & y) | (z & (x | y))) -#define K 0x8F1BBCDC - - P( A, B, C, D, E, R(40) ); - P( E, A, B, C, D, R(41) ); - P( D, E, A, B, C, R(42) ); - P( C, D, E, A, B, R(43) ); - P( B, C, D, E, A, R(44) ); - P( A, B, C, D, E, R(45) ); - P( E, A, B, C, D, R(46) ); - P( D, E, A, B, C, R(47) ); - P( C, D, E, A, B, R(48) ); - P( B, C, D, E, A, R(49) ); - P( A, B, C, D, E, R(50) ); - P( E, A, B, C, D, R(51) ); - P( D, E, A, B, C, R(52) ); - P( C, D, E, A, B, R(53) ); - P( B, C, D, E, A, R(54) ); - P( A, B, C, D, E, R(55) ); - P( E, A, B, C, D, R(56) ); - P( D, E, A, B, C, R(57) ); - P( C, D, E, A, B, R(58) ); - P( B, C, D, E, A, R(59) ); - -#undef K -#undef F - -#define F(x,y,z) (x ^ y ^ z) -#define K 0xCA62C1D6 - - P( A, B, C, D, E, R(60) ); - P( E, A, B, C, D, R(61) ); - P( D, E, A, B, C, R(62) ); - P( C, D, E, A, B, R(63) ); - P( B, C, D, E, A, R(64) ); - P( A, B, C, D, E, R(65) ); - P( E, A, B, C, D, R(66) ); - P( D, E, A, B, C, R(67) ); - P( C, D, E, A, B, R(68) ); - P( B, C, D, E, A, R(69) ); - P( A, B, C, D, E, R(70) ); - P( E, A, B, C, D, R(71) ); - P( D, E, A, B, C, R(72) ); - P( C, D, E, A, B, R(73) ); - P( B, C, D, E, A, R(74) ); - P( A, B, C, D, E, R(75) ); - P( E, A, B, C, D, R(76) ); - P( D, E, A, B, C, R(77) ); - P( C, D, E, A, B, R(78) ); - P( B, C, D, E, A, R(79) ); - -#undef K -#undef F - - ctx->state[0] += A; - ctx->state[1] += B; - ctx->state[2] += C; - ctx->state[3] += D; - ctx->state[4] += E; -} - -/*=========================================================================== - -FUNCTION SHA1_CPU_UPDATE - -DESCRIPTION - SHA1 update buffer - -DEPENDENCIES - None - -RETURN VALUE - None - -===========================================================================*/ -__device__ -void sha1_update( sha1_context *ctx, unsigned char *input, int ilen ) { - int fill; - unsigned long left; - - if( ilen <= 0 ) - return; - - left = ctx->total[0] & 0x3F; - fill = 64 - left; - - ctx->total[0] += ilen; - ctx->total[0] &= 0xFFFFFFFF; - - if ( ctx->total[0] < (unsigned long) ilen ) - ctx->total[1]++; - - if ( left && ilen >= fill ) { - /*memcpy( (void *) (ctx->buffer + left), - (void *) input, fill );*/ - for (int i = 0; i < fill; i++) { - ctx->buffer[i+left] = input[i]; - } - - - sha1_process( ctx, ctx->buffer ); - input += fill; - ilen -= fill; - left = 0; - } - - while ( ilen >= 64 ) { - sha1_process( ctx, input ); - input += 64; - ilen -= 64; - } - - if ( ilen > 0 ) { - /*memcpy( (void *) (ctx->buffer + left), - (void *) input, ilen );*/ - for (int i = 0; i < ilen; i++) { - ctx->buffer[i+left] = input[i]; - } - - } -} - - -/*=========================================================================== - -FUNCTION SHA1_CPU_FINISH - -DESCRIPTION - SHA1 final digest - -DEPENDENCIES - None - -RETURN VALUE - None - -===========================================================================*/ -__device__ -void sha1_finish( sha1_context *ctx, unsigned char *output ) { - unsigned long last, padn; - unsigned long high, low; - unsigned char msglen[8]; - - high = ( ctx->total[0] >> 29 ) - | ( ctx->total[1] << 3 ); - low = ( ctx->total[0] << 3 ); - - PUT_UINT32_BE( high, msglen, 0 ); - PUT_UINT32_BE( low, msglen, 4 ); - - last = ctx->total[0] & 0x3F; - padn = ( last < 56 ) ? ( 56 - last ) : ( 120 - last ); - - sha1_update( ctx, (unsigned char *) sha1_padding, padn ); - sha1_update( ctx, msglen, 8 ); - - PUT_UINT32_BE( ctx->state[0], output, 0 ); -#ifndef FEATURE_REDUCED_HASH_SIZE - PUT_UINT32_BE( ctx->state[1], output, 4 ); - PUT_UINT32_BE( ctx->state[2], output, 8 ); - PUT_UINT32_BE( ctx->state[3], output, 12 ); - PUT_UINT32_BE( ctx->state[4], output, 16 ); -#endif -} - -/*=========================================================================== - -FUNCTION SHA1_INTERNAL - -DESCRIPTION - Does the real sha1 algorithm - -DEPENDENCIES - None - -RETURN VALUE - output is the hash result - -===========================================================================*/ -__device__ -void sha1_internal( unsigned char *input, int ilen, - unsigned char *output ) { - sha1_context ctx; - - sha1_starts( &ctx ); - sha1_update( &ctx, input, ilen ); - sha1_finish( &ctx, output ); - - memset( &ctx, 0, sizeof( sha1_context ) ); -} - -#endif - -#ifdef FEATURE_SHARED_MEMORY -/*=========================================================================== - -FUNCTION SHA1_INTERNAL - -DESCRIPTION - Does the real sha1 algorithm. - -DEPENDENCIES - None - -RETURN VALUE - output is the hash result - -===========================================================================*/ - -__device__ -unsigned long macroRFunction(int t, unsigned int *sharedMemory) { - return sharedMemory[SHARED_MEMORY_INDEX((t - 3) & 0x0F)] ^ sharedMemory[SHARED_MEMORY_INDEX((t - 8) & 0x0F)] ^ - sharedMemory[SHARED_MEMORY_INDEX((t - 14) & 0x0F)] ^ sharedMemory[SHARED_MEMORY_INDEX( t & 0x0F)]; -} - - -__device__ -static void sha1_internal( unsigned int *input, unsigned int *sharedMemory, - unsigned int chunkSize, unsigned char *output ) { - - /* Number of passes (512 bit blocks) we have to do */ - int numberOfPasses = chunkSize / 64 + 1; - /* Used during the hashing process */ - unsigned long temp, A, B, C, D ,E; - //unsigned long shared14, shared15; - /* Needed to do the little endian stuff */ - unsigned char *data = (unsigned char *)sharedMemory; - - /* Will hold the hash value through the - intermediate stages of SHA1 algorithm */ - unsigned int state0 = 0x67452301; - unsigned int state1 = 0xEFCDAB89; - unsigned int state2 = 0x98BADCFE; - unsigned int state3 = 0x10325476; - unsigned int state4 = 0xC3D2E1F0; - - -/* int x0 = SHARED_MEMORY_INDEX(0); - int x1 = SHARED_MEMORY_INDEX(1); - int x2 = SHARED_MEMORY_INDEX(2); - int x3 = SHARED_MEMORY_INDEX(3); - int x4 = SHARED_MEMORY_INDEX(4); - int x5 = SHARED_MEMORY_INDEX(5); - int x6 = SHARED_MEMORY_INDEX(6); - int x7 = SHARED_MEMORY_INDEX(7); - int x8 = SHARED_MEMORY_INDEX(8); - int x9 = SHARED_MEMORY_INDEX(9); - int x10 = SHARED_MEMORY_INDEX(10); - int x11 = SHARED_MEMORY_INDEX(11); - int x12 = SHARED_MEMORY_INDEX(12); - int x13 = SHARED_MEMORY_INDEX(13); - int x14 = SHARED_MEMORY_INDEX(14); - int x15 = SHARED_MEMORY_INDEX(15); -*/ -#undef GET_CACHED_INDEX -#define GET_CACHED_INDEX(index) SHARED_MEMORY_INDEX(index)//(x##index) - - - for( int index = 0 ; index < (numberOfPasses) ; index++ ) { - - /* Move data to the thread's shared memory space */ - sharedMemory[GET_CACHED_INDEX(0)] = input[0 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(1)] = input[1 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(2)] = input[2 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(3)] = input[3 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(4)] = input[4 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(5)] = input[5 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(6)] = input[6 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(7)] = input[7 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(8)] = input[8 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(9)] = input[9 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(10)] = input[10 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(11)] = input[11 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(12)] = input[12 + 16 * index]; - - /* Testing the code with and without this if statement shows that - it has no effect on performance. */ - if(index == numberOfPasses -1 ) { - /* The last pass will contain the size of the chunk size (according to - official SHA1 algorithm). */ - sharedMemory[GET_CACHED_INDEX(13)] = 0x00000080; - - PUT_UINT32_BE( chunkSize >> 29, - data, GET_CACHED_INDEX(14) * 4 ); - PUT_UINT32_BE( chunkSize << 3, - data, GET_CACHED_INDEX(15) * 4 ); - - } - else { - sharedMemory[GET_CACHED_INDEX(13)] = input[13 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(14)] = input[14 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(15)] = input[15 + 16 * index]; - } - - /* Get the little endian stuff done. */ - GET_UINT32_BE( sharedMemory[ GET_CACHED_INDEX(0)], - data, GET_CACHED_INDEX(0) * 4 ); - GET_UINT32_BE( sharedMemory[ GET_CACHED_INDEX(1)], - data, GET_CACHED_INDEX(1) * 4 ); - GET_UINT32_BE( sharedMemory[ GET_CACHED_INDEX(2)], - data, GET_CACHED_INDEX(2) * 4 ); - GET_UINT32_BE( sharedMemory[ GET_CACHED_INDEX(3)], - data, GET_CACHED_INDEX(3) * 4 ); - GET_UINT32_BE( sharedMemory[ GET_CACHED_INDEX(4)], - data, GET_CACHED_INDEX(4) * 4 ); - GET_UINT32_BE( sharedMemory[ GET_CACHED_INDEX(5)], - data, GET_CACHED_INDEX(5) * 4 ); - GET_UINT32_BE( sharedMemory[ GET_CACHED_INDEX(6)], - data, GET_CACHED_INDEX(6) * 4 ); - GET_UINT32_BE( sharedMemory[ GET_CACHED_INDEX(7)], - data, GET_CACHED_INDEX(7) * 4 ); - GET_UINT32_BE( sharedMemory[ GET_CACHED_INDEX(8)], - data, GET_CACHED_INDEX(8) * 4 ); - GET_UINT32_BE( sharedMemory[ GET_CACHED_INDEX(9)], - data, GET_CACHED_INDEX(9) * 4 ); - GET_UINT32_BE( sharedMemory[GET_CACHED_INDEX(10)], - data, GET_CACHED_INDEX(10) * 4 ); - GET_UINT32_BE( sharedMemory[GET_CACHED_INDEX(11)], - data, GET_CACHED_INDEX(11) * 4 ); - GET_UINT32_BE( sharedMemory[GET_CACHED_INDEX(12)], - data, GET_CACHED_INDEX(12) * 4 ); - GET_UINT32_BE( sharedMemory[GET_CACHED_INDEX(13)], - data, GET_CACHED_INDEX(13) * 4 ); - GET_UINT32_BE( sharedMemory[GET_CACHED_INDEX(14)], - data, GET_CACHED_INDEX(14) * 4 ); - GET_UINT32_BE( sharedMemory[GET_CACHED_INDEX(15)], - data, GET_CACHED_INDEX(15) * 4 ); - - -#undef S -#define S(x,n) ((x << n) | ((x & 0xFFFFFFFF) >> (32 - n))) - - -#undef R -#define R(t) \ -( \ - temp = macroRFunction(t, sharedMemory) , \ - ( sharedMemory[SHARED_MEMORY_INDEX(t & 0x0F)] = S(temp,1) ) \ -) - -/* -#define R(t) \ -( \ - temp = sharedMemory[SHARED_MEMORY_INDEX((t - 3) & 0x0F)] ^ sharedMemory[SHARED_MEMORY_INDEX((t - 8) & 0x0F)] ^ \ - sharedMemory[SHARED_MEMORY_INDEX((t - 14) & 0x0F)] ^ sharedMemory[SHARED_MEMORY_INDEX( t & 0x0F)], \ - ( sharedMemory[SHARED_MEMORY_INDEX(t & 0x0F)] = S(temp,1) ) \ -) -*/ - -#undef P -#define P(a,b,c,d,e,x) \ -{ \ - e += S(a,5) + F(b,c,d) + K + x; b = S(b,30); \ -} - - A = state0; - B = state1; - C = state2; - D = state3; - E = state4; - - -#define F(x,y,z) (z ^ (x & (y ^ z))) -#define K 0x5A827999 - - P( A, B, C, D, E, sharedMemory[ GET_CACHED_INDEX(0)] ); - P( E, A, B, C, D, sharedMemory[ GET_CACHED_INDEX(1)] ); - P( D, E, A, B, C, sharedMemory[ GET_CACHED_INDEX(2)] ); - P( C, D, E, A, B, sharedMemory[ GET_CACHED_INDEX(3)] ); - P( B, C, D, E, A, sharedMemory[ GET_CACHED_INDEX(4)] ); - P( A, B, C, D, E, sharedMemory[ GET_CACHED_INDEX(5)] ); - P( E, A, B, C, D, sharedMemory[ GET_CACHED_INDEX(6)] ); - P( D, E, A, B, C, sharedMemory[ GET_CACHED_INDEX(7)] ); - P( C, D, E, A, B, sharedMemory[ GET_CACHED_INDEX(8)] ); - P( B, C, D, E, A, sharedMemory[ GET_CACHED_INDEX(9)] ); - P( A, B, C, D, E, sharedMemory[ GET_CACHED_INDEX(10)] ); - P( E, A, B, C, D, sharedMemory[ GET_CACHED_INDEX(11)] ); - P( D, E, A, B, C, sharedMemory[ GET_CACHED_INDEX(12)] ); - P( C, D, E, A, B, sharedMemory[ GET_CACHED_INDEX(13)] ); - P( B, C, D, E, A, sharedMemory[ GET_CACHED_INDEX(14)] ); - P( A, B, C, D, E, sharedMemory[ GET_CACHED_INDEX(15)] ); - P( E, A, B, C, D, R(16) ); - P( D, E, A, B, C, R(17) ); - P( C, D, E, A, B, R(18) ); - P( B, C, D, E, A, R(19) ); - - -#undef K -#undef F - -#define F(x,y,z) (x ^ y ^ z) -#define K 0x6ED9EBA1 - - P( A, B, C, D, E, R(20) ); - P( E, A, B, C, D, R(21) ); - P( D, E, A, B, C, R(22) ); - P( C, D, E, A, B, R(23) ); - P( B, C, D, E, A, R(24) ); - P( A, B, C, D, E, R(25) ); - P( E, A, B, C, D, R(26) ); - P( D, E, A, B, C, R(27) ); - P( C, D, E, A, B, R(28) ); - P( B, C, D, E, A, R(29) ); - P( A, B, C, D, E, R(30) ); - P( E, A, B, C, D, R(31) ); - P( D, E, A, B, C, R(32) ); - P( C, D, E, A, B, R(33) ); - P( B, C, D, E, A, R(34) ); - P( A, B, C, D, E, R(35) ); - P( E, A, B, C, D, R(36) ); - P( D, E, A, B, C, R(37) ); - P( C, D, E, A, B, R(38) ); - P( B, C, D, E, A, R(39) ); - -#undef K -#undef F - -#define F(x,y,z) ((x & y) | (z & (x | y))) -#define K 0x8F1BBCDC - - P( A, B, C, D, E, R(40) ); - P( E, A, B, C, D, R(41) ); - P( D, E, A, B, C, R(42) ); - P( C, D, E, A, B, R(43) ); - P( B, C, D, E, A, R(44) ); - P( A, B, C, D, E, R(45) ); - P( E, A, B, C, D, R(46) ); - P( D, E, A, B, C, R(47) ); - P( C, D, E, A, B, R(48) ); - P( B, C, D, E, A, R(49) ); - P( A, B, C, D, E, R(50) ); - P( E, A, B, C, D, R(51) ); - P( D, E, A, B, C, R(52) ); - P( C, D, E, A, B, R(53) ); - P( B, C, D, E, A, R(54) ); - P( A, B, C, D, E, R(55) ); - P( E, A, B, C, D, R(56) ); - P( D, E, A, B, C, R(57) ); - P( C, D, E, A, B, R(58) ); - P( B, C, D, E, A, R(59) ); - -#undef K -#undef F - -#define F(x,y,z) (x ^ y ^ z) -#define K 0xCA62C1D6 - - P( A, B, C, D, E, R(60) ); - P( E, A, B, C, D, R(61) ); - P( D, E, A, B, C, R(62) ); - P( C, D, E, A, B, R(63) ); - P( B, C, D, E, A, R(64) ); - P( A, B, C, D, E, R(65) ); - P( E, A, B, C, D, R(66) ); - P( D, E, A, B, C, R(67) ); - P( C, D, E, A, B, R(68) ); - P( B, C, D, E, A, R(69) ); - P( A, B, C, D, E, R(70) ); - P( E, A, B, C, D, R(71) ); - P( D, E, A, B, C, R(72) ); - P( C, D, E, A, B, R(73) ); - P( B, C, D, E, A, R(74) ); - P( A, B, C, D, E, R(75) ); - P( E, A, B, C, D, R(76) ); - P( D, E, A, B, C, R(77) ); - P( C, D, E, A, B, R(78) ); - P( B, C, D, E, A, R(79) ); - -#undef K -#undef F - - state0 += A; - state1 += B; - state2 += C; - state3 += D; - state4 += E; - } - - /* Got the hash, store it in the output buffer. */ - PUT_UINT32_BE( state0, output, 0 ); -#ifndef FEATURE_REDUCED_HASH_SIZE - PUT_UINT32_BE( state1, output, 4 ); - PUT_UINT32_BE( state2, output, 8 ); - PUT_UINT32_BE( state3, output, 12 ); - PUT_UINT32_BE( state4, output, 16 ); -#endif - -} - -__device__ -static void sha1_internal_overlap( unsigned int *input, unsigned int *sharedMemory, - unsigned int chunkSize, unsigned char *output ) { - - /* Number of passes (512 bit blocks) we have to do */ - int numberOfPasses = chunkSize / 64 + 1; - /* Used during the hashing process */ - unsigned long temp, A, B, C, D ,E; - //unsigned long shared14, shared15; - /* Needed to do the big endian stuff */ - unsigned char *data = (unsigned char *)sharedMemory; - // number of padding bytes. - int numPadBytes = 0; - int numPadInt = 0; - //int numPadRemain = 0; - - /* Will hold the hash value through the - intermediate stages of SHA1 algorithm */ - unsigned int state0 = 0x67452301; - unsigned int state1 = 0xEFCDAB89; - unsigned int state2 = 0x98BADCFE; - unsigned int state3 = 0x10325476; - unsigned int state4 = 0xC3D2E1F0; - - - int x0 = SHARED_MEMORY_INDEX(0); - int x1 = SHARED_MEMORY_INDEX(1); - int x2 = SHARED_MEMORY_INDEX(2); - int x3 = SHARED_MEMORY_INDEX(3); - int x4 = SHARED_MEMORY_INDEX(4); - int x5 = SHARED_MEMORY_INDEX(5); - int x6 = SHARED_MEMORY_INDEX(6); - int x7 = SHARED_MEMORY_INDEX(7); - int x8 = SHARED_MEMORY_INDEX(8); - int x9 = SHARED_MEMORY_INDEX(9); - int x10 = SHARED_MEMORY_INDEX(10); - int x11 = SHARED_MEMORY_INDEX(11); - int x12 = SHARED_MEMORY_INDEX(12); - int x13 = SHARED_MEMORY_INDEX(13); - int x14 = SHARED_MEMORY_INDEX(14); - int x15 = SHARED_MEMORY_INDEX(15); - -#undef GET_CACHED_INDEX -#define GET_CACHED_INDEX(index) (x##index) - - - for( int index = 0 ; index < (numberOfPasses) ; index++ ) { - - if(index == numberOfPasses -1 ){ - - numPadBytes = (64-12) - (chunkSize - (numberOfPasses-1)*64); - numPadInt = numPadBytes/sizeof(int); - /*numPadRemain = numPadBytes-numPadInt*sizeof(int); - printf("\nLast loop chunkSize = %d, numberOfPasses= %d and \nnumPadBytes = %d, numPadInt =%d, numPadRemain = %d\n", - chunkSize,numberOfPasses,numPadBytes,numPadInt,numPadRemain);*/ - - int i=0; - for(i=0;i> 29, - data, GET_CACHED_INDEX(14) * 4 ); - PUT_UINT32_BE( chunkSize << 3, - data, GET_CACHED_INDEX(15) * 4 ); - } - else{ - /* Move data to the thread's shared memory space */ - //printf("Not last loop\n"); - sharedMemory[GET_CACHED_INDEX(0)] = input[0 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(1)] = input[1 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(2)] = input[2 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(3)] = input[3 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(4)] = input[4 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(5)] = input[5 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(6)] = input[6 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(7)] = input[7 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(8)] = input[8 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(9)] = input[9 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(10)] = input[10 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(11)] = input[11 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(12)] = input[12 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(13)] = input[13 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(14)] = input[14 + 16 * index]; - sharedMemory[GET_CACHED_INDEX(15)] = input[15 + 16 * index]; - } - - /* int k=0; - printf("\nGPU DATA\n"); - for(k=0;k<16;k++){ - printf("%d\t",sharedMemory[SHARED_MEMORY_INDEX(k)]); - } - printf("\n\n");*/ - - /* Get the little endian stuff done. */ - GET_UINT32_BE( sharedMemory[ GET_CACHED_INDEX(0)], - data, GET_CACHED_INDEX(0) * 4 ); - GET_UINT32_BE( sharedMemory[ GET_CACHED_INDEX(1)], - data, GET_CACHED_INDEX(1) * 4 ); - GET_UINT32_BE( sharedMemory[ GET_CACHED_INDEX(2)], - data, GET_CACHED_INDEX(2) * 4 ); - GET_UINT32_BE( sharedMemory[ GET_CACHED_INDEX(3)], - data, GET_CACHED_INDEX(3) * 4 ); - GET_UINT32_BE( sharedMemory[ GET_CACHED_INDEX(4)], - data, GET_CACHED_INDEX(4) * 4 ); - GET_UINT32_BE( sharedMemory[ GET_CACHED_INDEX(5)], - data, GET_CACHED_INDEX(5) * 4 ); - GET_UINT32_BE( sharedMemory[ GET_CACHED_INDEX(6)], - data, GET_CACHED_INDEX(6) * 4 ); - GET_UINT32_BE( sharedMemory[ GET_CACHED_INDEX(7)], - data, GET_CACHED_INDEX(7) * 4 ); - GET_UINT32_BE( sharedMemory[ GET_CACHED_INDEX(8)], - data, GET_CACHED_INDEX(8) * 4 ); - GET_UINT32_BE( sharedMemory[ GET_CACHED_INDEX(9)], - data, GET_CACHED_INDEX(9) * 4 ); - GET_UINT32_BE( sharedMemory[GET_CACHED_INDEX(10)], - data, GET_CACHED_INDEX(10) * 4 ); - GET_UINT32_BE( sharedMemory[GET_CACHED_INDEX(11)], - data, GET_CACHED_INDEX(11) * 4 ); - GET_UINT32_BE( sharedMemory[GET_CACHED_INDEX(12)], - data, GET_CACHED_INDEX(12) * 4 ); - GET_UINT32_BE( sharedMemory[GET_CACHED_INDEX(13)], - data, GET_CACHED_INDEX(13) * 4 ); - GET_UINT32_BE( sharedMemory[GET_CACHED_INDEX(14)], - data, GET_CACHED_INDEX(14) * 4 ); - GET_UINT32_BE( sharedMemory[GET_CACHED_INDEX(15)], - data, GET_CACHED_INDEX(15) * 4 ); - -#undef S -#define S(x,n) ((x << n) | ((x & 0xFFFFFFFF) >> (32 - n))) - - -#undef R -#define R(t) \ -( \ - temp = macroRFunction(t, sharedMemory) , \ - ( sharedMemory[SHARED_MEMORY_INDEX(t & 0x0F)] = S(temp,1) ) \ -) - -/* -#define R(t) \ -( \ - temp = sharedMemory[SHARED_MEMORY_INDEX((t - 3) & 0x0F)] ^ sharedMemory[SHARED_MEMORY_INDEX((t - 8) & 0x0F)] ^ \ - sharedMemory[SHARED_MEMORY_INDEX((t - 14) & 0x0F)] ^ sharedMemory[SHARED_MEMORY_INDEX( t & 0x0F)], \ - ( sharedMemory[SHARED_MEMORY_INDEX(t & 0x0F)] = S(temp,1) ) \ -) -*/ - -#undef P -#define P(a,b,c,d,e,x) \ -{ \ - e += S(a,5) + F(b,c,d) + K + x; b = S(b,30); \ -} - - A = state0; - B = state1; - C = state2; - D = state3; - E = state4; - - -#define F(x,y,z) (z ^ (x & (y ^ z))) -#define K 0x5A827999 - - P( A, B, C, D, E, sharedMemory[ GET_CACHED_INDEX(0)] ); - P( E, A, B, C, D, sharedMemory[ GET_CACHED_INDEX(1)] ); - P( D, E, A, B, C, sharedMemory[ GET_CACHED_INDEX(2)] ); - P( C, D, E, A, B, sharedMemory[ GET_CACHED_INDEX(3)] ); - P( B, C, D, E, A, sharedMemory[ GET_CACHED_INDEX(4)] ); - P( A, B, C, D, E, sharedMemory[ GET_CACHED_INDEX(5)] ); - P( E, A, B, C, D, sharedMemory[ GET_CACHED_INDEX(6)] ); - P( D, E, A, B, C, sharedMemory[ GET_CACHED_INDEX(7)] ); - P( C, D, E, A, B, sharedMemory[ GET_CACHED_INDEX(8)] ); - P( B, C, D, E, A, sharedMemory[ GET_CACHED_INDEX(9)] ); - P( A, B, C, D, E, sharedMemory[ GET_CACHED_INDEX(10)] ); - P( E, A, B, C, D, sharedMemory[ GET_CACHED_INDEX(11)] ); - P( D, E, A, B, C, sharedMemory[ GET_CACHED_INDEX(12)] ); - P( C, D, E, A, B, sharedMemory[ GET_CACHED_INDEX(13)] ); - P( B, C, D, E, A, sharedMemory[ GET_CACHED_INDEX(14)] ); - P( A, B, C, D, E, sharedMemory[ GET_CACHED_INDEX(15)] ); - P( E, A, B, C, D, R(16) ); - P( D, E, A, B, C, R(17) ); - P( C, D, E, A, B, R(18) ); - P( B, C, D, E, A, R(19) ); - - -#undef K -#undef F - -#define F(x,y,z) (x ^ y ^ z) -#define K 0x6ED9EBA1 - - P( A, B, C, D, E, R(20) ); - P( E, A, B, C, D, R(21) ); - P( D, E, A, B, C, R(22) ); - P( C, D, E, A, B, R(23) ); - P( B, C, D, E, A, R(24) ); - P( A, B, C, D, E, R(25) ); - P( E, A, B, C, D, R(26) ); - P( D, E, A, B, C, R(27) ); - P( C, D, E, A, B, R(28) ); - P( B, C, D, E, A, R(29) ); - P( A, B, C, D, E, R(30) ); - P( E, A, B, C, D, R(31) ); - P( D, E, A, B, C, R(32) ); - P( C, D, E, A, B, R(33) ); - P( B, C, D, E, A, R(34) ); - P( A, B, C, D, E, R(35) ); - P( E, A, B, C, D, R(36) ); - P( D, E, A, B, C, R(37) ); - P( C, D, E, A, B, R(38) ); - P( B, C, D, E, A, R(39) ); - -#undef K -#undef F - -#define F(x,y,z) ((x & y) | (z & (x | y))) -#define K 0x8F1BBCDC - - P( A, B, C, D, E, R(40) ); - P( E, A, B, C, D, R(41) ); - P( D, E, A, B, C, R(42) ); - P( C, D, E, A, B, R(43) ); - P( B, C, D, E, A, R(44) ); - P( A, B, C, D, E, R(45) ); - P( E, A, B, C, D, R(46) ); - P( D, E, A, B, C, R(47) ); - P( C, D, E, A, B, R(48) ); - P( B, C, D, E, A, R(49) ); - P( A, B, C, D, E, R(50) ); - P( E, A, B, C, D, R(51) ); - P( D, E, A, B, C, R(52) ); - P( C, D, E, A, B, R(53) ); - P( B, C, D, E, A, R(54) ); - P( A, B, C, D, E, R(55) ); - P( E, A, B, C, D, R(56) ); - P( D, E, A, B, C, R(57) ); - P( C, D, E, A, B, R(58) ); - P( B, C, D, E, A, R(59) ); - -#undef K -#undef F - -#define F(x,y,z) (x ^ y ^ z) -#define K 0xCA62C1D6 - - P( A, B, C, D, E, R(60) ); - P( E, A, B, C, D, R(61) ); - P( D, E, A, B, C, R(62) ); - P( C, D, E, A, B, R(63) ); - P( B, C, D, E, A, R(64) ); - P( A, B, C, D, E, R(65) ); - P( E, A, B, C, D, R(66) ); - P( D, E, A, B, C, R(67) ); - P( C, D, E, A, B, R(68) ); - P( B, C, D, E, A, R(69) ); - P( A, B, C, D, E, R(70) ); - P( E, A, B, C, D, R(71) ); - P( D, E, A, B, C, R(72) ); - P( C, D, E, A, B, R(73) ); - P( B, C, D, E, A, R(74) ); - P( A, B, C, D, E, R(75) ); - P( E, A, B, C, D, R(76) ); - P( D, E, A, B, C, R(77) ); - P( C, D, E, A, B, R(78) ); - P( B, C, D, E, A, R(79) ); - -#undef K -#undef F - - state0 += A; - state1 += B; - state2 += C; - state3 += D; - state4 += E; - } - - /* Got the hash, store it in the output buffer. */ - PUT_UINT32_BE( state0, output, 0 ); -#ifndef FEATURE_REDUCED_HASH_SIZE - PUT_UINT32_BE( state1, output, 4 ); - PUT_UINT32_BE( state2, output, 8 ); - PUT_UINT32_BE( state3, output, 12 ); - PUT_UINT32_BE( state4, output, 16 ); -#endif - -} -#endif - -/*-------------------------------------------------------------------------- - - GLOBAL FUNCTIONS ---------------------------------------------------------------------------*/ -/*=========================================================================== - -FUNCTION SHA1 - -DESCRIPTION - Main sha1 hash function - -DEPENDENCIES - GPU must be initialized - -RETURN VALUE - output: the hash result - -===========================================================================*/ -__global__ -void sha1( unsigned char *input, int chunkSize, int totalThreads, - int padSize, unsigned char *scratch ) { - - // get the current thread index - int threadIndex = threadIdx.x + blockDim.x * blockIdx.x; - int chunkIndex = threadIndex * chunkSize; - int hashIndex = threadIndex * SHA1_HASH_SIZE; - - if(threadIndex >= totalThreads) - return; - - if ((threadIndex == (totalThreads - 1)) && (padSize > 0)) { - for(int i = 0 ; i < padSize ; i++) - input[chunkIndex + chunkSize - padSize + i] = 0; - } - -#ifdef FEATURE_SHARED_MEMORY - - __shared__ unsigned int sharedMemory[4 * 1024 - 32]; - - unsigned int *sharedMemoryIndex = sharedMemory + ((threadIdx.x >> 5) * 512); - unsigned char *tempInput = input + chunkIndex; - unsigned int *inputIndex = (unsigned int *)(tempInput); - - sha1_internal(inputIndex, sharedMemoryIndex, chunkSize, - scratch + hashIndex ); - -#else - sha1_internal(input + chunkIndex, chunkSize, scratch + hashIndex ); -#endif /* FEATURE_SHARED_MEMORY */ - -} - -__global__ -void sha1_overlap( unsigned char *input, int chunkSize, int offset, - int totalThreads, int padSize, unsigned char *output ) { - - int threadIndex = threadIdx.x + blockDim.x * blockIdx.x; - int chunkIndex = threadIndex * offset; - int hashIndex = threadIndex * SHA1_HASH_SIZE; - - if(threadIndex >= totalThreads) - return; - - if ((threadIndex == (totalThreads - 1))) { - chunkSize-= padSize; - } - -#ifdef FEATURE_SHARED_MEMORY - - __shared__ unsigned int sharedMemory[4 * 1024 - 32]; - - //NOTE : SAMER : this can exceed the size of the shared memory - unsigned int *sharedMemoryIndex = sharedMemory + ((threadIdx.x >> 5) * 512); - unsigned int *inputIndex = (unsigned int *)(input + chunkIndex); - - sha1_internal_overlap(inputIndex, sharedMemoryIndex, chunkSize, - output + hashIndex ); - -#else - sha1_internal(input + chunkIndex, chunkSize, output + hashIndex ); -#endif /* FEATURE_SHARED_MEMORY */ - - -} diff --git a/benchmarks/CUDA/STO/storeCPU.c b/benchmarks/CUDA/STO/storeCPU.c deleted file mode 100644 index 09d1997..0000000 --- a/benchmarks/CUDA/STO/storeCPU.c +++ /dev/null @@ -1,1114 +0,0 @@ -/*========================================================================== - S T O R E C P U - -* Copyright (c) 2008, NetSysLab at the University of British Columbia -* All rights reserved. -* -* Redistribution and use in source and binary forms, with or without -* modification, are permitted provided that the following conditions are met: -* * Redistributions of source code must retain the above copyright -* notice, this list of conditions and the following disclaimer. -* * Redistributions in binary form must reproduce the above copyright -* notice, this list of conditions and the following disclaimer in the -* documentation and/or other materials provided with the distribution. -* * Neither the name of the University nor the -* names of its contributors may be used to endorse or promote products -* derived from this software without specific prior written permission. -* -* THIS SOFTWARE IS PROVIDED BY NetSysLab ``AS IS'' AND ANY -* EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED -* WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE -* DISCLAIMED. IN NO EVENT SHALL NetSysLab BE LIABLE FOR ANY -* DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES -* (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; -* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND -* ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT -* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS -* SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. - - -DESCRIPTION - CPU version of the storeGPU library. - - -==========================================================================*/ - -/*========================================================================== - - INCLUDES - -==========================================================================*/ -#include -#include -#include - - -#include "cust.h" -#include "md5_cpu.h" -#include "sha1_cpu.h" - - -#ifdef FEATURE_WIN32_THREADS -#include -#endif - -/*========================================================================== - - DATA DECLARATIONS - -==========================================================================*/ - -/*-------------------------------------------------------------------------- - TYPE DEFINITIONS ---------------------------------------------------------------------------*/ -// defines an execution context -typedef struct sc_exec_context { - int threads_per_block; - int blocks_per_grid; - int total_threads; - int total_size; - int chunk_size; - int pad_size; -} sc_exec_context_type; - -#ifdef FEATURE_WIN32_THREADS -typedef struct thread_data_struct { - - unsigned char *input; - unsigned char *output; - int ilen; - -} thread_data_type, *pt_thread_data_type; -#endif /* FEATURE_WIN32_THREADS */ - - -/*-------------------------------------------------------------------------- - FUNCTION PROTOTYPES ---------------------------------------------------------------------------*/ - -/*-------------------------------------------------------------------------- - CONSTANTS ---------------------------------------------------------------------------*/ - -/*-------------------------------------------------------------------------- - GLOBAL VARIABLES ---------------------------------------------------------------------------*/ - - -/*-------------------------------------------------------------------------- - MACROS ---------------------------------------------------------------------------*/ - -#define GET_REAL_CHUNK_SIZE(chunk_size) ((chunk_size) - 12) - -/*========================================================================== - - FUNCTIONS - -==========================================================================*/ - -/*-------------------------------------------------------------------------- - LOCAL FUNCTIONS ---------------------------------------------------------------------------*/ - -#ifdef FEATURE_DYNAMIC_EXEC_CONTEXT -#ifdef FEATURE_MAXIMIZE_NUM_OF_THREADS -/*=========================================================================== - -FUNCTION SC_GET_EXEC_CONTEXT - -DESCRIPTION - sets the execution context the algorithm will run within: chunk size, - thread per block, blocks, padding and total number of threads according - to client buffer size. - -DEPENDENCIES - None - -RETURN VALUE - execution context - -===========================================================================*/ -static void sc_get_exec_context(int size, sc_exec_context_type* ctx){ - - int threads_per_block; - int blocks_per_grid; - int total_threads; - int chunk_size; - int pad_size; - - int total_chunks = 0; - int found = 0; - int index = 1; - - - //**** Determine the execution context ****// - /* The algorithm will try to determine the context by minimizing chunk - * size and maximizing total number of threads - * TODO: May be we can do better here - */ - while ( !found ) { - // Set chunk size - chunk_size = GET_REAL_CHUNK_SIZE(BASIC_CHUNK_SIZE * index); - - if ( chunk_size > MAX_CHUNK_SIZE ) - break; - - // Calculate the required padding for this chunk size - pad_size = ((size % chunk_size) == 0) ? 0 : - chunk_size - (size % chunk_size); - - // total number of chunks required if we are going to use this chunk size - total_chunks = (pad_size == 0) ? size / chunk_size : - (size / chunk_size) + 1; - - if ( total_chunks <= MAX_NUM_OF_THREADS ) { - // Got it, this is the minimum chunk size we can use. Now determine the - // threads and blocks numbers. - total_threads = total_chunks; - - // Get block and grid sizes - if (total_chunks <= MAX_THREADS_PER_BLOCK ) { - threads_per_block = total_chunks; - blocks_per_grid = 1; - - } else { - threads_per_block = MAX_THREADS_PER_BLOCK; - blocks_per_grid = ((total_threads % threads_per_block) == 0) ? - (total_threads/threads_per_block) : - (total_threads/threads_per_block) + 1; - - } - - found = 1; - - } - index++; - - } - - - //**** Fill the struct with the solution ****// - ctx->threads_per_block = threads_per_block; - ctx->blocks_per_grid = blocks_per_grid; - ctx->total_threads = total_threads; - ctx->total_size = size + pad_size; - ctx->chunk_size = chunk_size; - ctx->pad_size = pad_size; - -} - -#else /* FEATURE_MAXIMIZE_NUM_OF_THREADS */ -/*=========================================================================== - -FUNCTION SC_GET_EXEC_CONTEXT - -DESCRIPTION - sets the required chunk size, thread per block and number of blocks - needed for kernel execution according to client buffer size. - -DEPENDENCIES - None - -RETURN VALUE - execution context - -===========================================================================*/ -static void sc_get_exec_context(int size, sc_exec_context_type* ctx){ - - int threads_per_block; - int blocks_per_grid; - int total_threads; - int chunk_size; - int pad_size; - - int total_chunks = 0; - int found = 0; - - int index = MAX_CHUNK_SIZE / BASIC_CHUNK_SIZE; - - - - //**** Determine the execution context ****// - /* The algorithm will try to determine the context by minimizing chunk - * size and maximizing total number of threads - * TODO: May be we can do better here - */ - while ( 1 ) { - // Set chunk size - chunk_size = GET_REAL_CHUNK_SIZE(BASIC_CHUNK_SIZE * index); - - // don't go less than minimum chunk size - if ( chunk_size < BASIC_CHUNK_SIZE ) - break; - - // Calculate the required padding for this chunk size - pad_size = ((size % chunk_size) == 0) ? 0 : - chunk_size - (size % chunk_size); - - // total number of chunks required if we are going to use this chunk size - total_chunks = (pad_size == 0) ? size / chunk_size : - (size / chunk_size) + 1; - - - // don't go beyond the maximum number of threads or maximum global memory - // TODO: it seems that the kernel breaks way before reaching the maximum - // global memory size (around 94MByte input plus the required - // scratch space) - if (total_chunks > MAX_NUM_OF_THREADS) - break; - - // each thread will take care of one chunk - total_threads = total_chunks; - - - // Get block and grid sizes - if (total_chunks <= MAX_THREADS_PER_BLOCK ) { - threads_per_block = total_chunks; - blocks_per_grid = 1; - - } else { - threads_per_block = MAX_THREADS_PER_BLOCK; - blocks_per_grid = ((total_threads % threads_per_block) == 0) ? - (total_threads/threads_per_block) : - (total_threads/threads_per_block) + 1; - - } - found = 1; - - if( total_threads > NUM_OF_MULTIPROCESSORS * 32) - break; - - index--; - - } - - - //**** Fill the struct with the solution ****// - ctx->threads_per_block = threads_per_block; - ctx->blocks_per_grid = blocks_per_grid; - ctx->total_threads = total_threads; - ctx->total_size = size + pad_size; - ctx->chunk_size = chunk_size; - ctx->pad_size = pad_size; - -} -#endif /* FEATURE_MAXIMIZE_NUM_OF_THREADS */ - - -/*=========================================================================== - -FUNCTION SC_GET_OVERLAP_EXEC_CONTEXT - -DESCRIPTION - sets the required chunk size, thread per block and number of blocks - needed for kernel execution according to client buffer size, offset - and block size. - -DEPENDENCIES - None - -RETURN VALUE - execution context - -===========================================================================*/ -static void sc_get_overlap_exec_context( int size, int offset, - int block_size, - sc_exec_context_type* ctx ) { - - int threads_per_block; - int blocks_per_grid; - int total_threads; - int total_size; - int pad_size; - - - //**** Get the total number of threads required ****// - total_threads = (size + offset - block_size) / offset; - total_threads = ((size + offset - block_size) % offset) != 0 ? - total_threads + 1 : total_threads; - - - //**** Get the required padding for the last block ****// - pad_size = ((total_threads - 1) * offset + block_size) - size; - - - //**** threads and blocks ****// - if( total_threads > MAX_THREADS_PER_BLOCK ) { - - threads_per_block = MAX_THREADS_PER_BLOCK; - blocks_per_grid = (total_threads % MAX_THREADS_PER_BLOCK) == 0 ? - (total_threads / MAX_THREADS_PER_BLOCK) : - (total_threads / MAX_THREADS_PER_BLOCK) + 1; - } else { - - threads_per_block = total_threads; - blocks_per_grid = 1; - } - - total_size = size + pad_size; - - - //**** Fill the struct with the solution ****// - ctx->threads_per_block = threads_per_block; - ctx->blocks_per_grid = blocks_per_grid; - ctx->total_threads = total_threads; - ctx->total_size = total_size; - ctx->chunk_size = block_size; - ctx->pad_size = pad_size; - -} -#endif /* FEATURE_DYNAMIC_EXEC_CONTEXT */ - -/*=========================================================================== - -FUNCTION SC_PRINT_EXEC_CONTEXT - -DESCRIPTION - Prints out the passed execution context structure - -DEPENDENCIES - None - -RETURN VALUE - None - -===========================================================================*/ -static void sc_print_exec_context( sc_exec_context_type* ctx ) { - printf("\n== CPU Execution Context ==\n"); - printf("Threads : %d\n", ctx->threads_per_block); - printf("Blocks : %d\n", ctx->blocks_per_grid); - printf("Total Threads : %d\n", ctx->total_threads); - printf("Total Size : %d\n", ctx->total_size); - printf("Chunk Size : %d\n", ctx->chunk_size); - printf("Padding : %d\n\n", ctx->pad_size); -} - -#ifdef FEATURE_WIN32_THREADS -/*=========================================================================== - -FUNCTION MD5_CPU_MT - -DESCRIPTION - The multithread CPU implementation of the MD5 algorithm - -DEPENDENCIES - None - -RETURN VALUE - Hash - -===========================================================================*/ -static DWORD WINAPI md5_cpu_mt( LPVOID data ){ - - pt_thread_data_type thread_data; - - //cast to the correct data type - thread_data = (pt_thread_data_type)data; - - md5_cpu_internal(thread_data->input, thread_data->ilen, thread_data->output); - - return 0; -} - - -/*=========================================================================== - -FUNCTION MD5_CPU_MT - -DESCRIPTION - The multithread CPU implementation of the MD5 algorithm - -DEPENDENCIES - None - -RETURN VALUE - Hash - -===========================================================================*/ -static DWORD WINAPI sha1_cpu_mt( LPVOID data ){ - - pt_thread_data_type thread_data; - - //cast to the correct data type - thread_data = (pt_thread_data_type)data; - - sha1_cpu_internal(thread_data->input, thread_data->ilen, thread_data->output); - - return 0; -} -#endif /* FEATURE_WIN32_THREADS */ - - -/*-------------------------------------------------------------------------- - GLOBAL FUNCTIONS ---------------------------------------------------------------------------*/ -/*=========================================================================== - -FUNCTION SC_MD5_STANDARD - -DESCRIPTION - The standard MD5 algorithm - -DEPENDENCIES - None - -RETURN VALUE - Hash - -===========================================================================*/ -void sc_md5_standard( unsigned char* buffer, int size, unsigned char** output) { - - unsigned char * result; - - result = (unsigned char*)malloc( MD5_HASH_SIZE ); - - md5_cpu_internal( buffer, size, result ); - - *output = result; -} - -/*=========================================================================== - -FUNCTION SC_SHA1_STANDARD - -DESCRIPTION - The standard SHA1 algorithm - -DEPENDENCIES - None - -RETURN VALUE - Hash - -===========================================================================*/ -void sc_sha1_standard(unsigned char* buffer, int size, unsigned char** output) { - - unsigned char * result; - - result = (unsigned char*)malloc( SHA1_HASH_SIZE ); - - sha1_cpu_internal( buffer, size, result ); - - *output = result; - -} - -/*=========================================================================== - -FUNCTION SC_MD5 - -DESCRIPTION - CPU version of the MD5 algorithm - -DEPENDENCIES - None - -RETURN VALUE - Hash - -===========================================================================*/ -void sc_md5( unsigned char* buffer, int size, - unsigned char** output, int* output_size) { - - - //**** Variable Declarations ****// - sc_exec_context_type exec_context; - unsigned char* scratch_data; - int chunk_index; - int hash_index; - int k; - -#ifdef FEATURE_WIN32_THREADS - - /* This structure contains the input for a particular thread */ - pt_thread_data_type thread_data; - - /* Thread identifiers */ - DWORD *thread_id; - - /* Thread handlers */ - HANDLE *thread_handle; - -#endif /* FEATURE_WIN32_THREADS */ - -#ifdef FEATURE_DYNAMIC_EXEC_CONTEXT - //**** Calculate pad size and needed block and grid sizes ****// - sc_get_exec_context(size, &exec_context); -#else - //**** Fill the execution context structure ****// - exec_context.threads_per_block = THREADS_PER_BLOCK; - exec_context.blocks_per_grid = BLOCKS_PER_GRID; - exec_context.chunk_size = CHUNK_SIZE; - exec_context.total_size = size; - exec_context.total_threads = TOTAL_NUM_OF_THREADS; - exec_context.pad_size = 0; - -#endif /* FEATURE_DYNAMIC_EXEC_CONTEXT */ - - sc_print_exec_context( &exec_context ); - - scratch_data = (unsigned char *)malloc(MD5_HASH_SIZE * exec_context.total_threads); - -#ifdef FEATURE_WIN32_THREADS - - //allocate memory for the thread ids - thread_id = (DWORD *)malloc(sizeof(DWORD)*exec_context.total_threads - 1); - - //allocate memory for the thread handle - thread_handle = (HANDLE *)malloc(sizeof(HANDLE)*exec_context.total_threads-1); - - //create structures for thread ids - for( k = 0; k < exec_context.total_threads-1; k++ ) { - - // set indices - chunk_index = k * exec_context.chunk_size; - hash_index = k * MD5_HASH_SIZE; - - // Allocate memory for thread data. - // thread_data_type is a type that contains the input and output buffers - // wrapped up into a structure - // this is used by the thread to compute and store the hashed values. - thread_data = (pt_thread_data_type) HeapAlloc(GetProcessHeap(), - HEAP_ZERO_MEMORY, - sizeof(thread_data_type)); - - // In case something wrong happen. That is: if memory cannot be - // allocated in the Heap of the current process. - if( thread_data == NULL ) { - fprintf(stderr,"\n[FATAL ERROR] Unable to allocate memory in the heap for Thread[%d]",k); - ExitProcess(2); - } - - // Generate unique data for each thread. - thread_data->input = buffer + chunk_index; - thread_data->ilen = exec_context.chunk_size; - thread_data->output = scratch_data + hash_index; - - // Create a thread - thread_handle[k] = CreateThread(NULL, 0, md5_cpu_mt, thread_data, - 0, &thread_id[k]); - - int i; - - // Check whether the thread was created correctly. If it was not, close the - // handlers and release memory - if (thread_handle[k] == NULL) { - - fprintf(stderr,"\n[FATAL ERROR] Unable to spawn thread[%d].\n\t Releasing resources and saying goodbye!\n",k); - - for( i=0; i < exec_context.total_threads-1; i++) { - - if ( thread_handle[i] != NULL ) { - CloseHandle(thread_handle[i]); - } - } - - HeapFree(GetProcessHeap(), 0, thread_data); - - ExitProcess(k); - } - } - - // wait for each thread to finish - WaitForMultipleObjects(exec_context.total_threads-1, thread_handle, - TRUE, INFINITE); - - // Close all thread handles and free memory allocation. - for(k=0; k < exec_context.total_threads-1; k++) { - CloseHandle(thread_handle[k]); - } - - HeapFree(GetProcessHeap(), 0, thread_data); - -#else - - for( k = 0; k < exec_context.total_threads - 1; k++) { - chunk_index = k * exec_context.chunk_size; - hash_index = k * MD5_HASH_SIZE; - md5_cpu_internal(buffer + chunk_index, exec_context.chunk_size, - scratch_data + hash_index ); - } -#endif /* FEATURE_WIN32_THREADS */ - - chunk_index = k * exec_context.chunk_size; - hash_index = k * MD5_HASH_SIZE; - - if(exec_context.pad_size != 0) { - - unsigned char *last_chunk = (unsigned char*)malloc(exec_context.chunk_size); - - memset(last_chunk, 0, exec_context.chunk_size); - memcpy(last_chunk, buffer + chunk_index, - exec_context.chunk_size - exec_context.pad_size); - md5_cpu_internal(last_chunk, exec_context.chunk_size, - scratch_data + hash_index ); - } else { - - md5_cpu_internal(buffer + chunk_index, exec_context.chunk_size, - scratch_data + hash_index ); - } - - //**** will do the last hshing stage ****// - sc_md5_standard( scratch_data, MD5_HASH_SIZE * exec_context.total_threads, - output ); - - *output_size = MD5_HASH_SIZE; - -} - -/*=========================================================================== - -FUNCTION SC_MD5_OVERLAP - -DESCRIPTION - Returns the MD5 hash of each block for the provided buffer - -DEPENDENCIES - None - -RETURN VALUE - Hash value - -===========================================================================*/ -void sc_md5_overlap(unsigned char* buffer, int size, int block_size, - int offset, unsigned char** output, int* output_size) { - - //**** Variable Declarations ****// - sc_exec_context_type exec_context; - unsigned char* result; - int chunk_index; - int hash_index; - int k; - -#ifdef FEATURE_WIN32_THREADS - - /* This structure contains the input for a particular thread */ - pt_thread_data_type thread_data; - - /* Thread identifiers */ - DWORD *thread_id; - - /* Thread handlers */ - HANDLE *thread_handle; - -#endif /* FEATURE_WIN32_THREADS */ - -#ifdef FEATURE_DYNAMIC_EXEC_CONTEXT - //**** Calculate pad size and needed block and grid sizes ****// - sc_get_overlap_exec_context(size, offset, block_size, &exec_context); -#else - //**** Fill the execution context structure ****// - exec_context.threads_per_block = THREADS_PER_BLOCK; - exec_context.blocks_per_grid = BLOCKS_PER_GRID; - exec_context.chunk_size = CHUNK_SIZE; - exec_context.total_size = size; - exec_context.total_threads = TOTAL_NUM_OF_THREADS; - exec_context.pad_size = 0; -#endif /* FEATURE_DYNAMIC_EXEC_CONTEXT */ - - sc_print_exec_context( &exec_context ); - - result = (unsigned char*)malloc(MD5_HASH_SIZE * exec_context.total_threads); - -#ifdef FEATURE_WIN32_THREADS - - int i; - - //allocate memory for the thread ids - thread_id = (DWORD *)malloc(sizeof(DWORD)*exec_context.total_threads-1); - - //allocate memory for the thread handle - thread_handle = (HANDLE *)malloc(sizeof(HANDLE)*exec_context.total_threads-1); - - //create structures for thread ids - for( k = 0; k < exec_context.total_threads-1; k++ ) { - - // set indices - chunk_index = k * offset; - hash_index = k * MD5_HASH_SIZE; - - // Allocate memory for thread data. - // thread_data_type is a type that contains the input and output buffers - // wrapped up into a structure - // this is used by the thread to compute and store the hashed values. - thread_data = (pt_thread_data_type) HeapAlloc(GetProcessHeap(), - HEAP_ZERO_MEMORY, - sizeof(thread_data_type)); - - // In case something wrong happen. That is: if memory cannot be - // allocated in the Heap of the current process. - if( thread_data == NULL ) { - fprintf(stderr,"\n[FATAL ERROR] Unable to allocate memory in the heap for Thread[%d]",k); - ExitProcess(2); - } - - // Generate unique data for each thread. - thread_data->input = buffer + chunk_index; - thread_data->ilen = block_size; - thread_data->output = result + hash_index; - - // Create a thread - thread_handle[k] = CreateThread(NULL, 0, md5_cpu_mt, thread_data, - 0, &thread_id[k]); - - // Check whether the thread was created correctly. If it was not, close the - // handlers and release memory - if (thread_handle[k] == NULL) { - - fprintf(stderr,"\n[FATAL ERROR] Unable to spawn thread[%d].\n\t Releasing resources and saying goodbye!\n",k); - - for( i=0; i < exec_context.total_threads-1; i++) { - - if ( thread_handle[i] != NULL ) { - CloseHandle(thread_handle[i]); - } - } - - HeapFree(GetProcessHeap(), 0, thread_data); - - ExitProcess(k); - } - } - - // wait for each thread to finish - WaitForMultipleObjects(exec_context.total_threads-1, thread_handle, - TRUE, INFINITE); - - // Close all thread handles and free memory allocation. - for(k=0; k < exec_context.total_threads-1; k++) { - CloseHandle(thread_handle[k]); - } - - HeapFree(GetProcessHeap(), 0, thread_data); - -#else - - for(k = 0 ; k < exec_context.total_threads - 1; k++) { - chunk_index = k * offset; - hash_index = k * MD5_HASH_SIZE; - md5_cpu_internal(buffer + chunk_index, block_size, result + hash_index ); - } - -#endif /* FEATURE_WIN32_THREADS */ - - chunk_index = k * offset; - hash_index = k * MD5_HASH_SIZE; - md5_cpu_internal(buffer + chunk_index, block_size - exec_context.pad_size, - result + hash_index ); - - - *output = result; - *output_size = MD5_HASH_SIZE * exec_context.total_threads; - -} - -/*=========================================================================== - -FUNCTION SC_SHA1 - -DESCRIPTION - CPU version of the SHA1 algorithm - -DEPENDENCIES - None - -RETURN VALUE - Hash - -===========================================================================*/ -void sc_sha1( unsigned char* buffer, int size, - unsigned char** output, int* output_size) { - - //**** Variable Declarations ****// - sc_exec_context_type exec_context; - unsigned char* scratch_data; - int chunk_index; - int hash_index; - int k; - -#ifdef FEATURE_WIN32_THREADS - - /* This structure contains the input for a particular thread */ - pt_thread_data_type thread_data; - - /* Thread identifiers */ - DWORD *thread_id; - - /* Thread handlers */ - HANDLE *thread_handle; - -#endif /* FEATURE_WIN32_THREADS */ - -#ifdef FEATURE_DYNAMIC_EXEC_CONTEXT - //**** Calculate pad size and needed block and grid sizes ****// - sc_get_exec_context(size, &exec_context); -#else - //**** Fill the execution context structure ****// - exec_context.threads_per_block = THREADS_PER_BLOCK; - exec_context.blocks_per_grid = BLOCKS_PER_GRID; - exec_context.chunk_size = CHUNK_SIZE; - exec_context.total_size = size; - exec_context.total_threads = TOTAL_NUM_OF_THREADS; - exec_context.pad_size = 0; -#endif /* FEATURE_DYNAMIC_EXEC_CONTEXT */ - - sc_print_exec_context( &exec_context ); - - scratch_data = (unsigned char*)malloc(SHA1_HASH_SIZE * exec_context.total_threads); - -#ifdef FEATURE_WIN32_THREADS - - int i; - - //allocate memory for the thread ids - thread_id = (DWORD *)malloc(sizeof(DWORD)*exec_context.total_threads-1); - - //allocate memory for the thread handle - thread_handle = (HANDLE *)malloc(sizeof(HANDLE)*exec_context.total_threads-1); - - //create structures for thread ids - for( k = 0; k < exec_context.total_threads-1; k++ ) { - - // set indices - chunk_index = k * exec_context.chunk_size; - hash_index = k * SHA1_HASH_SIZE; - - // Allocate memory for thread data. - // thread_data_type is a type that contains the input and output buffers - // wrapped up into a structure - // this is used by the thread to compute and store the hashed values. - thread_data = (pt_thread_data_type) HeapAlloc(GetProcessHeap(), - HEAP_ZERO_MEMORY, - sizeof(thread_data_type)); - - // In case something wrong happen. That is: if memory cannot be - // allocated in the Heap of the current process. - if( thread_data == NULL ) { - fprintf(stderr,"\n[FATAL ERROR] Unable to allocate memory in the heap for Thread[%d]",k); - ExitProcess(2); - } - - // Generate unique data for each thread. - thread_data->input = buffer + chunk_index; - thread_data->ilen = exec_context.chunk_size; - thread_data->output = scratch_data + hash_index; - - // Create a thread - thread_handle[k] = CreateThread(NULL, 0, sha1_cpu_mt, thread_data, - 0, &thread_id[k]); - - // Check whether the thread was created correctly. If it was not, close the - // handlers and release memory - if (thread_handle[k] == NULL) { - - fprintf(stderr,"\n[FATAL ERROR] Unable to spawn thread[%d].\n\t Releasing resources and saying goodbye!\n",k); - - for( i=0; i < exec_context.total_threads-1; i++) { - - if ( thread_handle[i] != NULL ) { - CloseHandle(thread_handle[i]); - } - } - - HeapFree(GetProcessHeap(), 0, thread_data); - - ExitProcess(k); - } - } - - // wait for each thread to finish - WaitForMultipleObjects(exec_context.total_threads-1, thread_handle, - TRUE, INFINITE); - - // Close all thread handles and free memory allocation. - for(k=0; k < exec_context.total_threads-1; k++) { - CloseHandle(thread_handle[k]); - } - - HeapFree(GetProcessHeap(), 0, thread_data); - -#else - - for( k = 0; k < exec_context.total_threads - 1; k++) { - chunk_index = k * exec_context.chunk_size; - hash_index = k * SHA1_HASH_SIZE; - sha1_cpu_internal(buffer + chunk_index, exec_context.chunk_size, - scratch_data + hash_index ); - } - -#endif /* FEATURE_WIN32_THREADS */ - - chunk_index = k * exec_context.chunk_size; - hash_index = k * SHA1_HASH_SIZE; - - if(exec_context.pad_size != 0) { - - unsigned char *last_chunk = (unsigned char*)malloc(exec_context.chunk_size); - - memset(last_chunk, 0, exec_context.chunk_size); - memcpy(last_chunk, buffer + chunk_index, - exec_context.chunk_size - exec_context.pad_size); - sha1_cpu_internal(last_chunk, exec_context.chunk_size, - scratch_data + hash_index ); - } else { - - sha1_cpu_internal(buffer + chunk_index, exec_context.chunk_size, - scratch_data + hash_index ); - } - - //**** will do the last hshing stage ****// - sc_sha1_standard( scratch_data, SHA1_HASH_SIZE * exec_context.total_threads, - output ); - - *output_size = SHA1_HASH_SIZE; - -} - -/*=========================================================================== - -FUNCTION SC_SHA1_OVERLAP - -DESCRIPTION - CPU version of the SHA1 overlap algorithm - -DEPENDENCIES - None - -RETURN VALUE - Hash - -===========================================================================*/ -void sc_sha1_overlap(unsigned char* buffer, int size, int block_size, - int offset, unsigned char** output, int* output_size) { - - //**** Variable Declarations ****// - sc_exec_context_type exec_context; - unsigned char* result; - int chunk_index; - int hash_index; - int k; - -#ifdef FEATURE_WIN32_THREADS - - /* This structure contains the input for a particular thread */ - pt_thread_data_type thread_data; - - /* Thread identifiers */ - DWORD *thread_id; - - /* Thread handlers */ - HANDLE *thread_handle; - -#endif /* FEATURE_WIN32_THREADS */ - -#ifdef FEATURE_DYNAMIC_EXEC_CONTEXT - - //**** Calculate pad size and needed block and grid sizes ****// - sc_get_overlap_exec_context(size, offset, block_size, &exec_context); - -#else - //**** Fill the execution context structure ****// - exec_context.threads_per_block = THREADS_PER_BLOCK; - exec_context.blocks_per_grid = BLOCKS_PER_GRID; - exec_context.chunk_size = CHUNK_SIZE; - exec_context.total_size = size; - exec_context.total_threads = TOTAL_NUM_OF_THREADS; - exec_context.pad_size = 0; -#endif /* FEATURE_DYNAMIC_EXEC_CONTEXT */ - - sc_print_exec_context( &exec_context ); - - result = (unsigned char*)malloc(SHA1_HASH_SIZE * exec_context.total_threads); - -#ifdef FEATURE_WIN32_THREADS - - int i; - - //allocate memory for the thread ids - thread_id = (DWORD *)malloc(sizeof(DWORD)*exec_context.total_threads-1); - - //allocate memory for the thread handle - thread_handle = (HANDLE *)malloc(sizeof(HANDLE)*exec_context.total_threads-1); - - //create structures for thread ids - for( k = 0; k < exec_context.total_threads-1; k++ ) { - - // set indices - chunk_index = k * offset; - hash_index = k * SHA1_HASH_SIZE; - - // Allocate memory for thread data. - // thread_data_type is a type that contains the input and output buffers - // wrapped up into a structure - // this is used by the thread to compute and store the hashed values. - thread_data = (pt_thread_data_type) HeapAlloc(GetProcessHeap(), - HEAP_ZERO_MEMORY, - sizeof(thread_data_type)); - - // In case something wrong happen. That is: if memory cannot be - // allocated in the Heap of the current process. - if( thread_data == NULL ) { - fprintf(stderr,"\n[FATAL ERROR] Unable to allocate memory in the heap for Thread[%d]",k); - ExitProcess(2); - } - - // Generate unique data for each thread. - thread_data->input = buffer + chunk_index; - thread_data->ilen = block_size; - thread_data->output = result + hash_index; - - // Create a thread - thread_handle[k] = CreateThread(NULL, 0, sha1_cpu_mt, thread_data, - 0, &thread_id[k]); - - // Check whether the thread was created correctly. If it was not, close the - // handlers and release memory - if (thread_handle[k] == NULL) { - - fprintf(stderr,"\n[FATAL ERROR] Unable to spawn thread[%d].\n\t Releasing resources and saying goodbye!\n",k); - - for( i=0; i < exec_context.total_threads-1; i++) { - - if ( thread_handle[i] != NULL ) { - CloseHandle(thread_handle[i]); - } - } - - HeapFree(GetProcessHeap(), 0, thread_data); - - ExitProcess(k); - } - } - - // wait for each thread to finish - WaitForMultipleObjects(exec_context.total_threads-1, thread_handle, - TRUE, INFINITE); - - // Close all thread handles and free memory allocation. - for(k=0; k < exec_context.total_threads-1; k++) { - CloseHandle(thread_handle[k]); - } - - HeapFree(GetProcessHeap(), 0, thread_data); - -#else - - for(k = 0 ; k < exec_context.total_threads - 1; k++) { - chunk_index = k * offset; - hash_index = k * SHA1_HASH_SIZE; - sha1_cpu_internal(buffer + chunk_index, block_size, result + hash_index ); - } - -#endif /* FEATURE_WIN32_THREADS */ - - chunk_index = k * offset; - hash_index = k * SHA1_HASH_SIZE; - sha1_cpu_internal(buffer + chunk_index, block_size - exec_context.pad_size, - result + hash_index ); - - - *output = result; - *output_size = SHA1_HASH_SIZE * exec_context.total_threads; - -} diff --git a/benchmarks/CUDA/STO/storeCPU.h b/benchmarks/CUDA/STO/storeCPU.h deleted file mode 100644 index 6aada9b..0000000 --- a/benchmarks/CUDA/STO/storeCPU.h +++ /dev/null @@ -1,195 +0,0 @@ -#ifndef STORECPU_H -#define STORECPU_H -/*========================================================================== - S T O R E C P U - -* Copyright (c) 2008, NetSysLab at the University of British Columbia -* All rights reserved. -* -* Redistribution and use in source and binary forms, with or without -* modification, are permitted provided that the following conditions are met: -* * Redistributions of source code must retain the above copyright -* notice, this list of conditions and the following disclaimer. -* * Redistributions in binary form must reproduce the above copyright -* notice, this list of conditions and the following disclaimer in the -* documentation and/or other materials provided with the distribution. -* * Neither the name of the University nor the -* names of its contributors may be used to endorse or promote products -* derived from this software without specific prior written permission. -* -* THIS SOFTWARE IS PROVIDED BY NetSysLab ``AS IS'' AND ANY -* EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED -* WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE -* DISCLAIMED. IN NO EVENT SHALL NetSysLab BE LIABLE FOR ANY -* DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES -* (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; -* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND -* ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT -* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS -* SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. - -DESCRIPTION - CPU version of the storeGPU library. - - -==========================================================================*/ - -/*========================================================================== - - INCLUDES - -==========================================================================*/ - - - -/*========================================================================== - - DATA DECLARATIONS - -==========================================================================*/ - -/*-------------------------------------------------------------------------- - TYPE DEFINITIONS ---------------------------------------------------------------------------*/ - - -/*-------------------------------------------------------------------------- - FUNCTION PROTOTYPES ---------------------------------------------------------------------------*/ - -/*-------------------------------------------------------------------------- - CONSTANTS ---------------------------------------------------------------------------*/ - -/*-------------------------------------------------------------------------- - GLOBAL VARIABLES ---------------------------------------------------------------------------*/ - - -/*-------------------------------------------------------------------------- - MACROS ---------------------------------------------------------------------------*/ - -/*========================================================================== - - FUNCTIONS - -==========================================================================*/ - -/*-------------------------------------------------------------------------- - LOCAL FUNCTIONS ---------------------------------------------------------------------------*/ - - - -/*-------------------------------------------------------------------------- - GLOBAL FUNCTIONS ---------------------------------------------------------------------------*/ -/*=========================================================================== - -FUNCTION SC_MD5 - -DESCRIPTION - CPU version of the MD5 algorithm - -DEPENDENCIES - None - -RETURN VALUE - Hash - -===========================================================================*/ -//extern "C" -void sc_md5( unsigned char* buffer, int size, - unsigned char** output, int* output_size); - - -/*=========================================================================== - -FUNCTION SC_MD5_OVERLAP - -DESCRIPTION - CPU version of the MD5 overlap algorithm - -DEPENDENCIES - None - -RETURN VALUE - Hash - -===========================================================================*/ -//extern "C" -void sc_md5_overlap(unsigned char* buffer, int size, int block_size, int offset, - unsigned char** output, int* output_size); - -/*=========================================================================== - -FUNCTION SC_SHA1 - -DESCRIPTION - CPU version of the SHA1 algorithm - -DEPENDENCIES - None - -RETURN VALUE - Hash - -===========================================================================*/ -//extern "C" -void sc_sha1( unsigned char* buffer, int size, - unsigned char** output, int* output_size); - -/*=========================================================================== - -FUNCTION SC_SHA1_OVERLAP - -DESCRIPTION - CPU version of the SHA1 overlap algorithm - -DEPENDENCIES - None - -RETURN VALUE - Hash - -===========================================================================*/ -//extern "C" -void sc_sha1_overlap(unsigned char* buffer, int size, int block_size, - int offset, unsigned char** output, int* output_size); - -/*=========================================================================== - -FUNCTION SC_MD5_STANDARD - -DESCRIPTION - The standard MD5 algorithm - -DEPENDENCIES - None - -RETURN VALUE - Hash - -===========================================================================*/ -//extern "C" -void sc_md5_standard( unsigned char* buffer, int size, unsigned char** output); - -/*=========================================================================== - -FUNCTION SC_SHA1_STANDARD - -DESCRIPTION - The standard SHA1 algorithm - -DEPENDENCIES - None - -RETURN VALUE - Hash - -===========================================================================*/ -//extern "C" -void sc_sha1_standard( unsigned char* buffer, int size, unsigned char** output); - -#endif /* STORECPU_H */ diff --git a/benchmarks/CUDA/STO/storeGPU.cu b/benchmarks/CUDA/STO/storeGPU.cu deleted file mode 100644 index 70e23ab..0000000 --- a/benchmarks/CUDA/STO/storeGPU.cu +++ /dev/null @@ -1,1211 +0,0 @@ -/*========================================================================== - S T O R E G P U - -* Copyright (c) 2008, NetSysLab at the University of British Columbia -* All rights reserved. -* -* Redistribution and use in source and binary forms, with or without -* modification, are permitted provided that the following conditions are met: -* * Redistributions of source code must retain the above copyright -* notice, this list of conditions and the following disclaimer. -* * Redistributions in binary form must reproduce the above copyright -* notice, this list of conditions and the following disclaimer in the -* documentation and/or other materials provided with the distribution. -* * Neither the name of the University nor the -* names of its contributors may be used to endorse or promote products -* derived from this software without specific prior written permission. -* -* THIS SOFTWARE IS PROVIDED BY NetSysLab ``AS IS'' AND ANY -* EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED -* WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE -* DISCLAIMED. IN NO EVENT SHALL NetSysLab BE LIABLE FOR ANY -* DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES -* (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; -* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND -* ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT -* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS -* SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. - -DESCRIPTION - Main entry of the library. - - -==========================================================================*/ - -/*========================================================================== - - INCLUDES - -==========================================================================*/ -// system -#include -#include -#include - -// project -#include -#include -#include -#include -#include -#include - -// kernels -#include -#include - -/*========================================================================== - - DATA DECLARATIONS - -==========================================================================*/ - -/*-------------------------------------------------------------------------- - TYPE DEFINITIONS ---------------------------------------------------------------------------*/ - -// defines a GPU device properties -typedef struct sg_dev_prop { - int max_thread_per_block; - int max_grid_size; - int global_mem_size; - int warp_size; -} sg_dev_prop_type; - -// defines an execution context used to lunch a kernel. -typedef struct sg_exec_context { - int threads_per_block; - int blocks_per_grid; - int total_threads; - int chunk_size; - int total_size; - int pad_size; -} sg_exec_context_type; - - -/*-------------------------------------------------------------------------- - FUNCTION PROTOTYPES ---------------------------------------------------------------------------*/ - -/*-------------------------------------------------------------------------- - CONSTANTS ---------------------------------------------------------------------------*/ - - - -/*-------------------------------------------------------------------------- - GLOBAL VARIABLES ---------------------------------------------------------------------------*/ - - -/*-------------------------------------------------------------------------- - MACROS ---------------------------------------------------------------------------*/ - -#define GET_REAL_CHUNK_SIZE(chunk_size) ((chunk_size) - 12) - -/*========================================================================== - - FUNCTIONS - -==========================================================================*/ - -/*-------------------------------------------------------------------------- - LOCAL FUNCTIONS ---------------------------------------------------------------------------*/ -#ifdef FEATURE_DYNAMIC_EXEC_CONTEXT -/*=========================================================================== - -FUNCTION SG_GET_DEV_PROP - -DESCRIPTION - Probes the device for its properties - -DEPENDENCIES - None - -RETURN VALUE - device information - -===========================================================================*/ -static void sg_get_dev_prop(sg_dev_prop_type* dev_prop) { - - struct cudaDeviceProp prop; - int dev; - - cudaGetDevice(&dev); - cudaGetDeviceProperties(&prop, dev); - - printf("\n== Device Properties ==\n"); - printf("Max global memory : %d\n", prop.totalGlobalMem); - printf("Registers per block : %d\n", prop.regsPerBlock); - printf("Warp size : %d\n", prop.warpSize); - printf("Max threads per block: %d\n", prop.maxThreadsPerBlock); - printf("Block Dimensions : %d, %d, %d\n", - prop.maxThreadsDim[0], - prop.maxThreadsDim[1], - prop.maxThreadsDim[2]); - printf("Grid Dimensions : %d, %d, %d\n", - prop.maxGridSize[0], - prop.maxGridSize[1], - prop.maxGridSize[2]); - - dev_prop->max_thread_per_block = prop.maxThreadsDim[0]; - dev_prop->max_grid_size = prop.maxGridSize[0]; - dev_prop->global_mem_size = prop.totalGlobalMem; - dev_prop->warp_size = prop.warpSize; - -} - -#ifdef FEATURE_MAXIMIZE_NUM_OF_THREADS -/*=========================================================================== - -FUNCTION SG_GET_EXEC_CONTEXT - -DESCRIPTION - sets the required chunk size, thread per block and number of blocks - needed for kernel execution according to client buffer size. - -DEPENDENCIES - None - -RETURN VALUE - execution context - -===========================================================================*/ -static sg_status_type sg_get_exec_context(int size, int hash_size, - sg_exec_context_type* ctx){ - - sg_dev_prop_type dev_prop; - int threads_per_block; - int blocks_per_grid; - int total_threads; - int chunk_size; - int pad_size; - - int total_chunks = 0; - int found = 0; - int index = 1; - - - //**** Get device information ****// - sg_get_dev_prop(&dev_prop); - - //**** Determine the execution context ****// - /* The algorithm will try to determine the context by minimizing chunk - * size and maximizing total number of threads - * TODO: May be we can do better here - */ - while ( !found ) { - // Set chunk size - chunk_size = GET_REAL_CHUNK_SIZE(BASIC_CHUNK_SIZE * index); - - if ( chunk_size > MAX_CHUNK_SIZE ) - break; - - // Calculate the required padding for this chunk size - pad_size = ((size % chunk_size) == 0) ? 0 : - chunk_size - (size % chunk_size); - - // total number of chunks required if we are going to use this chunk size - total_chunks = (pad_size == 0) ? size / chunk_size : - (size / chunk_size) + 1; - - if ( total_chunks <= MAX_NUM_OF_THREADS ) { - // Got it, this is the minimum chunk size we can use. Now determine the - // threads and blocks numbers. - total_threads = total_chunks; - - // Get block and grid sizes - if (total_chunks <= MAX_THREADS_PER_BLOCK ) { - threads_per_block = total_chunks; - blocks_per_grid = 1; - - } else { - threads_per_block = MAX_THREADS_PER_BLOCK; - blocks_per_grid = ((total_threads % threads_per_block) == 0) ? - (total_threads/threads_per_block) : - (total_threads/threads_per_block) + 1; - - } - - if ( dev_prop.global_mem_size > - (total_threads * (chunk_size + hash_size) + hash_size) ) - found = 1; - - } - index++; - - } - - //**** Did we find a solution? ****// - if ( !found ) - return SG_ERR_DEV_MEM_OVERFLOW; - - //**** Fill the struct with the solution ****// - ctx->threads_per_block = threads_per_block; - ctx->blocks_per_grid = blocks_per_grid; - ctx->total_threads = total_threads; - ctx->total_size = size + pad_size; - ctx->chunk_size = chunk_size; - ctx->pad_size = pad_size; - - return SG_OK; - -} - -#else /* FEATURE_MAXIMIZE_NUM_OF_THREADS */ -/*=========================================================================== - -FUNCTION SG_GET_EXEC_CONTEXT - -DESCRIPTION - sets the required chunk size, thread per block and number of blocks - needed for kernel execution according to client buffer size. - -DEPENDENCIES - None - -RETURN VALUE - execution context - -===========================================================================*/ -static sg_status_type sg_get_exec_context(int size, int hash_size, - sg_exec_context_type* ctx){ - - sg_dev_prop_type dev_prop; - int threads_per_block; - int blocks_per_grid; - int total_threads; - int chunk_size; - int pad_size; - - int total_chunks = 0; - int found = 0; - - int index = MAX_CHUNK_SIZE / BASIC_CHUNK_SIZE; - - - //**** Get device information ****// - sg_get_dev_prop(&dev_prop); - - //**** Determine the execution context ****// - /* The algorithm will try to determine the context by minimizing chunk - * size and maximizing total number of threads - * TODO: May be we can do better here - */ - while ( 1 ) { - // Set chunk size - chunk_size = GET_REAL_CHUNK_SIZE(BASIC_CHUNK_SIZE * index); - - // don't go less than minimum chunk size - if ( chunk_size < GET_REAL_CHUNK_SIZE(BASIC_CHUNK_SIZE) ) - break; - - // Calculate the required padding for this chunk size - pad_size = ((size % chunk_size) == 0) ? 0 : - chunk_size - (size % chunk_size); - - // total number of chunks required if we are going to use this chunk size - total_chunks = (pad_size == 0) ? size / chunk_size : - (size / chunk_size) + 1; - - - // don't go beyond the maximum number of threads or maximum global memory - // TODO: it seems that the kernel breaks way before reaching the maximum - // global memory size (around 94MByte input plus the required - // scratch space) - if (( total_chunks > MAX_NUM_OF_THREADS) || - ( dev_prop.global_mem_size < (total_chunks * - (chunk_size + hash_size) + - hash_size))) - break; - - // each thread will take care of one chunk - total_threads = total_chunks; - - - // Get block and grid sizes - if (total_chunks <= MAX_THREADS_PER_BLOCK ) { - threads_per_block = total_chunks; - blocks_per_grid = 1; - - } else { - threads_per_block = MAX_THREADS_PER_BLOCK; - blocks_per_grid = ((total_threads % threads_per_block) == 0) ? - (total_threads/threads_per_block) : - (total_threads/threads_per_block) + 1; - - } - found = 1; - - if( total_threads > NUM_OF_MULTIPROCESSORS * dev_prop.warp_size) - break; - - index--; - - } - - //**** Did we find a solution? ****// - if ( !found ) - return SG_ERR_DEV_MEM_OVERFLOW; - - //**** Fill the struct with the solution ****// - ctx->threads_per_block = threads_per_block; - ctx->blocks_per_grid = blocks_per_grid; - ctx->total_threads = total_threads; - ctx->total_size = size + pad_size; - ctx->chunk_size = chunk_size; - ctx->pad_size = pad_size; - - return SG_OK; - -} -#endif /* FEATURE_MAXIMIZE_NUM_OF_THREADS */ - -/*=========================================================================== - -FUNCTION SG_GET_OVERLAP_EXEC_CONTEXT - -DESCRIPTION - sets the required chunk size, thread per block and number of blocks - needed for kernel execution according to client buffer size, offset - and block size. - -DEPENDENCIES - None - -RETURN VALUE - execution context - -===========================================================================*/ -static sg_status_type sg_get_overlap_exec_context(int size, int offset, - int block_size, int hash_size, - sg_exec_context_type* ctx) { - - sg_dev_prop_type dev_prop; - int threads_per_block; - int blocks_per_grid; - int total_threads; - int total_size; - int pad_size; - - - //**** Get device information ****// - sg_get_dev_prop(&dev_prop); - - - //**** Get the total number of threads required ****// - total_threads = (size + offset - block_size) / offset; - total_threads = ((size + offset - block_size) % offset) != 0 ? - total_threads + 1 : total_threads; - - - //**** Get the required padding for the last block ****// - pad_size = ((total_threads - 1) * offset + block_size) - size; - - - //**** threads and blocks ****// - if( total_threads > MAX_THREADS_PER_BLOCK ) { - - threads_per_block = MAX_THREADS_PER_BLOCK; - blocks_per_grid = (total_threads % MAX_THREADS_PER_BLOCK) == 0 ? - (total_threads / MAX_THREADS_PER_BLOCK) : - (total_threads / MAX_THREADS_PER_BLOCK) + 1; - } else { - - threads_per_block = total_threads; - blocks_per_grid = 1; - } - - total_size = size + pad_size; - - //**** Check for device memory overflow ****// - if( dev_prop.global_mem_size < (total_size + (hash_size * total_threads))) { - return SG_ERR_DEV_MEM_OVERFLOW; - } - - - //**** Fill the struct with the solution ****// - ctx->threads_per_block = threads_per_block; - ctx->blocks_per_grid = blocks_per_grid; - ctx->total_threads = total_threads; - ctx->total_size = total_size; - ctx->chunk_size = block_size; - ctx->pad_size = pad_size; - - return SG_OK; - -} -#endif /* FEATURE_DYNAMIC_EXEC_CONTEXT */ - -/*=========================================================================== - -FUNCTION SG_PRINT_EXEC_CONTEXT - -DESCRIPTION - Prints out the passed execution context structure - -DEPENDENCIES - None - -RETURN VALUE - None - -===========================================================================*/ -static void sg_print_exec_context( sg_exec_context_type* ctx ) { - printf("\n== GPU Execution Context ==\n"); - printf("Threads : %d\n", ctx->threads_per_block); - printf("Blocks : %d\n", ctx->blocks_per_grid); - printf("Total Threads : %d\n", ctx->total_threads); - printf("Total size : %d\n", ctx->total_size); - printf("Chunk Size : %d\n", ctx->chunk_size); - printf("Padding : %d\n\n", ctx->pad_size); -} - -/*-------------------------------------------------------------------------- - GLOBAL FUNCTIONS ---------------------------------------------------------------------------*/ - - - -/*=========================================================================== - -FUNCTION SG_INIT - -DESCRIPTION - Initialize the library - -DEPENDENCIES - None - -RETURN VALUE - None - -===========================================================================*/ -void sg_init( ) { - - char *buffer; - - //**** Utility library initialization ****// - // initialise card and timer - int deviceCount; - CUDA_SAFE_CALL_NO_SYNC(cudaGetDeviceCount(&deviceCount)); - if (deviceCount == 0) { - fprintf(stderr, "There is no device.\n"); - exit(EXIT_FAILURE); - } - int dev; - for (dev = 0; dev < deviceCount; ++dev) { - cudaDeviceProp deviceProp; - CUDA_SAFE_CALL_NO_SYNC(cudaGetDeviceProperties(&deviceProp, dev)); - if (deviceProp.major >= 1) - break; - } - if (dev == deviceCount) { - fprintf(stderr, "There is no device supporting CUDA.\n"); - exit(EXIT_FAILURE); - } - else - CUDA_SAFE_CALL(cudaSetDevice(dev)); - - - //**** force runtime initialization (CUDA ref. manual for more info.) ****// - cudaMallocHost( (void**) &buffer, 4 ); - cudaFreeHost( buffer ); - -} - -/*=========================================================================== - -FUNCTION SG_MALLOC - -DESCRIPTION - Allocate the required memory size. - -DEPENDENCIES - None - -RETURN VALUE - pointer to the reseved buffer - -===========================================================================*/ -void* sg_malloc(unsigned int size){ - - void* buffer; - -#ifdef FEATURE_PINNED_MODE - cudaMallocHost( (void**) &buffer, size ); -#else - - buffer = malloc( size ); -#endif /* FEATURE_PINNED_MODE */ - - return buffer; -} - -/*=========================================================================== - -FUNCTION SG_FREE - -DESCRIPTION - Free the allocated buffer. - -DEPENDENCIES - None - -RETURN VALUE - pointer to the reseved buffer - -===========================================================================*/ -void sg_free(void* buffer){ - - -#ifdef FEATURE_PINNED_MODE - cudaFreeHost(buffer ); -#else - free( buffer ); -#endif - -} - -/*=========================================================================== - -FUNCTION SG_MD5 - -DESCRIPTION - Returns the MD5 hash - -DEPENDENCIES - None - -RETURN VALUE - Hash value - -===========================================================================*/ -sg_status_type sg_md5(unsigned char* buffer, int size, - unsigned char** output, int* output_size, - sg_time_breakdown_type* time_breakdown) { - - //**** Variable Declarations ****// - sg_exec_context_type exec_context; - sg_status_type status = SG_OK; - unsigned char* d_scratchData; - unsigned char* h_scratchData; - unsigned char* d_input; - unsigned int timer; - - - //**** create the timer ****// - timer = 0; - CUT_SAFE_CALL( cutCreateTimer( &timer)); - - -#ifdef FEATURE_DYNAMIC_EXEC_CONTEXT - //**** Calculate pad size and needed block and grid sizes ****// - status = sg_get_exec_context(size, MD5_HASH_SIZE, &exec_context); - if ( status != SG_OK ) { - printf("Global memory overflow\n"); - return status; - } -#else - - //**** Fill the execution context structure ****// - exec_context.threads_per_block = THREADS_PER_BLOCK; - exec_context.blocks_per_grid = BLOCKS_PER_GRID; - exec_context.chunk_size = CHUNK_SIZE; - exec_context.total_size = size; - exec_context.total_threads = TOTAL_NUM_OF_THREADS; - exec_context.pad_size = 0; - -#endif /* FEATURE_DYNAMIC_EXEC_CONTEXT */ - sg_print_exec_context(&exec_context); - - - - //**** device memory allocation timing ****// - CUT_SAFE_CALL(cutResetTimer(timer)); - CUT_SAFE_CALL(cutStartTimer(timer)); - - /* allocate input data space */ - CUDA_SAFE_CALL(cudaMalloc((void**) &d_input, exec_context.total_size)); - - /* allocate scratch space */ - CUDA_SAFE_CALL(cudaMalloc((void**) &d_scratchData, - MD5_HASH_SIZE * exec_context.total_threads)); - - /* stop the timer (device memory allocation) */ - CUT_SAFE_CALL(cutStopTimer(timer)); - time_breakdown->device_mem_alloc_time = cutGetTimerValue(timer); - - - - //**** scratch buffer allocation timing ****// - CUT_SAFE_CALL(cutResetTimer(timer)); - CUT_SAFE_CALL(cutStartTimer(timer)); - - /* allocate buffer for the results */ - cudaMallocHost((void **)&h_scratchData, MD5_HASH_SIZE * - exec_context.total_threads); - - /* stop the timer (scratch buffer allocation) */ - CUT_SAFE_CALL(cutStopTimer(timer)); - time_breakdown->host_output_buffer_alloc_time = cutGetTimerValue(timer); - - - - //**** start timer for data copy in timing ****// - CUT_SAFE_CALL(cutResetTimer(timer)); - CUT_SAFE_CALL(cutStartTimer(timer)); - - /* move data to the device memory */ - CUDA_SAFE_CALL(cudaMemcpy( d_input, buffer, size, - cudaMemcpyHostToDevice)); - - /* stop the timer (copy in) */ - CUT_SAFE_CALL( cutStopTimer( timer)); - time_breakdown->copy_in_time = cutGetTimerValue( timer ); - - - - //**** setup execution parameters ****// - dim3 block( exec_context.threads_per_block ); - dim3 grid( exec_context.blocks_per_grid ); - - - - //**** start timer for kernel execution timing ****// - CUT_SAFE_CALL(cutResetTimer(timer)); - CUT_SAFE_CALL(cutStartTimer(timer)); - - - /* execute the kernel */ - md5<<< grid, block >>>(d_input, exec_context.chunk_size, - exec_context.total_threads, - exec_context.pad_size, - d_scratchData); - - // check if kernel execution generated an error - CUT_CHECK_ERROR("Kernel execution failed"); - - /* wait till the kernel finishes execution */ - CUDA_SAFE_CALL(cudaThreadSynchronize()); - - /* stop the timer (kernel execution) */ - CUT_SAFE_CALL(cutStopTimer(timer)); - time_breakdown->exec_time = cutGetTimerValue(timer); - - - - //**** start timer for output copy out timing ****// - CUT_SAFE_CALL(cutResetTimer(timer)); - CUT_SAFE_CALL(cutStartTimer(timer)); - - /* get the results from the device */ - CUDA_SAFE_CALL(cudaMemcpy(h_scratchData, - d_scratchData, - MD5_HASH_SIZE * exec_context.total_threads, - cudaMemcpyDeviceToHost)); - - /* stop the timer (output copy out) */ - CUT_SAFE_CALL(cutStopTimer(timer)); - time_breakdown->copy_out_time = cutGetTimerValue(timer); - - - - //**** start timer for last hasing stage timing ****// - CUT_SAFE_CALL(cutResetTimer(timer)); - CUT_SAFE_CALL(cutStartTimer(timer)); - - /* will do the last hshing stage on the CPU */ - sc_md5_standard(h_scratchData, MD5_HASH_SIZE * exec_context.total_threads, - output ); - - /* stop the timer (last stage) */ - CUT_SAFE_CALL(cutStopTimer(timer)); - time_breakdown->last_stage_time = cutGetTimerValue(timer); - - //**** free allocated memory ****// - CUDA_SAFE_CALL(cudaFree(d_input)); - CUDA_SAFE_CALL(cudaFree(d_scratchData)); - cudaFreeHost(h_scratchData); - - *output_size = MD5_HASH_SIZE; - - - return status; - -} - -/*=========================================================================== - -FUNCTION SG_MD5_OVERLAP - -DESCRIPTION - Returns the MD5 hash of each block for the provided buffer - -DEPENDENCIES - None - -RETURN VALUE - Hash value - -===========================================================================*/ -sg_status_type sg_md5_overlap(unsigned char* buffer, int size, - int block_size, int offset, - unsigned char** output, int* output_size, - sg_time_breakdown_type* time_breakdown) { - - - //**** Variable Declarations ****// - sg_exec_context_type exec_context; - sg_status_type status = SG_OK; - unsigned char* d_output; - unsigned char* d_input; - unsigned int timer; - - - - //**** create the timer ****// - timer = 0; - CUT_SAFE_CALL( cutCreateTimer( &timer)); - -#ifdef FEATURE_DYNAMIC_EXEC_CONTEXT - //**** Calculate pad size and needed block and grid sizes ****// - status = sg_get_overlap_exec_context(size, offset, block_size, - MD5_HASH_SIZE, &exec_context); - if ( status != SG_OK ) { - printf("Global memory overflow\n"); - return status; - } -#else - //**** Fill the execution context structure ****// - exec_context.threads_per_block = THREADS_PER_BLOCK; - exec_context.blocks_per_grid = BLOCKS_PER_GRID; - exec_context.chunk_size = CHUNK_SIZE; - exec_context.total_size = size; - exec_context.total_threads = TOTAL_NUM_OF_THREADS; - exec_context.pad_size = 0; -#endif /* FEATURE_DYNAMIC_EXEC_CONTEXT */ - sg_print_exec_context(&exec_context); - - - - //**** start timer for device memory allocation timing ****// - CUT_SAFE_CALL(cutResetTimer(timer)); - CUT_SAFE_CALL(cutStartTimer(timer)); - - /* allocate input space */ - CUDA_SAFE_CALL(cudaMalloc((void**) &d_input, exec_context.total_size)); - - /* allocate output space */ - CUDA_SAFE_CALL(cudaMalloc((void**) &d_output, - MD5_HASH_SIZE * exec_context.total_threads)); - - /* stop the timer (memory allocation) */ - CUT_SAFE_CALL(cutStopTimer(timer)); - time_breakdown->device_mem_alloc_time = cutGetTimerValue(timer); - - - - //**** start timer for output memory allocation timing ****// - CUT_SAFE_CALL(cutResetTimer(timer)); - CUT_SAFE_CALL(cutStartTimer(timer)); - - /**output = (unsigned char*) sg_malloc(MD5_HASH_SIZE * - exec_context.total_threads);*/ - cudaMallocHost( (void**) output, MD5_HASH_SIZE * - exec_context.total_threads ); - - /* stop the timer (output buffer allocation) */ - CUT_SAFE_CALL(cutStopTimer(timer)); - time_breakdown->host_output_buffer_alloc_time = cutGetTimerValue(timer); - - - - //**** start timer for data copy in timing ****// - CUT_SAFE_CALL(cutResetTimer(timer)); - CUT_SAFE_CALL(cutStartTimer(timer)); - - /* move data to the device memory */ - CUDA_SAFE_CALL(cudaMemcpy(d_input, buffer, size, - cudaMemcpyHostToDevice)); - - /* stop the timer (copy in) */ - CUT_SAFE_CALL(cutStopTimer(timer)); - time_breakdown->copy_in_time = cutGetTimerValue(timer); - - - - //**** setup execution parameters ****// - dim3 block( exec_context.threads_per_block ); - dim3 grid( exec_context.blocks_per_grid ); - - - - //**** start timer for kernel execution timing ****// - CUT_SAFE_CALL(cutResetTimer(timer)); - CUT_SAFE_CALL(cutStartTimer(timer)); - - /* execute the kernel */ - md5_overlap<<< grid, block >>>(d_input, exec_context.chunk_size, - offset, exec_context.total_threads, - exec_context.pad_size, d_output); - - // check if kernel execution generated an error - CUT_CHECK_ERROR("Kernel execution failed"); - - /* wait till the kernel finishes execution */ - CUDA_SAFE_CALL(cudaThreadSynchronize()); - - /* stop the timer (kernel execution) */ - CUT_SAFE_CALL(cutStopTimer(timer)); - time_breakdown->exec_time = cutGetTimerValue(timer); - - - - //**** start timer for output copy out timing ****// - CUT_SAFE_CALL(cutResetTimer(timer)); - CUT_SAFE_CALL(cutStartTimer(timer)); - - /* get the result from the device */ - CUDA_SAFE_CALL(cudaMemcpy(*output, - d_output, - MD5_HASH_SIZE * exec_context.total_threads, - cudaMemcpyDeviceToHost)); - - /* stop the timer (output copy out) */ - CUT_SAFE_CALL(cutStopTimer(timer)); - time_breakdown->copy_out_time = cutGetTimerValue(timer); - - - - //**** free allocated memory ****// - CUDA_SAFE_CALL(cudaFree(d_input)); - CUDA_SAFE_CALL(cudaFree(d_output)); - - *output_size = MD5_HASH_SIZE * exec_context.total_threads; - - return status; -} - -/*=========================================================================== - -FUNCTION SG_SHA1 - -DESCRIPTION - Returns the SHA1 hash of a the provided buffer - -DEPENDENCIES - None - -RETURN VALUE - Hash value - -===========================================================================*/ -sg_status_type sg_sha1(unsigned char* buffer, int size, - unsigned char** output, int* output_size, - sg_time_breakdown_type* time_breakdown) { - - //**** Variable Declarations ****// - sg_exec_context_type exec_context; - sg_status_type status = SG_OK; - unsigned char* d_scratchData; - unsigned char* h_scratchData; - unsigned char* d_input; - unsigned int timer; - - - //**** create the timer ****// - timer = 0; - CUT_SAFE_CALL( cutCreateTimer( &timer)); - - -#ifdef FEATURE_DYNAMIC_EXEC_CONTEXT - //**** Calculate pad size and needed block and grid sizes ****// - status = sg_get_exec_context(size, SHA1_HASH_SIZE, &exec_context); - if ( status != SG_OK ) { - printf("Global memory overflow\n"); - return status; - } -#else - - //**** Fill the execution context structure ****// - exec_context.threads_per_block = THREADS_PER_BLOCK; - exec_context.blocks_per_grid = BLOCKS_PER_GRID; - exec_context.chunk_size = CHUNK_SIZE; - exec_context.total_size = size; - exec_context.total_threads = TOTAL_NUM_OF_THREADS; - exec_context.pad_size = 0; - -#endif /* FEATURE_DYNAMIC_EXEC_CONTEXT */ - sg_print_exec_context(&exec_context); - - - - //**** device memory allocation timing ****// - CUT_SAFE_CALL(cutResetTimer(timer)); - CUT_SAFE_CALL(cutStartTimer(timer)); - - /* allocate input data space */ - CUDA_SAFE_CALL(cudaMalloc((void**) &d_input, exec_context.total_size)); - - /* allocate scratch space */ - CUDA_SAFE_CALL(cudaMalloc((void**) &d_scratchData, - SHA1_HASH_SIZE * exec_context.total_threads)); - - /* stop the timer (device memory allocation) */ - CUT_SAFE_CALL(cutStopTimer(timer)); - time_breakdown->device_mem_alloc_time = cutGetTimerValue(timer); - - - - //**** scratch buffer allocation timing ****// - CUT_SAFE_CALL(cutResetTimer(timer)); - CUT_SAFE_CALL(cutStartTimer(timer)); - - /* allocate buffer for the results */ - cudaMallocHost((void**)&h_scratchData, SHA1_HASH_SIZE * - exec_context.total_threads); - - /* stop the timer (scratch buffer allocation) */ - CUT_SAFE_CALL(cutStopTimer(timer)); - time_breakdown->host_output_buffer_alloc_time = cutGetTimerValue(timer); - - - - //**** start timer for data copy in timing ****// - CUT_SAFE_CALL(cutResetTimer(timer)); - CUT_SAFE_CALL(cutStartTimer(timer)); - - /* move data to the device memory */ - CUDA_SAFE_CALL(cudaMemcpy( d_input, buffer, size, - cudaMemcpyHostToDevice)); - - /* stop the timer (copy in) */ - CUT_SAFE_CALL( cutStopTimer( timer)); - time_breakdown->copy_in_time = cutGetTimerValue( timer ); - - - - //**** setup execution parameters ****// - dim3 block( exec_context.threads_per_block ); - dim3 grid( exec_context.blocks_per_grid ); - - - - //**** start timer for kernel execution timing ****// - CUT_SAFE_CALL(cutResetTimer(timer)); - CUT_SAFE_CALL(cutStartTimer(timer)); - - - /* execute the kernel */ - sha1<<< grid, block >>>(d_input, exec_context.chunk_size, - exec_context.total_threads, - exec_context.pad_size, - d_scratchData); - - // check if kernel execution generated an error - CUT_CHECK_ERROR("Kernel execution failed"); - - /* wait till the kernel finishes execution */ - CUDA_SAFE_CALL(cudaThreadSynchronize()); - - /* stop the timer (kernel execution) */ - CUT_SAFE_CALL(cutStopTimer(timer)); - time_breakdown->exec_time = cutGetTimerValue(timer); - - - - //**** start timer for output copy out timing ****// - CUT_SAFE_CALL(cutResetTimer(timer)); - CUT_SAFE_CALL(cutStartTimer(timer)); - - /* get the results from the device */ - CUDA_SAFE_CALL(cudaMemcpy(h_scratchData, - d_scratchData, - SHA1_HASH_SIZE * exec_context.total_threads, - cudaMemcpyDeviceToHost)); - - /* stop the timer (output copy out) */ - CUT_SAFE_CALL(cutStopTimer(timer)); - time_breakdown->copy_out_time = cutGetTimerValue(timer); - - - - //**** start timer for last hasing stage timing ****// - CUT_SAFE_CALL(cutResetTimer(timer)); - CUT_SAFE_CALL(cutStartTimer(timer)); - - /* will do the last hshing stage on the CPU */ - sc_sha1_standard(h_scratchData, SHA1_HASH_SIZE * exec_context.total_threads, - output ); - - /* stop the timer (last stage) */ - CUT_SAFE_CALL(cutStopTimer(timer)); - time_breakdown->last_stage_time = cutGetTimerValue(timer); - - - //**** free allocated memory ****// - CUDA_SAFE_CALL(cudaFree(d_input)); - CUDA_SAFE_CALL(cudaFree(d_scratchData)); - cudaFreeHost(h_scratchData); - - *output_size = SHA1_HASH_SIZE; - - return status; - -} - - -/*=========================================================================== - -FUNCTION SG_SHA1_OVERLAP - -DESCRIPTION - Returns the SHA1 hash of each block for the provided buffer - -DEPENDENCIES - None - -RETURN VALUE - Hash value - -===========================================================================*/ -sg_status_type sg_sha1_overlap(unsigned char* buffer, int size, - int block_size, int offset, - unsigned char** output, int* output_size, - sg_time_breakdown_type* time_breakdown) { - - - //**** Variable Declarations ****// - sg_exec_context_type exec_context; - sg_status_type status = SG_OK; - unsigned char* d_output; - unsigned char* d_input; - unsigned int timer; - - - - //**** create the timer ****// - timer = 0; - CUT_SAFE_CALL( cutCreateTimer( &timer)); - -#ifdef FEATURE_DYNAMIC_EXEC_CONTEXT - //**** Calculate pad size and needed block and grid sizes ****// - status = sg_get_overlap_exec_context(size, offset, block_size, - SHA1_HASH_SIZE, &exec_context); - if ( status != SG_OK ) { - printf("Global memory overflow\n"); - return status; - } -#else - //**** Fill the execution context structure ****// - exec_context.threads_per_block = THREADS_PER_BLOCK; - exec_context.blocks_per_grid = BLOCKS_PER_GRID; - exec_context.chunk_size = CHUNK_SIZE; - exec_context.total_size = size; - exec_context.total_threads = TOTAL_NUM_OF_THREADS; - exec_context.pad_size = 0; -#endif /* FEATURE_DYNAMIC_EXEC_CONTEXT */ - sg_print_exec_context(&exec_context); - - - - //**** start timer for device memory allocation timing ****// - CUT_SAFE_CALL(cutResetTimer(timer)); - CUT_SAFE_CALL(cutStartTimer(timer)); - - /* allocate input space */ - CUDA_SAFE_CALL(cudaMalloc((void**) &d_input, exec_context.total_size)); - - /* allocate output space */ - CUDA_SAFE_CALL(cudaMalloc((void**) &d_output, - SHA1_HASH_SIZE * exec_context.total_threads)); - - /* stop the timer (memory allocation) */ - CUT_SAFE_CALL(cutStopTimer(timer)); - time_breakdown->device_mem_alloc_time = cutGetTimerValue(timer); - - - - //**** start timer for output memory allocation timing ****// - CUT_SAFE_CALL(cutResetTimer(timer)); - CUT_SAFE_CALL(cutStartTimer(timer)); - - /**output = (unsigned char*) sg_malloc(SHA1_HASH_SIZE * - exec_context.total_threads);*/ - cudaMallocHost( (void**) output, SHA1_HASH_SIZE * - exec_context.total_threads ); - - /* stop the timer (output buffer allocation) */ - CUT_SAFE_CALL(cutStopTimer(timer)); - time_breakdown->host_output_buffer_alloc_time = cutGetTimerValue(timer); - - - - //**** start timer for data copy in timing ****// - CUT_SAFE_CALL(cutResetTimer(timer)); - CUT_SAFE_CALL(cutStartTimer(timer)); - - /* move data to the device memory */ - CUDA_SAFE_CALL(cudaMemcpy(d_input, buffer, size, - cudaMemcpyHostToDevice)); - - /* stop the timer (copy in) */ - CUT_SAFE_CALL(cutStopTimer(timer)); - time_breakdown->copy_in_time = cutGetTimerValue(timer); - - - - //**** setup execution parameters ****// - dim3 block( exec_context.threads_per_block ); - dim3 grid( exec_context.blocks_per_grid ); - - - - //**** start timer for kernel execution timing ****// - CUT_SAFE_CALL(cutResetTimer(timer)); - CUT_SAFE_CALL(cutStartTimer(timer)); - - /* execute the kernel */ - sha1_overlap<<< grid, block >>>(d_input, exec_context.chunk_size, - offset, exec_context.total_threads, - exec_context.pad_size, d_output); - - // check if kernel execution generated an error - CUT_CHECK_ERROR("Kernel execution failed"); - - /* wait till the kernel finishes execution */ - CUDA_SAFE_CALL(cudaThreadSynchronize()); - - /* stop the timer (kernel execution) */ - CUT_SAFE_CALL(cutStopTimer(timer)); - time_breakdown->exec_time = cutGetTimerValue(timer); - - - - //**** start timer for output copy out timing ****// - CUT_SAFE_CALL(cutResetTimer(timer)); - CUT_SAFE_CALL(cutStartTimer(timer)); - - /* get the result from the device */ - CUDA_SAFE_CALL(cudaMemcpy(*output, - d_output, - SHA1_HASH_SIZE * exec_context.total_threads, - cudaMemcpyDeviceToHost)); - - /* stop the timer (output copy out) */ - CUT_SAFE_CALL(cutStopTimer(timer)); - time_breakdown->copy_out_time = cutGetTimerValue(timer); - - - - //**** free allocated memory ****// - CUDA_SAFE_CALL(cudaFree(d_input)); - CUDA_SAFE_CALL(cudaFree(d_output)); - - *output_size = SHA1_HASH_SIZE * exec_context.total_threads; - - return status; -} diff --git a/benchmarks/CUDA/STO/storeGPU.h b/benchmarks/CUDA/STO/storeGPU.h deleted file mode 100644 index 3d8bd0b..0000000 --- a/benchmarks/CUDA/STO/storeGPU.h +++ /dev/null @@ -1,224 +0,0 @@ -#ifndef STOREGPU_H -#define STOREGPU_H -/*========================================================================== - S T O R E G P U - -* Copyright (c) 2008, NetSysLab at the University of British Columbia -* All rights reserved. -* -* Redistribution and use in source and binary forms, with or without -* modification, are permitted provided that the following conditions are met: -* * Redistributions of source code must retain the above copyright -* notice, this list of conditions and the following disclaimer. -* * Redistributions in binary form must reproduce the above copyright -* notice, this list of conditions and the following disclaimer in the -* documentation and/or other materials provided with the distribution. -* * Neither the name of the University nor the -* names of its contributors may be used to endorse or promote products -* derived from this software without specific prior written permission. -* -* THIS SOFTWARE IS PROVIDED BY NetSysLab ``AS IS'' AND ANY -* EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED -* WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE -* DISCLAIMED. IN NO EVENT SHALL NetSysLab BE LIABLE FOR ANY -* DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES -* (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; -* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND -* ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT -* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS -* SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. - -DESCRIPTION - StoreGPU library. - -==========================================================================*/ - -/*========================================================================== - - INCLUDES - -==========================================================================*/ -#include - -/*========================================================================== - - DATA DECLARATIONS - -==========================================================================*/ - - -/*-------------------------------------------------------------------------- - TYPE DEFINITIONS ---------------------------------------------------------------------------*/ -typedef enum { - - SG_OK = 0, - SG_ERR_DEV_MEM_OVERFLOW = -1, - - -} sg_status_type; - -// defines a set of elapsed time measurements taken at specified states while -// running the GPU version of the hashing algorithm -typedef struct sg_time_breakdown { - - float exec_time; /* kernel execution time */ - float device_mem_alloc_time; /* time elapsed to allocate device buffers */ - float host_output_buffer_alloc_time; /* time elapsed to allocate - host output buffer */ - float copy_in_time; /* time elapsed to push data into the GPU */ - float copy_out_time; /* time elapsed to get data from the GPU */ - float last_stage_time; /* time elapsed doing the last stage of hasing - on the CPU*/ -} sg_time_breakdown_type; - -/*-------------------------------------------------------------------------- - FUNCTION PROTOTYPES ---------------------------------------------------------------------------*/ - -/*-------------------------------------------------------------------------- - CONSTANTS ---------------------------------------------------------------------------*/ - - -/*-------------------------------------------------------------------------- - GLOBAL VARIABLES ---------------------------------------------------------------------------*/ - -/*-------------------------------------------------------------------------- - MACROS ---------------------------------------------------------------------------*/ - - - -/*========================================================================== - - FUNCTIONS - -==========================================================================*/ - -/*=========================================================================== - -FUNCTION SG_INIT - -DESCRIPTION - Library initialization - -DEPENDENCIES - None - -RETURN VALUE - None - -===========================================================================*/ -void sg_init(); - -/*=========================================================================== - -FUNCTION SG_MALLOC - -DESCRIPTION - Allocate the required memory size. - -DEPENDENCIES - None - -RETURN VALUE - pointer to the reseved buffer - -===========================================================================*/ -void* sg_malloc(unsigned int size); - -/*=========================================================================== - -FUNCTION SG_FREE - -DESCRIPTION - Free the allocated buffer. - -DEPENDENCIES - None - -RETURN VALUE - None - -===========================================================================*/ -void sg_free(void* buffer); - -/*=========================================================================== - -FUNCTION SG_MD5 - -DESCRIPTION - Returns the MD5 hash of a the supplied buffer - -DEPENDENCIES - None - -RETURN VALUE - Hash value - -===========================================================================*/ -sg_status_type sg_md5(unsigned char* buffer, int size, - unsigned char** output, int* output_size, - sg_time_breakdown_type* time_break_down ); - - -/*=========================================================================== - -FUNCTION SG_MD5_OVERLAP - -DESCRIPTION - Returns the MD5 hash of each block for the provided buffer - -DEPENDENCIES - None - -RETURN VALUE - Hash value - -===========================================================================*/ -sg_status_type sg_md5_overlap(unsigned char* buffer, int size, - int block_size, int offset, - unsigned char** output, int* output_size, - sg_time_breakdown_type* time_breakdown); - - -/*=========================================================================== - -FUNCTION SG_SHA1 - -DESCRIPTION - Returns the SHA1 hash of a the supplied buffer - -DEPENDENCIES - None - -RETURN VALUE - Hash value - -===========================================================================*/ -sg_status_type sg_sha1(unsigned char* buffer, int size, - unsigned char** output, int* output_size, - sg_time_breakdown_type* time_breakdown); - - -/*=========================================================================== - -FUNCTION SG_SHA1_OVERLAP - -DESCRIPTION - Returns the SHA1 hash of each block for the provided buffer - -DEPENDENCIES - None - -RETURN VALUE - Hash value - -===========================================================================*/ -sg_status_type sg_sha1_overlap(unsigned char* buffer, int size, - int block_size, int offset, - unsigned char** output, int* output_size, - sg_time_breakdown_type* time_breakdown); -#endif /* STOREGPU_H */ -- cgit v1.3