Move g_globals and g_constants

Signed-off-by: Mengchi Zhang <[email protected]>

5 files changed, 14 insertions, 17 deletions

diff --git a/libopencl/opencl_runtime_api.cc b/libopencl/opencl_runtime_api.cc
index aaaec4f..7f029c7 100644
--- a/libopencl/opencl_runtime_api.cc
+++ b/libopencl/opencl_runtime_api.cc
@@ -950,11 +950,11 @@ clEnqueueNDRangeKernel(cl_command_queue command_queue,
 
    gpgpu_t *gpu = command_queue->get_device()->the_device();
    if (kernel->get_implementation()->get_ptx_version().ver() <3.0){
-	   gpgpu_ptx_sim_memcpy_symbol( "%_global_size", _global_size, 3 * sizeof(int), 0, 1, gpu );
-	   gpgpu_ptx_sim_memcpy_symbol( "%_work_dim", &work_dim, 1 * sizeof(int), 0, 1, gpu  );
-	   gpgpu_ptx_sim_memcpy_symbol( "%_global_num_groups", &GridDim, 3 * sizeof(int), 0, 1, gpu );
-	   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 );
+	   ctx->func_sim->gpgpu_ptx_sim_memcpy_symbol( "%_global_size", _global_size, 3 * sizeof(int), 0, 1, gpu );
+	   ctx->func_sim->gpgpu_ptx_sim_memcpy_symbol( "%_work_dim", &work_dim, 1 * sizeof(int), 0, 1, gpu  );
+	   ctx->func_sim->gpgpu_ptx_sim_memcpy_symbol( "%_global_num_groups", &GridDim, 3 * sizeof(int), 0, 1, gpu );
+	   ctx->func_sim->gpgpu_ptx_sim_memcpy_symbol( "%_global_launch_offset", zeros, 3 * sizeof(int), 0, 1, gpu );
+	   ctx->func_sim->gpgpu_ptx_sim_memcpy_symbol( "%_global_block_offset", zeros, 3 * sizeof(int), 0, 1, gpu );
    }
    kernel_info_t *grid = ctx->func_sim->gpgpu_opencl_ptx_sim_init_grid(kernel->get_implementation(),params,GridDim,BlockDim,gpu);
    if ( ctx->func_sim->g_ptx_sim_mode )
diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc
index 7e8aab9..3d1da62 100644
--- a/src/cuda-sim/cuda-sim.cc
+++ b/src/cuda-sim/cuda-sim.cc
@@ -1942,9 +1942,6 @@ void print_splash()
    }
 }
 
-std::set<std::string>   g_globals;
-std::set<std::string>   g_constants;
-
 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 );
@@ -1957,7 +1954,7 @@ void cuda_sim::gpgpu_ptx_sim_register_global_variable(void *hostVar, const char
    g_global_name_lookup[hostVar] = deviceName;
 }
 
