Merge branch 'dev' of github.com:purdue-aalp/gpgpu-sim_distribution into dev

3 files changed, 23 insertions, 9 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..59d2a60 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__);
 }
 
@@ -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;
@@ -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 d819559..3c9f87c 100644
--- a/libcuda/gpgpu_context.h
+++ b/libcuda/gpgpu_context.h
@@ -4,15 +4,21 @@
 #include "../src/cuda-sim/ptx_loader.h"
 #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:
 	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);
+	    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;
@@ -22,6 +28,10 @@ class gpgpu_context {
 	ptxinfo_data* ptxinfo;
 	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 );
@@ -32,6 +42,7 @@ class gpgpu_context {
 	class symbol_table* init_parser(const char*);
 	class gpgpu_sim *gpgpu_ptx_sim_init_perf();
 	struct _cuda_device_id *GPGPUSim_Init();
+	void ptx_reg_options(option_parser_t opp);
 };
 gpgpu_context* GPGPU_Context();