From 69f2911e04ffb1b19eef1fafb8c040af271f656e Mon Sep 17 00:00:00 2001 From: Tor Aamodt Date: Thu, 15 Jul 2010 18:09:46 -0800 Subject: creating branch for adding support for CUDA 3.x and Fermi [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 6829] --- benchmarks/CUDA/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 insertions(+) create mode 100644 benchmarks/CUDA/STO/Makefile create mode 100644 benchmarks/CUDA/STO/README.GPGPU-Sim create mode 100644 benchmarks/CUDA/STO/cust.h create mode 100644 benchmarks/CUDA/STO/main.cu create mode 100644 benchmarks/CUDA/STO/md5_cpu.c create mode 100644 benchmarks/CUDA/STO/md5_cpu.h create mode 100644 benchmarks/CUDA/STO/md5_kernel.cu create mode 100644 benchmarks/CUDA/STO/sha1_cpu.c create mode 100644 benchmarks/CUDA/STO/sha1_cpu.h create mode 100644 benchmarks/CUDA/STO/sha1_kernel.cu create mode 100644 benchmarks/CUDA/STO/storeCPU.c create mode 100644 benchmarks/CUDA/STO/storeCPU.h create mode 100644 benchmarks/CUDA/STO/storeGPU.cu create mode 100644 benchmarks/CUDA/STO/storeGPU.h (limited to 'benchmarks/CUDA/STO') diff --git a/benchmarks/CUDA/STO/Makefile b/benchmarks/CUDA/STO/Makefile new file mode 100644 index 0000000..44e53f3 --- /dev/null +++ b/benchmarks/CUDA/STO/Makefile @@ -0,0 +1,52 @@ +################################################################################ +# +# 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 new file mode 100644 index 0000000..6056878 --- /dev/null +++ b/benchmarks/CUDA/STO/README.GPGPU-Sim @@ -0,0 +1,2 @@ +make +./gpgpu_ptx_sim__storegpu diff --git a/benchmarks/CUDA/STO/cust.h b/benchmarks/CUDA/STO/cust.h new file mode 100644 index 0000000..92219de --- /dev/null +++ b/benchmarks/CUDA/STO/cust.h @@ -0,0 +1,153 @@ +#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 new file mode 100644 index 0000000..31a234d --- /dev/null +++ b/benchmarks/CUDA/STO/main.cu @@ -0,0 +1,867 @@ +/*========================================================================== + 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 new file mode 100644 index 0000000..44fbef2 --- /dev/null +++ b/benchmarks/CUDA/STO/md5_cpu.c @@ -0,0 +1,397 @@ +/*========================================================================== + 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 new file mode 100644 index 0000000..eabdf5e --- /dev/null +++ b/benchmarks/CUDA/STO/md5_cpu.h @@ -0,0 +1,98 @@ +#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 new file mode 100644 index 0000000..419daa5 --- /dev/null +++ b/benchmarks/CUDA/STO/md5_kernel.cu @@ -0,0 +1,995 @@ +/*========================================================================== + 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 new file mode 100644 index 0000000..f4a6be5 --- /dev/null +++ b/benchmarks/CUDA/STO/sha1_cpu.c @@ -0,0 +1,429 @@ +/*========================================================================== + 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 new file mode 100644 index 0000000..637376a --- /dev/null +++ b/benchmarks/CUDA/STO/sha1_cpu.h @@ -0,0 +1,98 @@ +#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 new file mode 100644 index 0000000..2b70fd2 --- /dev/null +++ b/benchmarks/CUDA/STO/sha1_kernel.cu @@ -0,0 +1,1140 @@ +/*========================================================================== + 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 new file mode 100644 index 0000000..09d1997 --- /dev/null +++ b/benchmarks/CUDA/STO/storeCPU.c @@ -0,0 +1,1114 @@ +/*========================================================================== + 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 new file mode 100644 index 0000000..6aada9b --- /dev/null +++ b/benchmarks/CUDA/STO/storeCPU.h @@ -0,0 +1,195 @@ +#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 new file mode 100644 index 0000000..70e23ab --- /dev/null +++ b/benchmarks/CUDA/STO/storeGPU.cu @@ -0,0 +1,1211 @@ +/*========================================================================== + 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 new file mode 100644 index 0000000..3d8bd0b --- /dev/null +++ b/benchmarks/CUDA/STO/storeGPU.h @@ -0,0 +1,224 @@ +#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