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