From 73050b730b31e9f58f8c2a0bfc4b3a8cd92be7ff Mon Sep 17 00:00:00 2001 From: Suchita Pati Date: Wed, 10 Apr 2019 23:27:42 -0500 Subject: Adding texture Fix --- src/abstract_hardware_model.cc | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) (limited to 'src/abstract_hardware_model.cc') diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index 307df40..3db87c9 100644 --- a/src/abstract_hardware_model.cc +++ b/src/abstract_hardware_model.cc @@ -688,7 +688,7 @@ unsigned g_kernel_launch_latency; unsigned kernel_info_t::m_next_uid = 1; -kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry ) +kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry, std::map NameToCudaArray, std::map NameToTexureInfo) { m_kernel_entry=entry; m_grid_dim=gridDim; @@ -708,6 +708,8 @@ kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info * m_launch_latency = g_kernel_launch_latency; volta_cache_config_set=false; + m_NameToCudaArray = NameToCudaArray; + m_NameToTexureInfo = NameToTexureInfo; } kernel_info_t::~kernel_info_t() -- cgit v1.3 From 69af057daf1999f17b81d761c7a616f67f56adbb Mon Sep 17 00:00:00 2001 From: Suchita Pati Date: Tue, 16 Apr 2019 23:36:02 -0500 Subject: Adding additional fixes for the texture bugs --- libcuda/cuda_runtime_api.cc | 22 ++++++++++++---------- src/abstract_hardware_model.cc | 13 ++++++++++--- src/abstract_hardware_model.h | 36 +++++++++++++++++++++--------------- src/cuda-sim/cuda-sim.cc | 7 +++++++ src/cuda-sim/instructions.cc | 8 +++++++- 5 files changed, 57 insertions(+), 29 deletions(-) (limited to 'src/abstract_hardware_model.cc') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index 7569b20..f6e5e6f 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -1121,10 +1121,10 @@ __host__ cudaError_t CUDARTAPI cudaDeviceGetAttribute(int *value, enum cudaDevic *value= 0; break; case 75: - *value= 9 ; + *value= 7 ; break; case 76: - *value= 3 ; + *value= 0 ; break; case 78: *value= 0 ; //TODO: as of now, we dont support stream priorities. @@ -1199,22 +1199,23 @@ __host__ cudaError_t CUDARTAPI cudaGetDevice(int *device) *device = g_active_device; return g_last_cudaError = cudaSuccess; } + __host__ cudaError_t CUDARTAPI cudaDeviceGetLimit ( size_t* pValue, cudaLimit limit ) { if(g_debug_execution >= 3){ announce_call(__my_func__); } switch(limit) { - case 0: + case 0: // cudaLimitStackSize *pValue=1024; break; - case 2: + case 2: // cudaLimitMallocHeapSize *pValue=8388608; break; - case 3: + case 3: // cudaLimitDevRuntimeSyncDepth *pValue=2; break; - case 4: + case 4: // cudaLimitDevRuntimePendingLaunchCount *pValue=2048; break; default: @@ -1225,7 +1226,6 @@ __host__ cudaError_t CUDARTAPI cudaDeviceGetLimit ( size_t* pValue, cudaLimit li } - __host__ cudaError_t CUDARTAPI cudaStreamGetPriority ( cudaStream_t hStream, int* priority ) { if(g_debug_execution >= 3){ @@ -1585,7 +1585,7 @@ __host__ cudaError_t CUDARTAPI cudaStreamDestroy(cudaStream_t stream) announce_call(__my_func__); } #if (CUDART_VERSION >= 3000) - //synchronization required for application using external libraries without explicit synchronization in the code to + //per-stream synchronization required for application using external libraries without explicit synchronization in the code to //avoid the stream_manager from spinning forever to destroy non-empty streams without making any forward progress. stream->synchronize(); g_stream_manager->destroy_stream(stream); @@ -2714,7 +2714,6 @@ cudaError_t cudaDeviceReset ( void ) { return g_last_cudaError = cudaSuccess; } cudaError_t CUDARTAPI cudaDeviceSynchronize(void){ - // I don't know what this should do if(g_debug_execution >= 3){ announce_call(__my_func__); } @@ -2722,7 +2721,6 @@ cudaError_t CUDARTAPI cudaDeviceSynchronize(void){ return g_last_cudaError = cudaSuccess; } - void CUDARTAPI __cudaRegisterFunction( void **fatCubinHandle, const char *hostFun, @@ -3333,6 +3331,10 @@ kernel_info_t *gpgpu_cuda_ptx_sim_init_grid( const char *hostFun, } function_info *entry = context->get_kernel(hostFun); gpgpu_t* gpu= context->get_device()->get_gpgpu(); + /* + Passing a snapshot of the GPU's current texture mapping to the kernel's info + as kernels should use texture bindings present at the time of their launch. + */ kernel_info_t *result = new kernel_info_t(gridDim,blockDim,entry,gpu->getNameArrayMapping(),gpu->getNameInfoMapping()); if( entry == NULL ) { printf("GPGPU-Sim PTX: ERROR launching kernel -- no PTX implementation found for %p\n", hostFun); diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index 3db87c9..820f5fd 100644 --- a/src/abstract_hardware_model.cc +++ b/src/abstract_hardware_model.cc @@ -192,6 +192,9 @@ gpgpu_t::gpgpu_t( const gpgpu_functional_sim_config &config ) checkpoint_CTA_t = m_function_model_config.get_checkpoint_CTA_t(); checkpoint_insn_Y = m_function_model_config.get_checkpoint_insn_Y(); + // initialize texture mappings to empty + m_NameToTextureInfo.clear(); + m_NameToCudaArray.clear(); if(m_function_model_config.get_ptx_inst_debug_to_file() != 0) ptx_inst_debug_file = fopen(m_function_model_config.get_ptx_inst_debug_file(), "w"); @@ -688,7 +691,11 @@ unsigned g_kernel_launch_latency; unsigned kernel_info_t::m_next_uid = 1; -kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry, std::map NameToCudaArray, std::map NameToTexureInfo) +/*A snapshot of the texture mappings needs to be stored in the kernel's info as +kernels should use the texture bindings seen at the time of launch and textures + can be bound/unbound asynchronously with respect to streams. */ + +kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry, std::map nameToCudaArray, std::map nameToTexureInfo) { m_kernel_entry=entry; m_grid_dim=gridDim; @@ -708,8 +715,8 @@ kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info * m_launch_latency = g_kernel_launch_latency; volta_cache_config_set=false; - m_NameToCudaArray = NameToCudaArray; - m_NameToTexureInfo = NameToTexureInfo; + m_NameToCudaArray = nameToCudaArray; + m_NameToTexureInfo = nameToTexureInfo; } kernel_info_t::~kernel_info_t() diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index bbc31e5..201b21f 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -212,7 +212,7 @@ public: // m_num_cores_running=0; // m_param_mem=NULL; // } - kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry, std::map m_NameToCudaArray, std::map m_NameToTexureInfo); + kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry, std::map nameToCudaArray, std::map nameToTexureInfo); ~kernel_info_t(); void inc_running() { m_num_cores_running++; } @@ -275,19 +275,22 @@ public: std::list &active_threads() { return m_active_threads; } class memory_space *get_param_memory() { return m_param_mem; } + + //The following functions access texture bindings present at the kernel's launch + const struct cudaArray* get_texarray( const std::string &texname ) const - { - std::map::const_iterator t=m_NameToCudaArray.find(texname); - assert(t != m_NameToCudaArray.end()); - return t->second; - } + { + std::map::const_iterator t=m_NameToCudaArray.find(texname); + assert(t != m_NameToCudaArray.end()); + return t->second; + } - const struct textureInfo* get_texinfo( const std::string &texname ) const - { - std::map::const_iterator t=m_NameToTexureInfo.find(texname); - assert(t != m_NameToTexureInfo.end()); - return t->second; - } + const struct textureInfo* get_texinfo( const std::string &texname ) const + { + std::map::const_iterator t=m_NameToTexureInfo.find(texname); + assert(t != m_NameToTexureInfo.end()); + return t->second; + } private: kernel_info_t( const kernel_info_t & ); // disable copy constructor @@ -297,7 +300,8 @@ private: unsigned m_uid; static unsigned m_next_uid; - + + //These maps contain the snapshot of the texture mappings at kernel launch std::map m_NameToCudaArray; std::map m_NameToTexureInfo; @@ -611,7 +615,9 @@ public: const gpgpu_functional_sim_config &get_config() const { return m_function_model_config; } FILE* get_ptx_inst_debug_file() { return ptx_inst_debug_file; } - + + // These maps return the current texture mappings for the GPU at any given time. + std::map getNameArrayMapping() {return m_NameToCudaArray;} std::map getNameInfoMapping() {return m_NameToTexureInfo;} @@ -624,7 +630,7 @@ protected: class memory_space *m_surf_mem; unsigned long long m_dev_malloc; - + // These maps contain the current texture mappings for the GPU at any given time. std::map > m_NameToTextureRef; std::map m_TextureRefToName; std::map m_NameToCudaArray; diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index a51ea3e..9d59411 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -1507,6 +1507,13 @@ static unsigned get_tex_datasize( const ptx_instruction *pI, ptx_thread_info *th std::string texname = src1.name(); gpgpu_t *gpu = thread->get_gpu(); + /* + For programs with many streams, textures can be bound and unbound + asynchronously. This means we need to use the kernel's "snapshot" of + the state of the texture mappings when it was launched (so that we + don't try to access the incorrect texture mapping if it's been updated, + or that we don't access a mapping that has been unbound). + */ kernel_info_t& k = thread->get_kernel(); const struct textureInfo* texInfo = k.get_texinfo(texname); diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc index d29fc8b..11001d7 100644 --- a/src/cuda-sim/instructions.cc +++ b/src/cuda-sim/instructions.cc @@ -5258,7 +5258,13 @@ void tex_impl( const ptx_instruction *pI, ptx_thread_info *thread ) if (!ptx_tex_regs) ptx_tex_regs = new ptx_reg_t[4]; unsigned nelem = src2.get_vect_nelem(); thread->get_vector_operand_values(src2, ptx_tex_regs, nelem); //ptx_reg should be 4 entry vector type...coordinates into texture - + /* + For programs with many streams, textures can be bound and unbound + asynchronously. This means we need to use the kernel's "snapshot" of + the state of the texture mappings when it was launched (so that we + don't try to access the incorrect texture mapping if it's been updated, + or that we don't access a mapping that has been unbound). + */ gpgpu_t *gpu = thread->get_gpu(); kernel_info_t &k = thread->get_kernel(); const struct textureReference* texref = gpu->get_texref(texname); -- cgit v1.3 From 49c6eeff7752b5816fb08c0d6baf2b47c789bc18 Mon Sep 17 00:00:00 2001 From: Suchita Pati Date: Wed, 17 Apr 2019 22:09:17 -0500 Subject: Adding configs for device limits and fixed other misc bugs --- libcuda/cuda_runtime_api.cc | 11 +++++++---- src/abstract_hardware_model.cc | 9 ++++----- src/abstract_hardware_model.h | 17 ++++++++--------- src/cuda-sim/cuda-sim.cc | 4 ++-- src/gpgpu-sim/gpu-sim.cc | 8 ++++++++ src/gpgpu-sim/gpu-sim.h | 11 ++++++++++- 6 files changed, 39 insertions(+), 21 deletions(-) (limited to 'src/abstract_hardware_model.cc') diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index f6e5e6f..86435b9 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -1205,18 +1205,20 @@ __host__ cudaError_t CUDARTAPI cudaDeviceGetLimit ( size_t* pValue, cudaLimit li if(g_debug_execution >= 3){ announce_call(__my_func__); } + _cuda_device_id *dev = GPGPUSim_Init(); + const gpgpu_sim_config& config=dev->get_gpgpu()->get_config(); switch(limit) { case 0: // cudaLimitStackSize - *pValue=1024; + *pValue=config.stack_limit(); break; case 2: // cudaLimitMallocHeapSize - *pValue=8388608; + *pValue=config.heap_limit(); break; case 3: // cudaLimitDevRuntimeSyncDepth - *pValue=2; + *pValue=config.sync_depth_limit(); break; case 4: // cudaLimitDevRuntimePendingLaunchCount - *pValue=2048; + *pValue=config.pending_launch_count_limit(); break; default: printf("ERROR:Limit %s unimplemented \n",limit); @@ -2717,6 +2719,7 @@ cudaError_t CUDARTAPI cudaDeviceSynchronize(void){ if(g_debug_execution >= 3){ announce_call(__my_func__); } + //Blocks until the device has completed all preceding requested tasks synchronize(); return g_last_cudaError = cudaSuccess; } diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index 820f5fd..cebdb25 100644 --- a/src/abstract_hardware_model.cc +++ b/src/abstract_hardware_model.cc @@ -88,8 +88,6 @@ void checkpoint::load_global_mem(class memory_space *temp_mem, char * f1name) fclose ( fp2 ); } - - void checkpoint::store_global_mem(class memory_space * mem, char *fname, char * format) { @@ -195,6 +193,8 @@ gpgpu_t::gpgpu_t( const gpgpu_functional_sim_config &config ) // initialize texture mappings to empty m_NameToTextureInfo.clear(); m_NameToCudaArray.clear(); + m_TextureRefToName.clear(); + m_NameToAttribute.clear(); if(m_function_model_config.get_ptx_inst_debug_to_file() != 0) ptx_inst_debug_file = fopen(m_function_model_config.get_ptx_inst_debug_file(), "w"); @@ -694,8 +694,7 @@ unsigned kernel_info_t::m_next_uid = 1; /*A snapshot of the texture mappings needs to be stored in the kernel's info as kernels should use the texture bindings seen at the time of launch and textures can be bound/unbound asynchronously with respect to streams. */ - -kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry, std::map nameToCudaArray, std::map nameToTexureInfo) +kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry, std::map nameToCudaArray, std::map nameToTextureInfo) { m_kernel_entry=entry; m_grid_dim=gridDim; @@ -716,7 +715,7 @@ kernel_info_t::kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info * volta_cache_config_set=false; m_NameToCudaArray = nameToCudaArray; - m_NameToTexureInfo = nameToTexureInfo; + m_NameToTextureInfo = nameToTextureInfo; } kernel_info_t::~kernel_info_t() diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index 201b21f..e8716ab 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -212,7 +212,7 @@ public: // m_num_cores_running=0; // m_param_mem=NULL; // } - kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry, std::map nameToCudaArray, std::map nameToTexureInfo); + kernel_info_t( dim3 gridDim, dim3 blockDim, class function_info *entry, std::map nameToCudaArray, std::map nameToTextureInfo); ~kernel_info_t(); void inc_running() { m_num_cores_running++; } @@ -287,8 +287,8 @@ public: const struct textureInfo* get_texinfo( const std::string &texname ) const { - std::map::const_iterator t=m_NameToTexureInfo.find(texname); - assert(t != m_NameToTexureInfo.end()); + std::map::const_iterator t=m_NameToTextureInfo.find(texname); + assert(t != m_NameToTextureInfo.end()); return t->second; } @@ -303,7 +303,7 @@ private: //These maps contain the snapshot of the texture mappings at kernel launch std::map m_NameToCudaArray; - std::map m_NameToTexureInfo; + std::map m_NameToTextureInfo; dim3 m_grid_dim; dim3 m_block_dim; @@ -601,8 +601,8 @@ public: const struct textureInfo* get_texinfo( const std::string &texname ) const { - std::map::const_iterator t=m_NameToTexureInfo.find(texname); - assert(t != m_NameToTexureInfo.end()); + std::map::const_iterator t=m_NameToTextureInfo.find(texname); + assert(t != m_NameToTextureInfo.end()); return t->second; } @@ -617,9 +617,8 @@ public: FILE* get_ptx_inst_debug_file() { return ptx_inst_debug_file; } // These maps return the current texture mappings for the GPU at any given time. - std::map getNameArrayMapping() {return m_NameToCudaArray;} - std::map getNameInfoMapping() {return m_NameToTexureInfo;} + std::map getNameInfoMapping() {return m_NameToTextureInfo;} protected: const gpgpu_functional_sim_config &m_function_model_config; @@ -634,7 +633,7 @@ protected: std::map > m_NameToTextureRef; std::map m_TextureRefToName; std::map m_NameToCudaArray; - std::map m_NameToTexureInfo; + std::map m_NameToTextureInfo; std::map m_NameToAttribute; }; diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc index 9d59411..f7bb9cc 100644 --- a/src/cuda-sim/cuda-sim.cc +++ b/src/cuda-sim/cuda-sim.cc @@ -218,7 +218,7 @@ void gpgpu_t::gpgpu_ptx_sim_bindTextureToArray(const struct textureReference* te texInfo->Ty_numbits = intLOGB2(Ty); texInfo->texel_size = texel_size; texInfo->texel_size_numbits = intLOGB2(texel_size); - m_NameToTexureInfo[texname] = texInfo; + m_NameToTextureInfo[texname] = texInfo; } void gpgpu_t::gpgpu_ptx_sim_unbindTexture(const struct textureReference* texref) @@ -226,7 +226,7 @@ void gpgpu_t::gpgpu_ptx_sim_unbindTexture(const struct textureReference* texref) //assumes bind-use-unbind-bind-use-unbind pattern std::string texname = gpgpu_ptx_sim_findNamefromTexture(texref); m_NameToCudaArray.erase(texname); - m_NameToTexureInfo.erase(texname); + m_NameToTextureInfo.erase(texname); } unsigned g_assemble_code_next_pc=0; diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index ec570bf..2a862a4 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -532,6 +532,14 @@ void gpgpu_sim_config::reg_options(option_parser_t opp) option_parser_register(opp, "-visualizer_zlevel", OPT_INT32, &g_visualizer_zlevel, "Compression level of the visualizer output log (0=no comp, 9=highest)", "6"); + option_parser_register(opp, "-gpgpu_stack_size_limit", OPT_INT32, &stack_size_limit, + "GPU thread stack size", "1024" ); + option_parser_register(opp, "-gpgpu_heap_size_limit", OPT_INT32, &heap_size_limit, + "GPU malloc heap size ", "8388608" ); + option_parser_register(opp, "-gpgpu_runtime_sync_depth_limit", OPT_INT32, &runtime_sync_depth_limit, + "GPU device runtime synchronize depth", "2" ); + option_parser_register(opp, "-gpgpu_runtime_pending_launch_count_limit", OPT_INT32, &runtime_pending_launch_count_limit, + "GPU device runtime pending launch count", "2048" ); option_parser_register(opp, "-trace_enabled", OPT_BOOL, &Trace::enabled, "Turn on traces", "0"); diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h index 6ce5524..11ffe49 100644 --- a/src/gpgpu-sim/gpu-sim.h +++ b/src/gpgpu-sim/gpu-sim.h @@ -337,6 +337,11 @@ public: unsigned get_max_concurrent_kernel() const { return max_concurrent_kernel; } unsigned checkpoint_option; + size_t stack_limit() const {return stack_size_limit; } + size_t heap_limit() const {return heap_size_limit; } + size_t sync_depth_limit() const {return runtime_sync_depth_limit; } + size_t pending_launch_count_limit() const {return runtime_pending_launch_count_limit;} + private: void init_clock_domains(void ); @@ -377,7 +382,11 @@ private: int gpu_stat_sample_freq; int gpu_runtime_stat_flag; - + // Device Limits + size_t stack_size_limit; + size_t heap_size_limit; + size_t runtime_sync_depth_limit; + size_t runtime_pending_launch_count_limit; unsigned long long liveness_message_freq; -- cgit v1.3 From 0c16df3c8b108d8720846bb44b9abcc60ddf42f9 Mon Sep 17 00:00:00 2001 From: Mahmoud Date: Wed, 15 May 2019 20:16:54 -0400 Subject: make gpu_tot_cycle local variable not global variable --- src/abstract_hardware_model.cc | 12 ++++--- src/abstract_hardware_model.h | 15 ++++++-- src/cuda-sim/cuda_device_runtime.cc | 2 +- src/cuda-sim/ptx_sim.cc | 8 ++--- src/gpgpu-sim/dram.cc | 22 ++++++------ src/gpgpu-sim/dram.h | 6 ++-- src/gpgpu-sim/dram_sched.cc | 14 ++++---- src/gpgpu-sim/gpu-cache.cc | 15 ++++---- src/gpgpu-sim/gpu-cache.h | 21 ++++++----- src/gpgpu-sim/gpu-sim.cc | 35 ++++++++----------- src/gpgpu-sim/gpu-sim.h | 17 +++++++-- src/gpgpu-sim/l2cache.cc | 56 ++++++++++++++++-------------- src/gpgpu-sim/l2cache.h | 13 ++++--- src/gpgpu-sim/l2cache_trace.h | 4 +-- src/gpgpu-sim/mem_fetch.cc | 5 +-- src/gpgpu-sim/mem_fetch.h | 1 + src/gpgpu-sim/mem_latency_stat.cc | 9 ++--- src/gpgpu-sim/mem_latency_stat.h | 4 ++- src/gpgpu-sim/scoreboard.cc | 4 ++- src/gpgpu-sim/scoreboard.h | 4 ++- src/gpgpu-sim/shader.cc | 69 +++++++++++++++++++------------------ src/gpgpu-sim/shader.h | 10 +++--- src/gpgpu-sim/shader_trace.h | 4 +-- src/gpgpusim_entrypoint.cc | 2 +- src/stream_manager.cc | 2 +- src/trace.h | 14 +++++--- 26 files changed, 214 insertions(+), 154 deletions(-) (limited to 'src/abstract_hardware_model.cc') diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc index cebdb25..63b139e 100644 --- a/src/abstract_hardware_model.cc +++ b/src/abstract_hardware_model.cc @@ -198,6 +198,9 @@ gpgpu_t::gpgpu_t( const gpgpu_functional_sim_config &config ) if(m_function_model_config.get_ptx_inst_debug_to_file() != 0) ptx_inst_debug_file = fopen(m_function_model_config.get_ptx_inst_debug_file(), "w"); + + gpu_sim_cycle=0; + gpu_tot_sim_cycle=0; } address_type line_size_based_tag_func(new_addr_type address, new_addr_type line_size) @@ -830,10 +833,11 @@ void kernel_info_t::destroy_cta_streams() { m_cta_streams.clear(); } -simt_stack::simt_stack( unsigned wid, unsigned warpSize) +simt_stack::simt_stack( unsigned wid, unsigned warpSize, class gpgpu_sim * gpu) { m_warp_id=wid; m_warp_size = warpSize; + m_gpu=gpu; reset(); } @@ -1033,7 +1037,7 @@ void simt_stack::update( simt_mask_t &thread_done, addr_vector_t &next_pc, addre simt_stack_entry new_stack_entry; new_stack_entry.m_pc = tmp_next_pc; new_stack_entry.m_active_mask = tmp_active_mask; - new_stack_entry.m_branch_div_cycle = gpu_sim_cycle+gpu_tot_sim_cycle; + new_stack_entry.m_branch_div_cycle = m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle; new_stack_entry.m_type = STACK_ENTRY_TYPE_CALL; m_stack.push_back(new_stack_entry); return; @@ -1065,7 +1069,7 @@ void simt_stack::update( simt_mask_t &thread_done, addr_vector_t &next_pc, addre new_recvg_pc = recvg_pc; if (new_recvg_pc != top_recvg_pc) { m_stack.back().m_pc = new_recvg_pc; - m_stack.back().m_branch_div_cycle = gpu_sim_cycle+gpu_tot_sim_cycle; + m_stack.back().m_branch_div_cycle = m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle; m_stack.push_back(simt_stack_entry()); } @@ -1157,7 +1161,7 @@ void core_t::initilizeSIMTStack(unsigned warp_count, unsigned warp_size) { m_simt_stack = new simt_stack*[warp_count]; for (unsigned i = 0; i < warp_count; ++i) - m_simt_stack[i] = new simt_stack(i,warp_size); + m_simt_stack[i] = new simt_stack(i,warp_size,m_gpu); m_warp_size = warp_size; m_warp_count = warp_count; } diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h index 22ef509..1735c2f 100644 --- a/src/abstract_hardware_model.h +++ b/src/abstract_hardware_model.h @@ -399,7 +399,7 @@ typedef std::vector addr_vector_t; class simt_stack { public: - simt_stack( unsigned wid, unsigned warpSize); + simt_stack( unsigned wid, unsigned warpSize, class gpgpu_sim * gpu); void reset(); void launch( address_type start_pc, const simt_mask_t &active_mask ); @@ -416,6 +416,7 @@ protected: unsigned m_warp_id; unsigned m_warp_size; + enum stack_entry_type { STACK_ENTRY_TYPE_NORMAL = 0, STACK_ENTRY_TYPE_CALL @@ -433,6 +434,8 @@ protected: }; std::deque m_stack; + + class gpgpu_sim * m_gpu; }; #define GLOBAL_HEAP_START 0xC0000000 @@ -571,6 +574,12 @@ public: int resume_CTA; int checkpoint_CTA_t; int checkpoint_insn_Y; + + //Move some cycle core stats here instead of being global + unsigned long long gpu_sim_cycle; + unsigned long long gpu_tot_sim_cycle; + + void* gpu_malloc( size_t size ); void* gpu_mallocarray( size_t count ); void gpu_memset( size_t dst_start_addr, int c, size_t count ); @@ -835,8 +844,8 @@ public: class mem_fetch_allocator { public: - virtual mem_fetch *alloc( new_addr_type addr, mem_access_type type, unsigned size, bool wr ) const = 0; - virtual mem_fetch *alloc( const class warp_inst_t &inst, const mem_access_t &access ) const = 0; + virtual mem_fetch *alloc( new_addr_type addr, mem_access_type type, unsigned size, bool wr, unsigned long long cycle ) const = 0; + virtual mem_fetch *alloc( const class warp_inst_t &inst, const mem_access_t &access, unsigned long long cycle ) const = 0; }; // the maximum number of destination, source, or address uarch operands in a instruction diff --git a/src/cuda-sim/cuda_device_runtime.cc b/src/cuda-sim/cuda_device_runtime.cc index 917e7a8..86e8147 100644 --- a/src/cuda-sim/cuda_device_runtime.cc +++ b/src/cuda-sim/cuda_device_runtime.cc @@ -200,7 +200,7 @@ void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info * //create child kernel_info_t and index it with parameter_buffer address gpgpu_t* gpu=thread->get_gpu(); device_grid = new kernel_info_t(config.grid_dim, config.block_dim, device_kernel_entry, gpu->getNameArrayMapping(), gpu->getNameInfoMapping()); - device_grid->launch_cycle = gpu_sim_cycle + gpu_tot_sim_cycle; + device_grid->launch_cycle = gpu->gpu_sim_cycle + gpu->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/cuda-sim/ptx_sim.cc b/src/cuda-sim/ptx_sim.cc index 820287d..8ad651e 100644 --- a/src/cuda-sim/ptx_sim.cc +++ b/src/cuda-sim/ptx_sim.cc @@ -222,7 +222,7 @@ void ptx_thread_info::set_done() { assert( !m_at_barrier ); m_thread_done = true; - m_cycle_done = gpu_sim_cycle; + m_cycle_done = m_gpu->gpu_sim_cycle; } unsigned ptx_thread_info::get_builtin( int builtin_id, unsigned dim_mod ) @@ -230,15 +230,15 @@ unsigned ptx_thread_info::get_builtin( int builtin_id, unsigned dim_mod ) assert( m_valid ); switch ((builtin_id&0xFFFF)) { case CLOCK_REG: - return (unsigned)(gpu_sim_cycle + gpu_tot_sim_cycle); + return (unsigned)(m_gpu->gpu_sim_cycle + m_gpu->gpu_tot_sim_cycle); case CLOCK64_REG: abort(); // change return value to unsigned long long? // GPGPUSim clock is 4 times slower - multiply by 4 - return (gpu_sim_cycle + gpu_tot_sim_cycle)*4; + return (m_gpu->gpu_sim_cycle + m_gpu->gpu_tot_sim_cycle)*4; case HALFCLOCK_ID: // GPGPUSim clock is 4 times slower - multiply by 4 // Hardware clock counter is incremented at half the shader clock frequency - divide by 2 (Henry '10) - return (gpu_sim_cycle + gpu_tot_sim_cycle)*2; + return (m_gpu->gpu_sim_cycle + m_gpu->gpu_tot_sim_cycle)*2; case CTAID_REG: assert( dim_mod < 3 ); if( dim_mod == 0 ) return m_ctaid.x; diff --git a/src/gpgpu-sim/dram.cc b/src/gpgpu-sim/dram.cc index 192cb65..5e36d4b 100644 --- a/src/gpgpu-sim/dram.cc +++ b/src/gpgpu-sim/dram.cc @@ -42,12 +42,13 @@ template class fifo_pipeline; template class fifo_pipeline; dram_t::dram_t( unsigned int partition_id, const struct memory_config *config, memory_stats_t *stats, - memory_partition_unit *mp ) + memory_partition_unit *mp, gpgpu_sim* gpu ) { id = partition_id; m_memory_partition_unit = mp; m_stats = stats; m_config = config; + m_gpu = gpu; //rowblp access_num=0; @@ -191,11 +192,12 @@ unsigned int dram_t::queue_limit() const } -dram_req_t::dram_req_t( class mem_fetch *mf, unsigned banks, unsigned dram_bnk_indexing_policy) +dram_req_t::dram_req_t( class mem_fetch *mf, unsigned banks, unsigned dram_bnk_indexing_policy, class gpgpu_sim* gpu) { txbytes = 0; dqbytes = 0; data = mf; + m_gpu = gpu; const addrdec_t &tlx = mf->get_tlx_addr(); @@ -226,9 +228,9 @@ dram_req_t::dram_req_t( class mem_fetch *mf, unsigned banks, unsigned dram_bnk_i col = tlx.col; nbytes = mf->get_data_size(); - timestamp = gpu_tot_sim_cycle + gpu_sim_cycle; + timestamp = m_gpu->gpu_tot_sim_cycle + m_gpu->gpu_sim_cycle; addr = mf->get_addr(); - insertion_time = (unsigned) gpu_sim_cycle; + insertion_time = (unsigned) m_gpu->gpu_sim_cycle; rw = data->get_is_write()?WRITE:READ; } @@ -236,9 +238,9 @@ void dram_t::push( class mem_fetch *data ) { assert(id == data->get_tlx_addr().chip); // Ensure request is in correct memory partition - dram_req_t *mrq = new dram_req_t(data,m_config->nbk,m_config->dram_bnk_indexing_policy); + dram_req_t *mrq = new dram_req_t(data,m_config->nbk,m_config->dram_bnk_indexing_policy,m_memory_partition_unit->get_mgpu()); - data->set_status(IN_PARTITION_MC_INTERFACE_QUEUE,gpu_sim_cycle+gpu_tot_sim_cycle); + data->set_status(IN_PARTITION_MC_INTERFACE_QUEUE,m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle); mrqq->push(mrq); // stats... @@ -259,7 +261,7 @@ void dram_t::scheduler_fifo() if (!mrqq->empty()) { unsigned int bkn; dram_req_t *head_mrqq = mrqq->top(); - head_mrqq->data->set_status(IN_PARTITION_MC_BANK_ARB_QUEUE,gpu_sim_cycle+gpu_tot_sim_cycle); + head_mrqq->data->set_status(IN_PARTITION_MC_BANK_ARB_QUEUE,m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle); bkn = head_mrqq->bk; if (!bk[bkn]->mrq) bk[bkn]->mrq = mrqq->pop(); @@ -283,7 +285,7 @@ void dram_t::cycle() if (cmd->dqbytes >= cmd->nbytes) { mem_fetch *data = cmd->data; - data->set_status(IN_PARTITION_MC_RETURNQ,gpu_sim_cycle+gpu_tot_sim_cycle); + data->set_status(IN_PARTITION_MC_RETURNQ,m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle); if( data->get_access_type() != L1_WRBK_ACC && data->get_access_type() != L2_WRBK_ACC ) { data->set_reply(); returnq->push(data); @@ -566,7 +568,7 @@ bool dram_t::issue_col_command(int j) bool issued = false; unsigned grp = get_bankgrp_number(j); if (bk[j]->mrq) { //if currently servicing a memory request - bk[j]->mrq->data->set_status(IN_PARTITION_DRAM,gpu_sim_cycle+gpu_tot_sim_cycle); + bk[j]->mrq->data->set_status(IN_PARTITION_DRAM,m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle); // correct row activated for a READ if ( !issued && !CCDc && !bk[j]->RCDc && !(bkgrp[grp]->CCDLc) && @@ -654,7 +656,7 @@ bool dram_t::issue_row_command(int j) bool issued = false; unsigned grp = get_bankgrp_number(j); if (bk[j]->mrq) { //if currently servicing a memory request - bk[j]->mrq->data->set_status(IN_PARTITION_DRAM,gpu_sim_cycle+gpu_tot_sim_cycle); + bk[j]->mrq->data->set_status(IN_PARTITION_DRAM,m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle); // bank is idle //else if ( !issued && !RRDc && diff --git a/src/gpgpu-sim/dram.h b/src/gpgpu-sim/dram.h index 1ab0153..7a3a2da 100644 --- a/src/gpgpu-sim/dram.h +++ b/src/gpgpu-sim/dram.h @@ -48,7 +48,7 @@ class dram_req_t { public: - dram_req_t( class mem_fetch *data , unsigned banks, unsigned dram_bnk_indexing_policy); + dram_req_t( class mem_fetch *data , unsigned banks, unsigned dram_bnk_indexing_policy, class gpgpu_sim* gpu); unsigned int row; unsigned int col; @@ -62,6 +62,7 @@ public: unsigned long long int addr; unsigned int insertion_time; class mem_fetch * data; + class gpgpu_sim * m_gpu; }; struct bankgrp_t @@ -110,7 +111,7 @@ class dram_t { public: dram_t( unsigned int parition_id, const struct memory_config *config, class memory_stats_t *stats, - class memory_partition_unit *mp ); + class memory_partition_unit *mp, class gpgpu_sim* gpu ); bool full(bool is_write) const; void print( FILE* simFile ) const; @@ -129,6 +130,7 @@ public: void dram_log (int task); class memory_partition_unit *m_memory_partition_unit; + class gpgpu_sim* m_gpu; unsigned int id; // Power Model diff --git a/src/gpgpu-sim/dram_sched.cc b/src/gpgpu-sim/dram_sched.cc index ff50050..6ee6271 100644 --- a/src/gpgpu-sim/dram_sched.cc +++ b/src/gpgpu-sim/dram_sched.cc @@ -84,13 +84,13 @@ void frfcfs_scheduler::add_req( dram_req_t *req ) void frfcfs_scheduler::data_collection(unsigned int bank) { - if (gpu_sim_cycle > row_service_timestamp[bank]) { - curr_row_service_time[bank] = gpu_sim_cycle - row_service_timestamp[bank]; + if (m_dram->m_gpu->gpu_sim_cycle > row_service_timestamp[bank]) { + curr_row_service_time[bank] = m_dram->m_gpu->gpu_sim_cycle - row_service_timestamp[bank]; if (curr_row_service_time[bank] > m_stats->max_servicetime2samerow[m_dram->id][bank]) m_stats->max_servicetime2samerow[m_dram->id][bank] = curr_row_service_time[bank]; } curr_row_service_time[bank] = 0; - row_service_timestamp[bank] = gpu_sim_cycle; + row_service_timestamp[bank] = m_dram->m_gpu->gpu_sim_cycle; if (m_stats->concurrent_row_access[m_dram->id][bank] > m_stats->max_conc_access2samerow[m_dram->id][bank]) { m_stats->max_conc_access2samerow[m_dram->id][bank] = m_stats->concurrent_row_access[m_dram->id][bank]; } @@ -215,7 +215,7 @@ void dram_t::scheduler_frfcfs() m_stats->total_n_reads++; } - req->data->set_status(IN_PARTITION_MC_INPUT_QUEUE,gpu_sim_cycle+gpu_tot_sim_cycle); + req->data->set_status(IN_PARTITION_MC_INPUT_QUEUE,m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle); sched->add_req(req); } @@ -228,14 +228,14 @@ void dram_t::scheduler_frfcfs() req = sched->schedule(b, bk[b]->curr_row); if ( req ) { - req->data->set_status(IN_PARTITION_MC_BANK_ARB_QUEUE,gpu_sim_cycle+gpu_tot_sim_cycle); + req->data->set_status(IN_PARTITION_MC_BANK_ARB_QUEUE,m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle); prio = (prio+1)%m_config->nbk; bk[b]->mrq = req; if (m_config->gpgpu_memlatency_stat) { - mrq_latency = gpu_sim_cycle + gpu_tot_sim_cycle - bk[b]->mrq->timestamp; + mrq_latency = m_gpu->gpu_sim_cycle + m_gpu->gpu_tot_sim_cycle - bk[b]->mrq->timestamp; m_stats->tot_mrq_latency += mrq_latency; m_stats->tot_mrq_num++; - bk[b]->mrq->timestamp = gpu_tot_sim_cycle + gpu_sim_cycle; + bk[b]->mrq->timestamp =m_gpu->gpu_tot_sim_cycle + m_gpu->gpu_sim_cycle; m_stats->mrq_lat_table[LOGB2(mrq_latency)]++; if (mrq_latency > m_stats->max_mrq_latency) { m_stats->max_mrq_latency = mrq_latency; diff --git a/src/gpgpu-sim/gpu-cache.cc b/src/gpgpu-sim/gpu-cache.cc index 565fae1..62849f8 100644 --- a/src/gpgpu-sim/gpu-cache.cc +++ b/src/gpgpu-sim/gpu-cache.cc @@ -26,6 +26,7 @@ // OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include "gpu-cache.h" +#include "gpu-sim.h" #include "stat-tool.h" #include @@ -1183,7 +1184,8 @@ data_cache::wr_miss_wa_naive( new_addr_type addr, mf->get_wid(), mf->get_sid(), mf->get_tpc(), - mf->get_mem_config()); + mf->get_mem_config(), + m_gpu->gpu_tot_sim_cycle+m_gpu->gpu_sim_cycle); bool do_miss = false; bool wb = false; @@ -1201,7 +1203,7 @@ data_cache::wr_miss_wa_naive( new_addr_type addr, if( wb && (m_config.m_write_policy != WRITE_THROUGH) ) { assert(status == MISS); //SECTOR_MISS and HIT_RESERVED should not send write back mem_fetch *wb = m_memfetch_creator->alloc(evicted.m_block_addr, - m_wrbk_type,evicted.m_modified_size,true); + m_wrbk_type,evicted.m_modified_size,true,m_gpu->gpu_tot_sim_cycle+m_gpu->gpu_sim_cycle); send_write_request(wb, cache_event(WRITE_BACK_REQUEST_SENT, evicted), time, events); } return MISS; @@ -1245,7 +1247,7 @@ data_cache::wr_miss_wa_fetch_on_write( new_addr_type addr, // (already modified lower level) if( wb && (m_config.m_write_policy != WRITE_THROUGH) ) { mem_fetch *wb = m_memfetch_creator->alloc(evicted.m_block_addr, - m_wrbk_type,evicted.m_modified_size,true); + m_wrbk_type,evicted.m_modified_size,true,m_gpu->gpu_tot_sim_cycle+m_gpu->gpu_sim_cycle); send_write_request(wb, cache_event(WRITE_BACK_REQUEST_SENT, evicted), time, events); } return MISS; @@ -1297,6 +1299,7 @@ data_cache::wr_miss_wa_fetch_on_write( new_addr_type addr, mf->get_sid(), mf->get_tpc(), mf->get_mem_config(), + m_gpu->gpu_tot_sim_cycle+m_gpu->gpu_sim_cycle, NULL, mf); @@ -1320,7 +1323,7 @@ data_cache::wr_miss_wa_fetch_on_write( new_addr_type addr, // (already modified lower level) if(wb && (m_config.m_write_policy != WRITE_THROUGH) ){ mem_fetch *wb = m_memfetch_creator->alloc(evicted.m_block_addr, - m_wrbk_type,evicted.m_modified_size,true); + m_wrbk_type,evicted.m_modified_size,true,m_gpu->gpu_tot_sim_cycle+m_gpu->gpu_sim_cycle); send_write_request(wb, cache_event(WRITE_BACK_REQUEST_SENT, evicted), time, events); } return MISS; @@ -1373,7 +1376,7 @@ data_cache::wr_miss_wa_lazy_fetch_on_read( new_addr_type addr, // (already modified lower level) if( wb && (m_config.m_write_policy != WRITE_THROUGH) ) { mem_fetch *wb = m_memfetch_creator->alloc(evicted.m_block_addr, - m_wrbk_type,evicted.m_modified_size,true); + m_wrbk_type,evicted.m_modified_size,true,m_gpu->gpu_tot_sim_cycle+m_gpu->gpu_sim_cycle); send_write_request(wb, cache_event(WRITE_BACK_REQUEST_SENT, evicted), time, events); } return MISS; @@ -1458,7 +1461,7 @@ data_cache::rd_miss_base( new_addr_type addr, // (already modified lower level) if(wb && (m_config.m_write_policy != WRITE_THROUGH) ){ mem_fetch *wb = m_memfetch_creator->alloc(evicted.m_block_addr, - m_wrbk_type,evicted.m_modified_size,true); + m_wrbk_type,evicted.m_modified_size,true,m_gpu->gpu_tot_sim_cycle+m_gpu->gpu_sim_cycle); send_write_request(wb, WRITE_BACK_REQUEST_SENT, time, events); } return MISS; diff --git a/src/gpgpu-sim/gpu-cache.h b/src/gpgpu-sim/gpu-cache.h index 673e128..85e534e 100644 --- a/src/gpgpu-sim/gpu-cache.h +++ b/src/gpgpu-sim/gpu-cache.h @@ -1212,12 +1212,13 @@ public: data_cache( const char *name, cache_config &config, int core_id, int type_id, mem_fetch_interface *memport, mem_fetch_allocator *mfcreator, enum mem_fetch_status status, - mem_access_type wr_alloc_type, mem_access_type wrbk_type ) + mem_access_type wr_alloc_type, mem_access_type wrbk_type, class gpgpu_sim* gpu ) : baseline_cache(name,config,core_id,type_id,memport,status) { init( mfcreator ); m_wr_alloc_type = wr_alloc_type; m_wrbk_type = wrbk_type; + m_gpu=gpu; } virtual ~data_cache() {} @@ -1275,16 +1276,19 @@ protected: enum mem_fetch_status status, tag_array* new_tag_array, mem_access_type wr_alloc_type, - mem_access_type wrbk_type) + mem_access_type wrbk_type, + class gpgpu_sim* gpu ) : baseline_cache(name, config, core_id, type_id, memport,status, new_tag_array) { init( mfcreator ); m_wr_alloc_type = wr_alloc_type; m_wrbk_type = wrbk_type; + m_gpu=gpu; } mem_access_type m_wr_alloc_type; // Specifies type of write allocate request (e.g., L1 or L2) mem_access_type m_wrbk_type; // Specifies type of writeback request (e.g., L1 or L2) + class gpgpu_sim* m_gpu; //! A general function that takes the result of a tag_array probe // and performs the correspding functions based on the cache configuration @@ -1441,8 +1445,8 @@ class l1_cache : public data_cache { public: l1_cache(const char *name, cache_config &config, int core_id, int type_id, mem_fetch_interface *memport, - mem_fetch_allocator *mfcreator, enum mem_fetch_status status ) - : data_cache(name,config,core_id,type_id,memport,mfcreator,status, L1_WR_ALLOC_R, L1_WRBK_ACC){} + mem_fetch_allocator *mfcreator, enum mem_fetch_status status, class gpgpu_sim* gpu ) + : data_cache(name,config,core_id,type_id,memport,mfcreator,status, L1_WR_ALLOC_R, L1_WRBK_ACC, gpu){} virtual ~l1_cache(){} @@ -1460,10 +1464,11 @@ protected: mem_fetch_interface *memport, mem_fetch_allocator *mfcreator, enum mem_fetch_status status, - tag_array* new_tag_array ) + tag_array* new_tag_array, + class gpgpu_sim* gpu) : data_cache( name, config, - core_id,type_id,memport,mfcreator,status, new_tag_array, L1_WR_ALLOC_R, L1_WRBK_ACC ){} + core_id,type_id,memport,mfcreator,status, new_tag_array, L1_WR_ALLOC_R, L1_WRBK_ACC, gpu ){} }; @@ -1473,8 +1478,8 @@ class l2_cache : public data_cache { public: l2_cache(const char *name, cache_config &config, int core_id, int type_id, mem_fetch_interface *memport, - mem_fetch_allocator *mfcreator, enum mem_fetch_status status ) - : data_cache(name,config,core_id,type_id,memport,mfcreator,status, L2_WR_ALLOC_R, L2_WRBK_ACC){} + mem_fetch_allocator *mfcreator, enum mem_fetch_status status, class gpgpu_sim* gpu ) + : data_cache(name,config,core_id,type_id,memport,mfcreator,status, L2_WR_ALLOC_R, L2_WRBK_ACC, gpu){} virtual ~l2_cache() {} diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc index c1ba934..72bac92 100644 --- a/src/gpgpu-sim/gpu-sim.cc +++ b/src/gpgpu-sim/gpu-sim.cc @@ -83,23 +83,6 @@ class gpgpu_sim_wrapper {}; bool g_interactive_debugger_enabled=false; -unsigned long long gpu_sim_cycle = 0; -unsigned long long gpu_tot_sim_cycle = 0; - - -// performance counter for stalls due to congestion. -unsigned int gpu_stall_dramfull = 0; -unsigned int gpu_stall_icnt2sh = 0; -unsigned long long partiton_reqs_in_parallel = 0; -unsigned long long partiton_reqs_in_parallel_total = 0; -unsigned long long partiton_reqs_in_parallel_util = 0; -unsigned long long partiton_reqs_in_parallel_util_total = 0; -unsigned long long gpu_sim_cycle_parition_util = 0; -unsigned long long gpu_tot_sim_cycle_parition_util = 0; -unsigned long long partiton_replys_in_parallel = 0; -unsigned long long partiton_replys_in_parallel_total = 0; - -tr1_hash_map address_random_interleaving; /* Clock Domains */ @@ -731,7 +714,7 @@ gpgpu_sim::gpgpu_sim( const gpgpu_sim_config &config ) #endif m_shader_stats = new shader_core_stats(m_shader_config); - m_memory_stats = new memory_stats_t(m_config.num_shader(),m_shader_config,m_memory_config); + m_memory_stats = new memory_stats_t(m_config.num_shader(),m_shader_config,m_memory_config,this); average_pipeline_duty_cycle = (float *)malloc(sizeof(float)); active_sms=(float *)malloc(sizeof(float)); m_power_stats = new power_stat_t(m_shader_config,average_pipeline_duty_cycle,active_sms,m_shader_stats,m_memory_config,m_memory_stats); @@ -742,6 +725,16 @@ gpgpu_sim::gpgpu_sim( const gpgpu_sim_config &config ) m_total_cta_launched = 0; gpu_deadlock = false; + gpu_stall_dramfull = 0; + gpu_stall_icnt2sh = 0; + partiton_reqs_in_parallel = 0; + partiton_reqs_in_parallel_total = 0; + partiton_reqs_in_parallel_util = 0; + partiton_reqs_in_parallel_util_total = 0; + gpu_sim_cycle_parition_util = 0; + gpu_tot_sim_cycle_parition_util = 0; + partiton_replys_in_parallel = 0; + partiton_replys_in_parallel_total = 0; m_cluster = new simt_core_cluster*[m_shader_config->n_simt_clusters]; for (unsigned i=0;in_simt_clusters;i++) @@ -750,7 +743,7 @@ gpgpu_sim::gpgpu_sim( const gpgpu_sim_config &config ) m_memory_partition_unit = new memory_partition_unit*[m_memory_config->m_n_mem]; m_memory_sub_partition = new memory_sub_partition*[m_memory_config->m_n_mem_sub_partition]; for (unsigned i=0;im_n_mem;i++) { - m_memory_partition_unit[i] = new memory_partition_unit(i, m_memory_config, m_memory_stats); + m_memory_partition_unit[i] = new memory_partition_unit(i, m_memory_config, m_memory_stats, this); for (unsigned p = 0; p < m_memory_config->m_n_sub_partition_per_memory_channel; p++) { unsigned submpid = i * m_memory_config->m_n_sub_partition_per_memory_channel + p; m_memory_sub_partition[submpid] = m_memory_partition_unit[i]->get_sub_partition(p); @@ -1504,7 +1497,7 @@ void shader_core_ctx::issue_block2core( kernel_info_t &kernel ) shader_CTA_count_log(m_sid, 1); SHADER_DPRINTF(LIVENESS, "GPGPU-Sim uArch: cta:%2u, start_tid:%4u, end_tid:%4u, initialized @(%lld,%lld)\n", - free_cta_hw_id, start_thread, end_thread, gpu_sim_cycle, gpu_tot_sim_cycle ); + free_cta_hw_id, start_thread, end_thread, m_gpu->gpu_sim_cycle, m_gpu->gpu_tot_sim_cycle ); } @@ -1721,7 +1714,7 @@ void gpgpu_sim::cycle() for (unsigned i=0;in_simt_clusters;i++) { m_cluster[i]->get_current_occupancy(active, total); } - DPRINTF(LIVENESS, "uArch: inst.: %lld (ipc=%4.1f, occ=%0.4f\% [%llu / %llu]) sim_rate=%u (inst/sec) elapsed = %u:%u:%02u:%02u / %s", + DPRINTFG(LIVENESS, "uArch: inst.: %lld (ipc=%4.1f, occ=%0.4f\% [%llu / %llu]) sim_rate=%u (inst/sec) elapsed = %u:%u:%02u:%02u / %s", gpu_tot_sim_insn + gpu_sim_insn, (double)gpu_sim_insn/(double)gpu_sim_cycle, float(active)/float(total) * 100, active, total, diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h index c8dad89..8e3b6ee 100644 --- a/src/gpgpu-sim/gpu-sim.h +++ b/src/gpgpu-sim/gpu-sim.h @@ -292,9 +292,6 @@ struct memory_config { bool m_perf_sim_memcpy; }; -// global counters and flags (please try not to add to this list!!!) -extern unsigned long long gpu_sim_cycle; -extern unsigned long long gpu_tot_sim_cycle; extern bool g_interactive_debugger_enabled; class gpgpu_sim_config : public power_config, public gpgpu_functional_sim_config { @@ -566,6 +563,20 @@ public: occupancy_stats gpu_occupancy; occupancy_stats gpu_tot_occupancy; + // performance counter for stalls due to congestion. + unsigned int gpu_stall_dramfull; + unsigned int gpu_stall_icnt2sh; + unsigned long long partiton_reqs_in_parallel; + unsigned long long partiton_reqs_in_parallel_total; + unsigned long long partiton_reqs_in_parallel_util; + unsigned long long partiton_reqs_in_parallel_util_total; + unsigned long long gpu_sim_cycle_parition_util; + unsigned long long gpu_tot_sim_cycle_parition_util; + unsigned long long partiton_replys_in_parallel; + unsigned long long partiton_replys_in_parallel_total; + + tr1_hash_map address_random_interleaving; + FuncCache get_cache_config(std::string kernel_name); void set_cache_config(std::string kernel_name, FuncCache cacheConfig ); diff --git a/src/gpgpu-sim/l2cache.cc b/src/gpgpu-sim/l2cache.cc index 25da107..526e999 100644 --- a/src/gpgpu-sim/l2cache.cc +++ b/src/gpgpu-sim/l2cache.cc @@ -46,7 +46,7 @@ #include "l2cache_trace.h" -mem_fetch * partition_mf_allocator::alloc(new_addr_type addr, mem_access_type type, unsigned size, bool wr ) const +mem_fetch * partition_mf_allocator::alloc(new_addr_type addr, mem_access_type type, unsigned size, bool wr, unsigned long long cycle ) const { assert( wr ); mem_access_t access( type, addr, size, wr ); @@ -56,22 +56,25 @@ mem_fetch * partition_mf_allocator::alloc(new_addr_type addr, mem_access_type ty -1, -1, -1, - m_memory_config ); + m_memory_config, + cycle); return mf; } memory_partition_unit::memory_partition_unit( unsigned partition_id, const struct memory_config *config, - class memory_stats_t *stats ) -: m_id(partition_id), m_config(config), m_stats(stats), m_arbitration_metadata(config) + class memory_stats_t *stats, + class gpgpu_sim* gpu) +: m_id(partition_id), m_config(config), m_stats(stats), m_arbitration_metadata(config), m_gpu(gpu) { - m_dram = new dram_t(m_id,m_config,m_stats,this); + m_dram = new dram_t(m_id,m_config,m_stats,this,gpu); m_sub_partition = new memory_sub_partition*[m_config->m_n_sub_partition_per_memory_channel]; for (unsigned p = 0; p < m_config->m_n_sub_partition_per_memory_channel; p++) { unsigned sub_partition_id = m_id * m_config->m_n_sub_partition_per_memory_channel + p; - m_sub_partition[p] = new memory_sub_partition(sub_partition_id, m_config, stats); + m_sub_partition[p] = new memory_sub_partition(sub_partition_id, m_config, stats, gpu); } + } void memory_partition_unit::handle_memcpy_to_gpu( size_t addr, unsigned global_subpart_id, mem_access_sector_mask_t mask ) @@ -80,7 +83,7 @@ void memory_partition_unit::handle_memcpy_to_gpu( size_t addr, unsigned global_s std::string mystring = mask.to_string(); MEMPART_DPRINTF("Copy Engine Request Received For Address=%llx, local_subpart=%u, global_subpart=%u, sector_mask=%s \n", addr, p, global_subpart_id, mystring.c_str()); - m_sub_partition[p]->force_l2_tag_update(addr,gpu_sim_cycle+gpu_tot_sim_cycle, mask); + m_sub_partition[p]->force_l2_tag_update(addr,m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle, mask); } memory_partition_unit::~memory_partition_unit() @@ -218,7 +221,7 @@ void memory_partition_unit::dram_cycle() delete mf_return; } else { m_sub_partition[dest_spid]->dram_L2_queue_push(mf_return); - mf_return->set_status(IN_PARTITION_DRAM_TO_L2_QUEUE,gpu_sim_cycle+gpu_tot_sim_cycle); + mf_return->set_status(IN_PARTITION_DRAM_TO_L2_QUEUE,m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle); m_arbitration_metadata.return_credit(dest_spid); MEMPART_DPRINTF("mem_fetch request %p return from dram to sub partition %d\n", mf_return, dest_spid); } @@ -247,9 +250,9 @@ void memory_partition_unit::dram_cycle() MEMPART_DPRINTF("Issue mem_fetch request %p from sub partition %d to dram\n", mf, spid); dram_delay_t d; d.req = mf; - d.ready_cycle = gpu_sim_cycle+gpu_tot_sim_cycle + m_config->dram_latency; + d.ready_cycle = m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle + m_config->dram_latency; m_dram_latency_queue.push_back(d); - mf->set_status(IN_PARTITION_DRAM_LATENCY_QUEUE,gpu_sim_cycle+gpu_tot_sim_cycle); + mf->set_status(IN_PARTITION_DRAM_LATENCY_QUEUE,m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle); m_arbitration_metadata.borrow_credit(spid); break; // the DRAM should only accept one request per cycle } @@ -258,7 +261,7 @@ void memory_partition_unit::dram_cycle() // DRAM latency queue - if( !m_dram_latency_queue.empty() && ( (gpu_sim_cycle+gpu_tot_sim_cycle) >= m_dram_latency_queue.front().ready_cycle ) && !m_dram->full(m_dram_latency_queue.front().req->is_write()) ) { + if( !m_dram_latency_queue.empty() && ( (m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle) >= m_dram_latency_queue.front().ready_cycle ) && !m_dram->full(m_dram_latency_queue.front().req->is_write()) ) { mem_fetch* mf = m_dram_latency_queue.front().req; m_dram_latency_queue.pop_front(); m_dram->push(mf); @@ -310,11 +313,13 @@ void memory_partition_unit::print( FILE *fp ) const memory_sub_partition::memory_sub_partition( unsigned sub_partition_id, const struct memory_config *config, - class memory_stats_t *stats ) + class memory_stats_t *stats, + class gpgpu_sim* gpu) { m_id = sub_partition_id; m_config=config; m_stats=stats; + m_gpu = gpu; m_memcpy_cycle_offset = 0; assert(m_id < m_config->m_n_mem_sub_partition); @@ -325,7 +330,7 @@ memory_sub_partition::memory_sub_partition( unsigned sub_partition_id, m_mf_allocator = new partition_mf_allocator(config); if(!m_config->m_L2_config.disabled()) - m_L2cache = new l2_cache(L2c_name,m_config->m_L2_config,-1,-1,m_L2interface,m_mf_allocator,IN_PARTITION_L2_MISS_QUEUE); + m_L2cache = new l2_cache(L2c_name,m_config->m_L2_config,-1,-1,m_L2interface,m_mf_allocator,IN_PARTITION_L2_MISS_QUEUE, gpu); unsigned int icnt_L2; unsigned int L2_dram; @@ -357,7 +362,7 @@ void memory_sub_partition::cache_cycle( unsigned cycle ) mem_fetch *mf = m_L2cache->next_access(); if(mf->get_access_type() != L2_WR_ALLOC_R){ // Don't pass write allocate read request back to upper level cache mf->set_reply(); - mf->set_status(IN_PARTITION_L2_TO_ICNT_QUEUE,gpu_sim_cycle+gpu_tot_sim_cycle); + mf->set_status(IN_PARTITION_L2_TO_ICNT_QUEUE,m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle); m_L2_icnt_queue->push(mf); }else{ if(m_config->m_L2_config.m_write_alloc_policy == FETCH_ON_WRITE) @@ -365,7 +370,7 @@ void memory_sub_partition::cache_cycle( unsigned cycle ) mem_fetch* original_wr_mf = mf->get_original_wr_mf(); assert(original_wr_mf); original_wr_mf->set_reply(); - original_wr_mf->set_status(IN_PARTITION_L2_TO_ICNT_QUEUE,gpu_sim_cycle+gpu_tot_sim_cycle); + original_wr_mf->set_status(IN_PARTITION_L2_TO_ICNT_QUEUE,m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle); m_L2_icnt_queue->push(original_wr_mf); } m_request_tracker.erase(mf); @@ -379,13 +384,13 @@ void memory_sub_partition::cache_cycle( unsigned cycle ) mem_fetch *mf = m_dram_L2_queue->top(); if ( !m_config->m_L2_config.disabled() && m_L2cache->waiting_for_fill(mf) ) { if (m_L2cache->fill_port_free()) { - mf->set_status(IN_PARTITION_L2_FILL_QUEUE,gpu_sim_cycle+gpu_tot_sim_cycle); - m_L2cache->fill(mf,gpu_sim_cycle+gpu_tot_sim_cycle+m_memcpy_cycle_offset); + mf->set_status(IN_PARTITION_L2_FILL_QUEUE,m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle); + m_L2cache->fill(mf,m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle+m_memcpy_cycle_offset); m_dram_L2_queue->pop(); } } else if ( !m_L2_icnt_queue->full() ) { if(mf->is_write() && mf->get_type() == WRITE_ACK) - mf->set_status(IN_PARTITION_L2_TO_ICNT_QUEUE,gpu_sim_cycle+gpu_tot_sim_cycle); + mf->set_status(IN_PARTITION_L2_TO_ICNT_QUEUE,m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle); m_L2_icnt_queue->push(mf); m_dram_L2_queue->pop(); } @@ -406,7 +411,7 @@ void memory_sub_partition::cache_cycle( unsigned cycle ) bool port_free = m_L2cache->data_port_free(); if ( !output_full && port_free ) { std::list events; - enum cache_request_status status = m_L2cache->access(mf->get_addr(),mf,gpu_sim_cycle+gpu_tot_sim_cycle+m_memcpy_cycle_offset,events); + enum cache_request_status status = m_L2cache->access(mf->get_addr(),mf,m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle+m_memcpy_cycle_offset,events); bool write_sent = was_write_sent(events); bool read_sent = was_read_sent(events); MEM_SUBPART_DPRINTF("Probing L2 cache Address=%llx, status=%u\n", mf->get_addr(), status); @@ -420,7 +425,7 @@ void memory_sub_partition::cache_cycle( unsigned cycle ) delete mf; } else { mf->set_reply(); - mf->set_status(IN_PARTITION_L2_TO_ICNT_QUEUE,gpu_sim_cycle+gpu_tot_sim_cycle); + mf->set_status(IN_PARTITION_L2_TO_ICNT_QUEUE,m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle); m_L2_icnt_queue->push(mf); } m_icnt_L2_queue->pop(); @@ -431,7 +436,7 @@ void memory_sub_partition::cache_cycle( unsigned cycle ) } else if ( status != RESERVATION_FAIL ) { if(mf->is_write() && (m_config->m_L2_config.m_write_alloc_policy == FETCH_ON_WRITE || m_config->m_L2_config.m_write_alloc_policy == LAZY_FETCH_ON_READ) && !was_writeallocate_sent(events)) { mf->set_reply(); - mf->set_status(IN_PARTITION_L2_TO_ICNT_QUEUE,gpu_sim_cycle+gpu_tot_sim_cycle); + mf->set_status(IN_PARTITION_L2_TO_ICNT_QUEUE,m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle); m_L2_icnt_queue->push(mf); } // L2 cache accepted request @@ -444,7 +449,7 @@ void memory_sub_partition::cache_cycle( unsigned cycle ) } } else { // L2 is disabled or non-texture access to texture-only L2 - mf->set_status(IN_PARTITION_L2_TO_DRAM_QUEUE,gpu_sim_cycle+gpu_tot_sim_cycle); + mf->set_status(IN_PARTITION_L2_TO_DRAM_QUEUE,m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle); m_L2_dram_queue->push(mf); m_icnt_L2_queue->pop(); } @@ -455,7 +460,7 @@ void memory_sub_partition::cache_cycle( unsigned cycle ) mem_fetch* mf = m_rop.front().req; m_rop.pop(); m_icnt_L2_queue->push(mf); - mf->set_status(IN_PARTITION_ICNT_TO_L2_QUEUE,gpu_sim_cycle+gpu_tot_sim_cycle); + mf->set_status(IN_PARTITION_ICNT_TO_L2_QUEUE,m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle); } } @@ -635,6 +640,7 @@ std::vector memory_sub_partition::breakdown_request_to_sector_reques mf->get_sid(), mf->get_tpc(), mf->get_mem_config(), + m_gpu->gpu_tot_sim_cycle+m_gpu->gpu_sim_cycle, mf); result.push_back(n_mf); @@ -664,13 +670,13 @@ void memory_sub_partition::push( mem_fetch* m_req, unsigned long long cycle ) m_request_tracker.insert(req); if( req->istexture() ) { m_icnt_L2_queue->push(req); - req->set_status(IN_PARTITION_ICNT_TO_L2_QUEUE,gpu_sim_cycle+gpu_tot_sim_cycle); + req->set_status(IN_PARTITION_ICNT_TO_L2_QUEUE,m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle); } else { rop_delay_t r; r.req = req; r.ready_cycle = cycle + m_config->rop_latency; m_rop.push(r); - req->set_status(IN_PARTITION_ROP_DELAY,gpu_sim_cycle+gpu_tot_sim_cycle); + req->set_status(IN_PARTITION_ROP_DELAY,m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle); } } } diff --git a/src/gpgpu-sim/l2cache.h b/src/gpgpu-sim/l2cache.h index 18c0a8b..c8a213c 100644 --- a/src/gpgpu-sim/l2cache.h +++ b/src/gpgpu-sim/l2cache.h @@ -42,12 +42,12 @@ public: { m_memory_config = config; } - virtual mem_fetch * alloc(const class warp_inst_t &inst, const mem_access_t &access) const + virtual mem_fetch * alloc(const class warp_inst_t &inst, const mem_access_t &access, unsigned long long cycle) const { abort(); return NULL; } - virtual mem_fetch * alloc(new_addr_type addr, mem_access_type type, unsigned size, bool wr) const; + virtual mem_fetch * alloc(new_addr_type addr, mem_access_type type, unsigned size, bool wr, unsigned long long cycle) const; private: const memory_config *m_memory_config; }; @@ -58,7 +58,7 @@ private: class memory_partition_unit { public: - memory_partition_unit( unsigned partition_id, const struct memory_config *config, class memory_stats_t *stats ); + memory_partition_unit( unsigned partition_id, const struct memory_config *config, class memory_stats_t *stats, class gpgpu_sim* gpu ); ~memory_partition_unit(); bool busy() const; @@ -93,6 +93,8 @@ public: unsigned get_mpid() const { return m_id; } + class gpgpu_sim* get_mgpu() const { return m_gpu; } + private: unsigned m_id; @@ -140,12 +142,14 @@ private: class mem_fetch* req; }; std::list m_dram_latency_queue; + + class gpgpu_sim* m_gpu; }; class memory_sub_partition { public: - memory_sub_partition( unsigned sub_partition_id, const struct memory_config *config, class memory_stats_t *stats ); + memory_sub_partition( unsigned sub_partition_id, const struct memory_config *config, class memory_stats_t *stats, class gpgpu_sim* gpu ); ~memory_sub_partition(); unsigned get_id() const { return m_id; } @@ -192,6 +196,7 @@ private: const struct memory_config *m_config; class l2_cache *m_L2cache; class L2interface *m_L2interface; + class gpgpu_sim* m_gpu; partition_mf_allocator *m_mf_allocator; // model delay of ROP units with a fixed latency diff --git a/src/gpgpu-sim/l2cache_trace.h b/src/gpgpu-sim/l2cache_trace.h index 2235cdc..d2dd948 100644 --- a/src/gpgpu-sim/l2cache_trace.h +++ b/src/gpgpu-sim/l2cache_trace.h @@ -42,7 +42,7 @@ #define MEMPART_DPRINTF(...) do {\ if (MEMPART_DTRACE(MEMORY_PARTITION_UNIT)) {\ printf( MEMPART_PRINT_STR,\ - gpu_sim_cycle + gpu_tot_sim_cycle,\ + m_gpu->gpu_sim_cycle + m_gpu->gpu_tot_sim_cycle,\ Trace::trace_streams_str[Trace::MEMORY_PARTITION_UNIT],\ get_mpid() );\ printf(__VA_ARGS__);\ @@ -52,7 +52,7 @@ #define MEM_SUBPART_DPRINTF(...) do {\ if (MEM_SUBPART_DTRACE(MEMORY_PARTITION_UNIT)) {\ printf( MEM_SUBPART_PRINT_STR,\ - gpu_sim_cycle + gpu_tot_sim_cycle,\ + m_gpu->gpu_sim_cycle + m_gpu->gpu_tot_sim_cycle,\ Trace::trace_streams_str[Trace::MEMORY_SUBPARTITION_UNIT],\ m_id );\ printf(__VA_ARGS__);\ diff --git a/src/gpgpu-sim/mem_fetch.cc b/src/gpgpu-sim/mem_fetch.cc index a260a35..c9b0484 100644 --- a/src/gpgpu-sim/mem_fetch.cc +++ b/src/gpgpu-sim/mem_fetch.cc @@ -40,6 +40,7 @@ mem_fetch::mem_fetch( const mem_access_t &access, unsigned sid, unsigned tpc, const struct memory_config *config, + unsigned long long cycle, mem_fetch *m_original_mf, mem_fetch *m_original_wr_mf) @@ -58,10 +59,10 @@ mem_fetch::mem_fetch( const mem_access_t &access, config->m_address_mapping.addrdec_tlx(access.get_addr(),&m_raw_addr); m_partition_addr = config->m_address_mapping.partition_address(access.get_addr()); m_type = m_access.is_write()?WRITE_REQUEST:READ_REQUEST; - m_timestamp = gpu_sim_cycle + gpu_tot_sim_cycle; + m_timestamp = cycle; m_timestamp2 = 0; m_status = MEM_FETCH_INITIALIZED; - m_status_change = gpu_sim_cycle + gpu_tot_sim_cycle; + m_status_change = cycle; m_mem_config = config; icnt_flit_size = config->icnt_flit_size; original_mf = m_original_mf; diff --git a/src/gpgpu-sim/mem_fetch.h b/src/gpgpu-sim/mem_fetch.h index e5efffd..4eb3a52 100644 --- a/src/gpgpu-sim/mem_fetch.h +++ b/src/gpgpu-sim/mem_fetch.h @@ -56,6 +56,7 @@ public: unsigned sid, unsigned tpc, const struct memory_config *config, + unsigned long long cycle, mem_fetch *original_mf = NULL, mem_fetch *original_wr_mf = NULL); ~mem_fetch(); diff --git a/src/gpgpu-sim/mem_latency_stat.cc b/src/gpgpu-sim/mem_latency_stat.cc index 7f6cde9..c7d20d1 100644 --- a/src/gpgpu-sim/mem_latency_stat.cc +++ b/src/gpgpu-sim/mem_latency_stat.cc @@ -42,7 +42,7 @@ #include #include -memory_stats_t::memory_stats_t( unsigned n_shader, const struct shader_core_config *shader_config, const struct memory_config *mem_config ) +memory_stats_t::memory_stats_t( unsigned n_shader, const struct shader_core_config *shader_config, const struct memory_config *mem_config, const class gpgpu_sim* gpu ) { assert( mem_config->m_valid ); assert( shader_config->m_valid ); @@ -67,6 +67,7 @@ memory_stats_t::memory_stats_t( unsigned n_shader, const struct shader_core_conf m_n_shader=n_shader; m_memory_config=mem_config; + m_gpu=gpu; total_n_access=0; total_n_reads=0; total_n_writes=0; @@ -141,7 +142,7 @@ memory_stats_t::memory_stats_t( unsigned n_shader, const struct shader_core_conf unsigned memory_stats_t::memlatstat_done(mem_fetch *mf ) { unsigned mf_latency; - mf_latency = (gpu_sim_cycle+gpu_tot_sim_cycle) - mf->get_timestamp(); + mf_latency = (m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle) - mf->get_timestamp(); mf_num_lat_pw++; mf_tot_lat_pw += mf_latency; unsigned idx = LOGB2(mf_latency); @@ -161,7 +162,7 @@ void memory_stats_t::memlatstat_read_done(mem_fetch *mf) if (mf_latency > mf_max_lat_table[mf->get_tlx_addr().chip][mf->get_tlx_addr().bk]) mf_max_lat_table[mf->get_tlx_addr().chip][mf->get_tlx_addr().bk] = mf_latency; unsigned icnt2sh_latency; - icnt2sh_latency = (gpu_tot_sim_cycle+gpu_sim_cycle) - mf->get_return_timestamp(); + icnt2sh_latency = (m_gpu->gpu_tot_sim_cycle+m_gpu->gpu_sim_cycle) - mf->get_return_timestamp(); tot_icnt2sh_latency += icnt2sh_latency; icnt2sh_lat_table[LOGB2(icnt2sh_latency)]++; if (icnt2sh_latency > max_icnt2sh_latency) @@ -195,7 +196,7 @@ void memory_stats_t::memlatstat_icnt2mem_pop(mem_fetch *mf) { if (m_memory_config->gpgpu_memlatency_stat) { unsigned icnt2mem_latency; - icnt2mem_latency = (gpu_tot_sim_cycle+gpu_sim_cycle) - mf->get_timestamp(); + icnt2mem_latency = (m_gpu->gpu_tot_sim_cycle+m_gpu->gpu_sim_cycle) - mf->get_timestamp(); tot_icnt2mem_latency += icnt2mem_latency; icnt2mem_lat_table[LOGB2(icnt2mem_latency)]++; if (icnt2mem_latency > max_icnt2mem_latency) diff --git a/src/gpgpu-sim/mem_latency_stat.h b/src/gpgpu-sim/mem_latency_stat.h index 5b89202..b86740d 100644 --- a/src/gpgpu-sim/mem_latency_stat.h +++ b/src/gpgpu-sim/mem_latency_stat.h @@ -36,7 +36,8 @@ class memory_stats_t { public: memory_stats_t( unsigned n_shader, const struct shader_core_config *shader_config, - const struct memory_config *mem_config ); + const struct memory_config *mem_config, + const class gpgpu_sim* gpu); unsigned memlatstat_done( class mem_fetch *mf ); void memlatstat_read_done( class mem_fetch *mf ); @@ -51,6 +52,7 @@ public: const struct shader_core_config *m_shader_config; const struct memory_config *m_memory_config; + const class gpgpu_sim* m_gpu; unsigned max_mrq_latency; unsigned max_dq_latency; diff --git a/src/gpgpu-sim/scoreboard.cc b/src/gpgpu-sim/scoreboard.cc index ebec891..80f95c6 100644 --- a/src/gpgpu-sim/scoreboard.cc +++ b/src/gpgpu-sim/scoreboard.cc @@ -32,13 +32,15 @@ //Constructor -Scoreboard::Scoreboard( unsigned sid, unsigned n_warps ) +Scoreboard::Scoreboard( unsigned sid, unsigned n_warps, class gpgpu_t* gpu ) : longopregs() { m_sid = sid; //Initialize size of table reg_table.resize(n_warps); longopregs.resize(n_warps); + + m_gpu = gpu; } // Print scoreboard contents diff --git a/src/gpgpu-sim/scoreboard.h b/src/gpgpu-sim/scoreboard.h index 4a76ea3..a4baa19 100644 --- a/src/gpgpu-sim/scoreboard.h +++ b/src/gpgpu-sim/scoreboard.h @@ -38,7 +38,7 @@ class Scoreboard { public: - Scoreboard( unsigned sid, unsigned n_warps ); + Scoreboard( unsigned sid, unsigned n_warps, class gpgpu_t* gpu ); void reserveRegisters(const warp_inst_t *inst); void releaseRegisters(const warp_inst_t *inst); @@ -59,6 +59,8 @@ private: std::vector< std::set > reg_table; //Register that depend on a long operation (global, local or tex memory) std::vector< std::set > longopregs; + + class gpgpu_t* m_gpu; }; diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc index 007ad42..69b619a 100644 --- a/src/gpgpu-sim/shader.cc +++ b/src/gpgpu-sim/shader.cc @@ -133,7 +133,7 @@ shader_core_ctx::shader_core_ctx( class gpgpu_sim *gpu, m_L1I = new read_only_cache( name,m_config->m_L1I_config,m_sid,get_shader_instruction_cache_id(),m_icnt,IN_L1I_MISS_QUEUE); m_warp.resize(m_config->max_warps_per_shader, shd_warp_t(this, warp_size)); - m_scoreboard = new Scoreboard(m_sid, m_config->max_warps_per_shader); + m_scoreboard = new Scoreboard(m_sid, m_config->max_warps_per_shader, gpu); //scedulers //must currently occur after all inputs have been initialized. @@ -767,7 +767,7 @@ void shader_core_ctx::fetch() m_inst_fetch_buffer = ifetch_buffer_t(m_warp[mf->get_wid()].get_pc(), mf->get_access_size(), mf->get_wid()); assert( m_warp[mf->get_wid()].get_pc() == (mf->get_addr()-PROGRAM_MEM_START)); // Verify that we got the instruction we were expecting. m_inst_fetch_buffer.m_valid = true; - m_warp[mf->get_wid()].set_last_fetch(gpu_sim_cycle); + m_warp[mf->get_wid()].set_last_fetch(m_gpu->gpu_sim_cycle); delete mf; } else { @@ -815,17 +815,19 @@ void shader_core_ctx::fetch() warp_id, m_sid, m_tpc, - m_memory_config ); + m_memory_config, + m_gpu->gpu_tot_sim_cycle+m_gpu->gpu_sim_cycle + ); std::list events; - enum cache_request_status status = m_L1I->access( (new_addr_type)ppc, mf, gpu_sim_cycle+gpu_tot_sim_cycle,events); + enum cache_request_status status = m_L1I->access( (new_addr_type)ppc, mf, m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle,events); if( status == MISS ) { m_last_warp_fetched=warp_id; m_warp[warp_id].set_imiss_pending(); - m_warp[warp_id].set_last_fetch(gpu_sim_cycle); + m_warp[warp_id].set_last_fetch(m_gpu->gpu_sim_cycle); } else if( status == HIT ) { m_last_warp_fetched=warp_id; m_inst_fetch_buffer = ifetch_buffer_t(pc,nbytes,warp_id); - m_warp[warp_id].set_last_fetch(gpu_sim_cycle); + m_warp[warp_id].set_last_fetch(m_gpu->gpu_sim_cycle); delete mf; } else { m_last_warp_fetched=warp_id; @@ -859,7 +861,7 @@ void shader_core_ctx::issue_warp( register_set& pipe_reg_set, const warp_inst_t* m_warp[warp_id].ibuffer_free(); assert(next_inst->valid()); **pipe_reg = *next_inst; // static instruction information - (*pipe_reg)->issue( active_mask, warp_id, gpu_tot_sim_cycle + gpu_sim_cycle, m_warp[warp_id].get_dynamic_warp_id(), sch_id ); // dynamic instruction information + (*pipe_reg)->issue( active_mask, warp_id, m_gpu->gpu_tot_sim_cycle + m_gpu->gpu_sim_cycle, m_warp[warp_id].get_dynamic_warp_id(), sch_id ); // dynamic instruction information m_stats->shader_cycle_distro[2+(*pipe_reg)->active_count()]++; func_exec_inst( **pipe_reg ); if( next_inst->op == BARRIER_OP ){ @@ -1514,7 +1516,7 @@ void shader_core_ctx::warp_inst_complete(const warp_inst_t &inst) m_stats->m_num_sim_winsn[m_sid]++; m_gpu->gpu_sim_insn += inst.active_count(); - inst.completed(gpu_tot_sim_cycle + gpu_sim_cycle); + inst.completed(m_gpu->gpu_tot_sim_cycle + m_gpu->gpu_sim_cycle); } void shader_core_ctx::writeback() @@ -1552,9 +1554,9 @@ void shader_core_ctx::writeback() m_warp[warp_id].dec_inst_in_pipeline(); warp_inst_complete(*pipe_reg); m_gpu->gpu_sim_insn_last_update_sid = m_sid; - m_gpu->gpu_sim_insn_last_update = gpu_sim_cycle; - m_last_inst_gpu_sim_cycle = gpu_sim_cycle; - m_last_inst_gpu_tot_sim_cycle = gpu_tot_sim_cycle; + m_gpu->gpu_sim_insn_last_update = m_gpu->gpu_sim_cycle; + m_last_inst_gpu_sim_cycle = m_gpu->gpu_sim_cycle; + m_last_inst_gpu_tot_sim_cycle = m_gpu->gpu_tot_sim_cycle; pipe_reg->clear(); preg = m_pipeline_reg[EX_WB].get_ready(); pipe_reg = (preg==NULL)? NULL:*preg; @@ -1633,9 +1635,9 @@ mem_stage_stall_type ldst_unit::process_memory_access_queue( cache_t *cache, war return DATA_PORT_STALL; //const mem_access_t &access = inst.accessq_back(); - mem_fetch *mf = m_mf_allocator->alloc(inst,inst.accessq_back()); + mem_fetch *mf = m_mf_allocator->alloc(inst,inst.accessq_back(),m_core->get_gpu()->gpu_sim_cycle+m_core->get_gpu()->gpu_tot_sim_cycle); std::list events; - enum cache_request_status status = cache->access(mf->get_addr(),mf,gpu_sim_cycle+gpu_tot_sim_cycle,events); + enum cache_request_status status = cache->access(mf->get_addr(),mf,m_core->get_gpu()->gpu_sim_cycle+m_core->get_gpu()->gpu_tot_sim_cycle,events); return process_cache_access( cache, mf->get_addr(), inst, events, mf, status ); } @@ -1645,7 +1647,7 @@ mem_stage_stall_type ldst_unit::process_memory_access_queue_l1cache( l1_cache *c if( inst.accessq_empty() ) return result; - mem_fetch *mf = m_mf_allocator->alloc(inst,inst.accessq_back()); + mem_fetch *mf = m_mf_allocator->alloc(inst,inst.accessq_back(),m_core->get_gpu()->gpu_sim_cycle+m_core->get_gpu()->gpu_tot_sim_cycle); if(m_config->m_L1D_config.l1_latency > 0) { @@ -1675,7 +1677,7 @@ mem_stage_stall_type ldst_unit::process_memory_access_queue_l1cache( l1_cache *c else { std::list events; - enum cache_request_status status = cache->access(mf->get_addr(),mf,gpu_sim_cycle+gpu_tot_sim_cycle,events); + enum cache_request_status status = cache->access(mf->get_addr(),mf,m_core->get_gpu()->gpu_sim_cycle+m_core->get_gpu()->gpu_tot_sim_cycle,events); return process_cache_access( cache, mf->get_addr(), inst, events, mf, status ); } } @@ -1687,7 +1689,7 @@ void ldst_unit::L1_latency_queue_cycle() { mem_fetch* mf_next = l1_latency_queue[0]; std::list events; - enum cache_request_status status = m_L1D->access(mf_next->get_addr(),mf_next,gpu_sim_cycle+gpu_tot_sim_cycle,events); + enum cache_request_status status = m_L1D->access(mf_next->get_addr(),mf_next,m_core->get_gpu()->gpu_sim_cycle+m_core->get_gpu()->gpu_tot_sim_cycle,events); bool write_sent = was_write_sent(events); bool read_sent = was_read_sent(events); @@ -1804,7 +1806,7 @@ bool ldst_unit::memory_cycle( warp_inst_t &inst, mem_stage_stall_type &stall_rea if( m_icnt->full(size, inst.is_store() || inst.isatomic()) ) { stall_cond = ICNT_RC_FAIL; } else { - mem_fetch *mf = m_mf_allocator->alloc(inst,access); + mem_fetch *mf = m_mf_allocator->alloc(inst,access,m_core->get_gpu()->gpu_sim_cycle+m_core->get_gpu()->gpu_tot_sim_cycle); m_icnt->push(mf); inst.accessq_pop_back(); //inst.clear_active( access.get_warp_mask() ); @@ -1840,7 +1842,7 @@ bool ldst_unit::response_buffer_full() const void ldst_unit::fill( mem_fetch *mf ) { - mf->set_status(IN_SHADER_LDST_RESPONSE_FIFO,gpu_sim_cycle+gpu_tot_sim_cycle); + mf->set_status(IN_SHADER_LDST_RESPONSE_FIFO,m_core->get_gpu()->gpu_sim_cycle+m_core->get_gpu()->gpu_tot_sim_cycle); m_response_fifo.push_back(mf); } @@ -2115,7 +2117,8 @@ ldst_unit::ldst_unit( mem_fetch_interface *icnt, get_shader_normal_cache_id(), m_icnt, m_mf_allocator, - IN_L1D_MISS_QUEUE ); + IN_L1D_MISS_QUEUE, + core->get_gpu()); if(m_config->m_L1D_config.l1_latency > 0) { @@ -2202,8 +2205,8 @@ void ldst_unit::writeback() m_core->warp_inst_complete(m_next_wb); } m_next_wb.clear(); - m_last_inst_gpu_sim_cycle = gpu_sim_cycle; - m_last_inst_gpu_tot_sim_cycle = gpu_tot_sim_cycle; + m_last_inst_gpu_sim_cycle = m_core->get_gpu()->gpu_sim_cycle; + m_last_inst_gpu_tot_sim_cycle = m_core->get_gpu()->gpu_tot_sim_cycle; } } @@ -2311,13 +2314,13 @@ void ldst_unit::cycle() mem_fetch *mf = m_response_fifo.front(); if (mf->get_access_type() == TEXTURE_ACC_R) { if (m_L1T->fill_port_free()) { - m_L1T->fill(mf,gpu_sim_cycle+gpu_tot_sim_cycle); + m_L1T->fill(mf,m_core->get_gpu()->gpu_sim_cycle+m_core->get_gpu()->gpu_tot_sim_cycle); m_response_fifo.pop_front(); } } else if (mf->get_access_type() == CONST_ACC_R) { if (m_L1C->fill_port_free()) { - mf->set_status(IN_SHADER_FETCHED,gpu_sim_cycle+gpu_tot_sim_cycle); - m_L1C->fill(mf,gpu_sim_cycle+gpu_tot_sim_cycle); + mf->set_status(IN_SHADER_FETCHED,m_core->get_gpu()->gpu_sim_cycle+m_core->get_gpu()->gpu_tot_sim_cycle); + m_L1C->fill(mf,m_core->get_gpu()->gpu_sim_cycle+m_core->get_gpu()->gpu_tot_sim_cycle); m_response_fifo.pop_front(); } } else { @@ -2337,13 +2340,13 @@ void ldst_unit::cycle() } if( bypassL1D ) { if ( m_next_global == NULL ) { - mf->set_status(IN_SHADER_FETCHED,gpu_sim_cycle+gpu_tot_sim_cycle); + mf->set_status(IN_SHADER_FETCHED,m_core->get_gpu()->gpu_sim_cycle+m_core->get_gpu()->gpu_tot_sim_cycle); m_response_fifo.pop_front(); m_next_global = mf; } } else { if (m_L1D->fill_port_free()) { - m_L1D->fill(mf,gpu_sim_cycle+gpu_tot_sim_cycle); + m_L1D->fill(mf,m_core->get_gpu()->gpu_sim_cycle+m_core->get_gpu()->gpu_tot_sim_cycle); m_response_fifo.pop_front(); } } @@ -2432,7 +2435,7 @@ void shader_core_ctx::register_cta_thread_exit( unsigned cta_num, kernel_info_t shader_CTA_count_unlog(m_sid, 1); SHADER_DPRINTF(LIVENESS, "GPGPU-Sim uArch: Finished CTA #%d (%lld,%lld), %u CTAs running\n", - cta_num, gpu_sim_cycle, gpu_tot_sim_cycle, m_n_active_cta); + cta_num, m_gpu->gpu_sim_cycle, m_gpu->gpu_tot_sim_cycle, m_n_active_cta); if( m_n_active_cta == 0 ) { SHADER_DPRINTF(LIVENESS, "GPGPU-Sim uArch: Empty (last released kernel %u \'%s\').\n", @@ -2823,7 +2826,7 @@ void shader_core_ctx::display_pipeline(FILE *fout, int print_mem, int mask ) con { fprintf(fout, "=================================================\n"); fprintf(fout, "shader %u at cycle %Lu+%Lu (%u threads running)\n", m_sid, - gpu_tot_sim_cycle, gpu_sim_cycle, m_not_completed); + m_gpu->gpu_tot_sim_cycle, m_gpu->gpu_sim_cycle, m_not_completed); fprintf(fout, "=================================================\n"); dump_warp_state(fout); @@ -3221,7 +3224,7 @@ void barrier_set_t::warp_reaches_barrier(unsigned cta_id,unsigned warp_id,warp_i cta_to_warp_t::iterator w=m_cta_to_warps.find(cta_id); if( w == m_cta_to_warps.end() ) { // cta is active - printf("ERROR ** cta_id %u not found in barrier set on cycle %llu+%llu...\n", cta_id, gpu_tot_sim_cycle, gpu_sim_cycle ); + printf("ERROR ** cta_id %u not found in barrier set on cycle %llu+%llu...\n", cta_id, m_shader->get_gpu()->gpu_tot_sim_cycle, m_shader->get_gpu()->gpu_sim_cycle ); dump(); abort(); } @@ -3394,8 +3397,8 @@ bool shader_core_ctx::fetch_unit_response_buffer_full() const void shader_core_ctx::accept_fetch_response( mem_fetch *mf ) { - mf->set_status(IN_SHADER_FETCHED,gpu_sim_cycle+gpu_tot_sim_cycle); - m_L1I->fill(mf,gpu_sim_cycle+gpu_tot_sim_cycle); + mf->set_status(IN_SHADER_FETCHED,m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle); + m_L1I->fill(mf,m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle); } bool shader_core_ctx::ldst_unit_response_buffer_full() const @@ -3957,7 +3960,7 @@ void simt_core_cluster::icnt_inject_request_packet(class mem_fetch *mf) } m_stats->m_outgoing_traffic_stats->record_traffic(mf, packet_size); unsigned destination = mf->get_sub_partition_id(); - mf->set_status(IN_ICNT_TO_MEM,gpu_sim_cycle+gpu_tot_sim_cycle); + mf->set_status(IN_ICNT_TO_MEM,m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle); if (!mf->get_is_write() && !mf->isatomic()) ::icnt_push(m_cluster_id, m_config->mem2device(destination), (void*)mf, mf->get_ctrl_size() ); else @@ -3996,7 +3999,7 @@ void simt_core_cluster::icnt_cycle() // - For write-ack, the packet only has control metadata unsigned int packet_size = (mf->get_is_write())? mf->get_ctrl_size() : mf->size(); m_stats->m_incoming_traffic_stats->record_traffic(mf, packet_size); - mf->set_status(IN_CLUSTER_TO_SHADER_QUEUE,gpu_sim_cycle+gpu_tot_sim_cycle); + mf->set_status(IN_CLUSTER_TO_SHADER_QUEUE,m_gpu->gpu_sim_cycle+m_gpu->gpu_tot_sim_cycle); //m_memory_stats->memlatstat_read_done(mf,m_shader_config->max_warps_per_shader); m_response_fifo.push_back(mf); m_stats->n_mem_to_simt[m_cluster_id] += mf->get_num_flits(false); diff --git a/src/gpgpu-sim/shader.h b/src/gpgpu-sim/shader.h index a0c2b63..fde87b6 100644 --- a/src/gpgpu-sim/shader.h +++ b/src/gpgpu-sim/shader.h @@ -1730,7 +1730,7 @@ public: m_cluster_id = cluster_id; m_memory_config = config; } - mem_fetch *alloc( new_addr_type addr, mem_access_type type, unsigned size, bool wr ) const + mem_fetch *alloc( new_addr_type addr, mem_access_type type, unsigned size, bool wr, unsigned long long cycle ) const { mem_access_t access( type, addr, size, wr ); mem_fetch *mf = new mem_fetch( access, @@ -1739,11 +1739,12 @@ public: -1, m_core_id, m_cluster_id, - m_memory_config ); + m_memory_config, + cycle); return mf; } - mem_fetch *alloc( const warp_inst_t &inst, const mem_access_t &access ) const + mem_fetch *alloc( const warp_inst_t &inst, const mem_access_t &access, unsigned long long cycle ) const { warp_inst_t inst_copy = inst; mem_fetch *mf = new mem_fetch(access, @@ -1752,7 +1753,8 @@ public: inst.warp_id(), m_core_id, m_cluster_id, - m_memory_config); + m_memory_config, + cycle); return mf; } diff --git a/src/gpgpu-sim/shader_trace.h b/src/gpgpu-sim/shader_trace.h index de3e059..ac4e894 100644 --- a/src/gpgpu-sim/shader_trace.h +++ b/src/gpgpu-sim/shader_trace.h @@ -44,7 +44,7 @@ #define SHADER_DPRINTF(x, ...) do {\ if (SHADER_DTRACE(x)) {\ printf( SHADER_PRINT_STR,\ - gpu_sim_cycle + gpu_tot_sim_cycle,\ + m_gpu->gpu_sim_cycle + m_gpu->gpu_tot_sim_cycle,\ Trace::trace_streams_str[Trace::x],\ get_sid() );\ printf(__VA_ARGS__);\ @@ -56,7 +56,7 @@ #define SCHED_DPRINTF(...) do {\ if (SHADER_DTRACE(WARP_SCHEDULER)) {\ printf( SCHED_PRINT_STR,\ - gpu_sim_cycle + gpu_tot_sim_cycle,\ + m_shader->get_gpu()->gpu_sim_cycle + m_shader->get_gpu()->gpu_tot_sim_cycle,\ Trace::trace_streams_str[Trace::WARP_SCHEDULER],\ get_sid(),\ m_id );\ diff --git a/src/gpgpusim_entrypoint.cc b/src/gpgpusim_entrypoint.cc index 9e2bfa2..777f9bb 100644 --- a/src/gpgpusim_entrypoint.cc +++ b/src/gpgpusim_entrypoint.cc @@ -275,7 +275,7 @@ void print_simulation_time() printf("\n\ngpgpu_simulation_time = %u days, %u hrs, %u min, %u sec (%u sec)\n", (unsigned)d, (unsigned)h, (unsigned)m, (unsigned)s, (unsigned)difference ); printf("gpgpu_simulation_rate = %u (inst/sec)\n", (unsigned)(g_the_gpu->gpu_tot_sim_insn / difference) ); - printf("gpgpu_simulation_rate = %u (cycle/sec)\n", (unsigned)(gpu_tot_sim_cycle / difference) ); + printf("gpgpu_simulation_rate = %u (cycle/sec)\n", (unsigned)(g_the_gpu->gpu_tot_sim_cycle / difference) ); fflush(stdout); } diff --git a/src/stream_manager.cc b/src/stream_manager.cc index 6cd62a2..e07f4e4 100644 --- a/src/stream_manager.cc +++ b/src/stream_manager.cc @@ -190,7 +190,7 @@ bool stream_operation::do_operation( gpgpu_sim *gpu ) case stream_event: { printf("event update\n"); time_t wallclock = time((time_t *)NULL); - m_event->update( gpu_tot_sim_cycle, wallclock ); + m_event->update( gpu->gpu_tot_sim_cycle, wallclock ); m_stream->record_next_done(); } break; diff --git a/src/trace.h b/src/trace.h index a79b4a0..0b96dcf 100644 --- a/src/trace.h +++ b/src/trace.h @@ -31,9 +31,6 @@ #ifndef __TRACE_H__ #define __TRACE_H__ -extern unsigned long long gpu_sim_cycle; -extern unsigned long long gpu_tot_sim_cycle; - namespace Trace { #define TS_TUP_BEGIN(X) enum X { @@ -63,17 +60,26 @@ namespace Trace { #define DPRINTF(x, ...) do {\ if (DTRACE(x)) {\ printf( SIM_PRINT_STR,\ - gpu_sim_cycle + gpu_tot_sim_cycle,\ + m_gpu->gpu_sim_cycle + m_gpu->gpu_tot_sim_cycle,\ Trace::trace_streams_str[Trace::x] );\ printf(__VA_ARGS__);\ }\ } while (0) +#define DPRINTFG(x, ...) do {\ + if (DTRACE(x)) {\ + printf( SIM_PRINT_STR,\ + gpu_sim_cycle + gpu_tot_sim_cycle,\ + Trace::trace_streams_str[Trace::x] );\ + printf(__VA_ARGS__);\ + }\ +} while (0) #else #define DTRACE(x) (false) #define DPRINTF(x, ...) do {} while (0) +#define DPRINTFG(x, ...) do {} while (0) #endif -- cgit v1.3