Merge pull request #23 from echoedit/dev

Dev

17 files changed, 100 insertions, 86 deletions

diff --git a/libcuda/cuda_api_object.h b/libcuda/cuda_api_object.h
index 0054697..db5e6a4 100644
--- a/libcuda/cuda_api_object.h
+++ b/libcuda/cuda_api_object.h
@@ -169,9 +169,10 @@ private:
 
 class cuda_runtime_api {
     public:
-	cuda_runtime_api() {
+	cuda_runtime_api( gpgpu_context* ctx ) {
 	    g_glbmap = NULL;
 	    g_active_device = 0; //active gpu that runs the code
+	    gpgpu_ctx = ctx;
 	}
 	// global list
 	std::list<cuobjdumpSection*> cuobjdumpSectionList;
@@ -187,6 +188,8 @@ class cuda_runtime_api {
 	std::map<void *, size_t> pinned_memory_size;
 	glbmap_entry_t* g_glbmap;
 	int g_active_device; //active gpu that runs the code
+	// backward pointer
+	class gpgpu_context* gpgpu_ctx;
 	// member function list
 	void cuobjdumpInit();
 	void extract_code_using_cuobjdump();
diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index eb645ce..bbbaf23 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -767,9 +767,9 @@ void cudaRegisterVarInternal(
 		ctx->cuobjdumpParseBinary((unsigned)(unsigned long long)fatCubinHandle);
 	fflush(stdout);
 	if ( constant && !global && !ext ) {
-		gpgpu_ptx_sim_register_const_variable(hostVar,deviceName,size);
+		ctx->func_sim->gpgpu_ptx_sim_register_const_variable(hostVar,deviceName,size);
 	} else if ( !constant && !global && !ext ) {
-		gpgpu_ptx_sim_register_global_variable(hostVar,deviceName,size);
+		ctx->func_sim->gpgpu_ptx_sim_register_global_variable(hostVar,deviceName,size);
 	} else cuda_not_implemented(__my_func__,__LINE__);
 }
 
@@ -2936,7 +2936,7 @@ void gpgpu_context::cuobjdumpParseBinary(unsigned int handle){
 	if(context->get_device()->get_gpgpu()->get_config().convert_to_ptxplus() ) {
 		cuobjdumpELFSection* elfsection = api->findELFSection(ptx->getIdentifier());
 		assert (elfsection!= NULL);
-		char *ptxplus_str = gpgpu_ptx_sim_convert_ptx_and_sass_to_ptxplus(
+		char *ptxplus_str = ptxinfo->gpgpu_ptx_sim_convert_ptx_and_sass_to_ptxplus(
 				ptx->getPTXfilename(),
 				elfsection->getELFfilename(),
 				elfsection->getSASSfilename());
@@ -3545,7 +3545,7 @@ kernel_info_t * cuda_runtime_api::gpgpu_cuda_ptx_sim_init_grid( const char *host
 	}
 
 	entry->finalize(result->get_param_memory());
-	g_ptx_kernel_count++;
+	gpgpu_ctx->func_sim->g_ptx_kernel_count++;
 	fflush(stdout);
 	
 	if(g_debug_execution >= 4){
diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h
index 2e21009..a2ae7b6 100644
--- a/libcuda/gpgpu_context.h
+++ b/libcuda/gpgpu_context.h
@@ -10,7 +10,7 @@ class gpgpu_context {
     public:
 	gpgpu_context() {
 	    g_global_allfiles_symbol_table = NULL;
-	    api = new cuda_runtime_api();
+	    api = new cuda_runtime_api(this);
 	    ptxinfo = new ptxinfo_data(this);
 	    ptx_parser = new ptx_recognizer(this);
 	    the_gpgpusim = new GPGPUsim_ctx(this);
diff --git a/libopencl/opencl_runtime_api.cc b/libopencl/opencl_runtime_api.cc
index 03ec80c..0a6eb3e 100644
--- a/libopencl/opencl_runtime_api.cc
+++ b/libopencl/opencl_runtime_api.cc
@@ -877,6 +877,8 @@ clEnqueueNDRangeKernel(cl_command_queue command_queue,
                        const cl_event * event_wait_list,
                        cl_event *       event) CL_API_SUFFIX__VERSION_1_0
 {
+    gpgpu_context *ctx;
+    ctx = GPGPU_Context();
    int _global_size[3];
    int zeros[3] = { 0, 0, 0};
    printf("\n\n\n");
@@ -954,9 +956,9 @@ clEnqueueNDRangeKernel(cl_command_queue command_queue,
 	   gpgpu_ptx_sim_memcpy_symbol( "%_global_launch_offset", zeros, 3 * sizeof(int), 0, 1, gpu );
 	   gpgpu_ptx_sim_memcpy_symbol( "%_global_block_offset", zeros, 3 * sizeof(int), 0, 1, gpu );
    }
-   kernel_info_t *grid = gpgpu_opencl_ptx_sim_init_grid(kernel->get_implementation(),params,GridDim,BlockDim,gpu);
+   kernel_info_t *grid = ctx->func_sim->gpgpu_opencl_ptx_sim_init_grid(kernel->get_implementation(),params,GridDim,BlockDim,gpu);
    if ( g_ptx_sim_mode )
-      gpgpu_opencl_ptx_sim_main_func( grid );
+      ctx->func_sim->gpgpu_opencl_ptx_sim_main_func( grid );
    else
       gpgpu_opencl_ptx_sim_main_perf( grid );
    return CL_SUCCESS;
diff --git a/src/abstract_hardware_model.cc b/src/abstract_hardware_model.cc
index 7755477..248e7a5 100644
--- a/src/abstract_hardware_model.cc
+++ b/src/abstract_hardware_model.cc
@@ -173,9 +173,10 @@ void gpgpu_functional_sim_config::ptx_set_tex_cache_linesize(unsigned linesize)
    m_texcache_linesize = linesize;
 }
 
-gpgpu_t::gpgpu_t( const gpgpu_functional_sim_config &config )
+gpgpu_t::gpgpu_t( const gpgpu_functional_sim_config &config, gpgpu_context* ctx )
     : m_function_model_config(config)
 {
+   gpgpu_ctx = ctx;
    m_global_mem = new memory_space_impl<8192>("global",64*1024);
    
    m_tex_mem = new memory_space_impl<8192>("tex",64*1024);
diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h
index 68cb693..da29a11 100644
--- a/src/abstract_hardware_model.h
+++ b/src/abstract_hardware_model.h
@@ -31,6 +31,8 @@
 // Forward declarations
 class gpgpu_sim;
 class kernel_info_t;
+class gpgpu_context;
+
 
 //Set a hard limit of 32 CTAs per shader [cuda only has 8]
 #define MAX_CTA_PER_SHADER 32
@@ -529,7 +531,9 @@ private:
 
 class gpgpu_t {
 public:
-    gpgpu_t( const gpgpu_functional_sim_config &config );
+    gpgpu_t( const gpgpu_functional_sim_config &config, gpgpu_context* ctx );
+    // backward pointer
+    class gpgpu_context* gpgpu_ctx;
     int checkpoint_option;
     int checkpoint_kernel;
     int checkpoint_CTA;
diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc
index df0bbd7..b3e2965 100644
--- a/src/cuda-sim/cuda-sim.cc
+++ b/src/cuda-sim/cuda-sim.cc
@@ -56,20 +56,13 @@ typedef void * yyscan_t;
 int gpgpu_ptx_instruction_classification;
 void ** g_inst_classification_stat = NULL;
 void ** g_inst_op_classification_stat= NULL;
-int g_ptx_kernel_count = -1; // used for classification stat collection purposes 
 int g_debug_execution = 0;
 int g_debug_thread_uid = 0;
 addr_t g_debug_pc = 0xBEEF1518;
 // Output debug information to file options
-int cp_count;
-int cp_cta_resume;
 
-unsigned g_ptx_sim_num_insn = 0;
 unsigned gpgpu_param_num_shaders = 0;
 
-char *opcode_latency_fp, *opcode_latency_dp,*opcode_latency_sfu,*opcode_latency_tensor;
-char *opcode_initiation_int, *opcode_initiation_fp, *opcode_initiation_dp,*opcode_initiation_sfu,*opcode_initiation_tensor;
-char *cdp_latency_str;
 unsigned cdp_latency[5];
 
 void cuda_sim::ptx_opcocde_latency_options (option_parser_t opp) {
@@ -671,30 +664,30 @@ void ptx_instruction::set_opcode_and_latency()
 	sscanf(gpgpu_ctx->func_sim->opcode_latency_int, "%u,%u,%u,%u,%u",
 			&int_latency[0],&int_latency[1],&int_latency[2],
 			&int_latency[3],&int_latency[4]);
-	sscanf(opcode_latency_fp, "%u,%u,%u,%u,%u",
+	sscanf(gpgpu_ctx->func_sim->opcode_latency_fp, "%u,%u,%u,%u,%u",
 			&fp_latency[0],&fp_latency[1],&fp_latency[2],
 			&fp_latency[3],&fp_latency[4]);
-	sscanf(opcode_latency_dp, "%u,%u,%u,%u,%u",
+	sscanf(gpgpu_ctx->func_sim->opcode_latency_dp, "%u,%u,%u,%u,%u",
 			&dp_latency[0],&dp_latency[1],&dp_latency[2],
 			&dp_latency[3],&dp_latency[4]);
-	sscanf(opcode_latency_sfu, "%u",
+	sscanf(gpgpu_ctx->func_sim->opcode_latency_sfu, "%u",
 			&sfu_latency);
-	sscanf(opcode_latency_tensor, "%u",
+	sscanf(gpgpu_ctx->func_sim->opcode_latency_tensor, "%u",
 			&tensor_latency);
-	sscanf(opcode_initiation_int, "%u,%u,%u,%u,%u",
+	sscanf(gpgpu_ctx->func_sim->opcode_initiation_int, "%u,%u,%u,%u,%u",
 			&int_init[0],&int_init[1],&int_init[2],
 			&int_init[3],&int_init[4]);
-	sscanf(opcode_initiation_fp, "%u,%u,%u,%u,%u",
+	sscanf(gpgpu_ctx->func_sim->opcode_initiation_fp, "%u,%u,%u,%u,%u",
 			&fp_init[0],&fp_init[1],&fp_init[2],
 			&fp_init[3],&fp_init[4]);
-	sscanf(opcode_initiation_dp, "%u,%u,%u,%u,%u",
+	sscanf(gpgpu_ctx->func_sim->opcode_initiation_dp, "%u,%u,%u,%u,%u",
 			&dp_init[0],&dp_init[1],&dp_init[2],
 			&dp_init[3],&dp_init[4]);
-	sscanf(opcode_initiation_sfu, "%u",
+	sscanf(gpgpu_ctx->func_sim->opcode_initiation_sfu, "%u",
 			&sfu_init);
-	sscanf(opcode_initiation_tensor, "%u",
+	sscanf(gpgpu_ctx->func_sim->opcode_initiation_tensor, "%u",
 			&tensor_init);
-	sscanf(cdp_latency_str, "%u,%u,%u,%u,%u",
+	sscanf(gpgpu_ctx->func_sim->cdp_latency_str, "%u,%u,%u,%u,%u",
 			&cdp_latency[0],&cdp_latency[1],&cdp_latency[2],
             &cdp_latency[3],&cdp_latency[4]);
 
@@ -1486,7 +1479,7 @@ bool ptx_debug_exec_dump_cond(int thd_uid, addr_t pc)
    return false;
 }
 
-void init_inst_classification_stat() 
+void cuda_sim::init_inst_classification_stat()
 {
    static std::set<unsigned> init;
    if( init.find(g_ptx_kernel_count) != init.end() ) 
@@ -1631,7 +1624,7 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id)
       dim3 ctaid = get_ctaid();
       dim3 tid = get_tid();
       printf("%u [thd=%u][i=%u] : ctaid=(%u,%u,%u) tid=(%u,%u,%u) icount=%u [pc=%u] (%s:%u - %s)  [0x%llx]\n", 
-             g_ptx_sim_num_insn, 
+             m_gpu->gpgpu_ctx->func_sim->g_ptx_sim_num_insn,
              get_uid(),
              pI->uid(), ctaid.x,ctaid.y,ctaid.z,tid.x,tid.y,tid.z,
              get_icount(),
@@ -1689,14 +1682,14 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id)
          dump_regs(stdout);
    }
    update_pc();
-   g_ptx_sim_num_insn++;
+   m_gpu->gpgpu_ctx->func_sim->g_ptx_sim_num_insn++;
    
    //not using it with functional simulation mode
    if(!(this->m_functionalSimulationMode))
        ptx_file_line_stats_add_exec_count(pI);
    
    if ( gpgpu_ptx_instruction_classification ) {
-      init_inst_classification_stat();
+      m_gpu->gpgpu_ctx->func_sim->init_inst_classification_stat();
       unsigned space_type=0;
       switch ( pI->get_space().get_type() ) {
       case global_space: space_type = 10; break;
@@ -1712,15 +1705,15 @@ void ptx_thread_info::ptx_exec_inst( warp_inst_t &inst, unsigned lane_id)
          space_type = 0 ;
          break;
       }
-      StatAddSample( g_inst_classification_stat[g_ptx_kernel_count],  op_classification);
-      if (space_type) StatAddSample( g_inst_classification_stat[g_ptx_kernel_count], ( int )space_type);
-      StatAddSample( g_inst_op_classification_stat[g_ptx_kernel_count], (int)  pI->get_opcode() );
+      StatAddSample( g_inst_classification_stat[m_gpu->gpgpu_ctx->func_sim->g_ptx_kernel_count],  op_classification);
+      if (space_type) StatAddSample( g_inst_classification_stat[m_gpu->gpgpu_ctx->func_sim->g_ptx_kernel_count], ( int )space_type);
+      StatAddSample( g_inst_op_classification_stat[m_gpu->gpgpu_ctx->func_sim->g_ptx_kernel_count], (int)  pI->get_opcode() );
    }
-   if ( (g_ptx_sim_num_insn % 100000) == 0 ) {
+   if ( (m_gpu->gpgpu_ctx->func_sim->g_ptx_sim_num_insn % 100000) == 0 ) {
       dim3 ctaid = get_ctaid();
       dim3 tid = get_tid();
       DPRINTF(LIVENESS, "GPGPU-Sim PTX: %u instructions simulated : ctaid=(%u,%u,%u) tid=(%u,%u,%u)\n",
-             g_ptx_sim_num_insn, ctaid.x,ctaid.y,ctaid.z,tid.x,tid.y,tid.z );
+             m_gpu->gpgpu_ctx->func_sim->g_ptx_sim_num_insn, ctaid.x,ctaid.y,ctaid.z,tid.x,tid.y,tid.z );
       fflush(stdout);
    }
    
@@ -1923,7 +1916,7 @@ size_t get_kernel_code_size( class function_info *entry )
 }
 
 
-kernel_info_t *gpgpu_opencl_ptx_sim_init_grid(class function_info *entry,
+kernel_info_t *cuda_sim::gpgpu_opencl_ptx_sim_init_grid(class function_info *entry,
                                              gpgpu_ptx_sim_arg_list_t args, 
                                              struct dim3 gridDim,
                                              struct dim3 blockDim,
@@ -1955,18 +1948,16 @@ void print_splash()
    }
 }
 
-std::map<const void*,std::string>   g_const_name_lookup; // indexed by hostVar
-std::map<const void*,std::string>   g_global_name_lookup; // indexed by hostVar
 std::set<std::string>   g_globals;
 std::set<std::string>   g_constants;
 
-void gpgpu_ptx_sim_register_const_variable(void *hostVar, const char *deviceName, size_t size )
+void cuda_sim::gpgpu_ptx_sim_register_const_variable(void *hostVar, const char *deviceName, size_t size )
 {
    printf("GPGPU-Sim PTX registering constant %s (%zu bytes) to name mapping\n", deviceName, size );
    g_const_name_lookup[hostVar] = deviceName;
 }
 
-void gpgpu_ptx_sim_register_global_variable(void *hostVar, const char *deviceName, size_t size )
+void cuda_sim::gpgpu_ptx_sim_register_global_variable(void *hostVar, const char *deviceName, size_t size )
 {
    printf("GPGPU-Sim PTX registering global %s hostVar to name mapping\n", deviceName );
    g_global_name_lookup[hostVar] = deviceName;
@@ -1979,14 +1970,14 @@ void gpgpu_ptx_sim_memcpy_symbol(const char *hostVar, const void *src, size_t co
    memory_space_t mem_region = undefined_space;
    std::string sym_name;
 
-   std::map<const void*,std::string>::iterator c=g_const_name_lookup.find(hostVar);
-   if ( c!=g_const_name_lookup.end() ) {
+   std::map<const void*,std::string>::iterator c=gpu->gpgpu_ctx->func_sim->g_const_name_lookup.find(hostVar);
+   if ( c!=gpu->gpgpu_ctx->func_sim->g_const_name_lookup.end() ) {
       found_sym = true;
       sym_name = c->second;
       mem_region = const_space;
    }
-   std::map<const void*,std::string>::iterator g=g_global_name_lookup.find(hostVar);
-   if ( g!=g_global_name_lookup.end() ) {
+   std::map<const void*,std::string>::iterator g=gpu->gpgpu_ctx->func_sim->g_global_name_lookup.find(hostVar);
+   if ( g!=gpu->gpgpu_ctx->func_sim->g_global_name_lookup.end() ) {
       if ( found_sym ) {
          printf("Execution error: PTX symbol \"%s\" w/ hostVar=0x%Lx is declared both const and global?\n", 
                 sym_name.c_str(), (unsigned long long)hostVar );
@@ -2106,8 +2097,6 @@ void read_sim_environment_variables()
    }
 }
 
-ptx_cta_info *g_func_cta_info = NULL;
-
 #define MAX(a,b) (((a)>(b))?(a):(b))
 
 unsigned max_cta (const struct gpgpu_ptx_sim_info *kernel_info, unsigned threads_per_cta, unsigned int warp_size, unsigned int n_thread_per_shader, unsigned int gpgpu_shmem_size, unsigned int gpgpu_shader_registers, unsigned int max_cta_per_core)
@@ -2146,7 +2135,7 @@ unsigned max_cta (const struct gpgpu_ptx_sim_info *kernel_info, unsigned threads
 This function simulates the CUDA code functionally, it takes a kernel_info_t parameter 
 which holds the data for the CUDA kernel to be executed
 !*/
-void gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL )
+void cuda_sim::gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL )
 {
      printf("GPGPU-Sim: Performing Functional Simulation, executing kernel %s...\n",kernel.name().c_str());
 
@@ -2272,7 +2261,7 @@ void functionalCoreSim::initializeCTA(unsigned ctaid_cp)
         assert(m_thread[i]!=NULL && !m_thread[i]->is_done());
         char fname[2048];
         snprintf(fname,2048,"checkpoint_files/thread_%d_0_reg.txt",i );
-        if(cp_cta_resume==1)
+        if(m_gpu->gpgpu_ctx->func_sim->cp_cta_resume==1)
             m_thread[i]->resume_reg_thread(fname,symtab);
         ctaLiveThreads++;
     }
@@ -2296,7 +2285,7 @@ void  functionalCoreSim::createWarp(unsigned warpId)
    char fname[2048];
    snprintf(fname,2048,"checkpoint_files/warp_%d_0_simt.txt",warpId );
 
-   if(cp_cta_resume==1)
+   if(m_gpu->gpgpu_ctx->func_sim->cp_cta_resume==1)
    {
       unsigned pc,rpc;
       m_simt_stack[warpId]->resume(fname);
@@ -2312,8 +2301,8 @@ void  functionalCoreSim::createWarp(unsigned warpId)
 
 void functionalCoreSim::execute(int inst_count, unsigned ctaid_cp)
  {
-   cp_count= m_gpu->checkpoint_insn_Y;
-    cp_cta_resume= m_gpu->checkpoint_CTA_t;
+     m_gpu->gpgpu_ctx->func_sim->cp_count= m_gpu->checkpoint_insn_Y;
+     m_gpu->gpgpu_ctx->func_sim->cp_cta_resume= m_gpu->checkpoint_CTA_t;
     initializeCTA(ctaid_cp);
     
     int count=0;
@@ -2406,8 +2395,6 @@ unsigned translate_pc_to_ptxlineno(unsigned pc)
 
 extern std::map<unsigned,const char*> get_duplicate();
 
-int g_ptxinfo_error_detected;
-
 static char *g_ptxinfo_kname = NULL;
 static struct gpgpu_ptx_sim_info g_ptxinfo;
 static std::map<unsigned,const char*> g_duplicate;
diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h
index 96d34f5..25ebf7b 100644
--- a/src/cuda-sim/cuda-sim.h
+++ b/src/cuda-sim/cuda-sim.h
@@ -46,19 +46,8 @@ extern int g_debug_execution;
 extern int g_debug_thread_uid;
 extern void ** g_inst_classification_stat;
 extern void ** g_inst_op_classification_stat;
-extern int g_ptx_kernel_count; // used for classification stat collection purposes 
-extern char *opcode_latency_fp, *opcode_latency_dp,*opcode_latency_sfu,*opcode_latency_tensor;
 
-
-extern class kernel_info_t *gpgpu_opencl_ptx_sim_init_grid(class function_info *entry,
-                                            gpgpu_ptx_sim_arg_list_t args, 
-                                            struct dim3 gridDim, 
-                                            struct dim3 blockDim, 
-                                                          class gpgpu_t *gpu );
-extern void gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL = false );
 extern void   print_splash();
-extern void   gpgpu_ptx_sim_register_const_variable(void*, const char *deviceName, size_t size );
-extern void   gpgpu_ptx_sim_register_global_variable(void *hostVar, const char *deviceName, size_t size );
 extern void   gpgpu_ptx_sim_memcpy_symbol(const char *hostVar, const void *src, size_t count, size_t offset, int to, gpgpu_t *gpu );
 
 extern void read_sim_environment_variables();
@@ -135,10 +124,41 @@ struct gpgpu_ptx_sim_info get_ptxinfo();
 
 class cuda_sim {
     public:
+	cuda_sim() {
+	    g_ptx_sim_num_insn = 0;
+	    g_ptx_kernel_count = -1; // used for classification stat collection purposes
+	}
 	//global variables
 	char *opcode_latency_int;
+	char *opcode_latency_fp;
+	char *opcode_latency_dp;
+	char *opcode_latency_sfu;
+	char *opcode_latency_tensor;
+	char *opcode_initiation_int;
+	char *opcode_initiation_fp;
+	char *opcode_initiation_dp;
+	char *opcode_initiation_sfu;
+	char *opcode_initiation_tensor;
+	int cp_count;
+	int cp_cta_resume;
+	int g_ptxinfo_error_detected;
+	unsigned g_ptx_sim_num_insn;
+	char *cdp_latency_str;
+	int g_ptx_kernel_count; // used for classification stat collection purposes
+	std::map<const void*,std::string>   g_global_name_lookup; // indexed by hostVar
+	std::map<const void*,std::string>   g_const_name_lookup; // indexed by hostVar
 	//global functions
 	void ptx_opcocde_latency_options (option_parser_t opp);
+	void gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL = false );
+	int gpgpu_opencl_ptx_sim_main_func( kernel_info_t *grid );
+	void init_inst_classification_stat();
+	kernel_info_t *gpgpu_opencl_ptx_sim_init_grid(class function_info *entry,
+		gpgpu_ptx_sim_arg_list_t args,
+		struct dim3 gridDim,
+		struct dim3 blockDim,
+		gpgpu_t *gpu );
+	void   gpgpu_ptx_sim_register_global_variable(void *hostVar, const char *deviceName, size_t size );
+	void   gpgpu_ptx_sim_register_const_variable(void*, const char *deviceName, size_t size );
 };
 
 #endif
diff --git a/src/cuda-sim/ptx_loader.cc b/src/cuda-sim/ptx_loader.cc
index e0d9a11..2bf464c 100644
--- a/src/cuda-sim/ptx_loader.cc
+++ b/src/cuda-sim/ptx_loader.cc
@@ -58,7 +58,6 @@ extern int ptxinfo_lex_destroy(yyscan_t scanner);
 
 static bool g_save_embedded_ptx;
 static int g_occupancy_sm_number;
-bool m_ptx_save_converted_ptxplus;
 
 bool ptxinfo_data::keep_intermediate_files() {return g_keep_intermediate_files;}
 
@@ -71,7 +70,7 @@ void gpgpu_context::ptx_reg_options(option_parser_t opp)
                 "keep intermediate files created by GPGPU-Sim when interfacing with external programs",
                 "0");
    option_parser_register(opp, "-gpgpu_ptx_save_converted_ptxplus", OPT_BOOL,
-                &m_ptx_save_converted_ptxplus,
+                &(ptxinfo->m_ptx_save_converted_ptxplus),
                 "Saved converted ptxplus to a file",
                 "0");
    option_parser_register(opp, "-gpgpu_occupancy_sm_number", OPT_INT32, &g_occupancy_sm_number,
@@ -106,7 +105,7 @@ void gpgpu_context::print_ptx_file( const char *p, unsigned source_num, const ch
    fflush(stdout);
 }
 
-char* gpgpu_ptx_sim_convert_ptx_and_sass_to_ptxplus(const std::string ptxfilename, const std::string elffilename, const std::string sassfilename)
+char* ptxinfo_data::gpgpu_ptx_sim_convert_ptx_and_sass_to_ptxplus(const std::string ptxfilename, const std::string elffilename, const std::string sassfilename)
 {
 
 	printf("GPGPU-Sim PTX: converting EMBEDDED .ptx file to ptxplus \n");
diff --git a/src/cuda-sim/ptx_loader.h b/src/cuda-sim/ptx_loader.h
index aa3e2f3..decfdc8 100644
--- a/src/cuda-sim/ptx_loader.h
+++ b/src/cuda-sim/ptx_loader.h
@@ -43,14 +43,14 @@ class ptxinfo_data{
 	const char *g_ptxinfo_filename;
 	class gpgpu_context* gpgpu_ctx;
 	bool g_keep_intermediate_files;
+	bool m_ptx_save_converted_ptxplus;
 	void ptxinfo_addinfo();
 	bool keep_intermediate_files();
+	char* gpgpu_ptx_sim_convert_ptx_and_sass_to_ptxplus(const std::string ptx_str, const std::string sass_str, const std::string elf_str);
 };
 
 
 extern bool g_override_embedded_ptx;
 extern int no_of_ptx; //counter to track number of ptx files to be extracted in an application.
  
-char* gpgpu_ptx_sim_convert_ptx_and_sass_to_ptxplus(const std::string ptx_str, const std::string sass_str, const std::string elf_str);
-
 #endif
diff --git a/src/cuda-sim/ptx_parser.h b/src/cuda-sim/ptx_parser.h
index e09aab1..ef422e5 100644
--- a/src/cuda-sim/ptx_parser.h
+++ b/src/cuda-sim/ptx_parser.h
@@ -30,7 +30,6 @@
 
 #include "../abstract_hardware_model.h"
 #include "ptx_ir.h"
-extern int g_error_detected;
 
 #ifdef __cplusplus 
 const class ptx_instruction *ptx_instruction_lookup( const char *filename, unsigned linenumber );
diff --git a/src/cuda-sim/ptxinfo.l b/src/cuda-sim/ptxinfo.l
index 3a152b0..51371e3 100644
--- a/src/cuda-sim/ptxinfo.l
+++ b/src/cuda-sim/ptxinfo.l
@@ -88,13 +88,11 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 
 %%
 
-extern int g_ptxinfo_error_detected;
-
 int ptxinfo_error(yyscan_t yyscanner, ptxinfo_data* ptxinfo, const char* msg)
 {
     struct yyguts_t * yyg = (struct yyguts_t*)yyscanner;
 	int i;
-	g_ptxinfo_error_detected = 1;
+	ptxinfo->gpgpu_ctx->func_sim->g_ptxinfo_error_detected = 1;
 	fflush(stdout);
 	printf("GPGPU-Sim: ERROR while parsing output of ptxas (used to capture resource usage information)\n");
 	if( msg != NULL )
diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc
index 39acdd9..9f47067 100644
--- a/src/gpgpu-sim/gpu-sim.cc
+++ b/src/gpgpu-sim/gpu-sim.cc
@@ -699,7 +699,7 @@ void gpgpu_sim::stop_all_running_kernels(){
 }
 
 gpgpu_sim::gpgpu_sim( const gpgpu_sim_config &config, gpgpu_context* ctx )
-    : gpgpu_t(config), m_config(config)
+    : gpgpu_t(config, ctx), m_config(config)
 {
     gpgpu_ctx = ctx;
     m_shader_config = &m_config.m_shader_config;
@@ -1201,8 +1201,8 @@ void gpgpu_sim::gpu_print_stat()
       insn_warp_occ_print(stdout);
    }
    if ( gpgpu_ptx_instruction_classification ) {
-      StatDisp( g_inst_classification_stat[g_ptx_kernel_count]);
-      StatDisp( g_inst_op_classification_stat[g_ptx_kernel_count]);
+      StatDisp( g_inst_classification_stat[gpgpu_ctx->func_sim->g_ptx_kernel_count]);
+      StatDisp( g_inst_op_classification_stat[gpgpu_ctx->func_sim->g_ptx_kernel_count]);
    }
 
 #ifdef GPGPUSIM_POWER_MODEL
diff --git a/src/gpgpu-sim/gpu-sim.h b/src/gpgpu-sim/gpu-sim.h
index 7eeb7dd..119b934 100644
--- a/src/gpgpu-sim/gpu-sim.h
+++ b/src/gpgpu-sim/gpu-sim.h
@@ -62,6 +62,8 @@
 #define SAMPLELOG 222
 #define DUMPLOG 333
 
+class gpgpu_context;
+
 extern tr1_hash_map<new_addr_type,unsigned> address_random_interleaving;
 
 enum dram_ctrl_t {
@@ -494,6 +496,8 @@ public:
     */
     simt_core_cluster * getSIMTCluster();
 
+    // backward pointer
+    class gpgpu_context* gpgpu_ctx;
 
 private:
    // clocks
@@ -511,8 +515,6 @@ private:
    void gpgpu_debug();
 
 ///// data /////
-// backward pointer
-   class gpgpu_context* gpgpu_ctx;
 
    class simt_core_cluster **m_cluster;
    class memory_partition_unit **m_memory_partition_unit;
diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc
index 4d12068..69790fc 100644
--- a/src/gpgpu-sim/shader.cc
+++ b/src/gpgpu-sim/shader.cc
@@ -3022,15 +3022,15 @@ void shader_core_config::set_pipeline_latency() {
 			sscanf(gpgpu_ctx->func_sim->opcode_latency_int, "%u,%u,%u,%u,%u",
 					&int_latency[0],&int_latency[1],&int_latency[2],
 					&int_latency[3],&int_latency[4]);
-			sscanf(opcode_latency_fp, "%u,%u,%u,%u,%u",
+			sscanf(gpgpu_ctx->func_sim->opcode_latency_fp, "%u,%u,%u,%u,%u",
 					&fp_latency[0],&fp_latency[1],&fp_latency[2],
 					&fp_latency[3],&fp_latency[4]);
-			sscanf(opcode_latency_dp, "%u,%u,%u,%u,%u",
+			sscanf(gpgpu_ctx->func_sim->opcode_latency_dp, "%u,%u,%u,%u,%u",
 					&dp_latency[0],&dp_latency[1],&dp_latency[2],
 					&dp_latency[3],&dp_latency[4]);
-			sscanf(opcode_latency_sfu, "%u",
+			sscanf(gpgpu_ctx->func_sim->opcode_latency_sfu, "%u",
 					&sfu_latency);
-			sscanf(opcode_latency_tensor, "%u",
+			sscanf(gpgpu_ctx->func_sim->opcode_latency_tensor, "%u",
 					&tensor_latency);
 
 		//all div operation are executed on sfu
diff --git a/src/gpgpusim_entrypoint.cc b/src/gpgpusim_entrypoint.cc
index 5018305..d9d1023 100644
--- a/src/gpgpusim_entrypoint.cc
+++ b/src/gpgpusim_entrypoint.cc
@@ -134,7 +134,7 @@ void *gpgpu_sim_thread_concurrent(void*)
             if( GPGPUsim_ctx_ptr()->g_the_gpu->is_functional_sim()) {
                 kernel_info_t * kernel = GPGPUsim_ctx_ptr()->g_the_gpu->get_functional_kernel();
                 assert(kernel);
-                gpgpu_cuda_ptx_sim_main_func(*kernel);
+                GPGPUsim_ctx_ptr()->gpgpu_ctx->func_sim->gpgpu_cuda_ptx_sim_main_func(*kernel);
                 GPGPUsim_ctx_ptr()->g_the_gpu->finish_functional_sim(kernel);
             }
 
@@ -287,7 +287,7 @@ int gpgpu_opencl_ptx_sim_main_perf( kernel_info_t *grid )
 /*!
  * This function call the CUDA PTX functional simulator
  */
-int gpgpu_opencl_ptx_sim_main_func( kernel_info_t *grid )
+int cuda_sim::gpgpu_opencl_ptx_sim_main_func( kernel_info_t *grid )
 {
     //calling the CUDA PTX simulator, sending the kernel by reference and a flag set to true,
     //the flag used by the function to distinguish OpenCL calls from the CUDA simulation calls which
diff --git a/src/gpgpusim_entrypoint.h b/src/gpgpusim_entrypoint.h
index dfb82d0..887b3c8 100644
--- a/src/gpgpusim_entrypoint.h
+++ b/src/gpgpusim_entrypoint.h
@@ -83,6 +83,5 @@ struct GPGPUsim_ctx* GPGPUsim_ctx_ptr();
 class stream_manager* g_stream_manager();
 
 int gpgpu_opencl_ptx_sim_main_perf( kernel_info_t *grid );
-int gpgpu_opencl_ptx_sim_main_func( kernel_info_t *grid );
 
 #endif