refactoring: making g_current_symbol_table local to ptx_ir.cc

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

5 files changed, 289 insertions, 298 deletions

diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc
index 4acefe5..74b18bc 100644
--- a/src/cuda-sim/cuda-sim.cc
+++ b/src/cuda-sim/cuda-sim.cc
@@ -222,266 +222,10 @@ int gpgpu_ptx_sim_sizeofTexture(const char* name)
    return array->size;
 }
 
-extern unsigned int warp_size; 
-
-int g_warn_literal_operands_two_type_inst;
-
-std::list<operand_info> check_operands( int opcode,
-                                        const std::list<int> &scalar_type,
-                                        const std::list<operand_info> &operands )
-{
-    if( (opcode == CVT_OP) || (opcode == SET_OP) || (opcode == SLCT_OP) || (opcode == TEX_OP) ) {
-        // just make sure these do not have have const operands... 
-        if( !g_warn_literal_operands_two_type_inst ) {
-            std::list<operand_info>::const_iterator o;
-            for( o = operands.begin(); o != operands.end(); o++ ) {
-                const operand_info &op = *o;
-                if( op.is_literal() ) {
-                    printf("GPGPU-Sim PTX: PTX uses two scalar type intruction with literal operand.\n");
-                    g_warn_literal_operands_two_type_inst = 1;
-                }
-            }
-        }
-    } else {
-        assert( scalar_type.size() < 2 );
-        if( scalar_type.size() == 1 ) {
-            std::list<operand_info> result;
-            int inst_type = scalar_type.front();
-            std::list<operand_info>::const_iterator o;
-            for( o = operands.begin(); o != operands.end(); o++ ) {
-                const operand_info &op = *o;
-                if( op.is_literal() ) {
-                    if( (op.get_type() == double_op_t) && (inst_type == F32_TYPE) ) {
-                        ptx_reg_t v = op.get_literal_value();
-                        float u = (float)v.f64;
-                        operand_info n(u);
-                        result.push_back(n);
-                    } else {
-                        result.push_back(op);
-                    }
-                } else {
-                        result.push_back(op);
-                }
-            }
-            return result;
-        } 
-    }
-    return operands;
-}
-
-ptx_instruction::ptx_instruction( int opcode, 
-                                  const symbol *pred, 
-                                  int neg_pred, 
-                                  symbol *label,
-                                  const std::list<operand_info> &operands, 
-                                  const operand_info &return_var,
-                                  const std::list<int> &options, 
-                                  const std::list<int> &scalar_type,
-                                  memory_space_t space_spec,
-                                  const char *file, 
-                                  unsigned line,
-                                  const char *source ) 
-{
-   m_uid = ++g_num_ptx_inst_uid;
-   m_PC = 0;
-   m_opcode = opcode;
-   m_pred = pred;
-   m_neg_pred = neg_pred;
-   m_label = label;
-   const std::list<operand_info> checked_operands = check_operands(opcode,scalar_type,operands);
-   m_operands.insert(m_operands.begin(), checked_operands.begin(), checked_operands.end() );
-   m_return_var = return_var;
-   m_options = options;
-   m_wide = false;
-   m_hi = false;
-   m_lo = false;
-   m_uni = false;
-   m_to_option = false;
-   m_cache_option = 0;
-   m_rounding_mode = RN_OPTION;
-   m_compare_op = -1;
-   m_saturation_mode = 0;
-   m_geom_spec = 0;
-   m_vector_spec = 0;
-   m_atomic_spec = 0;
-   m_warp_size = ::warp_size;
-   m_membar_level = 0;
-
-   std::list<int>::const_iterator i;
-   unsigned n=1;
-   for ( i=options.begin(); i!= options.end(); i++, n++ ) {
-      int last_ptx_inst_option = *i;
-      switch ( last_ptx_inst_option ) {
-      case EQU_OPTION:
-      case NEU_OPTION:
-      case LTU_OPTION:
-      case LEU_OPTION:
-      case GTU_OPTION:
-      case GEU_OPTION:
-      case EQ_OPTION:
-      case NE_OPTION:
-      case LT_OPTION:
-      case LE_OPTION:
-      case GT_OPTION:
-      case GE_OPTION:
-      case LS_OPTION:
-      case HS_OPTION:
-         m_compare_op = last_ptx_inst_option;
-         break;
-      case NUM_OPTION:
-      case NAN_OPTION:
-    	  m_compare_op = last_ptx_inst_option;
-        // assert(0); // finish this
-         break;
-      case SAT_OPTION:
-         m_saturation_mode = 1;
-         break;
-      case RNI_OPTION:
-      case RZI_OPTION:
-      case RMI_OPTION:
-      case RPI_OPTION:
-      case RN_OPTION:
-      case RZ_OPTION:
-      case RM_OPTION:
-      case RP_OPTION:
-         m_rounding_mode = last_ptx_inst_option;
-         break;
-      case HI_OPTION:
-         m_compare_op = last_ptx_inst_option;
-         m_hi = true;
-         assert( !m_lo ); 
-         assert( !m_wide );
-         break;
-      case LO_OPTION:
-         m_compare_op = last_ptx_inst_option;
-         m_lo = true;
-         assert( !m_hi );
-         assert( !m_wide );
-         break;
-      case WIDE_OPTION:
-         m_wide = true;
-         assert( !m_lo ); 
-         assert( !m_hi ); 
-         break;
-      case UNI_OPTION:
-         m_uni = true; // don't care... < now we DO care when constructing flowgraph>
-         break;
-      case GEOM_MODIFIER_1D:
-      case GEOM_MODIFIER_2D:
-      case GEOM_MODIFIER_3D:
-         m_geom_spec = last_ptx_inst_option;
-         break;
-      case V2_TYPE:
-      case V3_TYPE:
-      case V4_TYPE:
-         m_vector_spec = last_ptx_inst_option;
-         break;
-      case ATOMIC_AND:
-      case ATOMIC_OR:
-      case ATOMIC_XOR:
-      case ATOMIC_CAS:
-      case ATOMIC_EXCH:
-      case ATOMIC_ADD:
-      case ATOMIC_INC:
-      case ATOMIC_DEC:
-      case ATOMIC_MIN:
-      case ATOMIC_MAX:
-         m_atomic_spec = last_ptx_inst_option;
-         break;
-      case APPROX_OPTION:
-         break;
-      case FULL_OPTION:
-         break;
-      case ANY_OPTION:
-         m_vote_mode = vote_any;
-         break;
-      case ALL_OPTION:
-         m_vote_mode = vote_all;
-         break;
-      case GLOBAL_OPTION:
-         m_membar_level = GLOBAL_OPTION;
-         break;
-      case CTA_OPTION:
-         m_membar_level = CTA_OPTION;
-         break;
-      case FTZ_OPTION:
-         break;
-      case TO_OPTION:
-         m_to_option = true;
-         break;
-      case CA_OPTION: case CG_OPTION: case CS_OPTION: case LU_OPTION: case CV_OPTION:
-         m_cache_option = last_ptx_inst_option;
-         break;
-
-      default:
-         assert(0);
-         break;
-      }
-   }
-   m_scalar_type = scalar_type;
-   m_space_spec = space_spec;
-   if( ( opcode == ST_OP || opcode == LD_OP ) && (space_spec == undefined_space) ) {
-      m_space_spec = generic_space;
-   }
-   m_source_file = file?file:"<unknown>";
-   m_source_line = line;
-   m_source = source;
-}
-
-void ptx_instruction::print_insn() const
-{
-   print_insn(stdout);
-   fflush(stdout);
-}
-
-void ptx_instruction::print_insn( FILE *fp ) const
-{
-   char buf[1024], *p;
-   snprintf(buf,1024,"%s", m_source.c_str());
-   p = strtok(buf,";");
-   if( !is_label() ) 
-      fprintf(fp,"PC=%3u [idx=%3u] ", m_PC, m_instr_mem_index );
-   else
-      fprintf(fp,"                " );
-   fprintf(fp,"(%s:%u) %s", m_source_file.c_str(), m_source_line, p );
-}
-
 unsigned g_assemble_code_next_pc=1; 
 std::map<unsigned,function_info*> g_pc_to_finfo;
 std::vector<ptx_instruction*> function_info::s_g_pc_to_insn;
 
