1. modify shader_core_ctx::execute_pipe() to model instruction throughput correctly (off-by-one error)

2. adding code to dump_pipeline to display reason for memory stage stalling
3. removing dead code resulting from prior changes

correlation vs. GT200 is 0.95 (need to add back modeling of memory writebacks, shared memory latency)

[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7831]

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();