diff options
| -rw-r--r-- | libcuda/cuda_runtime_api.cc | 35 | ||||
| -rw-r--r-- | src/gpgpu-sim/shader.cc | 26 | ||||
| -rw-r--r-- | src/gpgpu-sim/shader.h | 9 |
3 files changed, 19 insertions, 51 deletions
diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc index caf0ad2..df8ef4f 100644 --- a/libcuda/cuda_runtime_api.cc +++ b/libcuda/cuda_runtime_api.cc @@ -1479,8 +1479,6 @@ int CUDARTAPI __cudaSynchronizeThreads(void**, void*) //////// - - extern "C" int ptx_parse(); extern "C" int ptx__scan_string(const char*); extern "C" FILE *ptx_in; @@ -1492,39 +1490,6 @@ extern "C" FILE *ptxinfo_in; /// static functions - - -std::map<std::string,function_info*> *g_kernel_name_to_function_lookup=NULL; -std::map<const void*,std::string> *g_host_to_kernel_entrypoint_name_lookup=NULL; - -function_info *get_kernel(const char *kernel_key, std::string &kernel_func_name_mangled ) -{ - if ( g_host_to_kernel_entrypoint_name_lookup->find(kernel_key) == - g_host_to_kernel_entrypoint_name_lookup->end() ) { - printf("GPGPU-Sim PTX: ERROR ** cannot locate __global__ function from hostPtr\n" ); - printf("GPGPU-Sim PTX: registered PTX kernels: \n"); - std::map<const void*,std::string>::iterator i_eptr = g_host_to_kernel_entrypoint_name_lookup->begin(); - for (; i_eptr != g_host_to_kernel_entrypoint_name_lookup->end(); ++i_eptr) { - printf("GPGPU-Sim PTX: (%p,%s)\n", i_eptr->first, i_eptr->second.c_str()); - } - printf("\n"); - abort(); - } - kernel_func_name_mangled = (*g_host_to_kernel_entrypoint_name_lookup)[kernel_key]; - if ( g_kernel_name_to_function_lookup->find(kernel_func_name_mangled) == - g_kernel_name_to_function_lookup->end() ) { - printf("GPGPU-Sim PTX: ERROR ** function \'%s\' not found in ptx file\n", - kernel_func_name_mangled.c_str() ); - abort(); - } - return (*g_kernel_name_to_function_lookup)[kernel_func_name_mangled]; -} - - -void ptxinfo_cuda_addinfo( CUctx_st *context ) -{ -} - static int load_static_globals( symbol_table *symtab, unsigned min_gaddr, unsigned max_gaddr, gpgpu_t *gpu ) { printf( "GPGPU-Sim PTX: loading globals with explicit initializers... \n" ); diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc index 54e4ffa..1d51004 100644 --- a/src/gpgpu-sim/shader.cc +++ b/src/gpgpu-sim/shader.cc @@ -402,6 +402,7 @@ shader_core_ctx::shader_core_ctx( class gpgpu_sim *gpu, for (unsigned i = 0; i < config->max_warps_per_shader; ++i) m_pdom_warp[i] = new pdom_warp_ctx_t(i,this); m_shader_memory_new_instruction_processed = false; + m_mem_rc = NO_RC_FAIL, // Initialize scoreboard m_scoreboard = new Scoreboard(m_sid, m_config->max_warps_per_shader); @@ -419,11 +420,6 @@ shader_core_ctx::shader_core_ctx( class gpgpu_sim *gpu, m_memory_queue.texture.reserve(warp_size); m_memory_queue.local_global.reserve(warp_size); - // writeback - m_pl_tid = (int*) malloc(sizeof(int)*warp_size); - m_mshr_lat_info = (insn_latency_info*) malloc(sizeof(insn_latency_info) * warp_size); - m_pl_lat_info = (insn_latency_info*) malloc(sizeof(insn_latency_info) * warp_size); - // fetch m_last_warp_fetched = 0; m_last_warp_issued = 0; @@ -840,7 +836,7 @@ void shader_core_ctx::execute_pipe( unsigned pipeline, unsigned next_stage ) { if( !m_pipeline_reg[next_stage]->empty() ) return; - if( m_pipeline_reg[pipeline]->cycles ) { + if( m_pipeline_reg[pipeline]->cycles > 1 ) { m_pipeline_reg[pipeline]->cycles--; return; } @@ -1417,7 +1413,7 @@ void shader_core_ctx::memory() done &= memory_constant_cycle(rc_fail, type); done &= memory_texture_cycle(rc_fail, type); done &= memory_cycle(rc_fail, type); - + m_mem_rc = rc_fail; if (!done) { // log stall types and return assert(rc_fail != NO_RC_FAIL); m_stats->gpu_stall_shd_mem++; @@ -1425,7 +1421,7 @@ void shader_core_ctx::memory() return; } if( not m_pipeline_reg[MM_WB]->empty() ) - return; // writeback stalled + return; // writeback stalled move_warp(m_pipeline_reg[MM_WB],m_pipeline_reg[EX_MM]); } @@ -1631,6 +1627,20 @@ void shader_core_ctx::display_pipeline(FILE *fout, int print_mem, int mask ) } fprintf(fout, "EX/MEM = "); print_stage(EX_MM, fout, print_mem, mask); + if( m_mem_rc != NO_RC_FAIL ) { + fprintf(fout, "EX/MEM (stall condition: "); + switch ( m_mem_rc ) { + case BK_CONF: fprintf(fout,"BK_CONF"); break; + case MSHR_RC_FAIL: fprintf(fout,"MSHR_RC_FAIL"); break; + case ICNT_RC_FAIL: fprintf(fout,"ICNT_RC_FAIL"); break; + case COAL_STALL: fprintf(fout,"COAL_STALL"); break; + case WB_ICNT_RC_FAIL: fprintf(fout,"WB_ICNT_RC_FAIL"); break; + case WB_CACHE_RSRV_FAIL: fprintf(fout,"WB_CACHE_RSRV_FAIL"); break; + case N_MEM_STAGE_STALL_TYPE: fprintf(fout,"N_MEM_STAGE_STALL_TYPE"); break; + default: abort(); + } + fprintf(fout,")\n"); + } fprintf(fout, "MEM/WB = "); print_stage(MM_WB, fout, print_mem, mask); fprintf(fout, "\n"); diff --git a/src/gpgpu-sim/shader.h b/src/gpgpu-sim/shader.h index cdbaf37..a9e27b8 100644 --- a/src/gpgpu-sim/shader.h +++ b/src/gpgpu-sim/shader.h @@ -1130,8 +1130,6 @@ private: opndcoll_rfu_t m_operand_collector; mshr_shader_unit *m_mshr_unit; shader_queues_t m_memory_queue; - fifo_pipeline<std::vector<int> > *m_thd_commit_queue; - std::multiset<fixeddelay_queue_warp_t, fixeddelay_queue_warp_comp> m_fixeddelay_queue; // fetch int m_last_warp_fetched; @@ -1143,12 +1141,7 @@ private: cache_t *m_L1C; // constant cache bool m_shader_memory_new_instruction_processed; - int m_pending_mem_access; // number of memory access to be serviced (use for W0 classification) - - // used in writeback - int *m_pl_tid; - insn_latency_info *m_mshr_lat_info; - insn_latency_info *m_pl_lat_info; + enum mem_stage_stall_type m_mem_rc; }; void init_mshr_pool(); |