-unsigned function_info::sm_next_uid = 1;
-
-function_info::function_info(int entry_point ) 
-{
-   m_uid = sm_next_uid++;
-   m_entry_point = (entry_point==1)?true:false;
-   m_extern = (entry_point==2)?true:false;
-   num_reconvergence_pairs = 0;
-   m_symtab = NULL;
-   m_assembled = false;
-   m_return_var_sym = NULL; 
-   m_kernel_info.cmem = 0;
-   m_kernel_info.lmem = 0;
-   m_kernel_info.regs = 0;
-   m_kernel_info.smem = 0;
-   m_local_mem_framesize = 0;
-}
-
-void function_info::print_insn( unsigned pc, FILE * fp ) const
-{
-   unsigned index = pc - m_start_PC;
-   fprintf(fp,"FUNC[%s]",m_name.c_str() );
-   if ( index >= m_instr_mem_size ) {
-      fprintf(fp, "<past last instruction (max pc=%u)>", m_start_PC + m_instr_mem_size - 1 );
-   } else {
-      if ( m_instr_mem[index] != NULL )
-         m_instr_mem[index]->print_insn(fp);
-      else
-         fprintf(fp, "<no instruction at pc = %u>", pc );
-   }
-}
-
 void function_info::ptx_assemble()
 {
    if( m_assembled ) {
@@ -528,7 +272,7 @@ void function_info::ptx_assemble()
          }
          unsigned index = labels[ target.name() ]; //determine address from name
          unsigned PC = m_instr_mem[index]->get_PC();
-         g_current_symbol_table->set_label_address( target.get_symbol(), PC );
+         m_symtab->set_label_address( target.get_symbol(), PC );
          target.set_type(label_t);
       }
    }
