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, 0 insertions, 6975 deletions
diff --git a/benchmarks/CUDA/STO/Makefile b/benchmarks/CUDA/STO/Makefile
deleted file mode 100644
index 44e53f3..0000000
--- a/benchmarks/CUDA/STO/Makefile
+++ /dev/null
@@ -1,52 +0,0 @@
-################################################################################
-#
-# Copyright 1993-2006 NVIDIA Corporation. All rights reserved.
-#
-# NOTICE TO USER:
-#
-# This source code is subject to NVIDIA ownership rights under U.S. and
-# international Copyright laws.
-#
-# NVIDIA MAKES NO REPRESENTATION ABOUT THE SUITABILITY OF THIS SOURCE
-# CODE FOR ANY PURPOSE. IT IS PROVIDED "AS IS" WITHOUT EXPRESS OR
-# IMPLIED WARRANTY OF ANY KIND. NVIDIA DISCLAIMS ALL WARRANTIES WITH
-# REGARD TO THIS SOURCE CODE, INCLUDING ALL IMPLIED WARRANTIES OF
-# MERCHANTABILITY, NONINFRINGEMENT, AND FITNESS FOR A PARTICULAR PURPOSE.
-# IN NO EVENT SHALL NVIDIA BE LIABLE FOR ANY SPECIAL, INDIRECT, INCIDENTAL,
-# OR CONSEQUENTIAL DAMAGES, OR ANY DAMAGES WHATSOEVER RESULTING FROM LOSS
-# OF USE, DATA OR PROFITS, WHETHER IN AN ACTION OF CONTRACT, NEGLIGENCE
-# OR OTHER TORTIOUS ACTION, ARISING OUT OF OR IN CONNECTION WITH THE USE
-# OR PERFORMANCE OF THIS SOURCE CODE.
-#
-# U.S. Government End Users. This source code is a "commercial item" as
-# that term is defined at 48 C.F.R. 2.101 (OCT 1995), consisting of
-# "commercial computer software" and "commercial computer software
-# documentation" as such terms are used in 48 C.F.R. 12.212 (SEPT 1995)
-# and is provided to the U.S. Government only as a commercial end item.
-# Consistent with 48 C.F.R.12.212 and 48 C.F.R. 227.7202-1 through
-# 227.7202-4 (JUNE 1995), all U.S. Government End Users acquire the
-# source code with only those rights set forth herein.
-#
-################################################################################
-#
-# Build script for project
-#
-################################################################################
-
-# Add source files here
-EXECUTABLE := storegpu
-# CUDA source files (compiled with cudacc)
-CUFILES := main.cu storeGPU.cu
-# CUDA dependency files
-CU_DEPS := \
-md5_kernel.cu sha1_kernel.cu
-# C/C++ source files (compiled with gcc / c++)
-CCFILES := \
- storeCPU.c md5_cpu.c sha1_cpu.c
-#CUDACCFLAGS := -po maxrregcount=16
-
-################################################################################
-# Rules and targets
-
-GPGPUSIM_ROOT := ../../..
-include ../../../common/common.mk
diff --git a/benchmarks/CUDA/STO/README.GPGPU-Sim b/benchmarks/CUDA/STO/README.GPGPU-Sim
deleted file mode 100644
index 6056878..0000000
--- a/benchmarks/CUDA/STO/README.GPGPU-Sim
+++ /dev/null
@@ -1,2 +0,0 @@
-make
-./gpgpu_ptx_sim__storegpu
diff --git a/benchmarks/CUDA/STO/cust.h b/benchmarks/CUDA/STO/cust.h
deleted file mode 100644
index 92219de..0000000
--- a/benchmarks/CUDA/STO/cust.h
+++ /dev/null
@@ -1,153 +0,0 @@
-#ifndef CUST_H
-#define CUST_H
-/*===========================================================================
-
- CUSTOMIZATION
-
-* Copyright (c) 2008, NetSysLab at the University of British Columbia
-* All rights reserved.
-*
-* Redistribution and use in source and binary forms, with or without
-* modification, are permitted provided that the following conditions are met:
-* * Redistributions of source code must retain the above copyright
-* notice, this list of conditions and the following disclaimer.
-* * Redistributions in binary form must reproduce the above copyright
-* notice, this list of conditions and the following disclaimer in the
-* documentation and/or other materials provided with the distribution.
-* * Neither the name of the University nor the
-* names of its contributors may be used to endorse or promote products
-* derived from this software without specific prior written permission.
-*
-* THIS SOFTWARE IS PROVIDED BY NetSysLab ``AS IS'' AND ANY
-* EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
-* WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
-* DISCLAIMED. IN NO EVENT SHALL NetSysLab BE LIABLE FOR ANY
-* DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
-* (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
-* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
-* ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
-* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
-* SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
-
-DESCRIPTION
- Configuration file.
-
-===========================================================================*/
-
-
-/*===========================================================================
-
- PUBLIC DATA DECLARATIONS
-
-===========================================================================*/
-
-/* NOTE: Whenever you change anything in this file, you have to rebuild
- the project (Build->RebuildSolution) to take the effects. running directly
- using Debug->RunWithoutDebugging will not recompile the solution with the
- new parameters. */
-
-
-/* ------------------------------------------------------------------
-** GENERAL FEATURES
-** ------------------------------------------------------------------ */
-
-// Use shared memory implementation. Note that the total size of the test
-// must not exceed 96MB. This will be debugged later on.
-#define FEATURE_SHARED_MEMORY
-
-// Use pinned memory pages.
-#define FEATURE_PINNED_MODE
-
-// Hash size returned is trimmed to save bandwidth in the copy back.
-// Check md5.h for the actual hash size returned.
-#define FEATURE_REDUCED_HASH_SIZE
-
-// This feature guesses at run time the best execution context the GPU will
-// run within
-//#define FEATURE_DYNAMIC_EXEC_CONTEXT
-
-#ifdef FEATURE_DYNAMIC_EXEC_CONTEXT
-// when in dynamic execution context setting, this feature will enable
-// the algorithm that maximizes the number of threads rather than chunk size.
-//#define FEATURE_MAXIMIZE_NUM_OF_THREADS
-#endif
-
-// Tests to run
-// To run the overlap test instead of the original one.
-#define FEATURE_RUN_OVERLAP_TEST
-
-// Turn on to run SHA1, otherwise it will run MD5
-#define FEATURE_RUN_SHA1
-
-// Enable the multi-thread version for the CPU functions on windows
-//#define FEATURE_WIN32_THREADS
-
-// Enable the multi-thread version for the CPU functions on Linux (with pthreads)
-//#define FEATURE_PTHREADS
-
-/* ------------------------------------------------------------------
-** SANITY CHECKS
-** ------------------------------------------------------------------ */
-
-
-
-/* ------------------------------------------------------------------
-** CONSTANTS
-** ------------------------------------------------------------------ */
-#ifdef FEATURE_DYNAMIC_EXEC_CONTEXT
- #define MAX_THREADS_PER_BLOCK 192
- #define MAX_BLOCKS_PER_GRID (32 * 1024)
- #define MAX_NUM_OF_THREADS (MAX_THREADS_PER_BLOCK*MAX_BLOCKS_PER_GRID)
- #define BASIC_CHUNK_SIZE 64
- #define MAX_CHUNK_SIZE 2048
- #define NUM_OF_MULTIPROCESSORS 4
-#endif /* FEATURE_DYNAMIC_EXEC_CONTEXT */
-
-#ifdef FEATURE_REDUCED_HASH_SIZE
- #define MD5_HASH_SIZE 4
- #define SHA1_HASH_SIZE 4
-#else
- #define MD5_HASH_SIZE 16
- #define SHA1_HASH_SIZE 20
-#endif // FEATURE_REDUCED_HASH_SIZE
-
-
-#ifdef FEATURE_RUN_OVERLAP_TEST
-
-// THREADS_PER_BLOCK x CHUNK_SIZE should be < 16K if FEATURE_SHARED_MEMORY
- #define THREADS_PER_BLOCK 128
- #define BLOCKS_PER_GRID 384//32 //1024//16//384
- #define TOTAL_NUM_OF_THREADS ( THREADS_PER_BLOCK * BLOCKS_PER_GRID )
-
-// Chunk size must be multiple of 4.
- #define CHUNK_SIZE (52)
- //#define CHUNK_SIZE (1024-12)
-
-// Offset must be multiple of 4.
- #define OFFSET (4)
-
-// Size to be allocated for the original interface tests
- #define TEST_MEM_SIZE_OVERLAP (TOTAL_NUM_OF_THREADS * OFFSET + \
- CHUNK_SIZE - OFFSET)
-
-#else
-
- #define THREADS_PER_BLOCK 192
- #define BLOCKS_PER_GRID 512
- #define TOTAL_NUM_OF_THREADS ( THREADS_PER_BLOCK * BLOCKS_PER_GRID )
-
-// Currently chunk size must be of the form (64*i - 12). 12 is reserved for
-// padding issues.
- //#define CHUNK_SIZE (64-12)
- //#define CHUNK_SIZE (128-12)
- //#define CHUNK_SIZE (256-12)
- //#define CHUNK_SIZE (512-12)
- #define CHUNK_SIZE (1024-12)
- //#define CHUNK_SIZE (2048-12)
-
-// Size to be allocated for the tests
- #define TEST_MEM_SIZE (CHUNK_SIZE * TOTAL_NUM_OF_THREADS)
-
-#endif /* FEATURE_RUN_OVERLAP_TEST */
-
-#endif /* CUST_H */
diff --git a/benchmarks/CUDA/STO/main.cu b/benchmarks/CUDA/STO/main.cu
deleted file mode 100644
index 31a234d..0000000
--- a/benchmarks/CUDA/STO/main.cu
+++ /dev/null
@@ -1,867 +0,0 @@
-/*==========================================================================
- M A I N
-
-* Copyright (c) 2008, NetSysLab at the University of British Columbia
-* All rights reserved.
-*
-* Redistribution and use in source and binary forms, with or without
-* modification, are permitted provided that the following conditions are met:
-* * Redistributions of source code must retain the above copyright
-* notice, this list of conditions and the following disclaimer.
-* * Redistributions in binary form must reproduce the above copyright
-* notice, this list of conditions and the following disclaimer in the
-* documentation and/or other materials provided with the distribution.
-* * Neither the name of the University nor the
-* names of its contributors may be used to endorse or promote products
-* derived from this software without specific prior written permission.
-*
-* THIS SOFTWARE IS PROVIDED BY NetSysLab ``AS IS'' AND ANY
-* EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
-* WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
-* DISCLAIMED. IN NO EVENT SHALL NetSysLab BE LIABLE FOR ANY
-* DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
-* (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
-* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
-* ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
-* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
-* SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
-
-DESCRIPTION
- Main entry.
-
-
-==========================================================================*/
-
-/*==========================================================================
-
- INCLUDES
-
-==========================================================================*/
-// system
-#include <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
deleted file mode 100644
index 44fbef2..0000000
--- a/benchmarks/CUDA/STO/md5_cpu.c
+++ /dev/null
@@ -1,397 +0,0 @@
-/*==========================================================================
- M D 5 C P U
-
- CPU implementation of the md5 algorithm.
-
- *
- * FIPS-180-1 compliant SHA-1 implementation
- *
- * Copyright (C) 2006-2007 Christophe Devine
- *
- * This library is free software; you can redistribute it and/or
- * modify it under the terms of the GNU Lesser General Public
- * License, version 2.1 as published by the Free Software Foundation.
- *
- * This library is distributed in the hope that it will be useful,
- * but WITHOUT ANY WARRANTY; without even the implied warranty of
- * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU
- * Lesser General Public License for more details.
- *
- * You should have received a copy of the GNU Lesser General Public
- * License along with this library; if not, write to the Free Software
- * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston,
- * MA 02110-1301 USA
- *
- * The SHA-1 standard was published by NIST in 1993.
- *
- * http://www.itl.nist.gov/fipspubs/fip180-1.htm
-
-DESCRIPTION
- CPU implementation of the md5 algorithm.
-
-
-==========================================================================*/
-
-/*==========================================================================
-
- INCLUDES
-
-==========================================================================*/
-#include <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
deleted file mode 100644
index eabdf5e..0000000
--- a/benchmarks/CUDA/STO/md5_cpu.h
+++ /dev/null
@@ -1,98 +0,0 @@
-#ifndef MD5_CPU_H
-#define MD5_CPU_H
-
-/*==========================================================================
- M D 5 C P U
-
-* Copyright (c) 2008, NetSysLab at the University of British Columbia
-* All rights reserved.
-*
-* Redistribution and use in source and binary forms, with or without
-* modification, are permitted provided that the following conditions are met:
-* * Redistributions of source code must retain the above copyright
-* notice, this list of conditions and the following disclaimer.
-* * Redistributions in binary form must reproduce the above copyright
-* notice, this list of conditions and the following disclaimer in the
-* documentation and/or other materials provided with the distribution.
-* * Neither the name of the University nor the
-* names of its contributors may be used to endorse or promote products
-* derived from this software without specific prior written permission.
-*
-* THIS SOFTWARE IS PROVIDED BY NetSysLab ``AS IS'' AND ANY
-* EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
-* WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
-* DISCLAIMED. IN NO EVENT SHALL NetSysLab BE LIABLE FOR ANY
-* DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
-* (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
-* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
-* ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
-* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
-* SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
-
-DESCRIPTION
- CPU implementation of the md5 algorithm.
-
-
-==========================================================================*/
-
-/*==========================================================================
-
- INCLUDES
-
-==========================================================================*/
-
-
-/*==========================================================================
-
- DATA DECLARATIONS
-
-==========================================================================*/
-
-/*--------------------------------------------------------------------------
- TYPE DEFINITIONS
---------------------------------------------------------------------------*/
-
-
-/*--------------------------------------------------------------------------
- FUNCTION PROTOTYPES
---------------------------------------------------------------------------*/
-
-/*--------------------------------------------------------------------------
- CONSTANTS
---------------------------------------------------------------------------*/
-
-/*--------------------------------------------------------------------------
- GLOBAL VARIABLES
---------------------------------------------------------------------------*/
-
-
-/*--------------------------------------------------------------------------
- MACROS
---------------------------------------------------------------------------*/
-
-/*==========================================================================
-
- FUNCTIONS
-
-==========================================================================*/
-
-/*===========================================================================
-
-FUNCTION cpu_md5_internal
-
-DESCRIPTION
- CPU implementation of the MD5 algorithm
-
-DEPENDENCIES
- None
-
-RETURN VALUE
- Hash
-
-===========================================================================*/
-void md5_cpu_internal( unsigned char *input, int ilen,
- unsigned char *output );
-
-
-
-#endif /* MD5_CPU_H */
diff --git a/benchmarks/CUDA/STO/md5_kernel.cu b/benchmarks/CUDA/STO/md5_kernel.cu
deleted file mode 100644
index 419daa5..0000000
--- a/benchmarks/CUDA/STO/md5_kernel.cu
+++ /dev/null
@@ -1,995 +0,0 @@
-/*==========================================================================
- MD5 KERNEL
-
-* Copyright (c) 2008, NetSysLab at the University of British Columbia
-* All rights reserved.
-*
-* Redistribution and use in source and binary forms, with or without
-* modification, are permitted provided that the following conditions are met:
-* * Redistributions of source code must retain the above copyright
-* notice, this list of conditions and the following disclaimer.
-* * Redistributions in binary form must reproduce the above copyright
-* notice, this list of conditions and the following disclaimer in the
-* documentation and/or other materials provided with the distribution.
-* * Neither the name of the University nor the
-* names of its contributors may be used to endorse or promote products
-* derived from this software without specific prior written permission.
-*
-* THIS SOFTWARE IS PROVIDED BY NetSysLab ``AS IS'' AND ANY
-* EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
-* WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
-* DISCLAIMED. IN NO EVENT SHALL NetSysLab BE LIABLE FOR ANY
-* DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
-* (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
-* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
-* ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
-* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
-* SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
-
-DESCRIPTION
- CPU version of the storeGPU library.
-
-
-==========================================================================*/
-
-/*==========================================================================
-
- INCLUDES
-
-==========================================================================*/
-#include <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
deleted file mode 100644
index f4a6be5..0000000
--- a/benchmarks/CUDA/STO/sha1_cpu.c
+++ /dev/null
@@ -1,429 +0,0 @@
-/*==========================================================================
- S H A 1 C P U
-
-DESCRIPTION
- CPU implementation of the sha1 algorithm.
-
- *
- * FIPS-180-1 compliant SHA-1 implementation
- *
- * Copyright (C) 2006-2007 Christophe Devine
- *
- * This library is free software; you can redistribute it and/or
- * modify it under the terms of the GNU Lesser General Public
- * License, version 2.1 as published by the Free Software Foundation.
- *
- * This library is distributed in the hope that it will be useful,
- * but WITHOUT ANY WARRANTY; without even the implied warranty of
- * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU
- * Lesser General Public License for more details.
- *
- * You should have received a copy of the GNU Lesser General Public
- * License along with this library; if not, write to the Free Software
- * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston,
- * MA 02110-1301 USA
- *
- * The SHA-1 standard was published by NIST in 1993.
- *
- * http://www.itl.nist.gov/fipspubs/fip180-1.htm
- *
-==========================================================================*/
-
-/*==========================================================================
-
- INCLUDES
-
-==========================================================================*/
-#include <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
deleted file mode 100644
index 637376a..0000000
--- a/benchmarks/CUDA/STO/sha1_cpu.h
+++ /dev/null
@@ -1,98 +0,0 @@
-#ifndef SHA1_CPU_H
-#define SHA1_CPU_H
-
-/*==========================================================================
- S H A 1 C P U
-
-* Copyright (c) 2008, NetSysLab at the University of British Columbia
-* All rights reserved.
-*
-* Redistribution and use in source and binary forms, with or without
-* modification, are permitted provided that the following conditions are met:
-* * Redistributions of source code must retain the above copyright
-* notice, this list of conditions and the following disclaimer.
-* * Redistributions in binary form must reproduce the above copyright
-* notice, this list of conditions and the following disclaimer in the
-* documentation and/or other materials provided with the distribution.
-* * Neither the name of the University nor the
-* names of its contributors may be used to endorse or promote products
-* derived from this software without specific prior written permission.
-*
-* THIS SOFTWARE IS PROVIDED BY NetSysLab ``AS IS'' AND ANY
-* EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
-* WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
-* DISCLAIMED. IN NO EVENT SHALL NetSysLab BE LIABLE FOR ANY
-* DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
-* (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
-* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
-* ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
-* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
-* SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
-
-DESCRIPTION
- CPU implementation of the sha1 algorithm.
-
-
-==========================================================================*/
-
-/*==========================================================================
-
- INCLUDES
-
-==========================================================================*/
-
-
-/*==========================================================================
-
- DATA DECLARATIONS
-
-==========================================================================*/
-
-/*--------------------------------------------------------------------------
- TYPE DEFINITIONS
---------------------------------------------------------------------------*/
-
-
-/*--------------------------------------------------------------------------
- FUNCTION PROTOTYPES
---------------------------------------------------------------------------*/
-
-/*--------------------------------------------------------------------------
- CONSTANTS
---------------------------------------------------------------------------*/
-
-/*--------------------------------------------------------------------------
- GLOBAL VARIABLES
---------------------------------------------------------------------------*/
-
-
-/*--------------------------------------------------------------------------
- MACROS
---------------------------------------------------------------------------*/
-
-/*==========================================================================
-
- FUNCTIONS
-
-==========================================================================*/
-
-/*===========================================================================
-
-FUNCTION cpu_sha1_internal
-
-DESCRIPTION
- CPU implementation of the SHA1 algorithm
-
-DEPENDENCIES
- None
-
-RETURN VALUE
- Hash
-
-===========================================================================*/
-void sha1_cpu_internal( unsigned char *input, int ilen,
- unsigned char *output );
-
-
-
-#endif /* SHA1_CPU_H */
diff --git a/benchmarks/CUDA/STO/sha1_kernel.cu b/benchmarks/CUDA/STO/sha1_kernel.cu
deleted file mode 100644
index 2b70fd2..0000000
--- a/benchmarks/CUDA/STO/sha1_kernel.cu
+++ /dev/null
@@ -1,1140 +0,0 @@
-/*==========================================================================
- SHA1 KERNEL
-
-* Copyright (c) 2008, NetSysLab at the University of British Columbia
-* All rights reserved.
-*
-* Redistribution and use in source and binary forms, with or without
-* modification, are permitted provided that the following conditions are met:
-* * Redistributions of source code must retain the above copyright
-* notice, this list of conditions and the following disclaimer.
-* * Redistributions in binary form must reproduce the above copyright
-* notice, this list of conditions and the following disclaimer in the
-* documentation and/or other materials provided with the distribution.
-* * Neither the name of the University nor the
-* names of its contributors may be used to endorse or promote products
-* derived from this software without specific prior written permission.
-*
-* THIS SOFTWARE IS PROVIDED BY NetSysLab ``AS IS'' AND ANY
-* EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
-* WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
-* DISCLAIMED. IN NO EVENT SHALL NetSysLab BE LIABLE FOR ANY
-* DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
-* (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
-* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
-* ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
-* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
-* SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
-
-DESCRIPTION
- CPU version of the storeGPU library.
-
-
-==========================================================================*/
-
-/*==========================================================================
-
- INCLUDES
-
-==========================================================================*/
-#include <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
deleted file mode 100644
index 09d1997..0000000
--- a/benchmarks/CUDA/STO/storeCPU.c
+++ /dev/null
@@ -1,1114 +0,0 @@
-/*==========================================================================
- S T O R E C P U
-
-* Copyright (c) 2008, NetSysLab at the University of British Columbia
-* All rights reserved.
-*
-* Redistribution and use in source and binary forms, with or without
-* modification, are permitted provided that the following conditions are met:
-* * Redistributions of source code must retain the above copyright
-* notice, this list of conditions and the following disclaimer.
-* * Redistributions in binary form must reproduce the above copyright
-* notice, this list of conditions and the following disclaimer in the
-* documentation and/or other materials provided with the distribution.
-* * Neither the name of the University nor the
-* names of its contributors may be used to endorse or promote products
-* derived from this software without specific prior written permission.
-*
-* THIS SOFTWARE IS PROVIDED BY NetSysLab ``AS IS'' AND ANY
-* EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
-* WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
-* DISCLAIMED. IN NO EVENT SHALL NetSysLab BE LIABLE FOR ANY
-* DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
-* (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
-* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
-* ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
-* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
-* SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
-
-
-DESCRIPTION
- CPU version of the storeGPU library.
-
-
-==========================================================================*/
-
-/*==========================================================================
-
- INCLUDES
-
-==========================================================================*/
-#include <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
deleted file mode 100644
index 6aada9b..0000000
--- a/benchmarks/CUDA/STO/storeCPU.h
+++ /dev/null
@@ -1,195 +0,0 @@
-#ifndef STORECPU_H
-#define STORECPU_H
-/*==========================================================================
- S T O R E C P U
-
-* Copyright (c) 2008, NetSysLab at the University of British Columbia
-* All rights reserved.
-*
-* Redistribution and use in source and binary forms, with or without
-* modification, are permitted provided that the following conditions are met:
-* * Redistributions of source code must retain the above copyright
-* notice, this list of conditions and the following disclaimer.
-* * Redistributions in binary form must reproduce the above copyright
-* notice, this list of conditions and the following disclaimer in the
-* documentation and/or other materials provided with the distribution.
-* * Neither the name of the University nor the
-* names of its contributors may be used to endorse or promote products
-* derived from this software without specific prior written permission.
-*
-* THIS SOFTWARE IS PROVIDED BY NetSysLab ``AS IS'' AND ANY
-* EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
-* WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
-* DISCLAIMED. IN NO EVENT SHALL NetSysLab BE LIABLE FOR ANY
-* DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
-* (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
-* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
-* ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
-* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
-* SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
-
-DESCRIPTION
- CPU version of the storeGPU library.
-
-
-==========================================================================*/
-
-/*==========================================================================
-
- INCLUDES
-
-==========================================================================*/
-
-
-
-/*==========================================================================
-
- DATA DECLARATIONS
-
-==========================================================================*/
-
-/*--------------------------------------------------------------------------
- TYPE DEFINITIONS
---------------------------------------------------------------------------*/
-
-
-/*--------------------------------------------------------------------------
- FUNCTION PROTOTYPES
---------------------------------------------------------------------------*/
-
-/*--------------------------------------------------------------------------
- CONSTANTS
---------------------------------------------------------------------------*/
-
-/*--------------------------------------------------------------------------
- GLOBAL VARIABLES
---------------------------------------------------------------------------*/
-
-
-/*--------------------------------------------------------------------------
- MACROS
---------------------------------------------------------------------------*/
-
-/*==========================================================================
-
- FUNCTIONS
-
-==========================================================================*/
-
-/*--------------------------------------------------------------------------
- LOCAL FUNCTIONS
---------------------------------------------------------------------------*/
-
-
-
-/*--------------------------------------------------------------------------
- GLOBAL FUNCTIONS
---------------------------------------------------------------------------*/
-/*===========================================================================
-
-FUNCTION SC_MD5
-
-DESCRIPTION
- CPU version of the MD5 algorithm
-
-DEPENDENCIES
- None
-
-RETURN VALUE
- Hash
-
-===========================================================================*/
-//extern "C"
-void sc_md5( unsigned char* buffer, int size,
- unsigned char** output, int* output_size);
-
-
-/*===========================================================================
-
-FUNCTION SC_MD5_OVERLAP
-
-DESCRIPTION
- CPU version of the MD5 overlap algorithm
-
-DEPENDENCIES
- None
-
-RETURN VALUE
- Hash
-
-===========================================================================*/
-//extern "C"
-void sc_md5_overlap(unsigned char* buffer, int size, int block_size, int offset,
- unsigned char** output, int* output_size);
-
-/*===========================================================================
-
-FUNCTION SC_SHA1
-
-DESCRIPTION
- CPU version of the SHA1 algorithm
-
-DEPENDENCIES
- None
-
-RETURN VALUE
- Hash
-
-===========================================================================*/
-//extern "C"
-void sc_sha1( unsigned char* buffer, int size,
- unsigned char** output, int* output_size);
-
-/*===========================================================================
-
-FUNCTION SC_SHA1_OVERLAP
-
-DESCRIPTION
- CPU version of the SHA1 overlap algorithm
-
-DEPENDENCIES
- None
-
-RETURN VALUE
- Hash
-
-===========================================================================*/
-//extern "C"
-void sc_sha1_overlap(unsigned char* buffer, int size, int block_size,
- int offset, unsigned char** output, int* output_size);
-
-/*===========================================================================
-
-FUNCTION SC_MD5_STANDARD
-
-DESCRIPTION
- The standard MD5 algorithm
-
-DEPENDENCIES
- None
-
-RETURN VALUE
- Hash
-
-===========================================================================*/
-//extern "C"
-void sc_md5_standard( unsigned char* buffer, int size, unsigned char** output);
-
-/*===========================================================================
-
-FUNCTION SC_SHA1_STANDARD
-
-DESCRIPTION
- The standard SHA1 algorithm
-
-DEPENDENCIES
- None
-
-RETURN VALUE
- Hash
-
-===========================================================================*/
-//extern "C"
-void sc_sha1_standard( unsigned char* buffer, int size, unsigned char** output);
-
-#endif /* STORECPU_H */
diff --git a/benchmarks/CUDA/STO/storeGPU.cu b/benchmarks/CUDA/STO/storeGPU.cu
deleted file mode 100644
index 70e23ab..0000000
--- a/benchmarks/CUDA/STO/storeGPU.cu
+++ /dev/null
@@ -1,1211 +0,0 @@
-/*==========================================================================
- S T O R E G P U
-
-* Copyright (c) 2008, NetSysLab at the University of British Columbia
-* All rights reserved.
-*
-* Redistribution and use in source and binary forms, with or without
-* modification, are permitted provided that the following conditions are met:
-* * Redistributions of source code must retain the above copyright
-* notice, this list of conditions and the following disclaimer.
-* * Redistributions in binary form must reproduce the above copyright
-* notice, this list of conditions and the following disclaimer in the
-* documentation and/or other materials provided with the distribution.
-* * Neither the name of the University nor the
-* names of its contributors may be used to endorse or promote products
-* derived from this software without specific prior written permission.
-*
-* THIS SOFTWARE IS PROVIDED BY NetSysLab ``AS IS'' AND ANY
-* EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
-* WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
-* DISCLAIMED. IN NO EVENT SHALL NetSysLab BE LIABLE FOR ANY
-* DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
-* (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
-* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
-* ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
-* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
-* SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
-
-DESCRIPTION
- Main entry of the library.
-
-
-==========================================================================*/
-
-/*==========================================================================
-
- INCLUDES
-
-==========================================================================*/
-// system
-#include <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
deleted file mode 100644
index 3d8bd0b..0000000
--- a/benchmarks/CUDA/STO/storeGPU.h
+++ /dev/null
@@ -1,224 +0,0 @@
-#ifndef STOREGPU_H
-#define STOREGPU_H
-/*==========================================================================
- S T O R E G P U
-
-* Copyright (c) 2008, NetSysLab at the University of British Columbia
-* All rights reserved.
-*
-* Redistribution and use in source and binary forms, with or without
-* modification, are permitted provided that the following conditions are met:
-* * Redistributions of source code must retain the above copyright
-* notice, this list of conditions and the following disclaimer.
-* * Redistributions in binary form must reproduce the above copyright
-* notice, this list of conditions and the following disclaimer in the
-* documentation and/or other materials provided with the distribution.
-* * Neither the name of the University nor the
-* names of its contributors may be used to endorse or promote products
-* derived from this software without specific prior written permission.
-*
-* THIS SOFTWARE IS PROVIDED BY NetSysLab ``AS IS'' AND ANY
-* EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
-* WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
-* DISCLAIMED. IN NO EVENT SHALL NetSysLab BE LIABLE FOR ANY
-* DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
-* (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
-* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
-* ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
-* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
-* SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
-
-DESCRIPTION
- StoreGPU library.
-
-==========================================================================*/
-
-/*==========================================================================
-
- INCLUDES
-
-==========================================================================*/
-#include <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-11 00:41:41 +0000