From 69f2911e04ffb1b19eef1fafb8c040af271f656e Mon Sep 17 00:00:00 2001 From: Tor Aamodt Date: Thu, 15 Jul 2010 18:09:46 -0800 Subject: creating branch for adding support for CUDA 3.x and Fermi [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 6829] --- benchmarks/CUDA/STO/storeGPU.cu | 1211 +++++++++++++++++++++++++++++++++++++++ 1 file changed, 1211 insertions(+) create mode 100644 benchmarks/CUDA/STO/storeGPU.cu (limited to 'benchmarks/CUDA/STO/storeGPU.cu') diff --git a/benchmarks/CUDA/STO/storeGPU.cu b/benchmarks/CUDA/STO/storeGPU.cu new file mode 100644 index 0000000..70e23ab --- /dev/null +++ b/benchmarks/CUDA/STO/storeGPU.cu @@ -0,0 +1,1211 @@ +/*========================================================================== + S T O R E G P U + +* Copyright (c) 2008, NetSysLab at the University of British Columbia +* All rights reserved. +* +* Redistribution and use in source and binary forms, with or without +* modification, are permitted provided that the following conditions are met: +* * Redistributions of source code must retain the above copyright +* notice, this list of conditions and the following disclaimer. +* * Redistributions in binary form must reproduce the above copyright +* notice, this list of conditions and the following disclaimer in the +* documentation and/or other materials provided with the distribution. +* * Neither the name of the University nor the +* names of its contributors may be used to endorse or promote products +* derived from this software without specific prior written permission. +* +* THIS SOFTWARE IS PROVIDED BY NetSysLab ``AS IS'' AND ANY +* EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED +* WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE +* DISCLAIMED. IN NO EVENT SHALL NetSysLab BE LIABLE FOR ANY +* DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES +* (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; +* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND +* ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS +* SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + +DESCRIPTION + Main entry of the library. + + +==========================================================================*/ + +/*========================================================================== + + INCLUDES + +==========================================================================*/ +// system +#include +#include +#include + +// project +#include +#include +#include +#include +#include +#include + +// kernels +#include +#include + +/*========================================================================== + + DATA DECLARATIONS + +==========================================================================*/ + +/*-------------------------------------------------------------------------- + TYPE DEFINITIONS +--------------------------------------------------------------------------*/ + +// defines a GPU device properties +typedef struct sg_dev_prop { + int max_thread_per_block; + int max_grid_size; + int global_mem_size; + int warp_size; +} sg_dev_prop_type; + +// defines an execution context used to lunch a kernel. +typedef struct sg_exec_context { + int threads_per_block; + int blocks_per_grid; + int total_threads; + int chunk_size; + int total_size; + int pad_size; +} sg_exec_context_type; + + +/*-------------------------------------------------------------------------- + FUNCTION PROTOTYPES +--------------------------------------------------------------------------*/ + +/*-------------------------------------------------------------------------- + CONSTANTS +--------------------------------------------------------------------------*/ + + + +/*-------------------------------------------------------------------------- + GLOBAL VARIABLES +--------------------------------------------------------------------------*/ + + +/*-------------------------------------------------------------------------- + MACROS +--------------------------------------------------------------------------*/ + +#define GET_REAL_CHUNK_SIZE(chunk_size) ((chunk_size) - 12) + +/*========================================================================== + + FUNCTIONS + +==========================================================================*/ + +/*-------------------------------------------------------------------------- + LOCAL FUNCTIONS +--------------------------------------------------------------------------*/ +#ifdef FEATURE_DYNAMIC_EXEC_CONTEXT +/*=========================================================================== + +FUNCTION SG_GET_DEV_PROP + +DESCRIPTION + Probes the device for its properties + +DEPENDENCIES + None + +RETURN VALUE + device information + +===========================================================================*/ +static void sg_get_dev_prop(sg_dev_prop_type* dev_prop) { + + struct cudaDeviceProp prop; + int dev; + + cudaGetDevice(&dev); + cudaGetDeviceProperties(&prop, dev); + + printf("\n== Device Properties ==\n"); + printf("Max global memory : %d\n", prop.totalGlobalMem); + printf("Registers per block : %d\n", prop.regsPerBlock); + printf("Warp size : %d\n", prop.warpSize); + printf("Max threads per block: %d\n", prop.maxThreadsPerBlock); + printf("Block Dimensions : %d, %d, %d\n", + prop.maxThreadsDim[0], + prop.maxThreadsDim[1], + prop.maxThreadsDim[2]); + printf("Grid Dimensions : %d, %d, %d\n", + prop.maxGridSize[0], + prop.maxGridSize[1], + prop.maxGridSize[2]); + + dev_prop->max_thread_per_block = prop.maxThreadsDim[0]; + dev_prop->max_grid_size = prop.maxGridSize[0]; + dev_prop->global_mem_size = prop.totalGlobalMem; + dev_prop->warp_size = prop.warpSize; + +} + +#ifdef FEATURE_MAXIMIZE_NUM_OF_THREADS +/*=========================================================================== + +FUNCTION SG_GET_EXEC_CONTEXT + +DESCRIPTION + sets the required chunk size, thread per block and number of blocks + needed for kernel execution according to client buffer size. + +DEPENDENCIES + None + +RETURN VALUE + execution context + +===========================================================================*/ +static sg_status_type sg_get_exec_context(int size, int hash_size, + sg_exec_context_type* ctx){ + + sg_dev_prop_type dev_prop; + int threads_per_block; + int blocks_per_grid; + int total_threads; + int chunk_size; + int pad_size; + + int total_chunks = 0; + int found = 0; + int index = 1; + + + //**** Get device information ****// + sg_get_dev_prop(&dev_prop); + + //**** Determine the execution context ****// + /* The algorithm will try to determine the context by minimizing chunk + * size and maximizing total number of threads + * TODO: May be we can do better here + */ + while ( !found ) { + // Set chunk size + chunk_size = GET_REAL_CHUNK_SIZE(BASIC_CHUNK_SIZE * index); + + if ( chunk_size > MAX_CHUNK_SIZE ) + break; + + // Calculate the required padding for this chunk size + pad_size = ((size % chunk_size) == 0) ? 0 : + chunk_size - (size % chunk_size); + + // total number of chunks required if we are going to use this chunk size + total_chunks = (pad_size == 0) ? size / chunk_size : + (size / chunk_size) + 1; + + if ( total_chunks <= MAX_NUM_OF_THREADS ) { + // Got it, this is the minimum chunk size we can use. Now determine the + // threads and blocks numbers. + total_threads = total_chunks; + + // Get block and grid sizes + if (total_chunks <= MAX_THREADS_PER_BLOCK ) { + threads_per_block = total_chunks; + blocks_per_grid = 1; + + } else { + threads_per_block = MAX_THREADS_PER_BLOCK; + blocks_per_grid = ((total_threads % threads_per_block) == 0) ? + (total_threads/threads_per_block) : + (total_threads/threads_per_block) + 1; + + } + + if ( dev_prop.global_mem_size > + (total_threads * (chunk_size + hash_size) + hash_size) ) + found = 1; + + } + index++; + + } + + //**** Did we find a solution? ****// + if ( !found ) + return SG_ERR_DEV_MEM_OVERFLOW; + + //**** Fill the struct with the solution ****// + ctx->threads_per_block = threads_per_block; + ctx->blocks_per_grid = blocks_per_grid; + ctx->total_threads = total_threads; + ctx->total_size = size + pad_size; + ctx->chunk_size = chunk_size; + ctx->pad_size = pad_size; + + return SG_OK; + +} + +#else /* FEATURE_MAXIMIZE_NUM_OF_THREADS */ +/*=========================================================================== + +FUNCTION SG_GET_EXEC_CONTEXT + +DESCRIPTION + sets the required chunk size, thread per block and number of blocks + needed for kernel execution according to client buffer size. + +DEPENDENCIES + None + +RETURN VALUE + execution context + +===========================================================================*/ +static sg_status_type sg_get_exec_context(int size, int hash_size, + sg_exec_context_type* ctx){ + + sg_dev_prop_type dev_prop; + int threads_per_block; + int blocks_per_grid; + int total_threads; + int chunk_size; + int pad_size; + + int total_chunks = 0; + int found = 0; + + int index = MAX_CHUNK_SIZE / BASIC_CHUNK_SIZE; + + + //**** Get device information ****// + sg_get_dev_prop(&dev_prop); + + //**** Determine the execution context ****// + /* The algorithm will try to determine the context by minimizing chunk + * size and maximizing total number of threads + * TODO: May be we can do better here + */ + while ( 1 ) { + // Set chunk size + chunk_size = GET_REAL_CHUNK_SIZE(BASIC_CHUNK_SIZE * index); + + // don't go less than minimum chunk size + if ( chunk_size < GET_REAL_CHUNK_SIZE(BASIC_CHUNK_SIZE) ) + break; + + // Calculate the required padding for this chunk size + pad_size = ((size % chunk_size) == 0) ? 0 : + chunk_size - (size % chunk_size); + + // total number of chunks required if we are going to use this chunk size + total_chunks = (pad_size == 0) ? size / chunk_size : + (size / chunk_size) + 1; + + + // don't go beyond the maximum number of threads or maximum global memory + // TODO: it seems that the kernel breaks way before reaching the maximum + // global memory size (around 94MByte input plus the required + // scratch space) + if (( total_chunks > MAX_NUM_OF_THREADS) || + ( dev_prop.global_mem_size < (total_chunks * + (chunk_size + hash_size) + + hash_size))) + break; + + // each thread will take care of one chunk + total_threads = total_chunks; + + + // Get block and grid sizes + if (total_chunks <= MAX_THREADS_PER_BLOCK ) { + threads_per_block = total_chunks; + blocks_per_grid = 1; + + } else { + threads_per_block = MAX_THREADS_PER_BLOCK; + blocks_per_grid = ((total_threads % threads_per_block) == 0) ? + (total_threads/threads_per_block) : + (total_threads/threads_per_block) + 1; + + } + found = 1; + + if( total_threads > NUM_OF_MULTIPROCESSORS * dev_prop.warp_size) + break; + + index--; + + } + + //**** Did we find a solution? ****// + if ( !found ) + return SG_ERR_DEV_MEM_OVERFLOW; + + //**** Fill the struct with the solution ****// + ctx->threads_per_block = threads_per_block; + ctx->blocks_per_grid = blocks_per_grid; + ctx->total_threads = total_threads; + ctx->total_size = size + pad_size; + ctx->chunk_size = chunk_size; + ctx->pad_size = pad_size; + + return SG_OK; + +} +#endif /* FEATURE_MAXIMIZE_NUM_OF_THREADS */ + +/*=========================================================================== + +FUNCTION SG_GET_OVERLAP_EXEC_CONTEXT + +DESCRIPTION + sets the required chunk size, thread per block and number of blocks + needed for kernel execution according to client buffer size, offset + and block size. + +DEPENDENCIES + None + +RETURN VALUE + execution context + +===========================================================================*/ +static sg_status_type sg_get_overlap_exec_context(int size, int offset, + int block_size, int hash_size, + sg_exec_context_type* ctx) { + + sg_dev_prop_type dev_prop; + int threads_per_block; + int blocks_per_grid; + int total_threads; + int total_size; + int pad_size; + + + //**** Get device information ****// + sg_get_dev_prop(&dev_prop); + + + //**** Get the total number of threads required ****// + total_threads = (size + offset - block_size) / offset; + total_threads = ((size + offset - block_size) % offset) != 0 ? + total_threads + 1 : total_threads; + + + //**** Get the required padding for the last block ****// + pad_size = ((total_threads - 1) * offset + block_size) - size; + + + //**** threads and blocks ****// + if( total_threads > MAX_THREADS_PER_BLOCK ) { + + threads_per_block = MAX_THREADS_PER_BLOCK; + blocks_per_grid = (total_threads % MAX_THREADS_PER_BLOCK) == 0 ? + (total_threads / MAX_THREADS_PER_BLOCK) : + (total_threads / MAX_THREADS_PER_BLOCK) + 1; + } else { + + threads_per_block = total_threads; + blocks_per_grid = 1; + } + + total_size = size + pad_size; + + //**** Check for device memory overflow ****// + if( dev_prop.global_mem_size < (total_size + (hash_size * total_threads))) { + return SG_ERR_DEV_MEM_OVERFLOW; + } + + + //**** Fill the struct with the solution ****// + ctx->threads_per_block = threads_per_block; + ctx->blocks_per_grid = blocks_per_grid; + ctx->total_threads = total_threads; + ctx->total_size = total_size; + ctx->chunk_size = block_size; + ctx->pad_size = pad_size; + + return SG_OK; + +} +#endif /* FEATURE_DYNAMIC_EXEC_CONTEXT */ + +/*=========================================================================== + +FUNCTION SG_PRINT_EXEC_CONTEXT + +DESCRIPTION + Prints out the passed execution context structure + +DEPENDENCIES + None + +RETURN VALUE + None + +===========================================================================*/ +static void sg_print_exec_context( sg_exec_context_type* ctx ) { + printf("\n== GPU Execution Context ==\n"); + printf("Threads : %d\n", ctx->threads_per_block); + printf("Blocks : %d\n", ctx->blocks_per_grid); + printf("Total Threads : %d\n", ctx->total_threads); + printf("Total size : %d\n", ctx->total_size); + printf("Chunk Size : %d\n", ctx->chunk_size); + printf("Padding : %d\n\n", ctx->pad_size); +} + +/*-------------------------------------------------------------------------- + GLOBAL FUNCTIONS +--------------------------------------------------------------------------*/ + + + +/*=========================================================================== + +FUNCTION SG_INIT + +DESCRIPTION + Initialize the library + +DEPENDENCIES + None + +RETURN VALUE + None + +===========================================================================*/ +void sg_init( ) { + + char *buffer; + + //**** Utility library initialization ****// + // initialise card and timer + int deviceCount; + CUDA_SAFE_CALL_NO_SYNC(cudaGetDeviceCount(&deviceCount)); + if (deviceCount == 0) { + fprintf(stderr, "There is no device.\n"); + exit(EXIT_FAILURE); + } + int dev; + for (dev = 0; dev < deviceCount; ++dev) { + cudaDeviceProp deviceProp; + CUDA_SAFE_CALL_NO_SYNC(cudaGetDeviceProperties(&deviceProp, dev)); + if (deviceProp.major >= 1) + break; + } + if (dev == deviceCount) { + fprintf(stderr, "There is no device supporting CUDA.\n"); + exit(EXIT_FAILURE); + } + else + CUDA_SAFE_CALL(cudaSetDevice(dev)); + + + //**** force runtime initialization (CUDA ref. manual for more info.) ****// + cudaMallocHost( (void**) &buffer, 4 ); + cudaFreeHost( buffer ); + +} + +/*=========================================================================== + +FUNCTION SG_MALLOC + +DESCRIPTION + Allocate the required memory size. + +DEPENDENCIES + None + +RETURN VALUE + pointer to the reseved buffer + +===========================================================================*/ +void* sg_malloc(unsigned int size){ + + void* buffer; + +#ifdef FEATURE_PINNED_MODE + cudaMallocHost( (void**) &buffer, size ); +#else + + buffer = malloc( size ); +#endif /* FEATURE_PINNED_MODE */ + + return buffer; +} + +/*=========================================================================== + +FUNCTION SG_FREE + +DESCRIPTION + Free the allocated buffer. + +DEPENDENCIES + None + +RETURN VALUE + pointer to the reseved buffer + +===========================================================================*/ +void sg_free(void* buffer){ + + +#ifdef FEATURE_PINNED_MODE + cudaFreeHost(buffer ); +#else + free( buffer ); +#endif + +} + +/*=========================================================================== + +FUNCTION SG_MD5 + +DESCRIPTION + Returns the MD5 hash + +DEPENDENCIES + None + +RETURN VALUE + Hash value + +===========================================================================*/ +sg_status_type sg_md5(unsigned char* buffer, int size, + unsigned char** output, int* output_size, + sg_time_breakdown_type* time_breakdown) { + + //**** Variable Declarations ****// + sg_exec_context_type exec_context; + sg_status_type status = SG_OK; + unsigned char* d_scratchData; + unsigned char* h_scratchData; + unsigned char* d_input; + unsigned int timer; + + + //**** create the timer ****// + timer = 0; + CUT_SAFE_CALL( cutCreateTimer( &timer)); + + +#ifdef FEATURE_DYNAMIC_EXEC_CONTEXT + //**** Calculate pad size and needed block and grid sizes ****// + status = sg_get_exec_context(size, MD5_HASH_SIZE, &exec_context); + if ( status != SG_OK ) { + printf("Global memory overflow\n"); + return status; + } +#else + + //**** Fill the execution context structure ****// + exec_context.threads_per_block = THREADS_PER_BLOCK; + exec_context.blocks_per_grid = BLOCKS_PER_GRID; + exec_context.chunk_size = CHUNK_SIZE; + exec_context.total_size = size; + exec_context.total_threads = TOTAL_NUM_OF_THREADS; + exec_context.pad_size = 0; + +#endif /* FEATURE_DYNAMIC_EXEC_CONTEXT */ + sg_print_exec_context(&exec_context); + + + + //**** device memory allocation timing ****// + CUT_SAFE_CALL(cutResetTimer(timer)); + CUT_SAFE_CALL(cutStartTimer(timer)); + + /* allocate input data space */ + CUDA_SAFE_CALL(cudaMalloc((void**) &d_input, exec_context.total_size)); + + /* allocate scratch space */ + CUDA_SAFE_CALL(cudaMalloc((void**) &d_scratchData, + MD5_HASH_SIZE * exec_context.total_threads)); + + /* stop the timer (device memory allocation) */ + CUT_SAFE_CALL(cutStopTimer(timer)); + time_breakdown->device_mem_alloc_time = cutGetTimerValue(timer); + + + + //**** scratch buffer allocation timing ****// + CUT_SAFE_CALL(cutResetTimer(timer)); + CUT_SAFE_CALL(cutStartTimer(timer)); + + /* allocate buffer for the results */ + cudaMallocHost((void **)&h_scratchData, MD5_HASH_SIZE * + exec_context.total_threads); + + /* stop the timer (scratch buffer allocation) */ + CUT_SAFE_CALL(cutStopTimer(timer)); + time_breakdown->host_output_buffer_alloc_time = cutGetTimerValue(timer); + + + + //**** start timer for data copy in timing ****// + CUT_SAFE_CALL(cutResetTimer(timer)); + CUT_SAFE_CALL(cutStartTimer(timer)); + + /* move data to the device memory */ + CUDA_SAFE_CALL(cudaMemcpy( d_input, buffer, size, + cudaMemcpyHostToDevice)); + + /* stop the timer (copy in) */ + CUT_SAFE_CALL( cutStopTimer( timer)); + time_breakdown->copy_in_time = cutGetTimerValue( timer ); + + + + //**** setup execution parameters ****// + dim3 block( exec_context.threads_per_block ); + dim3 grid( exec_context.blocks_per_grid ); + + + + //**** start timer for kernel execution timing ****// + CUT_SAFE_CALL(cutResetTimer(timer)); + CUT_SAFE_CALL(cutStartTimer(timer)); + + + /* execute the kernel */ + md5<<< grid, block >>>(d_input, exec_context.chunk_size, + exec_context.total_threads, + exec_context.pad_size, + d_scratchData); + + // check if kernel execution generated an error + CUT_CHECK_ERROR("Kernel execution failed"); + + /* wait till the kernel finishes execution */ + CUDA_SAFE_CALL(cudaThreadSynchronize()); + + /* stop the timer (kernel execution) */ + CUT_SAFE_CALL(cutStopTimer(timer)); + time_breakdown->exec_time = cutGetTimerValue(timer); + + + + //**** start timer for output copy out timing ****// + CUT_SAFE_CALL(cutResetTimer(timer)); + CUT_SAFE_CALL(cutStartTimer(timer)); + + /* get the results from the device */ + CUDA_SAFE_CALL(cudaMemcpy(h_scratchData, + d_scratchData, + MD5_HASH_SIZE * exec_context.total_threads, + cudaMemcpyDeviceToHost)); + + /* stop the timer (output copy out) */ + CUT_SAFE_CALL(cutStopTimer(timer)); + time_breakdown->copy_out_time = cutGetTimerValue(timer); + + + + //**** start timer for last hasing stage timing ****// + CUT_SAFE_CALL(cutResetTimer(timer)); + CUT_SAFE_CALL(cutStartTimer(timer)); + + /* will do the last hshing stage on the CPU */ + sc_md5_standard(h_scratchData, MD5_HASH_SIZE * exec_context.total_threads, + output ); + + /* stop the timer (last stage) */ + CUT_SAFE_CALL(cutStopTimer(timer)); + time_breakdown->last_stage_time = cutGetTimerValue(timer); + + //**** free allocated memory ****// + CUDA_SAFE_CALL(cudaFree(d_input)); + CUDA_SAFE_CALL(cudaFree(d_scratchData)); + cudaFreeHost(h_scratchData); + + *output_size = MD5_HASH_SIZE; + + + return status; + +} + +/*=========================================================================== + +FUNCTION SG_MD5_OVERLAP + +DESCRIPTION + Returns the MD5 hash of each block for the provided buffer + +DEPENDENCIES + None + +RETURN VALUE + Hash value + +===========================================================================*/ +sg_status_type sg_md5_overlap(unsigned char* buffer, int size, + int block_size, int offset, + unsigned char** output, int* output_size, + sg_time_breakdown_type* time_breakdown) { + + + //**** Variable Declarations ****// + sg_exec_context_type exec_context; + sg_status_type status = SG_OK; + unsigned char* d_output; + unsigned char* d_input; + unsigned int timer; + + + + //**** create the timer ****// + timer = 0; + CUT_SAFE_CALL( cutCreateTimer( &timer)); + +#ifdef FEATURE_DYNAMIC_EXEC_CONTEXT + //**** Calculate pad size and needed block and grid sizes ****// + status = sg_get_overlap_exec_context(size, offset, block_size, + MD5_HASH_SIZE, &exec_context); + if ( status != SG_OK ) { + printf("Global memory overflow\n"); + return status; + } +#else + //**** Fill the execution context structure ****// + exec_context.threads_per_block = THREADS_PER_BLOCK; + exec_context.blocks_per_grid = BLOCKS_PER_GRID; + exec_context.chunk_size = CHUNK_SIZE; + exec_context.total_size = size; + exec_context.total_threads = TOTAL_NUM_OF_THREADS; + exec_context.pad_size = 0; +#endif /* FEATURE_DYNAMIC_EXEC_CONTEXT */ + sg_print_exec_context(&exec_context); + + + + //**** start timer for device memory allocation timing ****// + CUT_SAFE_CALL(cutResetTimer(timer)); + CUT_SAFE_CALL(cutStartTimer(timer)); + + /* allocate input space */ + CUDA_SAFE_CALL(cudaMalloc((void**) &d_input, exec_context.total_size)); + + /* allocate output space */ + CUDA_SAFE_CALL(cudaMalloc((void**) &d_output, + MD5_HASH_SIZE * exec_context.total_threads)); + + /* stop the timer (memory allocation) */ + CUT_SAFE_CALL(cutStopTimer(timer)); + time_breakdown->device_mem_alloc_time = cutGetTimerValue(timer); + + + + //**** start timer for output memory allocation timing ****// + CUT_SAFE_CALL(cutResetTimer(timer)); + CUT_SAFE_CALL(cutStartTimer(timer)); + + /**output = (unsigned char*) sg_malloc(MD5_HASH_SIZE * + exec_context.total_threads);*/ + cudaMallocHost( (void**) output, MD5_HASH_SIZE * + exec_context.total_threads ); + + /* stop the timer (output buffer allocation) */ + CUT_SAFE_CALL(cutStopTimer(timer)); + time_breakdown->host_output_buffer_alloc_time = cutGetTimerValue(timer); + + + + //**** start timer for data copy in timing ****// + CUT_SAFE_CALL(cutResetTimer(timer)); + CUT_SAFE_CALL(cutStartTimer(timer)); + + /* move data to the device memory */ + CUDA_SAFE_CALL(cudaMemcpy(d_input, buffer, size, + cudaMemcpyHostToDevice)); + + /* stop the timer (copy in) */ + CUT_SAFE_CALL(cutStopTimer(timer)); + time_breakdown->copy_in_time = cutGetTimerValue(timer); + + + + //**** setup execution parameters ****// + dim3 block( exec_context.threads_per_block ); + dim3 grid( exec_context.blocks_per_grid ); + + + + //**** start timer for kernel execution timing ****// + CUT_SAFE_CALL(cutResetTimer(timer)); + CUT_SAFE_CALL(cutStartTimer(timer)); + + /* execute the kernel */ + md5_overlap<<< grid, block >>>(d_input, exec_context.chunk_size, + offset, exec_context.total_threads, + exec_context.pad_size, d_output); + + // check if kernel execution generated an error + CUT_CHECK_ERROR("Kernel execution failed"); + + /* wait till the kernel finishes execution */ + CUDA_SAFE_CALL(cudaThreadSynchronize()); + + /* stop the timer (kernel execution) */ + CUT_SAFE_CALL(cutStopTimer(timer)); + time_breakdown->exec_time = cutGetTimerValue(timer); + + + + //**** start timer for output copy out timing ****// + CUT_SAFE_CALL(cutResetTimer(timer)); + CUT_SAFE_CALL(cutStartTimer(timer)); + + /* get the result from the device */ + CUDA_SAFE_CALL(cudaMemcpy(*output, + d_output, + MD5_HASH_SIZE * exec_context.total_threads, + cudaMemcpyDeviceToHost)); + + /* stop the timer (output copy out) */ + CUT_SAFE_CALL(cutStopTimer(timer)); + time_breakdown->copy_out_time = cutGetTimerValue(timer); + + + + //**** free allocated memory ****// + CUDA_SAFE_CALL(cudaFree(d_input)); + CUDA_SAFE_CALL(cudaFree(d_output)); + + *output_size = MD5_HASH_SIZE * exec_context.total_threads; + + return status; +} + +/*=========================================================================== + +FUNCTION SG_SHA1 + +DESCRIPTION + Returns the SHA1 hash of a the provided buffer + +DEPENDENCIES + None + +RETURN VALUE + Hash value + +===========================================================================*/ +sg_status_type sg_sha1(unsigned char* buffer, int size, + unsigned char** output, int* output_size, + sg_time_breakdown_type* time_breakdown) { + + //**** Variable Declarations ****// + sg_exec_context_type exec_context; + sg_status_type status = SG_OK; + unsigned char* d_scratchData; + unsigned char* h_scratchData; + unsigned char* d_input; + unsigned int timer; + + + //**** create the timer ****// + timer = 0; + CUT_SAFE_CALL( cutCreateTimer( &timer)); + + +#ifdef FEATURE_DYNAMIC_EXEC_CONTEXT + //**** Calculate pad size and needed block and grid sizes ****// + status = sg_get_exec_context(size, SHA1_HASH_SIZE, &exec_context); + if ( status != SG_OK ) { + printf("Global memory overflow\n"); + return status; + } +#else + + //**** Fill the execution context structure ****// + exec_context.threads_per_block = THREADS_PER_BLOCK; + exec_context.blocks_per_grid = BLOCKS_PER_GRID; + exec_context.chunk_size = CHUNK_SIZE; + exec_context.total_size = size; + exec_context.total_threads = TOTAL_NUM_OF_THREADS; + exec_context.pad_size = 0; + +#endif /* FEATURE_DYNAMIC_EXEC_CONTEXT */ + sg_print_exec_context(&exec_context); + + + + //**** device memory allocation timing ****// + CUT_SAFE_CALL(cutResetTimer(timer)); + CUT_SAFE_CALL(cutStartTimer(timer)); + + /* allocate input data space */ + CUDA_SAFE_CALL(cudaMalloc((void**) &d_input, exec_context.total_size)); + + /* allocate scratch space */ + CUDA_SAFE_CALL(cudaMalloc((void**) &d_scratchData, + SHA1_HASH_SIZE * exec_context.total_threads)); + + /* stop the timer (device memory allocation) */ + CUT_SAFE_CALL(cutStopTimer(timer)); + time_breakdown->device_mem_alloc_time = cutGetTimerValue(timer); + + + + //**** scratch buffer allocation timing ****// + CUT_SAFE_CALL(cutResetTimer(timer)); + CUT_SAFE_CALL(cutStartTimer(timer)); + + /* allocate buffer for the results */ + cudaMallocHost((void**)&h_scratchData, SHA1_HASH_SIZE * + exec_context.total_threads); + + /* stop the timer (scratch buffer allocation) */ + CUT_SAFE_CALL(cutStopTimer(timer)); + time_breakdown->host_output_buffer_alloc_time = cutGetTimerValue(timer); + + + + //**** start timer for data copy in timing ****// + CUT_SAFE_CALL(cutResetTimer(timer)); + CUT_SAFE_CALL(cutStartTimer(timer)); + + /* move data to the device memory */ + CUDA_SAFE_CALL(cudaMemcpy( d_input, buffer, size, + cudaMemcpyHostToDevice)); + + /* stop the timer (copy in) */ + CUT_SAFE_CALL( cutStopTimer( timer)); + time_breakdown->copy_in_time = cutGetTimerValue( timer ); + + + + //**** setup execution parameters ****// + dim3 block( exec_context.threads_per_block ); + dim3 grid( exec_context.blocks_per_grid ); + + + + //**** start timer for kernel execution timing ****// + CUT_SAFE_CALL(cutResetTimer(timer)); + CUT_SAFE_CALL(cutStartTimer(timer)); + + + /* execute the kernel */ + sha1<<< grid, block >>>(d_input, exec_context.chunk_size, + exec_context.total_threads, + exec_context.pad_size, + d_scratchData); + + // check if kernel execution generated an error + CUT_CHECK_ERROR("Kernel execution failed"); + + /* wait till the kernel finishes execution */ + CUDA_SAFE_CALL(cudaThreadSynchronize()); + + /* stop the timer (kernel execution) */ + CUT_SAFE_CALL(cutStopTimer(timer)); + time_breakdown->exec_time = cutGetTimerValue(timer); + + + + //**** start timer for output copy out timing ****// + CUT_SAFE_CALL(cutResetTimer(timer)); + CUT_SAFE_CALL(cutStartTimer(timer)); + + /* get the results from the device */ + CUDA_SAFE_CALL(cudaMemcpy(h_scratchData, + d_scratchData, + SHA1_HASH_SIZE * exec_context.total_threads, + cudaMemcpyDeviceToHost)); + + /* stop the timer (output copy out) */ + CUT_SAFE_CALL(cutStopTimer(timer)); + time_breakdown->copy_out_time = cutGetTimerValue(timer); + + + + //**** start timer for last hasing stage timing ****// + CUT_SAFE_CALL(cutResetTimer(timer)); + CUT_SAFE_CALL(cutStartTimer(timer)); + + /* will do the last hshing stage on the CPU */ + sc_sha1_standard(h_scratchData, SHA1_HASH_SIZE * exec_context.total_threads, + output ); + + /* stop the timer (last stage) */ + CUT_SAFE_CALL(cutStopTimer(timer)); + time_breakdown->last_stage_time = cutGetTimerValue(timer); + + + //**** free allocated memory ****// + CUDA_SAFE_CALL(cudaFree(d_input)); + CUDA_SAFE_CALL(cudaFree(d_scratchData)); + cudaFreeHost(h_scratchData); + + *output_size = SHA1_HASH_SIZE; + + return status; + +} + + +/*=========================================================================== + +FUNCTION SG_SHA1_OVERLAP + +DESCRIPTION + Returns the SHA1 hash of each block for the provided buffer + +DEPENDENCIES + None + +RETURN VALUE + Hash value + +===========================================================================*/ +sg_status_type sg_sha1_overlap(unsigned char* buffer, int size, + int block_size, int offset, + unsigned char** output, int* output_size, + sg_time_breakdown_type* time_breakdown) { + + + //**** Variable Declarations ****// + sg_exec_context_type exec_context; + sg_status_type status = SG_OK; + unsigned char* d_output; + unsigned char* d_input; + unsigned int timer; + + + + //**** create the timer ****// + timer = 0; + CUT_SAFE_CALL( cutCreateTimer( &timer)); + +#ifdef FEATURE_DYNAMIC_EXEC_CONTEXT + //**** Calculate pad size and needed block and grid sizes ****// + status = sg_get_overlap_exec_context(size, offset, block_size, + SHA1_HASH_SIZE, &exec_context); + if ( status != SG_OK ) { + printf("Global memory overflow\n"); + return status; + } +#else + //**** Fill the execution context structure ****// + exec_context.threads_per_block = THREADS_PER_BLOCK; + exec_context.blocks_per_grid = BLOCKS_PER_GRID; + exec_context.chunk_size = CHUNK_SIZE; + exec_context.total_size = size; + exec_context.total_threads = TOTAL_NUM_OF_THREADS; + exec_context.pad_size = 0; +#endif /* FEATURE_DYNAMIC_EXEC_CONTEXT */ + sg_print_exec_context(&exec_context); + + + + //**** start timer for device memory allocation timing ****// + CUT_SAFE_CALL(cutResetTimer(timer)); + CUT_SAFE_CALL(cutStartTimer(timer)); + + /* allocate input space */ + CUDA_SAFE_CALL(cudaMalloc((void**) &d_input, exec_context.total_size)); + + /* allocate output space */ + CUDA_SAFE_CALL(cudaMalloc((void**) &d_output, + SHA1_HASH_SIZE * exec_context.total_threads)); + + /* stop the timer (memory allocation) */ + CUT_SAFE_CALL(cutStopTimer(timer)); + time_breakdown->device_mem_alloc_time = cutGetTimerValue(timer); + + + + //**** start timer for output memory allocation timing ****// + CUT_SAFE_CALL(cutResetTimer(timer)); + CUT_SAFE_CALL(cutStartTimer(timer)); + + /**output = (unsigned char*) sg_malloc(SHA1_HASH_SIZE * + exec_context.total_threads);*/ + cudaMallocHost( (void**) output, SHA1_HASH_SIZE * + exec_context.total_threads ); + + /* stop the timer (output buffer allocation) */ + CUT_SAFE_CALL(cutStopTimer(timer)); + time_breakdown->host_output_buffer_alloc_time = cutGetTimerValue(timer); + + + + //**** start timer for data copy in timing ****// + CUT_SAFE_CALL(cutResetTimer(timer)); + CUT_SAFE_CALL(cutStartTimer(timer)); + + /* move data to the device memory */ + CUDA_SAFE_CALL(cudaMemcpy(d_input, buffer, size, + cudaMemcpyHostToDevice)); + + /* stop the timer (copy in) */ + CUT_SAFE_CALL(cutStopTimer(timer)); + time_breakdown->copy_in_time = cutGetTimerValue(timer); + + + + //**** setup execution parameters ****// + dim3 block( exec_context.threads_per_block ); + dim3 grid( exec_context.blocks_per_grid ); + + + + //**** start timer for kernel execution timing ****// + CUT_SAFE_CALL(cutResetTimer(timer)); + CUT_SAFE_CALL(cutStartTimer(timer)); + + /* execute the kernel */ + sha1_overlap<<< grid, block >>>(d_input, exec_context.chunk_size, + offset, exec_context.total_threads, + exec_context.pad_size, d_output); + + // check if kernel execution generated an error + CUT_CHECK_ERROR("Kernel execution failed"); + + /* wait till the kernel finishes execution */ + CUDA_SAFE_CALL(cudaThreadSynchronize()); + + /* stop the timer (kernel execution) */ + CUT_SAFE_CALL(cutStopTimer(timer)); + time_breakdown->exec_time = cutGetTimerValue(timer); + + + + //**** start timer for output copy out timing ****// + CUT_SAFE_CALL(cutResetTimer(timer)); + CUT_SAFE_CALL(cutStartTimer(timer)); + + /* get the result from the device */ + CUDA_SAFE_CALL(cudaMemcpy(*output, + d_output, + SHA1_HASH_SIZE * exec_context.total_threads, + cudaMemcpyDeviceToHost)); + + /* stop the timer (output copy out) */ + CUT_SAFE_CALL(cutStopTimer(timer)); + time_breakdown->copy_out_time = cutGetTimerValue(timer); + + + + //**** free allocated memory ****// + CUDA_SAFE_CALL(cudaFree(d_input)); + CUDA_SAFE_CALL(cudaFree(d_output)); + + *output_size = SHA1_HASH_SIZE * exec_context.total_threads; + + return status; +} -- cgit v1.3