summaryrefslogtreecommitdiff
path: root/benchmarks/CUDA/STO
diff options
context:
space:
mode:
Diffstat (limited to 'benchmarks/CUDA/STO')
-rw-r--r--benchmarks/CUDA/STO/Makefile52
-rw-r--r--benchmarks/CUDA/STO/README.GPGPU-Sim2
-rw-r--r--benchmarks/CUDA/STO/cust.h153
-rw-r--r--benchmarks/CUDA/STO/main.cu867
-rw-r--r--benchmarks/CUDA/STO/md5_cpu.c397
-rw-r--r--benchmarks/CUDA/STO/md5_cpu.h98
-rw-r--r--benchmarks/CUDA/STO/md5_kernel.cu995
-rw-r--r--benchmarks/CUDA/STO/sha1_cpu.c429
-rw-r--r--benchmarks/CUDA/STO/sha1_cpu.h98
-rw-r--r--benchmarks/CUDA/STO/sha1_kernel.cu1140
-rw-r--r--benchmarks/CUDA/STO/storeCPU.c1114
-rw-r--r--benchmarks/CUDA/STO/storeCPU.h195
-rw-r--r--benchmarks/CUDA/STO/storeGPU.cu1211
-rw-r--r--benchmarks/CUDA/STO/storeGPU.h224
14 files changed, 6975 insertions, 0 deletions
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 <stdlib.h>
+#include <stdio.h>
+#include <string.h>
+
+// project
+#include <cust.h>
+#include <storeGPU.h>
+#include <storeCPU.h>
+
+
+/*==========================================================================
+
+ 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 <string.h>
+#include <stdio.h>
+
+#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 <string.h>
+#include <stdio.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_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 <Name>
+
+DESCRIPTION
+ MD5 context setup
+
+DEPENDENCIES
+ <dep.>
+
+RETURN VALUE
+ <return>
+
+===========================================================================*/
+__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
+ <Desc.>
+
+DEPENDENCIES
+ <dep.>
+
+RETURN VALUE
+ <return>
+
+===========================================================================*/
+__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
+ <dep.>
+
+RETURN VALUE
+ <return>
+
+===========================================================================*/
+__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 ) {
+
+ //<ELSN>
+ /*memcpy( (void *) (ctx->buffer + left),
+ (void *) input, fill );*/
+ for (int i = 0; i < fill; i++) {
+ ctx->buffer[i+left] = input[i];
+ }
+ //</ELSN>
+
+ 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 ) {
+
+ //<ELSN>
+ /* memcpy( (void *) (ctx->buffer + left),
+ (void *) input, ilen );*/
+ for (int i = 0; i < ilen; i++) {
+ ctx->buffer[i+left] = input[i];
+ }
+ //</ELSN>
+
+ }
+}
+
+/*===========================================================================
+
+FUNCTION MD5_FINISH
+
+DESCRIPTION
+ MD5 final digest
+
+DEPENDENCIES
+ None.
+
+RETURN VALUE
+ <return>
+
+===========================================================================*/
+__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 <string.h>
+#include <stdio.h>
+
+#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 <string.h>
+#include <stdio.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_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<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 SHA1 algorithm). */
+ sharedMemory[SHARED_MEMORY_INDEX(13-i)] = 0x00000080;
+ //printf("the last one at %d\n",13-i);
+
+ PUT_UINT32_BE( chunkSize >> 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 <stdlib.h>
+#include <stdio.h>
+#include <string.h>
+
+
+#include "cust.h"
+#include "md5_cpu.h"
+#include "sha1_cpu.h"
+
+
+#ifdef FEATURE_WIN32_THREADS
+#include <windows.h>
+#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 <stdlib.h>
+#include <stdio.h>
+#include <string.h>
+
+// project
+#include <cutil.h>
+#include <cust.h>
+#include <md5_cpu.h>
+#include <sha1_cpu.h>
+#include <storeGPU.h>
+#include <storeCPU.h>
+
+// kernels
+#include <md5_kernel.cu>
+#include <sha1_kernel.cu>
+
+/*==========================================================================
+
+ 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 <cutil.h>
+
+/*==========================================================================
+
+ 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 */
generated by cgit v1.3 (git 2.53.0) at 2026-07-10 23:32:27 +0000