Move g_ptx_sim_mode

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

9 files changed, 46 insertions, 22 deletions

diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index bbbaf23..59d2a60 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -1076,7 +1076,7 @@ cudaError_t cudaLaunchInternal( const char *hostFun, gpgpu_context* gpgpu_ctx =
 	CUctx_st* context = GPGPUSim_Context();
 	char *mode = getenv("PTX_SIM_MODE_FUNC");
 	if( mode )
-		sscanf(mode,"%u", &g_ptx_sim_mode);
+		sscanf(mode,"%u", &(ctx->func_sim->g_ptx_sim_mode));
 	gpgpusim_ptx_assert( !ctx->api->g_cuda_launch_stack.empty(), "empty launch stack" );
 	kernel_config config = ctx->api->g_cuda_launch_stack.back();
 	{
@@ -1092,7 +1092,7 @@ cudaError_t cudaLaunchInternal( const char *hostFun, gpgpu_context* gpgpu_ctx =
 	}
 	struct CUstream_st *stream = config.get_stream();
 	printf("\nGPGPU-Sim PTX: cudaLaunch for 0x%p (mode=%s) on stream %u\n", hostFun,
-			g_ptx_sim_mode?"functional simulation":"performance simulation", stream?stream->get_uid():0 );
+			(ctx->func_sim->g_ptx_sim_mode)?"functional simulation":"performance simulation", stream?stream->get_uid():0 );
 	kernel_info_t *grid = ctx->api->gpgpu_cuda_ptx_sim_init_grid(hostFun,config.get_args(),config.grid_dim(),config.block_dim(),context);
         //do dynamic PDOM analysis for performance simulation scenario
 	std::string kname = grid->name();
@@ -1143,7 +1143,7 @@ cudaError_t cudaLaunchInternal( const char *hostFun, gpgpu_context* gpgpu_ctx =
 	}
 	printf("GPGPU-Sim PTX: pushing kernel \'%s\' to stream %u, gridDim= (%u,%u,%u) blockDim = (%u,%u,%u) \n",
 			kname.c_str(), stream?stream->get_uid():0, gridDim.x,gridDim.y,gridDim.z,blockDim.x,blockDim.y,blockDim.z );
-	stream_operation op(grid,g_ptx_sim_mode,stream);
+	stream_operation op(grid,ctx->func_sim->g_ptx_sim_mode,stream);
 	g_stream_manager()->push(op);
 	ctx->api->g_cuda_launch_stack.pop_back();
 	return g_last_cudaError = cudaSuccess;
diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h
index a2ae7b6..3c9f87c 100644
--- a/libcuda/gpgpu_context.h
+++ b/libcuda/gpgpu_context.h
@@ -5,6 +5,7 @@
 #include "../src/cuda-sim/ptx_parser.h"
 #include "../src/gpgpusim_entrypoint.h"
 #include "../src/cuda-sim/cuda-sim.h"
+#include "../src/cuda-sim/cuda_device_runtime.h"
 
 class gpgpu_context {
     public:
@@ -14,7 +15,10 @@ class gpgpu_context {
 	    ptxinfo = new ptxinfo_data(this);
 	    ptx_parser = new ptx_recognizer(this);
 	    the_gpgpusim = new GPGPUsim_ctx(this);
-	    func_sim = new cuda_sim();
+	    func_sim = new cuda_sim(this);
+#if (CUDART_VERSION >= 5000)
+	    device_runtime = new cuda_device_runtime(this);
+#endif
 	}
 	// global list
 	symbol_table *g_global_allfiles_symbol_table;
@@ -25,6 +29,9 @@ class gpgpu_context {
 	ptx_recognizer* ptx_parser;
 	GPGPUsim_ctx* the_gpgpusim;
 	cuda_sim* func_sim;
+#if (CUDART_VERSION >= 5000)
+	cuda_device_runtime* device_runtime;
+#endif
 	// member function list
 	void cuobjdumpParseBinary(unsigned int handle);
 	class symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source_num );
diff --git a/libopencl/opencl_runtime_api.cc b/libopencl/opencl_runtime_api.cc
index 0a6eb3e..aaaec4f 100644
--- a/libopencl/opencl_runtime_api.cc
+++ b/libopencl/opencl_runtime_api.cc
@@ -884,9 +884,9 @@ clEnqueueNDRangeKernel(cl_command_queue command_queue,
    printf("\n\n\n");
    char *mode = getenv("PTX_SIM_MODE_FUNC");
    if ( mode )
-      sscanf(mode,"%u", &g_ptx_sim_mode);
+      sscanf(mode,"%u", &(ctx->func_sim->g_ptx_sim_mode));
    printf("GPGPU-Sim OpenCL API: clEnqueueNDRangeKernel '%s' (mode=%s)\n", kernel->name().c_str(),
-          g_ptx_sim_mode?"functional simulation":"performance simulation");
+          (ctx->func_sim->g_ptx_sim_mode)?"functional simulation":"performance simulation");
    if ( !work_dim || work_dim > 3 ) return CL_INVALID_WORK_DIMENSION;
    size_t _local_size[3];
    if( local_work_size != NULL ) {
@@ -957,7 +957,7 @@ clEnqueueNDRangeKernel(cl_command_queue command_queue,
 	   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 ( g_ptx_sim_mode )
+   if ( ctx->func_sim->g_ptx_sim_mode )
       ctx->func_sim->gpgpu_opencl_ptx_sim_main_func( grid );
    else
       gpgpu_opencl_ptx_sim_main_perf( grid );
diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc
index b3e2965..7a7d205 100644
--- a/src/cuda-sim/cuda-sim.cc
+++ b/src/cuda-sim/cuda-sim.cc
@@ -2033,13 +2033,11 @@ void gpgpu_ptx_sim_memcpy_symbol(const char *hostVar, const void *src, size_t co
    fflush(stdout);
 }
 
-int g_ptx_sim_mode; // if non-zero run functional simulation only (i.e., no notion of a clock cycle)
-
 extern int ptx_debug;
 
 bool g_cuda_launch_blocking = false;
 
-void read_sim_environment_variables() 
+void cuda_sim::read_sim_environment_variables()
 {
    ptx_debug = 0;
    g_debug_execution = 0;
@@ -2185,7 +2183,7 @@ void cuda_sim::gpgpu_cuda_ptx_sim_main_func( kernel_info_t &kernel, bool openCL
            cta.execute(cp_count,temp);
 
             #if (CUDART_VERSION >= 5000)
-            	launch_all_device_kernels();
+               gpgpu_ctx->device_runtime->launch_all_device_kernels();
             #endif
          }
          else
diff --git a/src/cuda-sim/cuda-sim.h b/src/cuda-sim/cuda-sim.h
index 25ebf7b..3c4336d 100644
--- a/src/cuda-sim/cuda-sim.h
+++ b/src/cuda-sim/cuda-sim.h
@@ -36,12 +36,12 @@
 #include <string>
 #include"ptx_sim.h"
 
+class gpgpu_context;
 class memory_space;
 class function_info;
 class symbol_table;
 
 extern const char *g_gpgpusim_version_string;
-extern int g_ptx_sim_mode;
 extern int g_debug_execution;
 extern int g_debug_thread_uid;
 extern void ** g_inst_classification_stat;
@@ -50,7 +50,6 @@ extern void ** g_inst_op_classification_stat;
 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 read_sim_environment_variables();
 extern void ptxinfo_opencl_addinfo( std::map<std::string,function_info*> &kernels );
 unsigned ptx_sim_init_thread( kernel_info_t &kernel,
                               class ptx_thread_info** thread_info,
@@ -124,9 +123,10 @@ struct gpgpu_ptx_sim_info get_ptxinfo();
 
 class cuda_sim {
     public:
-	cuda_sim() {
+	cuda_sim( gpgpu_context* ctx ) {
 	    g_ptx_sim_num_insn = 0;
 	    g_ptx_kernel_count = -1; // used for classification stat collection purposes
+	    gpgpu_ctx = ctx;
 	}
 	//global variables
 	char *opcode_latency_int;
@@ -147,6 +147,9 @@ class cuda_sim {
 	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
+	int g_ptx_sim_mode; // if non-zero run functional simulation only (i.e., no notion of a clock cycle)
+	// backward pointer
+	class gpgpu_context* gpgpu_ctx;
 	//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 );
@@ -159,6 +162,7 @@ class cuda_sim {
 		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 );
+	void read_sim_environment_variables();
 };
 
 #endif
diff --git a/src/cuda-sim/cuda_device_runtime.cc b/src/cuda-sim/cuda_device_runtime.cc
index be8369f..354fa79 100644
--- a/src/cuda-sim/cuda_device_runtime.cc
+++ b/src/cuda-sim/cuda_device_runtime.cc
@@ -20,6 +20,7 @@ unsigned long long g_max_total_param_size = 0;
 #include "../stream_manager.h"
 #include "../gpgpusim_entrypoint.h"
 #include "cuda_device_runtime.h"
+#include "../../libcuda/gpgpu_context.h"
 
 #define DEV_RUNTIME_REPORT(a) \
    if( g_debug_execution ) { \
@@ -318,17 +319,17 @@ void gpgpusim_cuda_streamCreateWithFlags(const ptx_instruction * pI, ptx_thread_
 }
 
 
-void launch_one_device_kernel() {
+void cuda_device_runtime::launch_one_device_kernel() {
     if(!g_cuda_device_launch_op.empty()) {
         device_launch_operation_t &op = g_cuda_device_launch_op.front();
 
-        stream_operation stream_op = stream_operation(op.grid, g_ptx_sim_mode, op.stream);
+        stream_operation stream_op = stream_operation(op.grid, gpgpu_ctx->func_sim->g_ptx_sim_mode, op.stream);
         g_stream_manager()->push(stream_op);
         g_cuda_device_launch_op.pop_front();
     }
 }
 
-void launch_all_device_kernels() {
+void cuda_device_runtime::launch_all_device_kernels() {
     while(!g_cuda_device_launch_op.empty()) {
         launch_one_device_kernel();
     }
diff --git a/src/cuda-sim/cuda_device_runtime.h b/src/cuda-sim/cuda_device_runtime.h
index 6dbcd71..851fed2 100644
--- a/src/cuda-sim/cuda_device_runtime.h
+++ b/src/cuda-sim/cuda_device_runtime.h
@@ -6,6 +6,20 @@
 void gpgpusim_cuda_getParameterBufferV2(const ptx_instruction * pI, ptx_thread_info * thread, const function_info * target_func);
 void gpgpusim_cuda_launchDeviceV2(const ptx_instruction * pI, ptx_thread_info * thread, const function_info * target_func);
 void gpgpusim_cuda_streamCreateWithFlags(const ptx_instruction * pI, ptx_thread_info * thread, const function_info * target_func);
-void launch_all_device_kernels();
-void launch_one_device_kernel();
+#endif
+#if (CUDART_VERSION >= 5000)
+
+class gpgpu_context;
+
+class cuda_device_runtime {
+    public:
+	cuda_device_runtime( gpgpu_context* ctx ) {
+	    gpgpu_ctx = ctx;
+	}
+	// backward pointer
+	class gpgpu_context* gpgpu_ctx;
+	void launch_all_device_kernels();
+	void launch_one_device_kernel();
+};
+
 #endif
diff --git a/src/gpgpu-sim/gpu-sim.cc b/src/gpgpu-sim/gpu-sim.cc
index 9f47067..bbcc078 100644
--- a/src/gpgpu-sim/gpu-sim.cc
+++ b/src/gpgpu-sim/gpu-sim.cc
@@ -503,7 +503,7 @@ void gpgpu_sim_config::reg_options(option_parser_t opp)
                &gpgpu_ptx_instruction_classification, 
                "if enabled will classify ptx instruction types per kernel (Max 255 kernels now)", 
                "0");
-   option_parser_register(opp, "-gpgpu_ptx_sim_mode", OPT_INT32, &g_ptx_sim_mode, 
+   option_parser_register(opp, "-gpgpu_ptx_sim_mode", OPT_INT32, &(gpgpu_ctx->func_sim->g_ptx_sim_mode),
                "Select between Performance (default) or Functional simulation (1)", 
                "0");
    option_parser_register(opp, "-gpgpu_clock_domains", OPT_CSTR, &gpgpu_clock_domains, 
@@ -1753,7 +1753,7 @@ void gpgpu_sim::cycle()
 
 #if (CUDART_VERSION >= 5000)
       //launch device kernel
-      launch_one_device_kernel();
+      gpgpu_ctx->device_runtime->launch_one_device_kernel();
 #endif
    }
 }
diff --git a/src/gpgpusim_entrypoint.cc b/src/gpgpusim_entrypoint.cc
index d9d1023..683a695 100644
--- a/src/gpgpusim_entrypoint.cc
+++ b/src/gpgpusim_entrypoint.cc
@@ -213,7 +213,7 @@ gpgpu_sim *gpgpu_context::gpgpu_ptx_sim_init_perf()
 {
    srand(1);
    print_splash();
-   read_sim_environment_variables();
+   func_sim->read_sim_environment_variables();
    ptx_parser->read_parser_environment_variables();
    option_parser_t opp = option_parser_create();