From 582106ec595bb7745db092f2f4aa2b0fb9521b16 Mon Sep 17 00:00:00 2001 From: Jin Wang Date: Sat, 18 Oct 2014 23:07:53 -0400 Subject: MOD: add child kernel stream and scheduling support --- src/stream_manager.cc | 39 +++++++++++++++++++++------------------ 1 file changed, 21 insertions(+), 18 deletions(-) (limited to 'src/stream_manager.cc') diff --git a/src/stream_manager.cc b/src/stream_manager.cc index dd42f0a..1f93c03 100644 --- a/src/stream_manager.cc +++ b/src/stream_manager.cc @@ -155,7 +155,7 @@ void stream_operation::do_operation( gpgpu_sim *gpu ) gpu->set_cache_config(m_kernel->name()); 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 ); + gpu->functional_launch( m_kernel ); else gpu->launch( m_kernel ); } @@ -212,11 +212,9 @@ bool stream_manager::operation( bool * sim) bool stream_manager::check_finished_kernel() { - - unsigned grid_uid = m_gpu->finished_kernel(); - bool check=register_finished_kernel(grid_uid); - return check; - + unsigned grid_uid = m_gpu->finished_kernel(); + bool check=register_finished_kernel(grid_uid); + return check; } bool stream_manager::register_finished_kernel(unsigned grid_uid) @@ -226,13 +224,17 @@ bool stream_manager::register_finished_kernel(unsigned grid_uid) CUstream_st *stream = m_grid_id_to_stream[grid_uid]; kernel_info_t *kernel = stream->front().get_kernel(); assert( grid_uid == kernel->get_uid() ); - stream->record_next_done(); - m_grid_id_to_stream.erase(grid_uid); - delete kernel; - return true; - }else{ - return false; + + //Jin: should check children kernels for CDP + if(kernel->is_finished()) { + stream->record_next_done(); + m_grid_id_to_stream.erase(grid_uid); + kernel->notify_parent_finished(); + delete kernel; + return true; + } } + return false; } @@ -259,21 +261,22 @@ stream_operation stream_manager::front() { // called by gpu simulation thread stream_operation result; - if( concurrent_streams_empty() ) - m_service_stream_zero = true; +// if( concurrent_streams_empty() ) + m_service_stream_zero = true; if( m_service_stream_zero ) { - if( !m_stream_zero.empty() ) { - if( !m_stream_zero.busy() ) { + if( !m_stream_zero.empty() && !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 { + } + + if(!m_service_stream_zero) + { std::list::iterator s; for( s=m_streams.begin(); s != m_streams.end(); s++) { CUstream_st *stream = *s; -- cgit v1.3 From 6286547cfdc5d14c84568505db267f5e8dd9841f Mon Sep 17 00:00:00 2001 From: Jin Wang Date: Mon, 20 Oct 2014 23:23:00 -0400 Subject: BUG: multiple child kernels finish --- src/abstract_hardware_model.cc | 22 ++++++++++++++++++---- src/abstract_hardware_model.h | 2 ++ src/cuda-sim/cuda-sim.cc | 6 +----- src/stream_manager.cc | 3 ++- 4 files changed, 23 insertions(+), 10 deletions(-) (limited to 'src/stream_manager.cc') diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index b19b9d1..f3c5b21 100644 --- a/src/abstract_hardware_model.cc +++ b/src/abstract_hardware_model.cc @@ -595,6 +595,12 @@ void kernel_info_t::set_child(kernel_info_t * child) { m_child_kernels.push_back(child); } +void kernel_info_t::remove_child(kernel_info_t * child) { + assert(std::find(m_child_kernels.begin(), m_child_kernels.end(), child) + != m_child_kernels.end()); + m_child_kernels.remove(child); +} + bool kernel_info_t::is_finished() { if(done() && children_all_finished()) return true; @@ -603,16 +609,15 @@ bool kernel_info_t::is_finished() { } bool kernel_info_t::children_all_finished() { - for(auto child = m_child_kernels.begin(); child != m_child_kernels.end(); child++) - { - if(!(*child)->is_finished()) + if(!m_child_kernels.empty()) return false; - } + return true; } void kernel_info_t::notify_parent_finished() { if(m_parent_kernel) { + m_parent_kernel->remove_child(this); g_stream_manager->register_finished_kernel(m_parent_kernel->get_uid()); } } @@ -654,6 +659,15 @@ bool kernel_info_t::cta_has_stream(dim3 ctaid, CUstream_st* stream) { return true; } +void kernel_info_t::print_parent_info() { + if(m_parent_kernel) { + printf("Parent %d: \'%s\', Block (%d, %d, %d), Thread (%d, %d, %d)\n", + m_parent_kernel->get_uid(), m_parent_kernel->name().c_str(), + m_parent_ctaid.x, m_parent_ctaid.y, m_parent_ctaid.z, + m_parent_tid.x, m_parent_tid.y, m_parent_tid.z); + } +} + simt_stack::simt_stack( unsigned wid, unsigned warpSize) { diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index 11fee10..a0bf80d 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -275,12 +275,14 @@ public: //Jin: parent and child kernel management for CDP void set_parent(kernel_info_t * parent, dim3 parent_ctaid, dim3 parent_tid); void set_child(kernel_info_t * child); + void remove_child(kernel_info_t * child); bool is_finished(); bool children_all_finished(); void notify_parent_finished(); CUstream_st * create_stream_cta(dim3 ctaid); CUstream_st * get_default_stream_cta(dim3 ctaid); bool cta_has_stream(dim3 ctaid, CUstream_st* stream); + void print_parent_info(); private: kernel_info_t * m_parent_kernel; diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 58bd4e0..8dd1078 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1781,11 +1781,7 @@ void gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL ) //openCL kernel simulation calls don't register the kernel so we don't register its exit if(!openCL) { - - //Jin: for CDP, children should be finished first - if(kernel.is_finished()) { - g_stream_manager->register_finished_kernel(kernel.get_uid()); - } + g_stream_manager->register_finished_kernel(kernel.get_uid()); } //******PRINTING******* diff --git a/src/stream_manager.cc b/src/stream_manager.cc index 1f93c03..687d544 100644 --- a/src/stream_manager.cc +++ b/src/stream_manager.cc @@ -153,7 +153,8 @@ void stream_operation::do_operation( gpgpu_sim *gpu ) case stream_kernel_launch: if( gpu->can_start_kernel() ) { gpu->set_cache_config(m_kernel->name()); - printf("kernel \'%s\' transfer to GPU hardware scheduler\n", m_kernel->name().c_str() ); + printf("kernel %d: \'%s\' transfer to GPU hardware scheduler\n", m_kernel->get_uid(), m_kernel->name().c_str() ); + m_kernel->print_parent_info(); if( m_sim_mode ) gpu->functional_launch( m_kernel ); else -- cgit v1.3 From f372a4c641b9e6d38470ead6ae25743d26c5fed1 Mon Sep 17 00:00:00 2001 From: Jin Wang Date: Thu, 23 Oct 2014 15:33:59 -0400 Subject: BUG: kernels should return to stream if not pushed to concurrent kernel pool --- src/cuda-sim/cuda_device_runtime.cc | 8 +++++--- src/gpgpu-sim/shader.cc | 3 ++- src/stream_manager.cc | 37 ++++++++++++++++++++++++++++++++----- src/stream_manager.h | 3 ++- 4 files changed, 41 insertions(+), 10 deletions(-) (limited to 'src/stream_manager.cc') diff --git a/src/cuda-sim/cuda_device_runtime.cc b/src/cuda-sim/cuda_device_runtime.cc index be87f6a..c53ea04 100644 --- a/src/cuda-sim/cuda_device_runtime.cc +++ b/src/cuda-sim/cuda_device_runtime.cc @@ -122,7 +122,7 @@ void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info * addr_t from_addr = actual_param_op.get_symbol()->get_address(); if(arg == 0) {//paramter buffer for child kernel (in global memory) - //get parameter_buffer from the cudaDeviceLaunchV2_param0 + //get parameter_buffer from the cudaLaunchDeviceV2_param0 assert(size == sizeof(void *)); thread->m_local_mem->read(from_addr, size, ¶meter_buffer); assert((size_t)parameter_buffer >= GLOBAL_HEAP_START); @@ -152,11 +152,13 @@ void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info * kernel_info_t & parent_kernel = thread->get_kernel(); if(child_stream == 0) { //default stream on device for current CTA child_stream = parent_kernel.get_default_stream_cta(thread->get_ctaid()); - DEV_RUNTIME_REPORT("launching child kernel to default stream of the cta " << child_stream->get_uid() << ": " << child_stream); + DEV_RUNTIME_REPORT("launching child kernel " << child_grid->get_uid() << + " to default stream of the cta " << child_stream->get_uid() << ": " << child_stream); } else { assert(parent_kernel.cta_has_stream(thread->get_ctaid(), child_stream)); - DEV_RUNTIME_REPORT("launching child kernel to stream " << child_stream->get_uid() << ": " << child_stream); + DEV_RUNTIME_REPORT("launching child kernel " << child_grid->get_uid() << + " to stream " << child_stream->get_uid() << ": " << child_stream); } } diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc index ff2fac7..de5bcf6 100644 --- a/src/gpgpu-sim/shader.cc +++ b/src/gpgpu-sim/shader.cc @@ -1934,7 +1934,8 @@ void shader_core_ctx::register_cta_thread_exit( unsigned cta_num ) m_kernel->name().c_str() ); if( !m_gpu->kernel_more_cta_left(m_kernel) ) { if( !m_kernel->running() ) { - printf("GPGPU-Sim uArch: GPU detected kernel \'%s\' finished on shader %u.\n", m_kernel->name().c_str(), m_sid ); + printf("GPGPU-Sim uArch: GPU detected kernel %u \'%s\' finished on shader %u.\n", m_kernel->get_uid(), + m_kernel->name().c_str(), m_sid ); m_gpu->set_kernel_done( m_kernel ); } } diff --git a/src/stream_manager.cc b/src/stream_manager.cc index 687d544..546ab3b 100644 --- a/src/stream_manager.cc +++ b/src/stream_manager.cc @@ -95,6 +95,15 @@ stream_operation CUstream_st::next() return result; } +void CUstream_st::cancel_front() +{ + pthread_mutex_lock(&m_lock); + assert(m_pending); + m_pending = false; + pthread_mutex_unlock(&m_lock); + +} + void CUstream_st::print(FILE *fp) { pthread_mutex_lock(&m_lock); @@ -111,10 +120,10 @@ void CUstream_st::print(FILE *fp) } -void stream_operation::do_operation( gpgpu_sim *gpu ) +bool stream_operation::do_operation( gpgpu_sim *gpu ) { if( is_noop() ) - return; + return true; assert(!m_done && m_stream); if(g_debug_execution >= 3) @@ -153,16 +162,23 @@ void stream_operation::do_operation( gpgpu_sim *gpu ) case stream_kernel_launch: if( gpu->can_start_kernel() ) { gpu->set_cache_config(m_kernel->name()); - printf("kernel %d: \'%s\' transfer to GPU hardware scheduler\n", m_kernel->get_uid(), m_kernel->name().c_str() ); + if(g_debug_execution >= 3) + printf("kernel %d: \'%s\' transfer to GPU hardware scheduler\n", m_kernel->get_uid(), m_kernel->name().c_str() ); m_kernel->print_parent_info(); if( m_sim_mode ) gpu->functional_launch( m_kernel ); else gpu->launch( m_kernel ); } + else { + if(g_debug_execution >= 3) + printf("kernel %d: \'%s\' not ready to transfer to GPU hardware scheduler\n", m_kernel->get_uid(), m_kernel->name().c_str() ); + return false; + } break; case stream_event: { - printf("event update\n"); + if(g_debug_execution >= 3) + printf("event update\n"); time_t wallclock = time((time_t *)NULL); m_event->update( gpu_tot_sim_cycle, wallclock ); m_stream->record_next_done(); @@ -173,6 +189,7 @@ void stream_operation::do_operation( gpgpu_sim *gpu ) } m_done=true; fflush(stdout); + return true; } void stream_operation::print( FILE *fp ) const @@ -204,7 +221,16 @@ bool stream_manager::operation( bool * sim) bool check=check_finished_kernel(); if(check)m_gpu->print_stats(); stream_operation op =front(); - op.do_operation( m_gpu ); + if(!op.do_operation( m_gpu )) //not ready to execute + { + //cancel operation + if( op.is_kernel() ) { + unsigned grid_uid = op.get_kernel()->get_uid(); + m_grid_id_to_stream.erase(grid_uid); + } + op.get_stream()->cancel_front(); + + } pthread_mutex_unlock(&m_lock); //pthread_mutex_lock(&m_lock); // simulate a clock cycle on the GPU @@ -228,6 +254,7 @@ bool stream_manager::register_finished_kernel(unsigned grid_uid) //Jin: should check children kernels for CDP if(kernel->is_finished()) { + printf("kernel %d finishes, retires from stream %d\n", grid_uid, stream->get_uid()); stream->record_next_done(); m_grid_id_to_stream.erase(grid_uid); kernel->notify_parent_finished(); diff --git a/src/stream_manager.h b/src/stream_manager.h index 701b33c..222a1b2 100644 --- a/src/stream_manager.h +++ b/src/stream_manager.h @@ -150,7 +150,7 @@ public: 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 ); + bool 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; } @@ -218,6 +218,7 @@ public: void push( const stream_operation &op ); void record_next_done(); stream_operation next(); + void cancel_front(); //front operation fails, cancle the pending status stream_operation &front() { return m_operations.front(); } void print( FILE *fp ); unsigned get_uid() const { return m_uid; } -- cgit v1.3 From 3656376ccd83d5c514389c4c3818f1969bc30e0c Mon Sep 17 00:00:00 2001 From: Jin Wang Date: Fri, 1 Jul 2016 05:29:58 -0400 Subject: MOD: schedule one child kernel each cycle --- src/cuda-sim/cuda_device_runtime.cc | 214 ++++++++++++++++++++++++------------ src/cuda-sim/cuda_device_runtime.h | 2 + src/gpgpu-sim/gpu-sim.cc | 4 + src/stream_manager.cc | 2 +- 4 files changed, 148 insertions(+), 74 deletions(-) (limited to 'src/stream_manager.cc') diff --git a/src/cuda-sim/cuda_device_runtime.cc b/src/cuda-sim/cuda_device_runtime.cc index 2a90cba..df2653e 100644 --- a/src/cuda-sim/cuda_device_runtime.cc +++ b/src/cuda-sim/cuda_device_runtime.cc @@ -20,8 +20,45 @@ std::cout.flush(); \ } -std::map g_cuda_device_launch_map; -extern stream_manager * g_stream_manager; +class device_launch_config_t { + +public: + device_launch_config_t() {} + + device_launch_config_t(dim3 _grid_dim, + dim3 _block_dim, + unsigned int _shared_mem, + function_info * _entry): + grid_dim(_grid_dim), + block_dim(_block_dim), + shared_mem(_shared_mem), + entry(_entry) {} + + dim3 grid_dim; + dim3 block_dim; + unsigned int shared_mem; + function_info * entry; + +}; + +class device_launch_operation_t { + +public: + device_launch_operation_t() {} + device_launch_operation_t(kernel_info_t *_grid, + CUstream_st * _stream) : + grid(_grid), stream(_stream) {} + + kernel_info_t * grid; //a new child grid + + CUstream_st * stream; + +}; + + +std::map g_cuda_device_launch_param_map; +std::list g_cuda_device_launch_op; +extern stream_manager *g_stream_manager; //Handling device runtime api: //void * cudaGetParameterBufferV2(void *func, dim3 gridDimension, dim3 blockDimension, unsigned int sharedMemSize) @@ -35,8 +72,8 @@ void gpgpusim_cuda_getParameterBufferV2(const ptx_instruction * pI, ptx_thread_i assert( n_args == 4 ); function_info * child_kernel_entry; - struct dim3 gridDim, blockDim; - unsigned int sharedMem; + struct dim3 grid_dim, block_dim; + unsigned int shared_mem; for( unsigned arg=0; arg < n_args; arg ++ ) { const operand_info &actual_param_op = pI->operand_lookup(n_return+1+arg); //param# @@ -48,54 +85,48 @@ void gpgpusim_cuda_getParameterBufferV2(const ptx_instruction * pI, ptx_thread_i if(arg == 0) {//function_info* for the child kernel unsigned long long buf; - assert(size == sizeof(function_info *)); + assert(size == sizeof(function_info *)); thread->m_local_mem->read(from_addr, size, &buf); child_kernel_entry = (function_info *)buf; assert(child_kernel_entry); DEV_RUNTIME_REPORT("child kernel name " << child_kernel_entry->get_name()); } - else if(arg == 1) { //dim3 gridDim for the child kernel - assert(size == sizeof(struct dim3)); - thread->m_local_mem->read(from_addr, size, & gridDim); - DEV_RUNTIME_REPORT("grid (" << gridDim.x << ", " << gridDim.y << ", " << gridDim.z << ")"); + else if(arg == 1) { //dim3 grid_dim for the child kernel + assert(size == sizeof(struct dim3)); + thread->m_local_mem->read(from_addr, size, & grid_dim); + DEV_RUNTIME_REPORT("grid (" << grid_dim.x << ", " << grid_dim.y << ", " << grid_dim.z << ")"); } - else if(arg == 2) { //dim3 blockDim for the child kernel - assert(size == sizeof(struct dim3)); - thread->m_local_mem->read(from_addr, size, & blockDim); - DEV_RUNTIME_REPORT("block (" << blockDim.x << ", " << blockDim.y << ", " << blockDim.z << ")"); + else if(arg == 2) { //dim3 block_dim for the child kernel + assert(size == sizeof(struct dim3)); + thread->m_local_mem->read(from_addr, size, & block_dim); + DEV_RUNTIME_REPORT("block (" << block_dim.x << ", " << block_dim.y << ", " << block_dim.z << ")"); } - else if(arg == 3) { //unsigned int sharedMem - assert(size == sizeof(unsigned int)); - thread->m_local_mem->read(from_addr, size, & sharedMem); - DEV_RUNTIME_REPORT("shared memory " << sharedMem); + else if(arg == 3) { //unsigned int shared_mem + assert(size == sizeof(unsigned int)); + thread->m_local_mem->read(from_addr, size, & shared_mem); + DEV_RUNTIME_REPORT("shared memory " << shared_mem); } } - //get total child kernel argument size and malloc buffer in global memory - unsigned child_kernel_arg_size = child_kernel_entry->get_args_aligned_size(); - void * param_buffer = thread->get_gpu()->gpu_malloc(child_kernel_arg_size); - DEV_RUNTIME_REPORT("child kernel arg size total " << child_kernel_arg_size << ", parameter buffer allocated at " << param_buffer); - - //create child kernel_info_t and index it with parameter_buffer address - kernel_info_t * child_grid = new kernel_info_t(gridDim, blockDim, child_kernel_entry); - kernel_info_t & parent_grid = thread->get_kernel(); - DEV_RUNTIME_REPORT("child kernel launched by " << parent_grid.name() << ", cta (" << - thread->get_ctaid().x << ", " << thread->get_ctaid().y << ", " << thread->get_ctaid().z << - "), thread (" << thread->get_tid().x << ", " << thread->get_tid().y << ", " << thread->get_tid().z << - ")"); - child_grid->set_parent(&parent_grid, thread->get_ctaid(), thread->get_tid()); - assert(g_cuda_device_launch_map.find(param_buffer) == g_cuda_device_launch_map.end()); - g_cuda_device_launch_map[param_buffer] = child_grid; - - //copy the buffer address to retval0 + //get total child kernel argument size and malloc buffer in global memory + unsigned child_kernel_arg_size = child_kernel_entry->get_args_aligned_size(); + void * param_buffer = thread->get_gpu()->gpu_malloc(child_kernel_arg_size); + DEV_RUNTIME_REPORT("child kernel arg size total " << child_kernel_arg_size << ", parameter buffer allocated at " << param_buffer); + + //store param buffer address and launch config + device_launch_config_t device_launch_config(grid_dim, block_dim, shared_mem, child_kernel_entry); + assert(g_cuda_device_launch_param_map.find(param_buffer) == g_cuda_device_launch_param_map.end()); + g_cuda_device_launch_param_map[param_buffer] = device_launch_config; + + //copy the buffer address to retval0 const operand_info &actual_return_op = pI->operand_lookup(0); //retval0 const symbol *formal_return = target_func->get_return_var(); //void * - unsigned int return_size = formal_return->get_size_in_bytes(); - DEV_RUNTIME_REPORT("cudaGetParameterBufferV2 return value has size of " << return_size); - assert(actual_return_op.is_param_local()); - assert(actual_return_op.get_symbol()->get_size_in_bytes() == return_size && return_size == sizeof(void *)); + unsigned int return_size = formal_return->get_size_in_bytes(); + DEV_RUNTIME_REPORT("cudaGetParameterBufferV2 return value has size of " << return_size); + assert(actual_return_op.is_param_local()); + assert(actual_return_op.get_symbol()->get_size_in_bytes() == return_size && return_size == sizeof(void *)); addr_t ret_param_addr = actual_return_op.get_symbol()->get_address(); - thread->m_local_mem->write(ret_param_addr, return_size, ¶m_buffer, NULL, NULL); + thread->m_local_mem->write(ret_param_addr, return_size, ¶m_buffer, NULL, NULL); } @@ -109,10 +140,13 @@ void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info * unsigned n_args = target_func->num_args(); assert( n_args == 2 ); - kernel_info_t * child_grid = NULL; - function_info * child_kernel_entry = NULL; + kernel_info_t * device_grid = NULL; + function_info * device_kernel_entry = NULL; void * parameter_buffer; struct CUstream_st * child_stream; + device_launch_config_t config; + device_launch_operation_t device_launch_op; + for( unsigned arg=0; arg < n_args; arg ++ ) { const operand_info &actual_param_op = pI->operand_lookup(n_return+1+arg); //param# const symbol *formal_param = target_func->get_arg(arg); //cudaLaunchDeviceV2_param_# @@ -123,64 +157,80 @@ void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info * if(arg == 0) {//paramter buffer for child kernel (in global memory) //get parameter_buffer from the cudaLaunchDeviceV2_param0 - assert(size == sizeof(void *)); + assert(size == sizeof(void *)); thread->m_local_mem->read(from_addr, size, ¶meter_buffer); assert((size_t)parameter_buffer >= GLOBAL_HEAP_START); DEV_RUNTIME_REPORT("Parameter buffer locating at global memory " << parameter_buffer); //get child grid info through parameter_buffer address - assert(g_cuda_device_launch_map.find(parameter_buffer) != g_cuda_device_launch_map.end()); - child_grid = g_cuda_device_launch_map[parameter_buffer]; - child_kernel_entry = child_grid->entry(); - DEV_RUNTIME_REPORT("find child kernel " << child_kernel_entry->get_name()); - - //copy data in parameter_buffer to child kernel param memory - unsigned child_kernel_arg_size = child_kernel_entry->get_args_aligned_size(); - DEV_RUNTIME_REPORT("child_kernel_arg_size " << child_kernel_arg_size); - memory_space *child_kernel_param_mem = child_grid->get_param_memory(); + assert(g_cuda_device_launch_param_map.find(parameter_buffer) != g_cuda_device_launch_param_map.end()); + config = g_cuda_device_launch_param_map[parameter_buffer]; + //device_grid = op.grid; + device_kernel_entry = config.entry; + DEV_RUNTIME_REPORT("find device kernel " << device_kernel_entry->get_name()); + + //copy data in parameter_buffer to device kernel param memory + unsigned device_kernel_arg_size = device_kernel_entry->get_args_aligned_size(); + DEV_RUNTIME_REPORT("device_kernel_arg_size " << device_kernel_arg_size); + memory_space *device_kernel_param_mem; + + //create child kernel_info_t and index it with parameter_buffer address + device_grid = new kernel_info_t(config.grid_dim, config.block_dim, device_kernel_entry); + kernel_info_t & parent_grid = thread->get_kernel(); + DEV_RUNTIME_REPORT("child kernel launched by " << parent_grid.name() << ", cta (" << + thread->get_ctaid().x << ", " << thread->get_ctaid().y << ", " << thread->get_ctaid().z << + "), thread (" << thread->get_tid().x << ", " << thread->get_tid().y << ", " << thread->get_tid().z << + ")"); + device_grid->set_parent(&parent_grid, thread->get_ctaid(), thread->get_tid()); + device_launch_op = device_launch_operation_t(device_grid, NULL); + device_kernel_param_mem = device_grid->get_param_memory(); //kernel param + size_t param_start_address = 0; //copy in word - for(unsigned n = 0; n < child_kernel_arg_size; n += 4) { + for(unsigned n = 0; n < device_kernel_arg_size; n += 4) { unsigned int oneword; thread->get_gpu()->get_global_memory()->read((size_t)parameter_buffer + n, 4, &oneword); - child_kernel_param_mem->write(param_start_address + n, 4, &oneword, NULL, NULL); + device_kernel_param_mem->write(param_start_address + n, 4, &oneword, NULL, NULL); } } else if(arg == 1) { //cudaStream for the child kernel - assert(size == sizeof(cudaStream_t)); - thread->m_local_mem->read(from_addr, size, &child_stream); + assert(size == sizeof(cudaStream_t)); + thread->m_local_mem->read(from_addr, size, &child_stream); + kernel_info_t & parent_kernel = thread->get_kernel(); if(child_stream == 0) { //default stream on device for current CTA child_stream = parent_kernel.get_default_stream_cta(thread->get_ctaid()); - DEV_RUNTIME_REPORT("launching child kernel " << child_grid->get_uid() << + DEV_RUNTIME_REPORT("launching child kernel " << device_grid->get_uid() << " to default stream of the cta " << child_stream->get_uid() << ": " << child_stream); } else { assert(parent_kernel.cta_has_stream(thread->get_ctaid(), child_stream)); - DEV_RUNTIME_REPORT("launching child kernel " << child_grid->get_uid() << + DEV_RUNTIME_REPORT("launching child kernel " << device_grid->get_uid() << " to stream " << child_stream->get_uid() << ": " << child_stream); } + + device_launch_op.stream = child_stream; + } } - } + //launch child kernel - stream_operation op(child_grid, g_ptx_sim_mode, child_stream); - g_stream_manager->push(op); - g_cuda_device_launch_map.erase(parameter_buffer); + g_cuda_device_launch_op.push_back(device_launch_op); + g_cuda_device_launch_param_map.erase(parameter_buffer); //set retval0 const operand_info &actual_return_op = pI->operand_lookup(0); //retval0 const symbol *formal_return = target_func->get_return_var(); //cudaError_t - unsigned int return_size = formal_return->get_size_in_bytes(); - DEV_RUNTIME_REPORT("cudaLaunchDeviceV2 return value has size of " << return_size); - assert(actual_return_op.is_param_local()); - assert(actual_return_op.get_symbol()->get_size_in_bytes() == return_size + unsigned int return_size = formal_return->get_size_in_bytes(); + DEV_RUNTIME_REPORT("cudaLaunchDeviceV2 return value has size of " << return_size); + assert(actual_return_op.is_param_local()); + assert(actual_return_op.get_symbol()->get_size_in_bytes() == return_size && return_size == sizeof(cudaError_t)); cudaError_t error = cudaSuccess; addr_t ret_param_addr = actual_return_op.get_symbol()->get_address(); - thread->m_local_mem->write(ret_param_addr, return_size, &error, NULL, NULL); + thread->m_local_mem->write(ret_param_addr, return_size, &error, NULL, NULL); } @@ -208,7 +258,7 @@ void gpgpusim_cuda_streamCreateWithFlags(const ptx_instruction * pI, ptx_thread_ addr_t from_addr = actual_param_op.get_symbol()->get_address(); if(arg == 0) {//cudaStream_t * pStream, address of cudaStream_t - assert(size == sizeof(cudaStream_t *)); + assert(size == sizeof(cudaStream_t *)); thread->m_local_mem->read(from_addr, size, &generic_pStream_addr); //pStream should be non-zero address in local memory @@ -217,7 +267,7 @@ void gpgpusim_cuda_streamCreateWithFlags(const ptx_instruction * pI, ptx_thread_ DEV_RUNTIME_REPORT("pStream locating at local memory " << pStream_addr); } else if(arg == 1) { //unsigned int flags, should be cudaStreamNonBlocking - assert(size == sizeof(unsigned int)); + assert(size == sizeof(unsigned int)); thread->m_local_mem->read(from_addr, size, &flags); assert(flags == cudaStreamNonBlocking); } @@ -231,13 +281,31 @@ void gpgpusim_cuda_streamCreateWithFlags(const ptx_instruction * pI, ptx_thread_ //set retval0 const operand_info &actual_return_op = pI->operand_lookup(0); //retval0 const symbol *formal_return = target_func->get_return_var(); //cudaError_t - unsigned int return_size = formal_return->get_size_in_bytes(); - DEV_RUNTIME_REPORT("cudaStreamCreateWithFlags return value has size of " << return_size); - assert(actual_return_op.is_param_local()); - assert(actual_return_op.get_symbol()->get_size_in_bytes() == return_size + unsigned int return_size = formal_return->get_size_in_bytes(); + DEV_RUNTIME_REPORT("cudaStreamCreateWithFlags return value has size of " << return_size); + assert(actual_return_op.is_param_local()); + assert(actual_return_op.get_symbol()->get_size_in_bytes() == return_size && return_size == sizeof(cudaError_t)); cudaError_t error = cudaSuccess; addr_t ret_param_addr = actual_return_op.get_symbol()->get_address(); - thread->m_local_mem->write(ret_param_addr, return_size, &error, NULL, NULL); + thread->m_local_mem->write(ret_param_addr, return_size, &error, NULL, NULL); } + + +void launch_one_device_kernel() { + if(!g_cuda_device_launch_op.empty()) { + device_launch_operation_t &op = g_cuda_device_launch_op.front(); + + stream_operation stream_op = stream_operation(op.grid, g_ptx_sim_mode, op.stream); + g_stream_manager->push(stream_op); + g_cuda_device_launch_op.pop_front(); + } +} + +void launch_all_device_kernels() { + while(!g_cuda_device_launch_op.empty()) { + launch_one_device_kernel(); + } +} + diff --git a/src/cuda-sim/cuda_device_runtime.h b/src/cuda-sim/cuda_device_runtime.h index 7447d5a..385d605 100644 --- a/src/cuda-sim/cuda_device_runtime.h +++ b/src/cuda-sim/cuda_device_runtime.h @@ -6,3 +6,5 @@ void gpgpusim_cuda_getParameterBufferV2(const ptx_instruction * pI, ptx_thread_info * thread, const function_info * target_func); void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info * thread, const function_info * target_func); void gpgpusim_cuda_streamCreateWithFlags(const ptx_instruction * pI, ptx_thread_info * thread, const function_info * target_func); +void launch_all_device_kernels(); +void launch_one_device_kernel(); diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 47dbc89..126e007 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -61,6 +61,7 @@ #include "power_stat.h" #include "visualizer.h" #include "stats.h" +#include "../cuda-sim/cuda_device_runtime.h" #ifdef GPGPUSIM_POWER_MODEL #include "power_interface.h" @@ -1370,6 +1371,9 @@ void gpgpu_sim::cycle() } try_snap_shot(gpu_sim_cycle); spill_log_to_file (stdout, 0, gpu_sim_cycle); + + //launch device kernel + launch_one_device_kernel(); } } diff --git a/src/stream_manager.cc b/src/stream_manager.cc index 546ab3b..f90d9be 100644 --- a/src/stream_manager.cc +++ b/src/stream_manager.cc @@ -217,8 +217,8 @@ stream_manager::stream_manager( gpgpu_sim *gpu, bool cuda_launch_blocking ) bool stream_manager::operation( bool * sim) { - pthread_mutex_lock(&m_lock); bool check=check_finished_kernel(); + pthread_mutex_lock(&m_lock); if(check)m_gpu->print_stats(); stream_operation op =front(); if(!op.do_operation( m_gpu )) //not ready to execute -- cgit v1.3 From a31392370bc7b23e115160b170e95b117597ebc4 Mon Sep 17 00:00:00 2001 From: Jin Wang Date: Sat, 15 Nov 2014 02:09:24 -0500 Subject: ADD: add stats for kernel launching and complete cycle --- src/abstract_hardware_model.h | 6 ++++++ src/cuda-sim/cuda_device_runtime.cc | 2 +- src/gpgpu-sim/gpu-sim.cc | 2 ++ src/stream_manager.cc | 9 ++++++++- 4 files changed, 17 insertions(+), 2 deletions(-) (limited to 'src/stream_manager.cc') diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index 45334b6..e6ef521 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -284,6 +284,7 @@ public: bool cta_has_stream(dim3 ctaid, CUstream_st* stream); void destroy_cta_streams(); void print_parent_info(); + kernel_info_t * get_parent() { return m_parent_kernel; } private: kernel_info_t * m_parent_kernel; @@ -292,6 +293,11 @@ private: std::list m_child_kernels; //child kernel launched std::map< dim3, std::list, dim3comp > m_cta_streams; //streams created in each CTA +//Jin: kernel timing +public: + unsigned long long launch_cycle; + unsigned long long start_cycle; + unsigned long long end_cycle; }; struct core_config { diff --git a/src/cuda-sim/cuda_device_runtime.cc b/src/cuda-sim/cuda_device_runtime.cc index df2653e..1b8c8d9 100644 --- a/src/cuda-sim/cuda_device_runtime.cc +++ b/src/cuda-sim/cuda_device_runtime.cc @@ -176,6 +176,7 @@ void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info * //create child kernel_info_t and index it with parameter_buffer address device_grid = new kernel_info_t(config.grid_dim, config.block_dim, device_kernel_entry); + device_grid->launch_cycle = gpu_sim_cycle; kernel_info_t & parent_grid = thread->get_kernel(); DEV_RUNTIME_REPORT("child kernel launched by " << parent_grid.name() << ", cta (" << thread->get_ctaid().x << ", " << thread->get_ctaid().y << ", " << thread->get_ctaid().z << @@ -184,7 +185,6 @@ void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info * device_grid->set_parent(&parent_grid, thread->get_ctaid(), thread->get_tid()); device_launch_op = device_launch_operation_t(device_grid, NULL); device_kernel_param_mem = device_grid->get_param_memory(); //kernel param - size_t param_start_address = 0; //copy in word for(unsigned n = 0; n < device_kernel_arg_size; n += 4) { diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 81c9c9a..27362cf 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -531,6 +531,7 @@ kernel_info_t *gpgpu_sim::select_kernel() unsigned idx = (n+m_last_issued_kernel+1)%m_config.max_concurrent_kernel; if( kernel_more_cta_left(m_running_kernels[idx]) ){ m_last_issued_kernel=idx; + m_running_kernels[idx]->start_cycle = gpu_sim_cycle; // record this kernel for stat print if it is the first time this kernel is selected for execution unsigned launch_uid = m_running_kernels[idx]->get_uid(); if (std::find(m_executed_kernel_uids.begin(), m_executed_kernel_uids.end(), launch_uid) == m_executed_kernel_uids.end()) { @@ -560,6 +561,7 @@ void gpgpu_sim::set_kernel_done( kernel_info_t *kernel ) std::vector::iterator k; for( k=m_running_kernels.begin(); k!=m_running_kernels.end(); k++ ) { if( *k == kernel ) { + kernel->end_cycle = gpu_sim_cycle; *k = NULL; break; } diff --git a/src/stream_manager.cc b/src/stream_manager.cc index f90d9be..5bd7737 100644 --- a/src/stream_manager.cc +++ b/src/stream_manager.cc @@ -254,7 +254,14 @@ bool stream_manager::register_finished_kernel(unsigned grid_uid) //Jin: should check children kernels for CDP if(kernel->is_finished()) { - printf("kernel %d finishes, retires from stream %d\n", grid_uid, stream->get_uid()); +// std::ofstream kernel_stat("kernel_stat.txt", std::ofstream::out | std::ofstream::app); +// kernel_stat<< " kernel " << grid_uid; +// if(kernel->get_parent()) +// kernel_stat << ", parent " << kernel->get_parent()->get_uid() << +// ", launch " << kernel->launch_cycle; +// kernel_stat<< ", start " << kernel->start_cycle << +// ", end " << kernel->end_cycle << ", retire " << gpu_sim_cycle << "\n"; +// printf("kernel %d finishes, retires from stream %d\n", grid_uid, stream->get_uid()); stream->record_next_done(); m_grid_id_to_stream.erase(grid_uid); kernel->notify_parent_finished(); -- cgit v1.3 From b70a5a69fa14cc01f707b910f8f021e36067922f Mon Sep 17 00:00:00 2001 From: Jin Wang Date: Sat, 15 Nov 2014 05:44:03 -0500 Subject: ADD: print kernel parameter size footprint. BUG: concurrent kernels on same shader, should use hw_cta_id to store shared mem info --- src/abstract_hardware_model.cc | 2 ++ src/cuda-sim/cuda-sim.cc | 3 ++- src/cuda-sim/cuda_device_runtime.cc | 7 ++++++- src/gpgpu-sim/gpu-sim.cc | 23 +++++++++++++++-------- src/stream_manager.cc | 6 ++++-- 5 files changed, 29 insertions(+), 12 deletions(-) (limited to 'src/stream_manager.cc') diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index 4db5f2f..be5c5b9 100644 --- a/src/abstract_hardware_model.cc +++ b/src/abstract_hardware_model.cc @@ -618,6 +618,8 @@ bool kernel_info_t::children_all_finished() { void kernel_info_t::notify_parent_finished() { if(m_parent_kernel) { + extern unsigned long long g_total_param_size; + g_total_param_size -= ((m_kernel_entry->get_args_aligned_size() + 255)/256*256); m_parent_kernel->remove_child(this); g_stream_manager->register_finished_kernel(m_parent_kernel->get_uid()); } diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index c87e3e4..3f5af7e 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1499,7 +1499,8 @@ unsigned ptx_sim_init_thread( kernel_info_t &kernel, unsigned max_cta_per_sm = num_threads/cta_size; // e.g., 256 / 48 = 5 assert( max_cta_per_sm > 0 ); - unsigned sm_idx = (tid/cta_size)*gpgpu_param_num_shaders + sid; + //unsigned sm_idx = (tid/cta_size)*gpgpu_param_num_shaders + sid; + unsigned sm_idx = hw_cta_id*gpgpu_param_num_shaders + sid; if ( shared_memory_lookup.find(sm_idx) == shared_memory_lookup.end() ) { if ( g_debug_execution >= 1 ) { diff --git a/src/cuda-sim/cuda_device_runtime.cc b/src/cuda-sim/cuda_device_runtime.cc index 1b8c8d9..66cd063 100644 --- a/src/cuda-sim/cuda_device_runtime.cc +++ b/src/cuda-sim/cuda_device_runtime.cc @@ -59,6 +59,8 @@ public: std::map g_cuda_device_launch_param_map; std::list g_cuda_device_launch_op; extern stream_manager *g_stream_manager; +unsigned long long g_total_param_size = 0; +unsigned long long g_max_total_param_size = 0; //Handling device runtime api: //void * cudaGetParameterBufferV2(void *func, dim3 gridDimension, dim3 blockDimension, unsigned int sharedMemSize) @@ -111,7 +113,10 @@ void gpgpusim_cuda_getParameterBufferV2(const ptx_instruction * pI, ptx_thread_i //get total child kernel argument size and malloc buffer in global memory unsigned child_kernel_arg_size = child_kernel_entry->get_args_aligned_size(); void * param_buffer = thread->get_gpu()->gpu_malloc(child_kernel_arg_size); + g_total_param_size += ((child_kernel_arg_size + 255) / 256 * 256); DEV_RUNTIME_REPORT("child kernel arg size total " << child_kernel_arg_size << ", parameter buffer allocated at " << param_buffer); + if(g_total_param_size > g_max_total_param_size) + g_max_total_param_size = g_total_param_size; //store param buffer address and launch config device_launch_config_t device_launch_config(grid_dim, block_dim, shared_mem, child_kernel_entry); @@ -176,7 +181,7 @@ void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info * //create child kernel_info_t and index it with parameter_buffer address device_grid = new kernel_info_t(config.grid_dim, config.block_dim, device_kernel_entry); - device_grid->launch_cycle = gpu_sim_cycle; + device_grid->launch_cycle = gpu_sim_cycle + gpu_tot_sim_cycle; kernel_info_t & parent_grid = thread->get_kernel(); DEV_RUNTIME_REPORT("child kernel launched by " << parent_grid.name() << ", cta (" << thread->get_ctaid().x << ", " << thread->get_ctaid().y << ", " << thread->get_ctaid().z << diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index 27362cf..a9da1c9 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -524,20 +524,26 @@ bool gpgpu_sim::get_more_cta_left() const kernel_info_t *gpgpu_sim::select_kernel() { if(m_running_kernels[m_last_issued_kernel] && - !m_running_kernels[m_last_issued_kernel]->no_more_ctas_to_run()) + !m_running_kernels[m_last_issued_kernel]->no_more_ctas_to_run()) { + unsigned launch_uid = m_running_kernels[m_last_issued_kernel]->get_uid(); + if(std::find(m_executed_kernel_uids.begin(), m_executed_kernel_uids.end(), launch_uid) == m_executed_kernel_uids.end()) { + m_running_kernels[m_last_issued_kernel]->start_cycle = gpu_sim_cycle + gpu_tot_sim_cycle; + m_executed_kernel_uids.push_back(launch_uid); + m_executed_kernel_names.push_back(m_running_kernels[m_last_issued_kernel]->name()); + } return m_running_kernels[m_last_issued_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( kernel_more_cta_left(m_running_kernels[idx]) ){ m_last_issued_kernel=idx; - m_running_kernels[idx]->start_cycle = gpu_sim_cycle; + m_running_kernels[idx]->start_cycle = gpu_sim_cycle + gpu_tot_sim_cycle; // record this kernel for stat print if it is the first time this kernel is selected for execution unsigned launch_uid = m_running_kernels[idx]->get_uid(); - if (std::find(m_executed_kernel_uids.begin(), m_executed_kernel_uids.end(), launch_uid) == m_executed_kernel_uids.end()) { - m_executed_kernel_uids.push_back(launch_uid); - m_executed_kernel_names.push_back(m_running_kernels[idx]->name()); - } + assert(std::find(m_executed_kernel_uids.begin(), m_executed_kernel_uids.end(), launch_uid) == m_executed_kernel_uids.end()); + m_executed_kernel_uids.push_back(launch_uid); + m_executed_kernel_names.push_back(m_running_kernels[idx]->name()); return m_running_kernels[idx]; } @@ -561,7 +567,7 @@ void gpgpu_sim::set_kernel_done( kernel_info_t *kernel ) std::vector::iterator k; for( k=m_running_kernels.begin(); k!=m_running_kernels.end(); k++ ) { if( *k == kernel ) { - kernel->end_cycle = gpu_sim_cycle; + kernel->end_cycle = gpu_sim_cycle + gpu_tot_sim_cycle; *k = NULL; break; } @@ -942,7 +948,8 @@ void gpgpu_sim::gpu_print_stat() printf("gpu_tot_ipc = %12.4f\n", (float)(gpu_tot_sim_insn+gpu_sim_insn) / (gpu_tot_sim_cycle+gpu_sim_cycle)); printf("gpu_tot_issued_cta = %lld\n", gpu_tot_issued_cta + m_total_cta_launched); - + extern unsigned long long g_max_total_param_size; + fprintf(statfout, "max_total_param_size = %llu\n", g_max_total_param_size); // performance counter for stalls due to congestion. printf("gpu_stall_dramfull = %d\n", gpu_stall_dramfull); diff --git a/src/stream_manager.cc b/src/stream_manager.cc index 5bd7737..3459c6d 100644 --- a/src/stream_manager.cc +++ b/src/stream_manager.cc @@ -219,7 +219,7 @@ bool stream_manager::operation( bool * sim) { bool check=check_finished_kernel(); pthread_mutex_lock(&m_lock); - if(check)m_gpu->print_stats(); +// if(check)m_gpu->print_stats(); stream_operation op =front(); if(!op.do_operation( m_gpu )) //not ready to execute { @@ -260,8 +260,10 @@ bool stream_manager::register_finished_kernel(unsigned grid_uid) // kernel_stat << ", parent " << kernel->get_parent()->get_uid() << // ", launch " << kernel->launch_cycle; // kernel_stat<< ", start " << kernel->start_cycle << -// ", end " << kernel->end_cycle << ", retire " << gpu_sim_cycle << "\n"; +// ", end " << kernel->end_cycle << ", retire " << gpu_sim_cycle + gpu_tot_sim_cycle << "\n"; // printf("kernel %d finishes, retires from stream %d\n", grid_uid, stream->get_uid()); +// kernel_stat.flush(); +// kernel_stat.close(); stream->record_next_done(); m_grid_id_to_stream.erase(grid_uid); kernel->notify_parent_finished(); -- cgit v1.3 From 60fa05d4de0f3c926f9ab4f687b5d0748ec19285 Mon Sep 17 00:00:00 2001 From: Jin Wang Date: Sun, 16 Nov 2014 01:27:47 -0500 Subject: ADD: add kernel launching latency from stream to distributor --- src/abstract_hardware_model.cc | 4 ++++ src/abstract_hardware_model.h | 1 + src/gpgpu-sim/gpu-sim.cc | 6 ++++++ src/gpgpusim_entrypoint.cc | 1 + src/stream_manager.cc | 39 +++++++++++++++++++++++++-------------- 5 files changed, 37 insertions(+), 14 deletions(-) (limited to 'src/stream_manager.cc') diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index be5c5b9..819ad35 100644 --- a/src/abstract_hardware_model.cc +++ b/src/abstract_hardware_model.cc @@ -550,6 +550,8 @@ void warp_inst_t::completed( unsigned long long cycle ) const ptx_file_line_stats_add_latency(pc, latency * active_count()); } +//Jin: kernel launch latency +unsigned g_kernel_launch_latency; unsigned kernel_info_t::m_next_uid = 1; @@ -569,6 +571,8 @@ kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info * //Jin: parent and child kernel management for CDP m_parent_kernel = NULL; + //Jin: launch latency management + m_launch_latency = g_kernel_launch_latency; } kernel_info_t::~kernel_info_t() diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index e6ef521..ab77d63 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -298,6 +298,7 @@ public: unsigned long long launch_cycle; unsigned long long start_cycle; unsigned long long end_cycle; + unsigned m_launch_latency; }; struct core_config { diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index a9da1c9..7fb9ab3 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -443,6 +443,12 @@ void gpgpu_sim_config::reg_options(option_parser_t opp) &Trace::sampling_memory_partition, "The memory partition which is printed using MEMPART_DPRINTF. Default -1 (i.e. all)", "-1"); ptx_file_line_stats_options(opp); + + //Jin: kernel launch latency + extern unsigned g_kernel_launch_latency; + option_parser_register(opp, "-gpgpu_kernel_launch_latency", OPT_INT32, + &g_kernel_launch_latency, "Kernel launch latency in cycles. Default: 0", + "0"); } ///////////////////////////////////////////////////////////////////////////// diff --git a/src/gpgpusim_entrypoint.cc b/src/gpgpusim_entrypoint.cc index fb17eed..04845e7 100644 --- a/src/gpgpusim_entrypoint.cc +++ b/src/gpgpusim_entrypoint.cc @@ -153,6 +153,7 @@ void *gpgpu_sim_thread_concurrent(void*) printf("GPGPU-Sim: ** STOP simulation thread (no work) **\n"); fflush(stdout); } + g_the_gpu->print_stats(); if(sim_cycles) { g_the_gpu->update_stats(); print_simulation_time(); diff --git a/src/stream_manager.cc b/src/stream_manager.cc index 3459c6d..3b6cbd5 100644 --- a/src/stream_manager.cc +++ b/src/stream_manager.cc @@ -160,20 +160,31 @@ bool stream_operation::do_operation( gpgpu_sim *gpu ) m_stream->record_next_done(); break; case stream_kernel_launch: - if( gpu->can_start_kernel() ) { - gpu->set_cache_config(m_kernel->name()); - if(g_debug_execution >= 3) - printf("kernel %d: \'%s\' transfer to GPU hardware scheduler\n", m_kernel->get_uid(), m_kernel->name().c_str() ); - m_kernel->print_parent_info(); - if( m_sim_mode ) - gpu->functional_launch( m_kernel ); - else - gpu->launch( m_kernel ); + if( m_sim_mode ) { //Functional Sim + if(g_debug_execution >= 3) { + printf("kernel %d: \'%s\' transfer to GPU hardware scheduler\n", m_kernel->get_uid(), m_kernel->name().c_str() ); + m_kernel->print_parent_info(); + } + gpu->set_cache_config(m_kernel->name()); + gpu->functional_launch( m_kernel ); } - else { - if(g_debug_execution >= 3) - printf("kernel %d: \'%s\' not ready to transfer to GPU hardware scheduler\n", m_kernel->get_uid(), m_kernel->name().c_str() ); - return false; + else { //Performance Sim + if( gpu->can_start_kernel() && m_kernel->m_launch_latency == 0) { + if(g_debug_execution >= 3) { + printf("kernel %d: \'%s\' transfer to GPU hardware scheduler\n", m_kernel->get_uid(), m_kernel->name().c_str() ); + m_kernel->print_parent_info(); + } + gpu->set_cache_config(m_kernel->name()); + gpu->launch( m_kernel ); + } + else { + if(m_kernel->m_launch_latency) + m_kernel->m_launch_latency--; + if(g_debug_execution >= 3) + printf("kernel %d: \'%s\', latency %u not ready to transfer to GPU hardware scheduler\n", + m_kernel->get_uid(), m_kernel->name().c_str(), m_kernel->m_launch_latency); + return false; + } } break; case stream_event: { @@ -255,7 +266,7 @@ bool stream_manager::register_finished_kernel(unsigned grid_uid) //Jin: should check children kernels for CDP if(kernel->is_finished()) { // std::ofstream kernel_stat("kernel_stat.txt", std::ofstream::out | std::ofstream::app); -// kernel_stat<< " kernel " << grid_uid; +// kernel_stat<< " kernel " << grid_uid << ": " << kernel->name(); // if(kernel->get_parent()) // kernel_stat << ", parent " << kernel->get_parent()->get_uid() << // ", launch " << kernel->launch_cycle; -- cgit v1.3