-void gpgpu_ptx_sim_memcpy_symbol(const char *hostVar, const void *src, size_t count, size_t offset, int to, gpgpu_t *gpu )
+void cuda_sim::gpgpu_ptx_sim_memcpy_symbol(const char *hostVar, const void *src, size_t count, size_t offset, int to, gpgpu_t *gpu )
 {
    printf("GPGPU-Sim PTX: starting gpgpu_ptx_sim_memcpy_symbol with hostVar 0x%p\n", hostVar);
    bool found_sym = false;
diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h
index 628e434..5c95100 100644
--- a/src/cuda-sim/cuda-sim.h
+++ b/src/cuda-sim/cuda-sim.h
@@ -46,7 +46,6 @@ extern int g_debug_execution;
 extern int g_debug_thread_uid;
 
 extern void   print_splash();
-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 ptxinfo_opencl_addinfo( std::map<std::string,function_info*> &kernels );
 unsigned ptx_sim_init_thread( kernel_info_t &kernel,
@@ -161,6 +160,8 @@ class cuda_sim {
 	bool g_cuda_launch_blocking;
 	void ** g_inst_classification_stat;
 	void ** g_inst_op_classification_stat;
+	std::set<std::string>   g_globals;
+	std::set<std::string>   g_constants;
 	// backward pointer
 	class gpgpu_context* gpgpu_ctx;
 	//global functions
@@ -179,6 +180,7 @@ class cuda_sim {
 	void set_param_gpgpu_num_shaders(int num_shaders);
 	struct rec_pts find_reconvergence_points( function_info *finfo );
 	address_type get_converge_point( address_type pc );
+	void   gpgpu_ptx_sim_memcpy_symbol(const char *hostVar, const void *src, size_t count, size_t offset, int to, gpgpu_t *gpu );
 };
 
 #endif
diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc
index 0139534..269ec4d 100644
--- a/src/cuda-sim/ptx_parser.cc
+++ b/src/cuda-sim/ptx_parser.cc
@@ -322,9 +322,6 @@ bool ptx_recognizer::check_for_duplicates( const char *identifier )
    return ( s != NULL );
 }
 
-extern std::set<std::string>   g_globals;
-extern std::set<std::string>   g_constants;
-
 // Returns padding that needs to be inserted ahead of address to make it aligned to min(size, maxalign)
 /*
  * @param address the address in bytes
@@ -453,7 +450,7 @@ void ptx_recognizer::add_identifier( const char *identifier, int array_dim, unsi
          g_current_symbol_table->alloc_global( num_bits/8 + addr_pad ); 
       }
       if( g_current_symbol_table == g_global_symbol_table ) { 
-         g_constants.insert( identifier ); 
+         gpgpu_ctx->func_sim->g_constants.insert( identifier );
       }
       assert( g_current_symbol_table != NULL );
       g_sym_name_to_symbol_table[ identifier ] = g_current_symbol_table;
@@ -471,7 +468,7 @@ void ptx_recognizer::add_identifier( const char *identifier, int array_dim, unsi
       fflush(stdout);
       g_last_symbol->set_address( addr+addr_pad );
       g_current_symbol_table->alloc_global( num_bits/8 + addr_pad );
-      g_globals.insert( identifier );
+      gpgpu_ctx->func_sim->g_globals.insert( identifier );
       assert( g_current_symbol_table != NULL );
       g_sym_name_to_symbol_table[ identifier ] = g_current_symbol_table;
       break;
diff --git a/src/stream_manager.cc b/src/stream_manager.cc
index e07f4e4..be3dd71 100644
--- a/src/stream_manager.cc
+++ b/src/stream_manager.cc
@@ -29,6 +29,7 @@
 #include "gpgpusim_entrypoint.h"
 #include "cuda-sim/cuda-sim.h"
 #include "gpgpu-sim/gpu-sim.h"
+#include "../libcuda/gpgpu_context.h"
 
 unsigned CUstream_st::sm_next_stream_uid = 0;
 
@@ -150,13 +151,13 @@ bool stream_operation::do_operation( gpgpu_sim *gpu )
     case stream_memcpy_to_symbol:
         if(g_debug_execution >= 3)
             printf("memcpy to symbol\n");
-        gpgpu_ptx_sim_memcpy_symbol(m_symbol,m_host_address_src,m_cnt,m_offset,1,gpu);
+        gpu->gpgpu_ctx->func_sim->gpgpu_ptx_sim_memcpy_symbol(m_symbol,m_host_address_src,m_cnt,m_offset,1,gpu);
         m_stream->record_next_done();
         break;
     case stream_memcpy_from_symbol:
         if(g_debug_execution >= 3)
             printf("memcpy from symbol\n");
-        gpgpu_ptx_sim_memcpy_symbol(m_symbol,m_host_address_dst,m_cnt,m_offset,0,gpu);
+        gpu->gpgpu_ctx->func_sim->gpgpu_ptx_sim_memcpy_symbol(m_symbol,m_host_address_dst,m_cnt,m_offset,0,gpu);
         m_stream->record_next_done();
         break;
     case stream_kernel_launch: