diff options
| -rw-r--r-- | libcuda/cuda_runtime_api.cc | 203 | ||||
| -rw-r--r-- | libopencl/opencl_runtime_api.cc | 5 | ||||
| -rw-r--r-- | src/abstract_hardware_model.cc | 9 | ||||
| -rw-r--r-- | src/abstract_hardware_model.h | 23 | ||||
| -rw-r--r-- | src/cuda-sim/cuda-sim.cc | 9 | ||||
| -rw-r--r-- | src/cuda-sim/cuda-sim.h | 2 | ||||
| -rw-r--r-- | src/gpgpu-sim/gpu-sim.cc | 238 | ||||
| -rw-r--r-- | src/gpgpu-sim/gpu-sim.h | 39 | ||||
| -rw-r--r-- | src/gpgpu-sim/shader.cc | 104 | ||||
| -rw-r--r-- | src/gpgpu-sim/shader.h | 21 | ||||
| -rw-r--r-- | src/gpgpu-sim/stat-tool.cc | 23 | ||||
| -rw-r--r-- | src/gpgpu-sim/stat-tool.h | 17 | ||||
| -rw-r--r-- | src/gpgpusim_entrypoint.cc | 135 | ||||
| -rw-r--r-- | src/gpgpusim_entrypoint.h | 17 | ||||
| -rw-r--r-- | src/stream_manager.cc | 383 | ||||
| -rw-r--r-- | src/stream_manager.h | 294 |
16 files changed, 1205 insertions, 317 deletions
diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 1990076..f5e7810 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -128,6 +128,13 @@ #include "../src/cuda-sim/ptx_ir.h" #include "../src/cuda-sim/ptx_parser.h" #include "../src/gpgpusim_entrypoint.h" +#include "../src/stream_manager.h" + +#include <pthread.h> +#include <semaphore.h> + +extern void synchronize(); +extern void exit_simulation(); static int load_static_globals( symbol_table *symtab, unsigned min_gaddr, unsigned max_gaddr, gpgpu_t *gpu ); static int load_constants( symbol_table *symtab, addr_t min_gaddr, gpgpu_t *gpu ); @@ -161,6 +168,8 @@ struct cudaArray cudaError_t g_last_cudaError = cudaSuccess; +extern stream_manager *g_stream_manager; + void register_ptx_function( const char *name, function_info *impl ) { // no longer need this @@ -256,6 +265,32 @@ private: std::map<const void*,function_info*> m_kernel_lookup; // unique id (CUDA app function address) => kernel entry point }; +class kernel_config { +public: + kernel_config( dim3 GridDim, dim3 BlockDim, size_t sharedMem, struct CUstream_st *stream ) + { + m_GridDim=GridDim; + m_BlockDim=BlockDim; + m_sharedMem=sharedMem; + m_stream = stream; + } + void set_arg( const void *arg, size_t size, size_t offset ) + { + m_args.push_front( gpgpu_ptx_sim_arg(arg,size,offset) ); + } + dim3 grid_dim() const { return m_GridDim; } + dim3 block_dim() const { return m_BlockDim; } + gpgpu_ptx_sim_arg_list_t get_args() { return m_args; } + struct CUstream_st *get_stream() { return m_stream; } + +private: + dim3 m_GridDim; + dim3 m_BlockDim; + size_t m_sharedMem; + struct CUstream_st *m_stream; + gpgpu_ptx_sim_arg_list_t m_args; +}; + class _cuda_device_id *GPGPUSim_Init() { static _cuda_device_id *the_device = NULL; @@ -287,6 +322,7 @@ class _cuda_device_id *GPGPUSim_Init() the_gpu->set_prop(prop); the_device = new _cuda_device_id(the_gpu); } + start_sim_thread(1); return the_device; } @@ -353,67 +389,6 @@ void gpgpusim_ptx_assert_impl( int test_value, const char *func, const char *fil gpgpusim_ptx_error_impl(func, file, line, msg); } -class kernel_config { -public: - kernel_config( dim3 GridDim, dim3 BlockDim, size_t sharedMem, cudaStream_t stream ) - { - m_GridDim=GridDim; - m_BlockDim=BlockDim; - m_sharedMem=sharedMem; - m_stream=stream; - } - void set_arg(const void *arg, size_t size, size_t offset) - { - m_args.push_front( gpgpu_ptx_sim_arg(arg,size,offset) ); - } - dim3 grid_dim() const { return m_GridDim; } - dim3 block_dim() const { return m_BlockDim; } - gpgpu_ptx_sim_arg_list_t get_args() { return m_args; } - -private: - dim3 m_GridDim; - dim3 m_BlockDim; - size_t m_sharedMem; - cudaStream_t m_stream; - gpgpu_ptx_sim_arg_list_t m_args; -}; - -struct CUstream_st { - int _; -}; - -class CUevent_st { -public: - CUevent_st( bool blocking ) - { - m_uid = ++m_next_event_uid; - m_blocking = blocking; - m_updates = 0; - m_wallclock = 0; - m_gpu_tot_sim_cycle = 0; - m_done = true; - } - void update( double cycle, time_t clk ) - { - m_updates++; - m_wallclock=clk; - m_gpu_tot_sim_cycle=cycle; - } - //void set_done() { assert(!m_done); m_done=true; } - int get_uid() const { return m_uid; } - unsigned num_updates() const { return m_updates; } - bool done() const { return m_done; } - time_t clock() const { return m_wallclock; } -private: - int m_uid; - bool m_blocking; - bool m_done; - int m_updates; - time_t m_wallclock; - double m_gpu_tot_sim_cycle; - - static int m_next_event_uid; -}; typedef std::map<unsigned,CUevent_st*> event_tracker_t; @@ -516,17 +491,17 @@ __host__ cudaError_t CUDARTAPI cudaFree(void *devPtr) * * *******************************************************************************/ - __host__ cudaError_t CUDARTAPI cudaMemcpy(void *dst, const void *src, size_t count, enum cudaMemcpyKind kind) +__host__ cudaError_t CUDARTAPI cudaMemcpy(void *dst, const void *src, size_t count, enum cudaMemcpyKind kind) { - CUctx_st *context = GPGPUSim_Context(); - gpgpu_t *gpu = context->get_device()->get_gpgpu(); + //CUctx_st *context = GPGPUSim_Context(); + //gpgpu_t *gpu = context->get_device()->get_gpgpu(); printf("GPGPU-Sim PTX: cudaMemcpy(): devPtr = %p\n", dst); if( kind == cudaMemcpyHostToDevice ) - gpu->memcpy_to_gpu( (size_t)dst, src, count ); + g_stream_manager->push( stream_operation(src,(size_t)dst,count,0) ); else if( kind == cudaMemcpyDeviceToHost ) - gpu->memcpy_from_gpu( dst, (size_t)src, count ); + g_stream_manager->push( stream_operation((size_t)src,dst,count,0) ); else if( kind == cudaMemcpyDeviceToDevice ) - gpu->memcpy_gpu_to_gpu( (size_t)dst, (size_t)src, count ); + g_stream_manager->push( stream_operation((size_t)src,(size_t)dst,count,0) ); else { printf("GPGPU-Sim PTX: cudaMemcpy - ERROR : unsupported cudaMemcpyKind\n"); abort(); @@ -636,20 +611,23 @@ __host__ cudaError_t CUDARTAPI cudaFree(void *devPtr) __host__ cudaError_t CUDARTAPI cudaMemcpyToSymbol(const char *symbol, const void *src, size_t count, size_t offset __dv(0), enum cudaMemcpyKind kind __dv(cudaMemcpyHostToDevice)) { - CUctx_st *context = GPGPUSim_Context(); + //CUctx_st *context = GPGPUSim_Context(); assert(kind == cudaMemcpyHostToDevice); printf("GPGPU-Sim PTX: cudaMemcpyToSymbol: symbol = %p\n", symbol); - gpgpu_ptx_sim_memcpy_symbol(symbol,src,count,offset,1,context->get_device()->get_gpgpu()); + //stream_operation( const char *symbol, const void *src, size_t count, size_t offset ) + g_stream_manager->push( stream_operation(src,symbol,count,offset,0) ); + //gpgpu_ptx_sim_memcpy_symbol(symbol,src,count,offset,1,context->get_device()->get_gpgpu()); return g_last_cudaError = cudaSuccess; } __host__ cudaError_t CUDARTAPI cudaMemcpyFromSymbol(void *dst, const char *symbol, size_t count, size_t offset __dv(0), enum cudaMemcpyKind kind __dv(cudaMemcpyDeviceToHost)) { - CUctx_st *context = GPGPUSim_Context(); + //CUctx_st *context = GPGPUSim_Context(); assert(kind == cudaMemcpyDeviceToHost); printf("GPGPU-Sim PTX: cudaMemcpyFromSymbol: symbol = %p\n", symbol); - gpgpu_ptx_sim_memcpy_symbol(symbol,dst,count,offset,0,context->get_device()->get_gpgpu()); + g_stream_manager->push( stream_operation(symbol,dst,count,offset,0) ); + //gpgpu_ptx_sim_memcpy_symbol(symbol,dst,count,offset,0,context->get_device()->get_gpgpu()); return g_last_cudaError = cudaSuccess; } @@ -663,8 +641,14 @@ __host__ cudaError_t CUDARTAPI cudaFree(void *devPtr) __host__ cudaError_t CUDARTAPI cudaMemcpyAsync(void *dst, const void *src, size_t count, enum cudaMemcpyKind kind, cudaStream_t stream) { - printf("GPGPU-Sim PTX: warning cudaMemcpyAsync is implemented as blocking in this version of GPGPU-Sim...\n"); - return cudaMemcpy(dst,src,count,kind); + switch( kind ) { + case cudaMemcpyHostToDevice: g_stream_manager->push( stream_operation(src,(size_t)dst,count,stream) ); break; + case cudaMemcpyDeviceToHost: g_stream_manager->push( stream_operation((size_t)src,dst,count,stream) ); break; + case cudaMemcpyDeviceToDevice: g_stream_manager->push( stream_operation((size_t)src,(size_t)dst,count,stream) ); break; + default: + abort(); + } + return g_last_cudaError = cudaSuccess; } @@ -920,15 +904,20 @@ __host__ cudaError_t CUDARTAPI cudaLaunch( const char *hostFun ) char *mode = getenv("PTX_SIM_MODE_FUNC"); if( mode ) sscanf(mode,"%u", &g_ptx_sim_mode); - printf("\nGPGPU-Sim PTX: cudaLaunch for 0x%p (mode=%s)\n", hostFun, - g_ptx_sim_mode?"functional simulation":"performance simulation"); gpgpusim_ptx_assert( !g_cuda_launch_stack.empty(), "empty launch stack" ); - kernel_config &config = g_cuda_launch_stack.back(); + kernel_config config = g_cuda_launch_stack.back(); + cudaStream_t stream = config.get_stream(); + printf("\nGPGPU-Sim PTX: cudaLaunch for 0x%p (mode=%s) on stream %u\n", hostFun, + g_ptx_sim_mode?"functional simulation":"performance simulation", stream?stream->get_uid():0 ); kernel_info_t grid = gpgpu_cuda_ptx_sim_init_grid(hostFun,config.get_args(),config.grid_dim(),config.block_dim(),context); - if( g_ptx_sim_mode ) - gpgpu_cuda_ptx_sim_main_func( grid, config.grid_dim(), config.block_dim(), config.get_args() ); - else - gpgpu_cuda_ptx_sim_main_perf( grid, config.grid_dim(), config.block_dim(), config.get_args() ); + std::string kname = grid.name(); + dim3 gridDim = config.grid_dim(); + dim3 blockDim = config.block_dim(); + printf("GPGPU-Sim PTX: pushing kernel \'%s\' to stream %u, gridDim= (%u,%u,%u) blockDim = (%u,%u,%u) \n", + kname.c_str(), stream?stream->get_uid():0, gridDim.x,gridDim.y,gridDim.z,blockDim.x,blockDim.y,blockDim.z ); + stream_operation op(grid,g_ptx_sim_mode,stream); + g_stream_manager->push(op); + g_cuda_launch_stack.pop_back(); return g_last_cudaError = cudaSuccess; } @@ -943,8 +932,10 @@ __host__ cudaError_t CUDARTAPI cudaStreamCreate(cudaStream_t *stream) printf("GPGPU-Sim PTX: cudaStreamCreate\n"); #if (CUDART_VERSION >= 3000) *stream = new struct CUstream_st(); + g_stream_manager->add_stream(*stream); #else *stream = 0; + printf("GPGPU-Sim PTX: WARNING: Asynchronous kernel execution not supported (%s)\n", __my_func__); #endif return g_last_cudaError = cudaSuccess; } @@ -952,21 +943,33 @@ __host__ cudaError_t CUDARTAPI cudaStreamCreate(cudaStream_t *stream) __host__ cudaError_t CUDARTAPI cudaStreamDestroy(cudaStream_t stream) { #if (CUDART_VERSION >= 3000) - delete stream; + g_stream_manager->destroy_stream(stream); #endif return g_last_cudaError = cudaSuccess; } __host__ cudaError_t CUDARTAPI cudaStreamSynchronize(cudaStream_t stream) { - printf("GPGPU-Sim PTX: WARNING: This implementation of %s is only a stub! \n", __my_func__); - return g_last_cudaError = cudaSuccess; // it is stub because all cuda calls are synchronous +#if (CUDART_VERSION >= 3000) + if( stream == NULL ) + return g_last_cudaError = cudaErrorInvalidResourceHandle; + stream->synchronize(); +#else + printf("GPGPU-Sim PTX: WARNING: Asynchronous kernel execution not supported (%s)\n", __my_func__); +#endif + return g_last_cudaError = cudaSuccess; } __host__ cudaError_t CUDARTAPI cudaStreamQuery(cudaStream_t stream) { - printf("GPGPU-Sim PTX: WARNING: This implementation of %s is only a stub! \n", __my_func__); +#if (CUDART_VERSION >= 3000) + if( stream == NULL ) + return g_last_cudaError = cudaErrorInvalidResourceHandle; + return g_last_cudaError = stream->empty()?cudaSuccess:cudaErrorNotReady; +#else + printf("GPGPU-Sim PTX: WARNING: Asynchronous kernel execution not supported (%s)\n", __my_func__); return g_last_cudaError = cudaSuccess; // it is always success because all cuda calls are synchronous +#endif } /******************************************************************************* @@ -1005,15 +1008,15 @@ __host__ cudaError_t CUDARTAPI cudaEventRecord(cudaEvent_t event, cudaStream_t s { CUevent_st *e = get_event(event); if( !e ) return g_last_cudaError = cudaErrorUnknown; - time_t wallclock = time((time_t *)NULL); - e->update( gpu_tot_sim_cycle, wallclock ); + stream_operation op(e,stream); + g_stream_manager->push(op); return g_last_cudaError = cudaSuccess; } __host__ cudaError_t CUDARTAPI cudaEventQuery(cudaEvent_t event) { CUevent_st *e = get_event(event); - if( e == NULL || e->num_updates() == 0 ) { + if( e == NULL ) { return g_last_cudaError = cudaErrorInvalidValue; } else if( e->done() ) { return g_last_cudaError = cudaSuccess; @@ -1024,7 +1027,13 @@ __host__ cudaError_t CUDARTAPI cudaEventQuery(cudaEvent_t event) __host__ cudaError_t CUDARTAPI cudaEventSynchronize(cudaEvent_t event) { - return g_last_cudaError = cudaSuccess; + printf("GPGPU-Sim API: cudaEventSynchronize ** waiting for event\n"); + fflush(stdout); + while( !event->done() ) + ; + printf("GPGPU-Sim API: cudaEventSynchronize ** event detected\n"); + fflush(stdout); + return g_last_cudaError = cudaSuccess; } __host__ cudaError_t CUDARTAPI cudaEventDestroy(cudaEvent_t event) @@ -1061,22 +1070,20 @@ __host__ cudaError_t CUDARTAPI cudaEventElapsedTime(float *ms, cudaEvent_t start __host__ cudaError_t CUDARTAPI cudaThreadExit(void) { - // TODO... manage memory resources? + exit_simulation(); return g_last_cudaError = cudaSuccess; } - __host__ cudaError_t CUDARTAPI cudaThreadSynchronize(void) { - //Called on host side - //TODO This function should syncronize if we support Asyn kernel calls - return g_last_cudaError = cudaSuccess; + //Called on host side + synchronize(); + return g_last_cudaError = cudaSuccess; }; int CUDARTAPI __cudaSynchronizeThreads(void**, void*) { - //TODO This function should syncronize if we support Asyn kernel calls - return g_last_cudaError = cudaSuccess; + return cudaThreadExit(); } @@ -1563,12 +1570,6 @@ kernel_info_t gpgpu_cuda_ptx_sim_init_grid( const char *hostFun, printf("GPGPU-Sim PTX: ERROR launching kernel -- no PTX implementation found\n"); abort(); } - std::string kname = entry->get_name(); - - printf("GPGPU-Sim PTX: Launching kernel \'%s\' gridDim= (%u,%u,%u) blockDim = (%u,%u,%u); ntuid=%u\n", - kname.c_str(), gridDim.x,gridDim.y,gridDim.z,blockDim.x,blockDim.y,blockDim.z, - g_ptx_thread_info_uid_next ); - unsigned argcount=args.size(); unsigned argn=1; for( gpgpu_ptx_sim_arg_list_t::iterator a = args.begin(); a != args.end(); a++ ) { diff --git a/libopencl/opencl_runtime_api.cc b/libopencl/opencl_runtime_api.cc index c3f5e8a..8943061 100644 --- a/libopencl/opencl_runtime_api.cc +++ b/libopencl/opencl_runtime_api.cc @@ -623,6 +623,7 @@ class _cl_device_id *GPGPUSim_Init() gpgpu_sim *the_gpu = gpgpu_ptx_sim_init_perf(); the_device = new _cl_device_id(the_gpu); } + start_sim_thread(2); return the_device; } @@ -874,9 +875,9 @@ clEnqueueNDRangeKernel(cl_command_queue command_queue, kernel_info_t grid = gpgpu_opencl_ptx_sim_init_grid(kernel->get_implementation(),params,GridDim,BlockDim,gpu); if ( g_ptx_sim_mode ) - gpgpu_opencl_ptx_sim_main_func( grid, GridDim, BlockDim, params ); + gpgpu_opencl_ptx_sim_main_func( grid ); else - gpgpu_opencl_ptx_sim_main_perf( grid, GridDim, BlockDim, params ); + gpgpu_opencl_ptx_sim_main_perf( grid ); return CL_SUCCESS; } diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index 3f2883a..f779e3d 100644 --- a/src/abstract_hardware_model.cc +++ b/src/abstract_hardware_model.cc @@ -1,6 +1,7 @@ #include "abstract_hardware_model.h" #include "cuda-sim/memory.h" #include "option_parser.h" +#include "cuda-sim/ptx_ir.h" #include <algorithm> unsigned mem_access_t::sm_next_access_uid = 0; @@ -306,3 +307,11 @@ void warp_inst_t::generate_mem_accesses() m_mem_accesses_created=true; } + +unsigned kernel_info_t::m_next_uid = 1; + + +std::string kernel_info_t::name() const +{ + return m_kernel_entry->get_name(); +} diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index c519be1..77a35a7 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -91,6 +91,8 @@ public: { m_valid=false; m_kernel_entry=NULL; + m_uid=0; + m_num_cores_running=0; } kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry ) { @@ -102,8 +104,22 @@ public: m_next_cta.y=0; m_next_cta.z=0; m_next_tid=m_next_cta; + m_num_cores_running=0; + m_uid = m_next_uid++; } + void inc_running() { m_num_cores_running++; } + void dec_running() + { + assert( m_num_cores_running > 0 ); + m_num_cores_running--; + } + bool running() const { return m_num_cores_running>0; } + bool valid() const { return m_valid; } + bool done() const + { + return (!valid()) || (no_more_ctas_to_run() && !running()); + } class function_info *entry() { return m_kernel_entry; } const class function_info *entry() const { return m_kernel_entry; } @@ -143,15 +159,22 @@ public: { return m_next_tid.z < m_block_dim.z && m_next_tid.y < m_block_dim.y && m_next_tid.z < m_block_dim.x; } + unsigned get_uid() const { return m_uid; } + std::string name() const; private: bool m_valid; class function_info *m_kernel_entry; + unsigned m_uid; + static unsigned m_next_uid; + dim3 m_grid_dim; dim3 m_block_dim; dim3 m_next_cta; dim3 m_next_tid; + + unsigned m_num_cores_running; }; struct core_config { diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 9b65c25..b6a15e8 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1331,6 +1331,8 @@ int g_ptx_sim_mode; // if non-zero run functional simulation only (i.e., no noti extern "C" int ptx_debug; +bool g_cuda_launch_blocking = false; + void read_sim_environment_variables() { ptx_debug = 0; @@ -1375,7 +1377,12 @@ void read_sim_environment_variables() fflush(stdout); g_override_embedded_ptx = true; } + char *blocking = getenv("CUDA_LAUNCH_BLOCKING"); + if( blocking && !strcmp(blocking,"1") ) { + g_cuda_launch_blocking = true; + } #else + g_cuda_launch_blocking = true; g_override_embedded_ptx = true; #endif @@ -1388,7 +1395,7 @@ ptx_cta_info *g_func_cta_info = NULL; #define MAX(a,b) (((a)>(b))?(a):(b)) -void gpgpu_cuda_ptx_sim_main_func( kernel_info_t kernel, dim3 gridDim, dim3 blockDim, gpgpu_ptx_sim_arg_list_t args) +void gpgpu_cuda_ptx_sim_main_func( kernel_info_t kernel ) { printf("GPGPU-Sim: Performing Functional Simulation...\n"); diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h index e7c1965..ef045fb 100644 --- a/src/cuda-sim/cuda-sim.h +++ b/src/cuda-sim/cuda-sim.h @@ -25,7 +25,7 @@ extern class kernel_info_t gpgpu_opencl_ptx_sim_init_grid(class function_info *e struct dim3 gridDim, struct dim3 blockDim, class gpgpu_t *gpu ); -extern void gpgpu_cuda_ptx_sim_main_func( kernel_info_t kernel, dim3 gridDim, dim3 blockDim, gpgpu_ptx_sim_arg_list_t args); +extern void gpgpu_cuda_ptx_sim_main_func( kernel_info_t kernel ); extern void print_splash(); extern void gpgpu_ptx_sim_register_const_variable(void*, const char *deviceName, size_t size ); extern void gpgpu_ptx_sim_register_global_variable(void *hostVar, const char *deviceName, size_t size ); diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 9bb6a23..9ae05d8 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -276,6 +276,8 @@ void gpgpu_sim_config::reg_options(option_parser_t opp) option_parser_register(opp, "-gpgpu_clock_domains", OPT_CSTR, &gpgpu_clock_domains, "Clock Domain Frequencies in MhZ {<Core Clock>:<ICNT Clock>:<L2 Clock>:<DRAM Clock>}", "500.0:2000.0:2000.0:2000.0"); + option_parser_register(opp, "-gpgpu_max_concurrent_kernel", OPT_INT32, &max_concurrent_kernel, + "maximum kernels that can run concurrently on GPU", "8" ); option_parser_register(opp, "-gpgpu_cflog_interval", OPT_INT32, &gpgpu_cflog_interval, "Interval between each snapshot in control flow logger", "0"); @@ -307,7 +309,6 @@ void increment_x_then_y_then_z( dim3 &i, const dim3 &bound) } } - void gpgpu_sim::launch( kernel_info_t &kinfo ) { unsigned cta_size = kinfo.threads_per_cta(); @@ -319,15 +320,57 @@ void gpgpu_sim::launch( kernel_info_t &kinfo ) printf(" modify the CUDA source to decrease the kernel block size.\n"); abort(); } + unsigned n=0; + for(n=0; n < m_running_kernels.size(); n++ ) { + if( m_running_kernels[n].done() ) { + m_running_kernels[n] = kinfo; + break; + } + } + assert(n < m_running_kernels.size()); +} + +bool gpgpu_sim::can_start_kernel() +{ + for(unsigned n=0; n < m_running_kernels.size(); n++ ) { + if( m_running_kernels[n].done() ) + return true; + } + return false; +} - m_running_kernels.push_back(kinfo); +bool gpgpu_sim::get_more_cta_left() const +{ + if (m_config.gpu_max_cta_opt != 0) { + if( m_total_cta_launched >= m_config.gpu_max_cta_opt ) + return false; + } + for(unsigned n=0; n < m_running_kernels.size(); n++ ) { + if( m_running_kernels[n].valid() && !m_running_kernels[n].no_more_ctas_to_run() ) + return true; + } + return false; +} + +kernel_info_t *gpgpu_sim::select_kernel() +{ + for(unsigned n=0; n < m_running_kernels.size(); n++ ) { + unsigned idx = (n+m_last_issued_kernel+1)%m_config.max_concurrent_kernel; + if( m_running_kernels[idx].valid() && !m_running_kernels[idx].no_more_ctas_to_run() ) { + m_last_issued_kernel=idx; + return &m_running_kernels[idx]; + } + } + return NULL; } -kernel_info_t *gpgpu_sim::next_grid() +unsigned gpgpu_sim::finished_kernel() { - m_the_kernel = m_running_kernels.front(); - m_running_kernels.pop_front(); - return &m_the_kernel; + if( m_finished_kernel.empty() ) + return 0; + unsigned result = m_finished_kernel.front(); + m_finished_kernel.pop_front(); + return result; } void set_ptx_warp_size(const struct core_config * warp_size); @@ -360,6 +403,10 @@ gpgpu_sim::gpgpu_sim( const gpgpu_sim_config &config ) time_vector_create(NUM_MEM_REQ_STAT); fprintf(stdout, "GPGPU-Sim uArch: performance model initialization complete.\n"); + + m_running_kernels.resize( config.max_concurrent_kernel ); + m_last_issued_kernel = 0; + m_last_cluster_issue = 0; } int gpgpu_sim::shared_mem_size() const @@ -421,94 +468,76 @@ void gpgpu_sim::reinit_clock_domains(void) l2_time = 0; } -// return the number of cycle required to run all the trace on the gpu -unsigned int gpgpu_sim::run_gpu_sim() +bool gpgpu_sim::active() { - // run a CUDA grid on the GPU microarchitecture simulator - kernel_info_t &entry = m_the_kernel; - size_t program_size = get_kernel_code_size(entry.entry()); - - int not_completed; - int mem_busy; - int icnt2mem_busy; - - gpu_sim_cycle = 0; - not_completed = 1; - mem_busy = 1; - icnt2mem_busy = 1; - more_thread = true; - g_total_cta_left=0; - gpu_sim_insn = 0; - m_shader_stats->new_grid(); + if (m_config.gpu_max_cycle_opt && (gpu_tot_sim_cycle + gpu_sim_cycle) >= m_config.gpu_max_cycle_opt) + return false; + if (m_config.gpu_max_insn_opt && (gpu_tot_sim_insn + gpu_sim_insn) >= m_config.gpu_max_insn_opt) + return false; + if (m_config.gpu_max_cta_opt && (gpu_tot_issued_cta >= m_config.gpu_max_cta_opt) ) + return false; + if (m_config.gpu_deadlock_detect && gpu_deadlock) + return false; + for (unsigned i=0;i<m_shader_config->n_simt_clusters;i++) + if( m_cluster[i]->get_not_completed()>0 ) + return true;; + for (unsigned i=0;i<m_memory_config->m_n_mem;i++) + if( m_memory_partition_unit[i]->busy()>0 ) + return true;; + if( icnt_busy() ) + return true; + if( get_more_cta_left() ) + return true; + return false; +} - reinit_clock_domains(); - set_param_gpgpu_num_shaders(m_config.num_shader()); - for (unsigned i=0;i<m_shader_config->n_simt_clusters;i++) - m_cluster[i]->reinit(); - if (m_config.gpu_max_cta_opt != 0) { - g_total_cta_left = m_config.gpu_max_cta_opt; - } else { - g_total_cta_left = m_the_kernel.num_blocks(); - } - if (m_config.gpu_max_cta_opt != 0) { - // the maximum number of CTA has been reached, stop any further simulation - if (gpu_tot_issued_cta >= m_config.gpu_max_cta_opt) - return 0; - } +void gpgpu_sim::init() +{ + // run a CUDA grid on the GPU microarchitecture simulator + gpu_sim_cycle = 0; + gpu_sim_insn = 0; + m_total_cta_launched=0; - if (m_config.gpu_max_cycle_opt && (gpu_tot_sim_cycle + gpu_sim_cycle) >= m_config.gpu_max_cycle_opt) { - return gpu_sim_cycle; - } - if (m_config.gpu_max_insn_opt && (gpu_tot_sim_insn + gpu_sim_insn) >= m_config.gpu_max_insn_opt) { - return gpu_sim_cycle; - } + reinit_clock_domains(); + set_param_gpgpu_num_shaders(m_config.num_shader()); + for (unsigned i=0;i<m_shader_config->n_simt_clusters;i++) + m_cluster[i]->reinit(); + m_shader_stats->new_grid(); + // initialize the control-flow, memory access, memory latency logger + create_thread_CFlogger( m_config.num_shader(), m_shader_config->n_thread_per_shader, 0, m_config.gpgpu_cflog_interval ); + shader_CTA_count_create( m_config.num_shader(), m_config.gpgpu_cflog_interval); + if (m_config.gpgpu_cflog_interval != 0) { + insn_warp_occ_create( m_config.num_shader(), m_shader_config->warp_size ); + shader_warp_occ_create( m_config.num_shader(), m_shader_config->warp_size, m_config.gpgpu_cflog_interval); + shader_mem_acc_create( m_config.num_shader(), m_memory_config->m_n_mem, 4, m_config.gpgpu_cflog_interval); + shader_mem_lat_create( m_config.num_shader(), m_config.gpgpu_cflog_interval); + shader_cache_access_create( m_config.num_shader(), 3, m_config.gpgpu_cflog_interval); + set_spill_interval (m_config.gpgpu_cflog_interval * 40); + } - // initialize the control-flow, memory access, memory latency logger - create_thread_CFlogger( m_config.num_shader(), m_shader_config->n_thread_per_shader, program_size, 0, m_config.gpgpu_cflog_interval ); - shader_CTA_count_create( m_config.num_shader(), m_config.gpgpu_cflog_interval); - if (m_config.gpgpu_cflog_interval != 0) { - insn_warp_occ_create( m_config.num_shader(), m_shader_config->warp_size, program_size ); - shader_warp_occ_create( m_config.num_shader(), m_shader_config->warp_size, m_config.gpgpu_cflog_interval); - shader_mem_acc_create( m_config.num_shader(), m_memory_config->m_n_mem, 4, m_config.gpgpu_cflog_interval); - shader_mem_lat_create( m_config.num_shader(), m_config.gpgpu_cflog_interval); - shader_cache_access_create( m_config.num_shader(), 3, m_config.gpgpu_cflog_interval); - set_spill_interval (m_config.gpgpu_cflog_interval * 40); - } + if (g_network_mode) + icnt_init_grid(); +} - if (g_network_mode) - icnt_init_grid(); +void gpgpu_sim::print_stats() +{ + m_memory_stats->memlatstat_lat_pw(); + gpu_tot_sim_cycle += gpu_sim_cycle; + gpu_tot_sim_insn += gpu_sim_insn; - last_gpu_sim_insn = 0; - while (not_completed || mem_busy || icnt2mem_busy) { - cycle(); - not_completed = 0; - for (unsigned i=0;i<m_shader_config->n_simt_clusters;i++) - not_completed += m_cluster[i]->get_not_completed(); - mem_busy = 0; - for (unsigned i=0;i<m_memory_config->m_n_mem;i++) - mem_busy += m_memory_partition_unit[i]->busy(); - icnt2mem_busy = icnt_busy(); - if (m_config.gpu_max_cycle_opt && (gpu_tot_sim_cycle + gpu_sim_cycle) >= m_config.gpu_max_cycle_opt) - break; - if (m_config.gpu_max_insn_opt && (gpu_tot_sim_insn + gpu_sim_insn) >= m_config.gpu_max_insn_opt) - break; - if (m_config.gpu_deadlock_detect && gpu_deadlock) - break; - } - m_memory_stats->memlatstat_lat_pw(); - gpu_tot_sim_cycle += gpu_sim_cycle; - gpu_tot_sim_insn += gpu_sim_insn; - - ptx_file_line_stats_write_file(); + ptx_file_line_stats_write_file(); + gpu_print_stat(); - gpu_print_stat(); - if (g_network_mode) { - interconnect_stats(); - printf("----------------------------Interconnect-DETAILS---------------------------------" ); - icnt_overal_stat(); - printf("----------------------------END-of-Interconnect-DETAILS-------------------------" ); - } + if (g_network_mode) { + interconnect_stats(); + printf("----------------------------Interconnect-DETAILS---------------------------------" ); + icnt_overal_stat(); + printf("----------------------------END-of-Interconnect-DETAILS-------------------------" ); + } +} +void gpgpu_sim::deadlock_check() +{ if (m_config.gpu_deadlock_detect && gpu_deadlock) { fflush(stdout); printf("\n\nGPGPU-Sim uArch: ERROR ** deadlock detected: last writeback core %u @ gpu_sim_cycle %u (+ gpu_tot_sim_cycle %u) (%u cycles ago)\n", @@ -545,7 +574,6 @@ unsigned int gpgpu_sim::run_gpu_sim() fflush(stdout); abort(); } - return gpu_sim_cycle; } void gpgpu_sim::gpu_print_stat() const @@ -634,11 +662,11 @@ void shader_core_ctx::mem_instruction_stats(const warp_inst_t &inst) void shader_core_ctx::issue_block2core( kernel_info_t &kernel ) { + set_max_cta(kernel); + // find a free CTA context unsigned free_cta_hw_id=(unsigned)-1; - unsigned max_concurrent_cta_this_kernel = m_config->max_cta(kernel); - assert( max_concurrent_cta_this_kernel <= MAX_CTA_PER_SHADER ); - for (unsigned i=0;i<max_concurrent_cta_this_kernel;i++ ) { + for (unsigned i=0;i<kernel_max_cta_per_shader;i++ ) { if( m_cta_status[i]==0 ) { free_cta_hw_id=i; break; @@ -722,6 +750,18 @@ int gpgpu_sim::next_clock_domain(void) return mask; } +void gpgpu_sim::issue_block2core() +{ + for (unsigned i=0;i<m_shader_config->n_simt_clusters;i++) { + unsigned idx = (i+m_last_cluster_issue+1) % m_shader_config->n_simt_clusters; + unsigned num = m_cluster[idx]->issue_block2core(); + if( num ) { + m_last_cluster_issue=idx; + m_total_cta_launched += num; + } + } +} + unsigned long long g_single_step=0; // set this in gdb to single step the pipeline void gpgpu_sim::cycle() @@ -778,10 +818,11 @@ void gpgpu_sim::cycle() icnt_transfer(); } + last_gpu_sim_insn = 0; if (clock_mask & CORE) { // L1 cache + shader core pipeline stages for (unsigned i=0;i<m_shader_config->n_simt_clusters;i++) { - if (m_cluster[i]->get_not_completed() || more_thread) { + if (m_cluster[i]->get_not_completed() || get_more_cta_left() ) { m_cluster[i]->core_cycle(); } } @@ -791,18 +832,9 @@ void gpgpu_sim::cycle() gpu_sim_cycle++; if( g_interactive_debugger_enabled ) gpgpu_debug(); - - for (unsigned i=0;i<m_shader_config->n_simt_clusters && more_thread;i++) { - if ( ( (m_cluster[i]->get_n_active_cta()+1) <= m_cluster[i]->max_cta(m_the_kernel) ) && g_total_cta_left ) { - unsigned num = m_cluster[i]->issue_block2core( m_the_kernel ); - if (num >= g_total_cta_left) { - g_total_cta_left = 0; - more_thread = false; - } else - g_total_cta_left -= num; - } - } - + + issue_block2core(); + // Flush the caches once all of threads are completed. if (m_config.gpgpu_flush_cache) { int all_threads_complete = 1 ; diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h index db0b18f..5631ab9 100644 --- a/src/gpgpu-sim/gpu-sim.h +++ b/src/gpgpu-sim/gpu-sim.h @@ -192,20 +192,8 @@ public: m_valid=true; } - void set_max_cta( const kernel_info_t &kernel ) - { - // calcaulte the max cta count and cta size for local memory address mapping - m_shader_config.gpu_max_cta_per_shader = m_shader_config.max_cta(kernel); - //gpu_max_cta_per_shader is limited by number of CTAs if not enough - if( kernel.num_blocks() < m_shader_config.gpu_max_cta_per_shader*num_shader() ) { - m_shader_config.gpu_max_cta_per_shader = (kernel.num_blocks() / num_shader()); - if (kernel.num_blocks() % num_shader()) - m_shader_config.gpu_max_cta_per_shader++; - } - unsigned int gpu_cta_size = kernel.threads_per_cta(); - m_shader_config.gpu_padded_cta_size = (gpu_cta_size%32) ? 32*((gpu_cta_size/32)+1) : gpu_cta_size; - } unsigned num_shader() const { return m_shader_config.num_shader(); } + unsigned get_max_concurrent_kernel() const { return max_concurrent_kernel; } private: void init_clock_domains(void ); @@ -234,6 +222,7 @@ private: int gpgpu_dram_sched_queue_size; int gpgpu_cflog_interval; char * gpgpu_clock_domains; + unsigned max_concurrent_kernel; // visualizer bool g_visualizer_enabled; @@ -254,12 +243,17 @@ public: void set_prop( struct cudaDeviceProp *prop ); void launch( kernel_info_t &kinfo ); - kernel_info_t *next_grid(); + bool can_start_kernel(); + unsigned finished_kernel(); + void set_kernel_done( unsigned uid ) { m_finished_kernel.push_back(uid); } - unsigned run_gpu_sim(); + void init(); + void cycle(); + bool active(); + void print_stats(); + void deadlock_check(); void get_pdom_stack_top_info( unsigned sid, unsigned tid, unsigned *pc, unsigned *rpc ); - const kernel_info_t &the_kernel() const { return m_the_kernel; } int shared_mem_size() const; int num_registers_per_core() const; @@ -269,6 +263,8 @@ public: enum divergence_support_t simd_model() const; unsigned threads_per_core() const; + bool get_more_cta_left() const; + kernel_info_t *select_kernel(); const gpgpu_sim_config &get_config() const { return m_config; } void gpu_print_stat() const; @@ -278,8 +274,8 @@ private: // clocks void reinit_clock_domains(void); int next_clock_domain(void); + void issue_block2core(); - void cycle(); void L2c_print_cache_stat() const; void shader_print_runtime_stat( FILE *fout ); void shader_print_l1_miss_stat( FILE *fout ); @@ -293,11 +289,12 @@ private: class simt_core_cluster **m_cluster; class memory_partition_unit **m_memory_partition_unit; - kernel_info_t m_the_kernel; - std::list<kernel_info_t> m_running_kernels; + std::vector<kernel_info_t> m_running_kernels; + unsigned m_last_issued_kernel; - unsigned g_total_cta_left; - bool more_thread; + std::list<unsigned> m_finished_kernel; + unsigned m_total_cta_launched; + unsigned m_last_cluster_issue; // time of next rising edge double core_time; diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc index 7a0ad59..07aeb6d 100644 --- a/src/gpgpu-sim/shader.cc +++ b/src/gpgpu-sim/shader.cc @@ -110,6 +110,7 @@ shader_core_ctx::shader_core_ctx( class gpgpu_sim *gpu, shader_core_stats *stats ) : m_barriers( config->max_warps_per_shader, config->max_cta_per_core ) { + m_kernel = NULL; m_gpu = gpu; m_cluster = cluster; m_config = config; @@ -519,7 +520,8 @@ void shader_core_ctx::fetch() for( unsigned t=0; t<m_config->warp_size;t++) { unsigned tid=warp_id*m_config->warp_size+t; if( m_thread[tid].m_functional_model_thread_state ) { - register_cta_thread_exit(tid); + unsigned cta_id = m_warp[warp_id].get_cta_id(); + register_cta_thread_exit(cta_id); m_not_completed -= 1; m_active_threads.reset(tid); m_thread[tid].m_functional_model_thread_state=NULL; @@ -725,26 +727,37 @@ address_type shader_core_ctx::translate_local_memaddr( address_type localaddr, u { // During functional execution, each thread sees its own memory space for local memory, but these // need to be mapped to a shared address space for timing simulation. We do that mapping here. - localaddr /=4; + + address_type thread_base = 0; + unsigned max_concurrent_threads=0; if (m_config->gpgpu_local_mem_map) { - // Dnew = D*nTpC*nCpS*nS + nTpC*C + T%nTpC - // C = S + nS*(T/nTpC) - // D = data index; T = thread; C = CTA; S = shader core; p = per - // keep threads in a warp contiguous + // Dnew = D*N + T%nTpC + nTpC*C + // N = nTpC*nCpS*nS (max concurent threads) + // C = nS*K + S (hw cta number per gpu) + // K = T/nTpC (hw cta number per core) + // D = data index + // T = thread + // nTpC = number of threads per CTA + // nCpS = number of CTA per shader + // + // for a given local memory address threads in a CTA map to contiguous addresses, // then distribute across memory space by CTAs from successive shader cores first, // then by successive CTA in same shader core - localaddr *= m_config->gpu_padded_cta_size * m_config->gpu_max_cta_per_shader * num_shader; - localaddr += m_config->gpu_padded_cta_size * (m_sid + num_shader * (tid / m_config->gpu_padded_cta_size)); - localaddr += tid % m_config->gpu_padded_cta_size; + thread_base = 4*(kernel_padded_threads_per_cta * (m_sid + num_shader * (tid / kernel_padded_threads_per_cta)) + + tid % kernel_padded_threads_per_cta); + max_concurrent_threads = kernel_padded_threads_per_cta * kernel_max_cta_per_shader * num_shader; } else { // legacy mapping that maps the same address in the local memory space of all threads // to a single contiguous address region - localaddr *= num_shader * m_config->n_thread_per_shader; - localaddr += (m_config->n_thread_per_shader *m_sid) + tid; + thread_base = 4*(m_config->n_thread_per_shader * m_sid + tid); + max_concurrent_threads = num_shader * m_config->n_thread_per_shader; } - localaddr *= 4; + assert( thread_base < 4/*word size*/*max_concurrent_threads ); - return localaddr; + address_type local_word = localaddr/4; + address_type word_offset = localaddr%4; + address_type linear_address = local_word*max_concurrent_threads + thread_base + word_offset; + return linear_address; } ///////////////////////////////////////////////////////////////////////////////////////// @@ -1181,19 +1194,29 @@ void ldst_unit::cycle() } } -void shader_core_ctx::register_cta_thread_exit(int tid ) +void shader_core_ctx::register_cta_thread_exit( unsigned cta_num ) { - unsigned padded_cta_size = m_gpu->the_kernel().threads_per_cta(); - if (padded_cta_size%m_config->warp_size) - padded_cta_size = ((padded_cta_size/m_config->warp_size)+1)*(m_config->warp_size); - int cta_num = tid/padded_cta_size; assert( m_cta_status[cta_num] > 0 ); m_cta_status[cta_num]--; if (!m_cta_status[cta_num]) { m_n_active_cta--; m_barriers.deallocate_barrier(cta_num); shader_CTA_count_unlog(m_sid, 1); - printf("GPGPU-Sim uArch: Shader %d finished CTA #%d (%lld,%lld)\n", m_sid, cta_num, gpu_sim_cycle, gpu_tot_sim_cycle ); + printf("GPGPU-Sim uArch: Shader %d finished CTA #%d (%lld,%lld), %u CTAs running\n", m_sid, cta_num, gpu_sim_cycle, gpu_tot_sim_cycle, + m_n_active_cta ); + if( m_n_active_cta == 0 ) { + assert( m_kernel != NULL ); + m_kernel->dec_running(); + printf("GPGPU-Sim uArch: Shader %u empty.\n", m_sid ); + if( m_kernel->no_more_ctas_to_run() ) { + if( !m_kernel->running() ) { + m_gpu->set_kernel_done( m_kernel->get_uid() ); + printf("GPGPU-Sim uArch: GPU detected kernel \'%s\' finished on shader %u.\n", m_kernel->name().c_str(), m_sid ); + } + } + m_kernel=NULL; + fflush(stdout); + } } } @@ -1483,11 +1506,20 @@ unsigned int shader_core_config::max_cta( const kernel_info_t &k ) const printf ("\n"); } - if (result < 1) { - printf ("GPGPU-Sim uArch: ERROR ** Kernel requires more resources than shader has.\n"); - abort(); - } - return result; + //gpu_max_cta_per_shader is limited by number of CTAs if not enough to keep all cores busy + if( k.num_blocks() < result*num_shader() ) { + result = k.num_blocks() / num_shader(); + if (k.num_blocks() % num_shader()) + result++; + } + + assert( result <= MAX_CTA_PER_SHADER ); + if (result < 1) { + printf ("GPGPU-Sim uArch: ERROR ** Kernel requires more resources than shader has.\n"); + abort(); + } + + return result; } void shader_core_ctx::cycle() @@ -1712,6 +1744,16 @@ bool shader_core_ctx::warp_waiting_at_mem_barrier( unsigned warp_id ) return true; } +void shader_core_ctx::set_max_cta( const kernel_info_t &kernel ) +{ + // calculate the max cta count and cta size for local memory address mapping + kernel_max_cta_per_shader = m_config->max_cta(kernel); + unsigned int gpu_cta_size = kernel.threads_per_cta(); + kernel_padded_threads_per_cta = (gpu_cta_size%m_config->warp_size) ? + m_config->warp_size*((gpu_cta_size/m_config->warp_size)+1) : + gpu_cta_size; +} + gpgpu_sim *shader_core_ctx::get_gpu() { return m_gpu; @@ -2120,13 +2162,21 @@ unsigned simt_core_cluster::get_n_active_cta() const return n; } -unsigned simt_core_cluster::issue_block2core( class kernel_info_t &kernel ) +unsigned simt_core_cluster::issue_block2core() { unsigned num_blocks_issued=0; for( unsigned i=0; i < m_config->n_simt_cores_per_cluster; i++ ) { unsigned core = (i+m_cta_issue_next_core+1)%m_config->n_simt_cores_per_cluster; - if( m_core[core]->get_n_active_cta() < m_config->max_cta(kernel) ) { - m_core[core]->issue_block2core(kernel); + if( m_core[core]->get_not_completed() == 0 ) { + if( m_core[core]->get_kernel() == NULL ) { + kernel_info_t *k = m_gpu->select_kernel(); + if( k ) + m_core[core]->set_kernel(k); + } + } + kernel_info_t *kernel = m_core[core]->get_kernel(); + if( kernel && !kernel->no_more_ctas_to_run() && (m_core[core]->get_n_active_cta() < m_config->max_cta(*kernel)) ) { + m_core[core]->issue_block2core(*kernel); num_blocks_issued++; m_cta_issue_next_core=core; break; diff --git a/src/gpgpu-sim/shader.h b/src/gpgpu-sim/shader.h index 9200c99..7f41cfe 100644 --- a/src/gpgpu-sim/shader.h +++ b/src/gpgpu-sim/shader.h @@ -983,11 +983,8 @@ struct shader_core_config : public core_config unsigned gpgpu_shader_registers; int gpgpu_warpdistro_shader; unsigned gpgpu_num_reg_banks; - unsigned gpu_max_cta_per_shader; // TODO: modify this for fermi... computed based upon kernel - // resource usage; used in shader_core_ctx::translate_local_memaddr bool gpgpu_reg_bank_use_warp_id; bool gpgpu_local_mem_map; - int gpu_padded_cta_size; unsigned max_sp_latency; unsigned max_sfu_latency; @@ -1122,12 +1119,20 @@ public: void cache_flush(); void accept_fetch_response( mem_fetch *mf ); void accept_ldst_unit_response( class mem_fetch * mf ); + void set_kernel( kernel_info_t *k ) + { + assert(k); + m_kernel=k; + if(k) k->inc_running(); + printf("GPGPU-Sim uArch: Shader %d kernel = 0x%p\n", m_sid, m_kernel ); + } // accessors bool fetch_unit_response_buffer_full() const; bool ldst_unit_response_buffer_full() const; unsigned get_not_completed() const { return m_not_completed; } unsigned get_n_active_cta() const { return m_n_active_cta; } + kernel_info_t *get_kernel() { return m_kernel; } // used by functional simulation: // modifiers @@ -1148,6 +1153,7 @@ public: void dec_inst_in_pipeline( unsigned warp_id ) { m_warp[warp_id].dec_inst_in_pipeline(); } // also used in writeback() void store_ack( class mem_fetch *mf ); bool warp_waiting_at_mem_barrier( unsigned warp_id ); + void set_max_cta( const kernel_info_t &kernel ); // accessors std::list<unsigned> get_regs_written( const inst_t &fvt ) const; @@ -1163,7 +1169,7 @@ private: address_type next_pc( int tid ) const; void fetch(); - void register_cta_thread_exit(int cta_num ); + void register_cta_thread_exit( unsigned cta_num ); void decode(); void issue_warp( warp_inst_t *&warp, const warp_inst_t *pI, const active_mask_t &active_mask, unsigned warp_id ); @@ -1189,6 +1195,7 @@ private: const memory_config *m_memory_config; class simt_core_cluster *m_cluster; class gpgpu_sim *m_gpu; + kernel_info_t *m_kernel; // statistics shader_core_stats *m_stats; @@ -1228,6 +1235,10 @@ private: ldst_unit *m_ldst_unit; static const unsigned MAX_ALU_LATENCY = 64; std::bitset<MAX_ALU_LATENCY> m_result_bus; + + // used for local address mapping with single kernel launch + unsigned kernel_max_cta_per_shader; + unsigned kernel_padded_threads_per_cta; }; class simt_core_cluster { @@ -1243,7 +1254,7 @@ public: void icnt_cycle(); void reinit(); - unsigned issue_block2core( class kernel_info_t &kernel ); + unsigned issue_block2core(); void cache_flush(); bool icnt_injection_buffer_full(unsigned size, bool write); void icnt_inject_request_packet(class mem_fetch *mf); diff --git a/src/gpgpu-sim/stat-tool.cc b/src/gpgpu-sim/stat-tool.cc index 744a39d..576412c 100644 --- a/src/gpgpu-sim/stat-tool.cc +++ b/src/gpgpu-sim/stat-tool.cc @@ -152,7 +152,7 @@ unsigned translate_pc_to_ptxlineno(unsigned pc); static int n_thread_CFloggers = 0; static thread_CFlocality** thread_CFlogger = NULL; -void create_thread_CFlogger( int n_loggers, int n_threads, int n_insn, address_type start_pc, unsigned long long logging_interval) +void create_thread_CFlogger( int n_loggers, int n_threads, address_type start_pc, unsigned long long logging_interval) { destroy_thread_CFlogger(); @@ -163,7 +163,7 @@ void create_thread_CFlogger( int n_loggers, int n_threads, int n_insn, address_t char buffer[32]; for (int i = 0; i < n_thread_CFloggers; i++) { snprintf(buffer, 32, "%02d", i); - thread_CFlogger[i] = new thread_CFlocality( name_tpl + buffer, logging_interval, n_threads, n_insn, start_pc); + thread_CFlogger[i] = new thread_CFlocality( name_tpl + buffer, logging_interval, n_threads, start_pc); if (logging_interval != 0) { add_snap_shot_trigger(thread_CFlogger[i]); add_spill_log(thread_CFlogger[i]); @@ -222,10 +222,10 @@ int insn_warp_occ_logger::s_ids = 0; static std::vector<insn_warp_occ_logger> iwo_logger; -void insn_warp_occ_create( int n_loggers, int simd_width, int n_insn) +void insn_warp_occ_create( int n_loggers, int simd_width ) { iwo_logger.clear(); - iwo_logger.assign(n_loggers, insn_warp_occ_logger(simd_width, n_insn)); + iwo_logger.assign(n_loggers, insn_warp_occ_logger(simd_width)); for (unsigned i = 0; i < iwo_logger.size(); i++) { iwo_logger[i].set_id(i); } @@ -509,12 +509,12 @@ void shader_CTA_count_visualizer_gzprint( gzFile fout ) //////////////////////////////////////////////////////////////////////////////// //////////////////////////////////////////////////////////////////////////////// -thread_insn_span::thread_insn_span(unsigned long long cycle, int n_insn) - : m_cycle(cycle), m_n_insn(n_insn), +thread_insn_span::thread_insn_span(unsigned long long cycle) + : m_cycle(cycle), #ifdef USE_MAP m_insn_span_count() #else - m_insn_span_count(n_insn * 2) + m_insn_span_count(32*1024) #endif { } @@ -522,7 +522,7 @@ thread_insn_span::thread_insn_span(unsigned long long cycle, int n_insn) thread_insn_span::~thread_insn_span() { } thread_insn_span::thread_insn_span(const thread_insn_span& other) - : m_cycle(other.m_cycle), m_n_insn(other.m_n_insn), + : m_cycle(other.m_cycle), m_insn_span_count(other.m_insn_span_count) { } @@ -530,8 +530,7 @@ thread_insn_span::thread_insn_span(const thread_insn_span& other) thread_insn_span& thread_insn_span::operator=(const thread_insn_span& other) { printf("thread_insn_span& operator=\n"); - if (this != &other && m_n_insn != other.m_n_insn) { - m_n_insn = other.m_n_insn; + if (this != &other) { m_insn_span_count = other.m_insn_span_count; m_cycle = other.m_cycle; } @@ -540,7 +539,6 @@ thread_insn_span& thread_insn_span::operator=(const thread_insn_span& other) thread_insn_span& thread_insn_span::operator+=(const thread_insn_span& other) { - assert(m_n_insn == other.m_n_insn); // no way to aggregate if they are different programs span_count_map::const_iterator i_sc = other.m_insn_span_count.begin(); for (; i_sc != other.m_insn_span_count.end(); ++i_sc) { m_insn_span_count[i_sc->first] += i_sc->second; @@ -615,12 +613,11 @@ void thread_insn_span::print_sparse_histo(gzFile fout) const thread_CFlocality::thread_CFlocality(std::string name, unsigned long long snap_shot_interval, int nthreads, - int n_insn, address_type start_pc, unsigned long long start_cycle) : snap_shot_trigger(snap_shot_interval), m_name(name), m_nthreads(nthreads), m_thread_pc(nthreads, start_pc), m_cycle(start_cycle), - m_thd_span(start_cycle, n_insn) + m_thd_span(start_cycle) { std::fill(m_thread_pc.begin(), m_thread_pc.end(), -1); // so that hw thread with no work assigned will not clobber results } diff --git a/src/gpgpu-sim/stat-tool.h b/src/gpgpu-sim/stat-tool.h index 00b7771..31a1dcb 100644 --- a/src/gpgpu-sim/stat-tool.h +++ b/src/gpgpu-sim/stat-tool.h @@ -118,7 +118,7 @@ public: class thread_insn_span { public: - thread_insn_span(unsigned long long cycle, int n_insn); + thread_insn_span(unsigned long long cycle); thread_insn_span(const thread_insn_span& other); ~thread_insn_span(); @@ -135,14 +135,13 @@ public: private: typedef my_hash_map<address_type, int> span_count_map; unsigned long long m_cycle; - int m_n_insn; span_count_map m_insn_span_count; }; class thread_CFlocality : public snap_shot_trigger, public spill_log_interface { public: thread_CFlocality(std::string name, unsigned long long snap_shot_interval, - int nthreads, int n_insn, address_type start_pc, unsigned long long start_cycle = 0); + int nthreads, address_type start_pc, unsigned long long start_cycle = 0); ~thread_CFlocality(); void update_thread_pc( int thread_id, address_type pc ); @@ -170,9 +169,9 @@ private: class insn_warp_occ_logger { public: - insn_warp_occ_logger(int simd_width, int n_insn) + insn_warp_occ_logger(int simd_width) : m_simd_width(simd_width), - m_insn_warp_occ(n_insn, linear_histogram(1, "", m_simd_width)), + m_insn_warp_occ(1,linear_histogram(1, "", m_simd_width)), m_id(s_ids++) {} insn_warp_occ_logger(const insn_warp_occ_logger& other) @@ -191,7 +190,9 @@ public: void set_id(int id) { m_id = id; } void log(address_type pc, int warp_occ) { - m_insn_warp_occ[pc].add2bin(warp_occ - 1); + if( pc >= m_insn_warp_occ.size() ) + m_insn_warp_occ.resize(2*pc, linear_histogram(1, "", m_simd_width)); + m_insn_warp_occ[pc].add2bin(warp_occ - 1); } void print(FILE *fout) const @@ -307,7 +308,7 @@ void try_snap_shot (unsigned long long current_cycle); void set_spill_interval (unsigned long long interval); void spill_log_to_file (FILE *fout, int final, unsigned long long current_cycle); -void create_thread_CFlogger( int n_loggers, int n_threads, int n_insn, address_type start_pc, unsigned long long logging_interval); +void create_thread_CFlogger( int n_loggers, int n_threads, address_type start_pc, unsigned long long logging_interval); void destroy_thread_CFlogger( ); void cflog_update_thread_pc( int logger_id, int thread_id, address_type pc ); void cflog_snapshot( int logger_id, unsigned long long cycle ); @@ -316,7 +317,7 @@ void cflog_print_path_expression(FILE *fout); void cflog_visualizer_print(FILE *fout); void cflog_visualizer_gzprint(gzFile fout); -void insn_warp_occ_create( int n_loggers, int simd_width, int n_insn ); +void insn_warp_occ_create( int n_loggers, int simd_width ); void insn_warp_occ_log( int logger_id, address_type pc, int warp_occ ); void insn_warp_occ_print( FILE *fout ); diff --git a/src/gpgpusim_entrypoint.cc b/src/gpgpusim_entrypoint.cc index 8579f08..00ba5a2 100644 --- a/src/gpgpusim_entrypoint.cc +++ b/src/gpgpusim_entrypoint.cc @@ -71,6 +71,7 @@ #include "cuda-sim/ptx_parser.h" #include "gpgpu-sim/gpu-sim.h" #include "gpgpu-sim/icnt_wrapper.h" +#include "stream_manager.h" #include <pthread.h> #include <semaphore.h> @@ -84,30 +85,116 @@ struct gpgpu_ptx_sim_arg *grid_params; sem_t g_sim_signal_start; sem_t g_sim_signal_finish; +sem_t g_sim_signal_exit; time_t g_simulation_starttime; pthread_t g_simulation_thread; gpgpu_sim_config g_the_gpu_config; gpgpu_sim *g_the_gpu; +stream_manager *g_stream_manager; static void print_simulation_time(); -void *gpgpu_sim_thread(void*) +void *gpgpu_sim_thread_sequential(void*) { - kernel_info_t *kernel = NULL; + // at most one kernel running at a time + bool done; do { sem_wait(&g_sim_signal_start); - kernel = g_the_gpu->next_grid(); - if( kernel ) { - g_the_gpu_config.set_max_cta(*kernel); - g_the_gpu->run_gpu_sim(); + done = true; + if( g_the_gpu->get_more_cta_left() ) { + done = false; + g_the_gpu->init(); + while( g_the_gpu->active() ) + g_the_gpu->cycle(); + g_the_gpu->print_stats(); + g_the_gpu->deadlock_check(); print_simulation_time(); } sem_post(&g_sim_signal_finish); - } while(kernel); + } while(!done); + sem_post(&g_sim_signal_exit); return NULL; } +pthread_mutex_t g_sim_lock = PTHREAD_MUTEX_INITIALIZER; +bool g_sim_active = false; +bool g_sim_done = true; + +void *gpgpu_sim_thread_concurrent(void*) +{ + // concurrent kernel execution simulation thread + g_the_gpu->init(); + do { + printf("GPGPU-Sim: *** simulation thread starting and spinning waiting for work ***\n"); + fflush(stdout); + while( g_stream_manager->empty() && !g_sim_done ) + ; + printf("GPGPU-Sim: ** START simulation thread (detected work) **\n"); + g_stream_manager->print(stdout); + fflush(stdout); + pthread_mutex_lock(&g_sim_lock); + g_sim_active = true; + pthread_mutex_unlock(&g_sim_lock); + bool active = false; + do { + // check if a kernel has completed + unsigned grid_uid = g_the_gpu->finished_kernel(); + if( grid_uid ) + g_stream_manager->register_finished_kernel(grid_uid); + + // launch operation on device if one is pending and can be run + stream_operation op = g_stream_manager->front(); + op.do_operation(g_the_gpu); + + // simulate a clock cycle on the GPU + if( g_the_gpu->active() ) + g_the_gpu->cycle(); + g_the_gpu->deadlock_check(); + active = g_the_gpu->active() || !g_stream_manager->empty(); + } while( active ); + printf("GPGPU-Sim: ** STOP simulation thread (no work) **\n"); + fflush(stdout); + g_the_gpu->print_stats(); + pthread_mutex_lock(&g_sim_lock); + g_sim_active = false; + pthread_mutex_unlock(&g_sim_lock); + } while( !g_sim_done ); + printf("GPGPU-Sim: *** simulation thread exiting ***\n"); + fflush(stdout); + sem_post(&g_sim_signal_exit); + return NULL; +} + +void synchronize() +{ + printf("GPGPU-Sim: synchronize waiting for inactive GPU simulation\n"); + g_stream_manager->print(stdout); + fflush(stdout); +// sem_wait(&g_sim_signal_finish); + bool done = false; + do { + pthread_mutex_lock(&g_sim_lock); + done = g_stream_manager->empty() && !g_sim_active; + pthread_mutex_unlock(&g_sim_lock); + } while (!done); + printf("GPGPU-Sim: detected inactive GPU simulation thread\n"); + fflush(stdout); +// sem_post(&g_sim_signal_start); +} + +void exit_simulation() +{ + g_sim_done=true; + printf("GPGPU-Sim: exit_simulation called\n"); + fflush(stdout); + sem_wait(&g_sim_signal_exit); + printf("GPGPU-Sim: simulation thread signaled exit\n"); + fflush(stdout); +} + +extern bool g_cuda_launch_blocking; + gpgpu_sim *gpgpu_ptx_sim_init_perf() { srand(1); @@ -125,16 +212,29 @@ gpgpu_sim *gpgpu_ptx_sim_init_perf() g_the_gpu_config.init(); g_the_gpu = new gpgpu_sim(g_the_gpu_config); + g_stream_manager = new stream_manager(g_the_gpu,g_cuda_launch_blocking); g_simulation_starttime = time((time_t *)NULL); sem_init(&g_sim_signal_start,0,0); sem_init(&g_sim_signal_finish,0,0); - pthread_create(&g_simulation_thread,NULL,gpgpu_sim_thread,NULL); + sem_init(&g_sim_signal_exit,0,0); return g_the_gpu; } +void start_sim_thread(int api) +{ + if( g_sim_done ) { + g_sim_done = false; + if( g_the_gpu_config.get_max_concurrent_kernel() > 1 && api == 1 ) { + pthread_create(&g_simulation_thread,NULL,gpgpu_sim_thread_concurrent,NULL); + } else { + pthread_create(&g_simulation_thread,NULL,gpgpu_sim_thread_sequential,NULL); + } + } +} + void print_simulation_time() { time_t current_time, difference, d, h, m, s; @@ -154,21 +254,15 @@ void print_simulation_time() fflush(stdout); } -int gpgpu_cuda_ptx_sim_main_perf( kernel_info_t grid, - struct dim3 gridDim, - struct dim3 blockDim, - gpgpu_ptx_sim_arg_list_t grid_params ) +int gpgpu_cuda_ptx_sim_main_perf( kernel_info_t grid ) { g_the_gpu->launch(grid); - sem_post(&g_sim_signal_start); - sem_wait(&g_sim_signal_finish); + //sem_post(&g_sim_signal_start); + //sem_wait(&g_sim_signal_finish); return 0; } -int gpgpu_opencl_ptx_sim_main_perf( kernel_info_t grid, - struct dim3 gridDim, - struct dim3 blockDim, - gpgpu_ptx_sim_arg_list_t grid_params ) +int gpgpu_opencl_ptx_sim_main_perf( kernel_info_t grid ) { g_the_gpu->launch(grid); sem_post(&g_sim_signal_start); @@ -176,10 +270,7 @@ int gpgpu_opencl_ptx_sim_main_perf( kernel_info_t grid, return 0; } -int gpgpu_opencl_ptx_sim_main_func( kernel_info_t grid, - struct dim3 gridDim, - struct dim3 blockDim, - gpgpu_ptx_sim_arg_list_t grid_params ) +int gpgpu_opencl_ptx_sim_main_func( kernel_info_t grid ) { printf("GPGPU-Sim PTX API: OpenCL functional-only simulation not yet implemented (use performance simulation)\n"); exit(1); diff --git a/src/gpgpusim_entrypoint.h b/src/gpgpusim_entrypoint.h index 689dcf9..b310752 100644 --- a/src/gpgpusim_entrypoint.h +++ b/src/gpgpusim_entrypoint.h @@ -71,19 +71,10 @@ extern time_t g_simulation_starttime; class gpgpu_sim *gpgpu_ptx_sim_init_perf(); +void start_sim_thread(int api); -int gpgpu_cuda_ptx_sim_main_perf( kernel_info_t grid, - struct dim3 gridDim, - struct dim3 blockDim, - gpgpu_ptx_sim_arg_list_t grid_params ); +int gpgpu_cuda_ptx_sim_main_perf( kernel_info_t grid ); +int gpgpu_opencl_ptx_sim_main_perf( kernel_info_t grid ); +int gpgpu_opencl_ptx_sim_main_func( kernel_info_t grid ); -int gpgpu_opencl_ptx_sim_main_perf( kernel_info_t grid, - struct dim3 gridDim, - struct dim3 blockDim, - gpgpu_ptx_sim_arg_list_t grid_params ); - -int gpgpu_opencl_ptx_sim_main_func( kernel_info_t grid, - struct dim3 gridDim, - struct dim3 blockDim, - gpgpu_ptx_sim_arg_list_t grid_params ); #endif diff --git a/src/stream_manager.cc b/src/stream_manager.cc new file mode 100644 index 0000000..5d1cfe2 --- /dev/null +++ b/src/stream_manager.cc @@ -0,0 +1,383 @@ +/* + * stream_manager.cc + * + * Copyright © 2009 by Tor M. Aamodt, Wilson W. L. Fung, Ali Bakhoda, + * George L. Yuan and the University of British Columbia, Vancouver, + * BC V6T 1Z4, All Rights Reserved. + * + * THIS IS A LEGAL DOCUMENT BY DOWNLOADING GPGPU-SIM, YOU ARE AGREEING TO THESE + * TERMS AND CONDITIONS. + * + * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "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 THE COPYRIGHT OWNERS OR CONTRIBUTORS 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. + * + * NOTE: The files libcuda/cuda_runtime_api.c and src/cuda-sim/cuda-math.h + * are derived from the CUDA Toolset available from http://www.nvidia.com/cuda + * (property of NVIDIA). The files benchmarks/BlackScholes/ and + * benchmarks/template/ are derived from the CUDA SDK available from + * http://www.nvidia.com/cuda (also property of NVIDIA). The files from + * src/intersim/ are derived from Booksim (a simulator provided with the + * textbook "Principles and Practices of Interconnection Networks" available + * from http://cva.stanford.edu/books/ppin/). As such, those files are bound by + * the corresponding legal terms and conditions set forth separately (original + * copyright notices are left in files from these sources and where we have + * modified a file our copyright notice appears before the original copyright + * notice). + * + * Using this version of GPGPU-Sim requires a complete installation of CUDA + * which is distributed seperately by NVIDIA under separate terms and + * conditions. To use this version of GPGPU-Sim with OpenCL requires a + * recent version of NVIDIA's drivers which support OpenCL. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions are met: + * + * 1. Redistributions of source code must retain the above copyright notice, + * this list of conditions and the following disclaimer. + * + * 2. 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. + * + * 3. Neither the name of the University of British Columbia nor the names of + * its contributors may be used to endorse or promote products derived from + * this software without specific prior written permission. + * + * 4. This version of GPGPU-SIM is distributed freely for non-commercial use only. + * + * 5. No nonprofit user may place any restrictions on the use of this software, + * including as modified by the user, by any other authorized user. + * + * 6. GPGPU-SIM was developed primarily by Tor M. Aamodt, Wilson W. L. Fung, + * Ali Bakhoda, George L. Yuan, at the University of British Columbia, + * Vancouver, BC V6T 1Z4 + */ + +#include "stream_manager.h" +#include "gpgpusim_entrypoint.h" +#include "cuda-sim/cuda-sim.h" +#include "gpgpu-sim/gpu-sim.h" + +unsigned CUstream_st::sm_next_stream_uid = 0; + +CUstream_st::CUstream_st() +{ + m_pending = false; + m_uid = sm_next_stream_uid++; + pthread_mutex_init(&m_lock,NULL); +} + +bool CUstream_st::empty() +{ + pthread_mutex_lock(&m_lock); + bool empty = m_operations.empty(); + pthread_mutex_unlock(&m_lock); + return empty; +} + +bool CUstream_st::busy() +{ + pthread_mutex_lock(&m_lock); + bool pending = m_pending; + pthread_mutex_unlock(&m_lock); + return pending; +} + +void CUstream_st::synchronize() +{ + // called by host thread + bool done=false; + do{ + pthread_mutex_lock(&m_lock); + done = m_operations.empty(); + pthread_mutex_unlock(&m_lock); + } while ( !done ); +} + +void CUstream_st::push( const stream_operation &op ) +{ + // called by host thread + pthread_mutex_lock(&m_lock); + m_operations.push_back( op ); + pthread_mutex_unlock(&m_lock); +} + +void CUstream_st::record_next_done() +{ + // called by gpu thread + pthread_mutex_lock(&m_lock); + assert(m_pending); + m_operations.pop_front(); + m_pending=false; + pthread_mutex_unlock(&m_lock); +} + + +stream_operation CUstream_st::next() +{ + // called by gpu thread + pthread_mutex_lock(&m_lock); + m_pending = true; + stream_operation result = m_operations.front(); + pthread_mutex_unlock(&m_lock); + return result; +} + +void CUstream_st::print(FILE *fp) +{ + pthread_mutex_lock(&m_lock); + fprintf(fp,"GPGPU-Sim API: stream %u has %zu operations\n", m_uid, m_operations.size() ); + std::list<stream_operation>::iterator i; + unsigned n=0; + for( i=m_operations.begin(); i!=m_operations.end(); i++ ) { + stream_operation &op = *i; + fprintf(fp,"GPGPU-Sim API: %u : ", n++); + op.print(fp); + fprintf(fp,"\n"); + } + pthread_mutex_unlock(&m_lock); +} + +void stream_operation::do_operation( gpgpu_sim *gpu ) +{ + if( is_noop() ) + return; + + assert(!m_done && m_stream); + printf("GPGPU-Sim API: stream %u performing ", m_stream->get_uid() ); + switch( m_type ) { + case stream_memcpy_host_to_device: + printf("memcpy host-to-device\n"); + gpu->memcpy_to_gpu(m_device_address_dst,m_host_address_src,m_cnt); + m_stream->record_next_done(); + break; + case stream_memcpy_device_to_host: + printf("memcpy device-to-host\n"); + gpu->memcpy_from_gpu(m_host_address_dst,m_device_address_src,m_cnt); + m_stream->record_next_done(); + break; + case stream_memcpy_device_to_device: + printf("memcpy device-to-device\n"); + gpu->memcpy_gpu_to_gpu(m_device_address_dst,m_device_address_src,m_cnt); + m_stream->record_next_done(); + break; + case stream_memcpy_to_symbol: + printf("memcpy to symbol\n"); + gpgpu_ptx_sim_memcpy_symbol(m_symbol,m_host_address_src,m_cnt,m_offset,1,gpu); + m_stream->record_next_done(); + break; + case stream_memcpy_from_symbol: + printf("memcpy from symbol\n"); + gpgpu_ptx_sim_memcpy_symbol(m_symbol,m_host_address_dst,m_cnt,m_offset,0,gpu); + m_stream->record_next_done(); + break; + case stream_kernel_launch: + if( gpu->can_start_kernel() ) { + printf("kernel \'%s\' transfer to GPU hardware scheduler\n", m_kernel.name().c_str() ); + if( m_sim_mode ) + gpgpu_cuda_ptx_sim_main_func( m_kernel ); + else + gpu->launch( m_kernel ); + } + break; + case stream_event: { + printf("event update\n"); + time_t wallclock = time((time_t *)NULL); + m_event->update( gpu_tot_sim_cycle, wallclock ); + m_stream->record_next_done(); + } + break; + default: + abort(); + } + m_done=true; + fflush(stdout); +} + +void stream_operation::print( FILE *fp ) const +{ + fprintf(fp," stream operation " ); + switch( m_type ) { + case stream_event: fprintf(fp,"event"); break; + case stream_kernel_launch: fprintf(fp,"kernel"); break; + case stream_memcpy_device_to_device: fprintf(fp,"memcpy device-to-device"); break; + case stream_memcpy_device_to_host: fprintf(fp,"memcpy device-to-host"); break; + case stream_memcpy_host_to_device: fprintf(fp,"memcpy host-to-device"); break; + case stream_memcpy_to_symbol: fprintf(fp,"memcpy to symbol"); break; + case stream_memcpy_from_symbol: fprintf(fp,"memcpy from symbol"); break; + case stream_no_op: fprintf(fp,"no-op"); break; + } +} + +stream_manager::stream_manager( gpgpu_sim *gpu, bool cuda_launch_blocking ) +{ + m_gpu = gpu; + m_service_stream_zero = false; + m_cuda_launch_blocking = cuda_launch_blocking; + pthread_mutex_init(&m_lock,NULL); +} + +void stream_manager::register_finished_kernel( unsigned grid_uid ) +{ + // called by gpu simulation thread + pthread_mutex_lock(&m_lock); + CUstream_st *stream = m_grid_id_to_stream[grid_uid]; + assert( grid_uid == stream->front().get_kernel().get_uid() ); + stream->record_next_done(); + m_grid_id_to_stream.erase(grid_uid); + pthread_mutex_unlock(&m_lock); +} + +stream_operation stream_manager::front() +{ + // called by gpu simulation thread + stream_operation result; + pthread_mutex_lock(&m_lock); + if( concurrent_streams_empty() ) + m_service_stream_zero = true; + if( m_service_stream_zero ) { + if( !m_stream_zero.empty() ) { + if( !m_stream_zero.busy() ) { + result = m_stream_zero.next(); + if( result.is_kernel() ) { + unsigned grid_id = result.get_kernel().get_uid(); + m_grid_id_to_stream[grid_id] = &m_stream_zero; + } + } + } else { + m_service_stream_zero = false; + } + } else { + std::list<struct CUstream_st*>::iterator s; + for( s=m_streams.begin(); s != m_streams.end(); s++) { + CUstream_st *stream = *s; + if( !stream->busy() && !stream->empty() ) { + result = stream->next(); +// stream_barrier *b = result.get_barrier(); +// if( b && b->value() > 1 ) { +// b->dec(); +// result = stream_operation(); +// } else { +// assert( b->value() == 1 ); +// delete b; + if( result.is_kernel() ) { + unsigned grid_id = result.get_kernel().get_uid(); + m_grid_id_to_stream[grid_id] = stream; + } +// } + break; + } + } + } + pthread_mutex_unlock(&m_lock); + return result; +} + +void stream_manager::add_stream( struct CUstream_st *stream ) +{ + // called by host thread + pthread_mutex_lock(&m_lock); + m_streams.push_back(stream); + pthread_mutex_unlock(&m_lock); +} + +void stream_manager::destroy_stream( CUstream_st *stream ) +{ + // called by host thread + pthread_mutex_lock(&m_lock); + while( !stream->empty() ) + ; + std::list<CUstream_st *>::iterator s; + for( s=m_streams.begin(); s != m_streams.end(); s++ ) { + if( *s == stream ) { + m_streams.erase(s); + break; + } + } + delete stream; + pthread_mutex_unlock(&m_lock); +} + +bool stream_manager::concurrent_streams_empty() +{ + bool result = true; + // called by gpu simulation thread + std::list<struct CUstream_st *>::iterator s; + for( s=m_streams.begin(); s!=m_streams.end();++s ) { + struct CUstream_st *stream = *s; + if( !stream->empty() ) { + //stream->print(stdout); + result = false; + } + } + return result; +} + +bool stream_manager::empty() +{ + bool result = true; + pthread_mutex_lock(&m_lock); + if( !concurrent_streams_empty() ) + result = false; + if( !m_stream_zero.empty() ) + result = false; + pthread_mutex_unlock(&m_lock); + return result; +} + +void stream_manager::print( FILE *fp) +{ + pthread_mutex_lock(&m_lock); + print_impl(fp); + pthread_mutex_unlock(&m_lock); +} + +void stream_manager::print_impl( FILE *fp) +{ + fprintf(fp,"GPGPU-Sim API: Stream Manager State\n"); + std::list<struct CUstream_st *>::iterator s; + for( s=m_streams.begin(); s!=m_streams.end();++s ) { + struct CUstream_st *stream = *s; + if( !stream->empty() ) + stream->print(fp); + } + if( !m_stream_zero.empty() ) + m_stream_zero.print(fp); +} + +void stream_manager::push( stream_operation op ) +{ + struct CUstream_st *stream = op.get_stream(); + + // block if stream 0 (or concurrency disabled) and pending concurrent operations exist + bool block= !stream || m_cuda_launch_blocking; + while(block) { + pthread_mutex_lock(&m_lock); + block = !concurrent_streams_empty(); + pthread_mutex_unlock(&m_lock); + }; + + pthread_mutex_lock(&m_lock); + if( stream && !m_cuda_launch_blocking ) { + stream->push(op); + } else { + op.set_stream(&m_stream_zero); + m_stream_zero.push(op); + } + print_impl(stdout); + pthread_mutex_unlock(&m_lock); + if( m_cuda_launch_blocking || stream == NULL ) { + while( !empty() ) + ; + } +} + diff --git a/src/stream_manager.h b/src/stream_manager.h new file mode 100644 index 0000000..d47f0c7 --- /dev/null +++ b/src/stream_manager.h @@ -0,0 +1,294 @@ +/* + * stream_manager.h + * + * Copyright © 2009 by Tor M. Aamodt, Wilson W. L. Fung, Ali Bakhoda, + * George L. Yuan and the University of British Columbia, Vancouver, + * BC V6T 1Z4, All Rights Reserved. + * + * THIS IS A LEGAL DOCUMENT BY DOWNLOADING GPGPU-SIM, YOU ARE AGREEING TO THESE + * TERMS AND CONDITIONS. + * + * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "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 THE COPYRIGHT OWNERS OR CONTRIBUTORS 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. + * + * NOTE: The files libcuda/cuda_runtime_api.c and src/cuda-sim/cuda-math.h + * are derived from the CUDA Toolset available from http://www.nvidia.com/cuda + * (property of NVIDIA). The files benchmarks/BlackScholes/ and + * benchmarks/template/ are derived from the CUDA SDK available from + * http://www.nvidia.com/cuda (also property of NVIDIA). The files from + * src/intersim/ are derived from Booksim (a simulator provided with the + * textbook "Principles and Practices of Interconnection Networks" available + * from http://cva.stanford.edu/books/ppin/). As such, those files are bound by + * the corresponding legal terms and conditions set forth separately (original + * copyright notices are left in files from these sources and where we have + * modified a file our copyright notice appears before the original copyright + * notice). + * + * Using this version of GPGPU-Sim requires a complete installation of CUDA + * which is distributed seperately by NVIDIA under separate terms and + * conditions. To use this version of GPGPU-Sim with OpenCL requires a + * recent version of NVIDIA's drivers which support OpenCL. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions are met: + * + * 1. Redistributions of source code must retain the above copyright notice, + * this list of conditions and the following disclaimer. + * + * 2. 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. + * + * 3. Neither the name of the University of British Columbia nor the names of + * its contributors may be used to endorse or promote products derived from + * this software without specific prior written permission. + * + * 4. This version of GPGPU-SIM is distributed freely for non-commercial use only. + * + * 5. No nonprofit user may place any restrictions on the use of this software, + * including as modified by the user, by any other authorized user. + * + * 6. GPGPU-SIM was developed primarily by Tor M. Aamodt, Wilson W. L. Fung, + * Ali Bakhoda, George L. Yuan, at the University of British Columbia, + * Vancouver, BC V6T 1Z4 + */ + +#ifndef STREAM_MANAGER_H_INCLUDED +#define STREAM_MANAGER_H_INCLUDED + +#include "abstract_hardware_model.h" +#include <list> +#include <pthread.h> +#include <time.h> + +//class stream_barrier { +//public: +// stream_barrier() { m_pending_streams=0; } +// void inc() { m_pending_streams++; } +// void dec() { assert(m_pending_streams); m_pending_streams--; } +// unsigned value() const { return m_pending_streams; } +//private: +// unsigned m_pending_streams; +//}; + +enum stream_operation_type { + stream_no_op, + stream_memcpy_host_to_device, + stream_memcpy_device_to_host, + stream_memcpy_device_to_device, + stream_memcpy_to_symbol, + stream_memcpy_from_symbol, + stream_kernel_launch, + stream_event +}; + +class stream_operation { +public: + stream_operation() + { + m_type = stream_no_op; + m_stream = NULL; + m_done=true; + } + stream_operation( const void *src, const char *symbol, size_t count, size_t offset, struct CUstream_st *stream ) + { + m_stream = stream; + m_type=stream_memcpy_to_symbol; + m_host_address_src=src; + m_symbol=symbol; + m_cnt=count; + m_offset=offset; + + m_done=false; + } + stream_operation( const char *symbol, void *dst, size_t count, size_t offset, struct CUstream_st *stream ) + { + m_stream = stream; + m_type=stream_memcpy_from_symbol; + m_host_address_dst=dst; + m_symbol=symbol; + m_cnt=count; + m_offset=offset; + m_done=false; + } + stream_operation( const kernel_info_t &kernel, bool sim_mode, struct CUstream_st *stream ) + { + m_type=stream_kernel_launch; + m_kernel=kernel; + m_sim_mode=sim_mode; + m_stream=stream; + m_done=false; + } + stream_operation( class CUevent_st *e, struct CUstream_st *stream ) + { + m_type=stream_event; + m_event=e; + m_stream=stream; + m_done=false; + } + stream_operation( const void *host_address_src, size_t device_address_dst, size_t cnt, struct CUstream_st *stream ) + { + m_type=stream_memcpy_host_to_device; + m_host_address_src =host_address_src; + m_device_address_dst=device_address_dst; + m_host_address_dst=NULL; + m_device_address_src=0; + m_cnt=cnt; + m_stream=stream; + m_sim_mode=false; + m_done=false; + } + stream_operation( size_t device_address_src, void *host_address_dst, size_t cnt, struct CUstream_st *stream ) + { + m_type=stream_memcpy_device_to_host; + m_device_address_src=device_address_src; + m_host_address_dst=host_address_dst; + m_device_address_dst=0; + m_host_address_src=NULL; + m_cnt=cnt; + m_stream=stream; + m_sim_mode=false; + m_done=false; + } + stream_operation( size_t device_address_src, size_t device_address_dst, size_t cnt, struct CUstream_st *stream ) + { + m_type=stream_memcpy_device_to_device; + m_device_address_src=device_address_src; + m_device_address_dst=device_address_dst; + m_host_address_src=NULL; + m_host_address_dst=NULL; + m_cnt=cnt; + m_stream=stream; + m_sim_mode=false; + m_done=false; + } + + //void add_barrier( stream_barrier *b ) { m_barrier=b; } + //stream_barrier *get_barrier() { return m_barrier; } + bool is_kernel() const { return m_type == stream_kernel_launch; } + bool is_mem() const { + return m_type == stream_memcpy_host_to_device || + m_type == stream_memcpy_device_to_host || + m_type == stream_memcpy_host_to_device; + } + bool is_noop() const { return m_type == stream_no_op; } + bool is_done() const { return m_done; } + kernel_info_t &get_kernel() { return m_kernel; } + void do_operation( gpgpu_sim *gpu ); + void print( FILE *fp ) const; + struct CUstream_st *get_stream() { return m_stream; } + void set_stream( CUstream_st *stream ) { m_stream = stream; } + +private: + struct CUstream_st *m_stream; + + bool m_done; + + stream_operation_type m_type; + size_t m_device_address_dst; + size_t m_device_address_src; + void *m_host_address_dst; + const void *m_host_address_src; + size_t m_cnt; + + const char *m_symbol; + size_t m_offset; + + bool m_sim_mode; + kernel_info_t m_kernel; + class CUevent_st *m_event; + + //stream_barrier *m_barrier; +}; + +class CUevent_st { +public: + CUevent_st( bool blocking ) + { + m_uid = ++m_next_event_uid; + m_blocking = blocking; + m_updates = 0; + m_wallclock = 0; + m_gpu_tot_sim_cycle = 0; + m_done = false; + } + void update( double cycle, time_t clk ) + { + m_updates++; + m_wallclock=clk; + m_gpu_tot_sim_cycle=cycle; + m_done = true; + } + //void set_done() { assert(!m_done); m_done=true; } + int get_uid() const { return m_uid; } + unsigned num_updates() const { return m_updates; } + bool done() const { return m_done; } + time_t clock() const { return m_wallclock; } +private: + int m_uid; + bool m_blocking; + bool m_done; + int m_updates; + time_t m_wallclock; + double m_gpu_tot_sim_cycle; + + static int m_next_event_uid; +}; + +struct CUstream_st { +public: + CUstream_st(); + bool empty(); + bool busy(); + void synchronize(); + void push( const stream_operation &op ); + void record_next_done(); + stream_operation next(); + stream_operation &front() { return m_operations.front(); } + void print( FILE *fp ); + unsigned get_uid() const { return m_uid; } + +private: + unsigned m_uid; + static unsigned sm_next_stream_uid; + + std::list<stream_operation> m_operations; + bool m_pending; // front operation has started but not yet completed + + pthread_mutex_t m_lock; // ensure only one host or gpu manipulates stream operation at one time +}; + +class stream_manager { +public: + stream_manager( gpgpu_sim *gpu, bool cuda_launch_blocking ); + void register_finished_kernel( unsigned grid_uid ); + stream_operation front(); + void add_stream( CUstream_st *stream ); + void destroy_stream( CUstream_st *stream ); + bool concurrent_streams_empty(); + bool empty(); + void print( FILE *fp); + void push( stream_operation op ); + +private: + void print_impl( FILE *fp); + + bool m_cuda_launch_blocking; + gpgpu_sim *m_gpu; + std::list<CUstream_st *> m_streams; + std::map<unsigned,CUstream_st *> m_grid_id_to_stream; + CUstream_st m_stream_zero; + bool m_service_stream_zero; + pthread_mutex_t m_lock; +}; + +#endif |