@@ -1036,7 +780,7 @@ void function_info::add_param_data( unsigned argn, struct gpgpu_ptx_sim_arg *arg
    }
 }
 
-void function_info::finalize( memory_space *param_mem, symbol_table *symtab  ) 
+void function_info::finalize( memory_space *param_mem ) 
 {
    unsigned param_address = 0;
    for( std::map<unsigned,param_info>::iterator i=m_ptx_kernel_param_info.begin(); i!=m_ptx_kernel_param_info.end(); i++ ) {
@@ -1045,7 +789,7 @@ void function_info::finalize( memory_space *param_mem, symbol_table *symtab  )
       int type = p.get_type();
       param_t param_value = p.get_value();
       param_value.type = type;
-      symbol *param = symtab->lookup(name.c_str());
+      symbol *param = m_symtab->lookup(name.c_str());
       unsigned xtype = param->type()->get_key().scalar_type();
       assert(xtype==(unsigned)type);
       size_t size;
@@ -1069,14 +813,11 @@ void function_info::finalize( memory_space *param_mem, symbol_table *symtab  )
 
 void function_info::list_param( FILE *fout ) const
 {
-   symbol_table *symtab = g_current_symbol_table;
    for( std::map<unsigned,param_info>::const_iterator i=m_ptx_kernel_param_info.begin(); i!=m_ptx_kernel_param_info.end(); i++ ) {
       const param_info &p = i->second;
       std::string name = p.get_name();
-      symbol *param = symtab->lookup(name.c_str());
-
+      symbol *param = m_symtab->lookup(name.c_str());
       addr_t param_addr = param->get_address();
-
       fprintf(fout, "%s: %#08x\n", name.c_str(), param_addr);
    }
    fflush(fout);
@@ -1146,7 +887,6 @@ void function_info::ptx_exec_inst( ptx_thread_info *thread,
    ptx_instruction *pI = m_instr_mem[index];
    try {
 
-   g_current_symbol_table = thread->get_finfo()->get_symtab();
    thread->clearRPC();
    thread->m_last_set_operand_value.u64 = 0;
 
@@ -1415,7 +1155,7 @@ unsigned ptx_sim_init_thread( ptx_thread_info** thread_info,int sid,unsigned tid
                local_mem = new memory_space_impl<32>(buf,32);
                local_mem_lookup[new_tid] = local_mem;
             }
-            thd->set_info(g_entrypoint_symbol_table,g_entrypoint_func_info);
+            thd->set_info(g_entrypoint_func_info);
             thd->set_nctaid(g_cudaGridDim.x,g_cudaGridDim.y,g_cudaGridDim.z);
             thd->set_ntid(g_cudaBlockDim.x,g_cudaBlockDim.y,g_cudaBlockDim.z);
             thd->set_ctaid(g_gx,g_gy,g_gz);
@@ -1510,20 +1250,20 @@ void gpgpu_ptx_sim_init_grid( const char *kernel_key, struct gpgpu_ptx_sim_arg*
       }
       printf("\n");
       abort();
-   } else {
-      std::string kname = (*g_host_to_kernel_entrypoint_name_lookup)[kernel_key];
-      printf("GPGPU-Sim PTX: Launching kernel \'%s\' gridDim= (%u,%u,%u) blockDim = (%u,%u,%u); ntuid=%u\n",
-             kname.c_str(), g_cudaGridDim.x,g_cudaGridDim.y,g_cudaGridDim.z,g_cudaBlockDim.x,g_cudaBlockDim.y,g_cudaBlockDim.z, 
-             g_ptx_thread_info_uid_next );
+   } 
 
-      if ( g_kernel_name_to_function_lookup->find(kname) ==
-           g_kernel_name_to_function_lookup->end() ) {
-         printf("GPGPU-Sim PTX: ERROR ** function \'%s\' not found in ptx file\n", kname.c_str() );
-         abort();
-      }
-      g_entrypoint_func_info = g_func_info = (*g_kernel_name_to_function_lookup)[kname];
-      g_entrypoint_symbol_table = g_current_symbol_table = g_kernel_name_to_symtab_lookup[kname];
+   std::string kname = (*g_host_to_kernel_entrypoint_name_lookup)[kernel_key];
+   printf("GPGPU-Sim PTX: Launching kernel \'%s\' gridDim= (%u,%u,%u) blockDim = (%u,%u,%u); ntuid=%u\n",
+          kname.c_str(), g_cudaGridDim.x,g_cudaGridDim.y,g_cudaGridDim.z,g_cudaBlockDim.x,g_cudaBlockDim.y,g_cudaBlockDim.z, 
+          g_ptx_thread_info_uid_next );
+
+   if ( g_kernel_name_to_function_lookup->find(kname) ==
+        g_kernel_name_to_function_lookup->end() ) {
+      printf("GPGPU-Sim PTX: ERROR ** function \'%s\' not found in ptx file\n", kname.c_str() );
+      abort();
    }
+   g_entrypoint_func_info = g_func_info = (*g_kernel_name_to_function_lookup)[kname];
+   g_entrypoint_symbol_table = g_kernel_name_to_symtab_lookup[kname];
 
    unsigned argcount=0;
    struct gpgpu_ptx_sim_arg *tmparg = args;
@@ -1538,7 +1278,7 @@ void gpgpu_ptx_sim_init_grid( const char *kernel_key, struct gpgpu_ptx_sim_arg*
       args = args->m_next;
       argn++;
    }
-   g_func_info->finalize(g_param_mem,g_current_symbol_table);
+   g_func_info->finalize(g_param_mem);
    g_ptx_kernel_count++; 
    if ( gpgpu_ptx_instruction_classification ) {
       init_inst_classification_stat();
@@ -1661,11 +1401,10 @@ void gpgpu_ptx_sim_memcpy_symbol(const char *hostVar, const void *src, size_t co
    memory_space *mem = NULL;
 
    std::map<std::string,symbol_table*>::iterator st = g_sym_name_to_symbol_table.find(sym_name.c_str());
-   if( st != g_sym_name_to_symbol_table.end() ) {
-      g_current_symbol_table = st->second;
-   }
+   assert( st != g_sym_name_to_symbol_table.end() );
+   symbol_table *symtab = st->second;
 
-   symbol *sym = g_current_symbol_table->lookup(sym_name.c_str());
+   symbol *sym = symtab->lookup(sym_name.c_str());
    assert(sym);
    unsigned dst = sym->get_address() + offset; 
    switch (mem_region.get_type()) {
@@ -1931,15 +1670,13 @@ void gpgpu_ptx_assemble( std::string kname, void *kinfo )
        printf("GPGPU-Sim PTX: skipping assembly for extern declared function \'%s\'\n", func_info->get_name().c_str() );
        return;
     }
-    g_func_info = func_info;
-    g_current_symbol_table = g_kernel_name_to_symtab_lookup[ kname ];
-    if( g_current_symbol_table == NULL ) {
+    if( g_kernel_name_to_symtab_lookup[ kname ] == NULL ) {
        printf("\nGPGPU-Sim PTX: ERROR no information for kernel \'%s\'\n"
               "               this can happen for kernels contained in CUDA\n"
               "               libraries (such as CUBLAS, CUFFT, or CUDPP)\n", kname.c_str() );
        exit(1);
     }
-
+    g_func_info = func_info;
     func_info->ptx_assemble();
 }
 
@@ -2075,7 +1812,7 @@ void gpgpu_ptx_sim_main_func( const char *kernel_key, dim3 gridDim, dim3 blockDi
                      }
 
 
-                     thd->set_info(g_current_symbol_table,g_func_info);
+                     thd->set_info(g_func_info);
                      thd->set_nctaid(gridDim.x,gridDim.y,gridDim.z);
                      thd->set_ntid(blockDim.x, blockDim.y, blockDim.z);
                      thd->set_ctaid(gx,gy,gz);
diff --git a/src/cuda-sim/ptx_ir.cc b/src/cuda-sim/ptx_ir.cc
index b2c8552..d049a6a 100644
--- a/src/cuda-sim/ptx_ir.cc
+++ b/src/cuda-sim/ptx_ir.cc
@@ -78,7 +78,7 @@ extern "C" int ptx_error( const char *s );
 
 // the program intermediate representation...
 symbol_table *g_global_symbol_table = NULL;
-symbol_table *g_current_symbol_table = NULL;
+static symbol_table *g_current_symbol_table = NULL;
 std::list<ptx_instruction*> g_instructions;
 symbol *g_last_symbol = NULL;
 std::map<std::string,symbol_table*> g_sym_name_to_symbol_table;
@@ -1424,3 +1424,259 @@ void copy_buffer_list_into_frame(ptx_thread_info * thread, arg_buffer_list_t &ar
    }
 }
 
+
+
+static std::list<operand_info> check_operands( int opcode,
+                                        const std::list<int> &scalar_type,
+                                        const std::list<operand_info> &operands )
+{
+   static int g_warn_literal_operands_two_type_inst;
+    if( (opcode == CVT_OP) || (opcode == SET_OP) || (opcode == SLCT_OP) || (opcode == TEX_OP) ) {
+        // just make sure these do not have have const operands... 
+        if( !g_warn_literal_operands_two_type_inst ) {
+            std::list<operand_info>::const_iterator o;
+            for( o = operands.begin(); o != operands.end(); o++ ) {
+                const operand_info &op = *o;
+                if( op.is_literal() ) {
+                    printf("GPGPU-Sim PTX: PTX uses two scalar type intruction with literal operand.\n");
+                    g_warn_literal_operands_two_type_inst = 1;
+                }
+            }
+        }
+    } else {
+        assert( scalar_type.size() < 2 );
+        if( scalar_type.size() == 1 ) {
+            std::list<operand_info> result;
+            int inst_type = scalar_type.front();
+            std::list<operand_info>::const_iterator o;
+            for( o = operands.begin(); o != operands.end(); o++ ) {
+                const operand_info &op = *o;
+                if( op.is_literal() ) {
+                    if( (op.get_type() == double_op_t) && (inst_type == F32_TYPE) ) {
+                        ptx_reg_t v = op.get_literal_value();
+                        float u = (float)v.f64;
+                        operand_info n(u);
+                        result.push_back(n);
+                    } else {
+                        result.push_back(op);
+                    }
+                } else {
+                        result.push_back(op);
+                }
+            }
+            return result;
+        } 
+    }
+    return operands;
+}
+
+
+ptx_instruction::ptx_instruction( int opcode, 
+                                  const symbol *pred, 
+                                  int neg_pred, 
+                                  symbol *label,
+                                  const std::list<operand_info> &operands, 
+                                  const operand_info &return_var,
+                                  const std::list<int> &options, 
+                                  const std::list<int> &scalar_type,
+                                  memory_space_t space_spec,
+                                  const char *file, 
+                                  unsigned line,
+                                  const char *source ) 
+{
+   m_uid = ++g_num_ptx_inst_uid;
+   m_PC = 0;
+   m_opcode = opcode;
+   m_pred = pred;
+   m_neg_pred = neg_pred;
+   m_label = label;
+   const std::list<operand_info> checked_operands = check_operands(opcode,scalar_type,operands);
+   m_operands.insert(m_operands.begin(), checked_operands.begin(), checked_operands.end() );
+   m_return_var = return_var;
+   m_options = options;
+   m_wide = false;
+   m_hi = false;
+   m_lo = false;
+   m_uni = false;
+   m_to_option = false;
+   m_cache_option = 0;
+   m_rounding_mode = RN_OPTION;
+   m_compare_op = -1;
+   m_saturation_mode = 0;
+   m_geom_spec = 0;
+   m_vector_spec = 0;
+   m_atomic_spec = 0;
+   m_warp_size = ::warp_size;
+   m_membar_level = 0;
+
+   std::list<int>::const_iterator i;
+   unsigned n=1;
+   for ( i=options.begin(); i!= options.end(); i++, n++ ) {
+      int last_ptx_inst_option = *i;
+      switch ( last_ptx_inst_option ) {
+      case EQU_OPTION:
+      case NEU_OPTION:
+      case LTU_OPTION:
+      case LEU_OPTION:
+      case GTU_OPTION:
+      case GEU_OPTION:
+      case EQ_OPTION:
+      case NE_OPTION:
+      case LT_OPTION:
+      case LE_OPTION:
+      case GT_OPTION:
+      case GE_OPTION:
+      case LS_OPTION:
+      case HS_OPTION:
+         m_compare_op = last_ptx_inst_option;
+         break;
+      case NUM_OPTION:
+      case NAN_OPTION:
+    	  m_compare_op = last_ptx_inst_option;
+        // assert(0); // finish this
+         break;
+      case SAT_OPTION:
+         m_saturation_mode = 1;
+         break;
+      case RNI_OPTION:
+      case RZI_OPTION:
+      case RMI_OPTION:
+      case RPI_OPTION:
+      case RN_OPTION:
+      case RZ_OPTION:
+      case RM_OPTION:
+      case RP_OPTION:
+         m_rounding_mode = last_ptx_inst_option;
+         break;
+      case HI_OPTION:
+         m_compare_op = last_ptx_inst_option;
+         m_hi = true;
+         assert( !m_lo ); 
+         assert( !m_wide );
+         break;
+      case LO_OPTION:
+         m_compare_op = last_ptx_inst_option;
+         m_lo = true;
+         assert( !m_hi );
+         assert( !m_wide );
+         break;
+      case WIDE_OPTION:
+         m_wide = true;
+         assert( !m_lo ); 
+         assert( !m_hi ); 
+         break;
+      case UNI_OPTION:
+         m_uni = true; // don't care... < now we DO care when constructing flowgraph>
+         break;
+      case GEOM_MODIFIER_1D:
+      case GEOM_MODIFIER_2D:
+      case GEOM_MODIFIER_3D:
+         m_geom_spec = last_ptx_inst_option;
+         break;
+      case V2_TYPE:
+      case V3_TYPE:
+      case V4_TYPE:
+         m_vector_spec = last_ptx_inst_option;
+         break;
+      case ATOMIC_AND:
+      case ATOMIC_OR:
+      case ATOMIC_XOR:
+      case ATOMIC_CAS:
+      case ATOMIC_EXCH:
+      case ATOMIC_ADD:
+      case ATOMIC_INC:
+      case ATOMIC_DEC:
+      case ATOMIC_MIN:
+      case ATOMIC_MAX:
+         m_atomic_spec = last_ptx_inst_option;
+         break;
+      case APPROX_OPTION:
+         break;
+      case FULL_OPTION:
+         break;
+      case ANY_OPTION:
+         m_vote_mode = vote_any;
+         break;
+      case ALL_OPTION:
+         m_vote_mode = vote_all;
+         break;
+      case GLOBAL_OPTION:
+         m_membar_level = GLOBAL_OPTION;
+         break;
+      case CTA_OPTION:
+         m_membar_level = CTA_OPTION;
+         break;
+      case FTZ_OPTION:
+         break;
+      case TO_OPTION:
+         m_to_option = true;
+         break;
+      case CA_OPTION: case CG_OPTION: case CS_OPTION: case LU_OPTION: case CV_OPTION:
+         m_cache_option = last_ptx_inst_option;
+         break;
+
+      default:
+         assert(0);
+         break;
+      }
+   }
+   m_scalar_type = scalar_type;
+   m_space_spec = space_spec;
+   if( ( opcode == ST_OP || opcode == LD_OP ) && (space_spec == undefined_space) ) {
+      m_space_spec = generic_space;
+   }
+   m_source_file = file?file:"<unknown>";
+   m_source_line = line;
+   m_source = source;
+}
+
+void ptx_instruction::print_insn() const
+{
+   print_insn(stdout);
+   fflush(stdout);
+}
+
+void ptx_instruction::print_insn( FILE *fp ) const
+{
+   char buf[1024], *p;
+   snprintf(buf,1024,"%s", m_source.c_str());
+   p = strtok(buf,";");
+   if( !is_label() ) 
+      fprintf(fp,"PC=%3u [idx=%3u] ", m_PC, m_instr_mem_index );
+   else
+      fprintf(fp,"                " );
+   fprintf(fp,"(%s:%u) %s", m_source_file.c_str(), m_source_line, p );
+}
+
+unsigned function_info::sm_next_uid = 1;
+
+function_info::function_info(int entry_point ) 
+{
+   m_uid = sm_next_uid++;
+   m_entry_point = (entry_point==1)?true:false;
+   m_extern = (entry_point==2)?true:false;
+   num_reconvergence_pairs = 0;
+   m_symtab = NULL;
+   m_assembled = false;
+   m_return_var_sym = NULL; 
+   m_kernel_info.cmem = 0;
+   m_kernel_info.lmem = 0;
+   m_kernel_info.regs = 0;
+   m_kernel_info.smem = 0;
+   m_local_mem_framesize = 0;
+}
+
+void function_info::print_insn( unsigned pc, FILE * fp ) const
+{
+   unsigned index = pc - m_start_PC;
+   fprintf(fp,"FUNC[%s]",m_name.c_str() );
+   if ( index >= m_instr_mem_size ) {
+      fprintf(fp, "<past last instruction (max pc=%u)>", m_start_PC + m_instr_mem_size - 1 );
+   } else {
+      if ( m_instr_mem[index] != NULL )
+         m_instr_mem[index]->print_insn(fp);
+      else
+         fprintf(fp, "<no instruction at pc = %u>", pc );
+   }
+}
+
diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h
index 1daf9f1..839cceb 100644
--- a/src/cuda-sim/ptx_ir.h
+++ b/src/cuda-sim/ptx_ir.h
@@ -1010,7 +1010,7 @@ public:
        return m_start_PC;
    }
 
-   void finalize( memory_space *param_mem, symbol_table *symtab );
+   void finalize( memory_space *param_mem );
    void list_param( FILE *fout ) const;
 
    const struct gpgpu_ptx_sim_kernel_info* get_kernel_info () {
@@ -1259,7 +1259,6 @@ extern function_info *g_func_info;
 extern int g_error_detected;
 extern bool g_debug_ir_generation;
 extern std::list<ptx_instruction*> g_instructions;
-extern symbol_table *g_current_symbol_table;
 extern symbol_table *g_entrypoint_symbol_table;
 extern function_info *g_entrypoint_func_info;
 extern symbol_table *g_global_symbol_table;
diff --git a/src/cuda-sim/ptx_sim.cc b/src/cuda-sim/ptx_sim.cc
index 9a2acf6..31c8c81 100644
--- a/src/cuda-sim/ptx_sim.cc
+++ b/src/cuda-sim/ptx_sim.cc
@@ -301,9 +301,9 @@ unsigned ptx_thread_info::get_builtin( int builtin_id, unsigned dim_mod )
    return 0;
 }
 
-void ptx_thread_info::set_info( symbol_table *symtab, function_info *func ) 
+void ptx_thread_info::set_info( function_info *func ) 
 {
-  m_symbol_table = symtab;
+  m_symbol_table = func->get_symtab();
   m_func_info = func;
   m_PC = func->get_start_PC();
 }
@@ -313,9 +313,9 @@ void ptx_thread_info::print_insn( unsigned pc, FILE * fp ) const
    m_func_info->print_insn(pc,fp);
 }
 
-static void print_reg( std::string name, ptx_reg_t value )
+static void print_reg( std::string name, ptx_reg_t value, symbol_table *symtab )
 {
-   const symbol *sym = g_current_symbol_table->lookup(name.c_str());
+   const symbol *sym = symtab->lookup(name.c_str());
    printf("  %8s   ", name.c_str() );
    if( sym == NULL ) {
       printf("<unknown type> 0x%llx\n", (unsigned long long ) value.u64 );
@@ -451,8 +451,7 @@ void ptx_thread_info::dump_regs()
    for ( r=m_regs.back().begin(); r != m_regs.back().end(); ++r ) {
       std::string name = r->first->name();
       ptx_reg_t value = r->second;
-      print_reg(name,value);
-
+      print_reg(name,value,m_symbol_table);
    }
 }
 
@@ -465,7 +464,7 @@ void ptx_thread_info::dump_modifiedregs()
    for ( r=m_debug_trace_regs_modified.begin(); r != m_debug_trace_regs_modified.end(); ++r ) {
       std::string name = r->first->name();
       ptx_reg_t value = r->second;
-      print_reg(name,value);
+      print_reg(name,value,m_symbol_table);
    }
 }
 
diff --git a/src/cuda-sim/ptx_sim.h b/src/cuda-sim/ptx_sim.h
index 12075c8..0379686 100644
--- a/src/cuda-sim/ptx_sim.h
+++ b/src/cuda-sim/ptx_sim.h
@@ -270,7 +270,7 @@ public:
       return m_func_info; 
    }
    void print_insn( unsigned pc, FILE * fp ) const;
-   void set_info( symbol_table *symtab, function_info *func );
+   void set_info( function_info *func );
    unsigned get_uid() const
    {
       return m_uid;