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

7 files changed, 1623 insertions, 1140 deletions

diff --git a/libcuda/Makefile b/libcuda/Makefile
index 13932e2..c8ff2e3 100644
--- a/libcuda/Makefile
+++ b/libcuda/Makefile
@@ -111,10 +111,10 @@ $(OUTPUT_DIR)/%.o: %.cc
 	$(CPP) $(CXXFLAGS) -I./ -I$(OUTPUT_DIR) -I$(CUDA_INSTALL_PATH)/include  -c $< -o $@
 
 $(OUTPUT_DIR)/%.o: %.c
-	$(CC) $(CCFLAGS) -I./ -I$(OUTPUT_DIR) -I$(CUDA_INSTALL_PATH)/include  -c $< -o $@
+	$(CPP) $(CCFLAGS) -I./ -I$(OUTPUT_DIR) -I$(CUDA_INSTALL_PATH)/include  -c $< -o $@
 
 $(OUTPUT_DIR)/%.o: $(OUTPUT_DIR)/%.c
-	$(CC) $(CCFLAGS) -I./ -I$(OUTPUT_DIR) -I$(CUDA_INSTALL_PATH)/include  -c $< -o $@
+	$(CPP) $(CCFLAGS) -I./ -I$(OUTPUT_DIR) -I$(CUDA_INSTALL_PATH)/include  -c $< -o $@
 
 $(OUTPUT_DIR)/cuobjdump_parser.c: cuobjdump.y
 	$(YACC) $(YFLAGS) -p cuobjdump_  -o$@ $< --file-prefix=$(OUTPUT_DIR)/cuobjdump
diff --git a/libcuda/cuda_api_object.h b/libcuda/cuda_api_object.h
new file mode 100644
index 0000000..51416f2
--- /dev/null
+++ b/libcuda/cuda_api_object.h
@@ -0,0 +1,213 @@
+#ifndef __cuda_api_object_h__
+#define __cuda_api_object_h__
+
+#include <list>
+#include <map>
+#include <set>
+#include <string>
+
+#include "builtin_types.h"
+
+#include "../src/gpgpu-sim/gpu-sim.h"
+#include "../src/cuda-sim/ptx_ir.h"
+#include "../src/abstract_hardware_model.h"
+#include "cuobjdump.h"
+
+typedef std::list<gpgpu_ptx_sim_arg> gpgpu_ptx_sim_arg_list_t;
+
+#ifndef OPENGL_SUPPORT
+typedef unsigned long GLuint;
+#endif
+
+struct glbmap_entry {
+	GLuint m_bufferObj;
+	void *m_devPtr;
+	size_t m_size;
+	struct glbmap_entry *m_next;
+};
+
+typedef struct glbmap_entry glbmap_entry_t;
+
+struct _cuda_device_id {
+	_cuda_device_id(gpgpu_sim* gpu) {m_id = 0; m_next = NULL; m_gpgpu=gpu;}
+	struct _cuda_device_id *next() { return m_next; }
+	unsigned num_shader() const { return m_gpgpu->get_config().num_shader(); }
+	int num_devices() const {
+		if( m_next == NULL ) return 1;
+		else return 1 + m_next->num_devices();
+	}
+	struct _cuda_device_id *get_device( unsigned n )
+	{
+		assert( n < (unsigned)num_devices() );
+		struct _cuda_device_id *p=this;
+		for(unsigned i=0; i<n; i++)
+			p = p->m_next;
+		return p;
+	}
+	const struct cudaDeviceProp *get_prop() const
+	{
+		return m_gpgpu->get_prop();
+	}
+	unsigned get_id() const { return m_id; }
+
+	gpgpu_sim *get_gpgpu() { return m_gpgpu; }
+private:
+	unsigned m_id;
+	class gpgpu_sim *m_gpgpu;
+	struct _cuda_device_id *m_next;
+};
+
+struct CUctx_st {
+	CUctx_st( _cuda_device_id *gpu )
+	{
+		m_gpu = gpu;
+		m_binary_info.cmem = 0;
+		m_binary_info.gmem = 0;
+		no_of_ptx=0;
+	}
+
+	_cuda_device_id *get_device() { return m_gpu; }
+
+	void add_binary( symbol_table *symtab, unsigned fat_cubin_handle )
+	{
+		m_code[fat_cubin_handle] = symtab;
+		m_last_fat_cubin_handle = fat_cubin_handle;
+	}
+
+	void add_ptxinfo( const char *deviceFun, const struct gpgpu_ptx_sim_info &info )
+	{
+		symbol *s = m_code[m_last_fat_cubin_handle]->lookup(deviceFun);
+		assert( s != NULL );
+		function_info *f = s->get_pc();
+		assert( f != NULL );
+		f->set_kernel_info(info);
+	}
+
+	void add_ptxinfo( const struct gpgpu_ptx_sim_info &info )
+	{
+		m_binary_info = info;
+	}
+
+	void register_function( unsigned fat_cubin_handle, const char *hostFun, const char *deviceFun )
+	{
+		if( m_code.find(fat_cubin_handle) != m_code.end() ) {
+			symbol *s = m_code[fat_cubin_handle]->lookup(deviceFun);
+			if(s != NULL) {
+				function_info *f = s->get_pc();
+				assert( f != NULL );
+				m_kernel_lookup[hostFun] = f;
+			}
+			else {
+				printf("Warning: cannot find deviceFun %s\n", deviceFun);
+				m_kernel_lookup[hostFun] = NULL;
+			}
+	//		assert( s != NULL );
+	//		function_info *f = s->get_pc();
+	//		assert( f != NULL );
+	//		m_kernel_lookup[hostFun] = f;
+		} else {
+			m_kernel_lookup[hostFun] = NULL;
+		}
+	}
+
+    void register_hostFun_function( const char*hostFun, function_info* f){
+        m_kernel_lookup[hostFun] = f;
+    }
+
+	function_info *get_kernel(const char *hostFun)
+	{
+		std::map<const void*,function_info*>::iterator i=m_kernel_lookup.find(hostFun);
+		assert( i != m_kernel_lookup.end() );
+		return i->second;
+	}
+
+	int no_of_ptx;
+
+private:
+	_cuda_device_id *m_gpu; // selected gpu
+	std::map<unsigned,symbol_table*> m_code; // fat binary handle => global symbol table
+	unsigned m_last_fat_cubin_handle;
+	std::map<const void*,function_info*> m_kernel_lookup; // unique id (CUDA app function address) => kernel entry point
+	struct gpgpu_ptx_sim_info m_binary_info;
+
+};
+
+class kernel_config {
+public:
+	kernel_config( dim3 GridDim, dim3 BlockDim, size_t sharedMem, struct CUstream_st *stream )
+	{
+		m_GridDim=GridDim;
+		m_BlockDim=BlockDim;
+		m_sharedMem=sharedMem;
+		m_stream = stream;
+	}
+	kernel_config()
+	{
+		m_GridDim=dim3(-1,-1,-1);
+		m_BlockDim=dim3(-1,-1,-1);
+		m_sharedMem=0;
+		m_stream =NULL;
+	}
+	void set_arg( const void *arg, size_t size, size_t offset )
+	{
+		m_args.push_front( gpgpu_ptx_sim_arg(arg,size,offset) );
+	}
+	dim3 grid_dim() const { return m_GridDim; }
+	dim3 block_dim() const { return m_BlockDim; }
+	void set_grid_dim(dim3 *d) { m_GridDim = *d; }
+	void set_block_dim(dim3 *d) { m_BlockDim = *d; }
+	gpgpu_ptx_sim_arg_list_t get_args() { return m_args; }
+	struct CUstream_st *get_stream() { return m_stream; }
+
+private:
+	dim3 m_GridDim;
+	dim3 m_BlockDim;
+	size_t m_sharedMem;
+	struct CUstream_st *m_stream;
+	gpgpu_ptx_sim_arg_list_t m_args;
+};
+
+class cuda_runtime_api {
+    public:
+	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;
+	std::list<cuobjdumpSection*> libSectionList;
+	std::list<kernel_config> g_cuda_launch_stack;
+	std::map<int, bool>fatbin_registered;
+	std::map<int, std::string> fatbinmap;
+	std::map<std::string, symbol_table*> name_symtab;
+	std::map<unsigned long long, size_t> g_mallocPtr_Size;
+	//maps sm version number to set of filenames
+	std::map<unsigned, std::set<std::string> > version_filename;
+	std::map<void *,void **> pinned_memory; //support for pinned memories added
+	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();
+	void extract_ptx_files_using_cuobjdump(CUctx_st *context);
+	std::list<cuobjdumpSection*> pruneSectionList(CUctx_st *context);
+	std::list<cuobjdumpSection*> mergeMatchingSections(std::string identifier);
+	std::list<cuobjdumpSection*> mergeSections();
+	cuobjdumpELFSection* findELFSection(const std::string identifier);
+	cuobjdumpPTXSection* findPTXSection(const std::string identifier);
+	cuobjdumpPTXSection* findPTXSectionInList(std::list<cuobjdumpSection*> &sectionlist, const std::string identifier);
+	void cuobjdumpRegisterFatBinary(unsigned int handle, const char* filename, CUctx_st *context);
+	kernel_info_t *gpgpu_cuda_ptx_sim_init_grid( const char *kernel_key,
+		gpgpu_ptx_sim_arg_list_t args,
+		struct dim3 gridDim,
+		struct dim3 blockDim,
+		struct CUctx_st* context );
+	int load_static_globals( symbol_table *symtab, unsigned min_gaddr, unsigned max_gaddr, gpgpu_t *gpu );
+	int load_constants( symbol_table *symtab, addr_t min_gaddr, gpgpu_t *gpu );
+
+};
+#endif /* __cuda_api_object_h__ */
diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index 45511d4..9d685e6 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -131,6 +131,8 @@
 #if (CUDART_VERSION < 8000)
 #include "__cudaFatFormat.h"
 #endif
+#include "gpgpu_context.h"
+#include "cuda_api_object.h"
 #include "../src/gpgpu-sim/gpu-sim.h"
 #include "../src/cuda-sim/ptx_loader.h"
 #include "../src/cuda-sim/cuda-sim.h"
@@ -147,25 +149,10 @@
 #include <mach-o/dyld.h>
 #endif
 
-std::map<void *,void **> pinned_memory; //support for pinned memories added
-std::map<void *, size_t> pinned_memory_size;
-std::map<unsigned long long, size_t> g_mallocPtr_Size;
-int no_of_ptx=0;
-//maps sm version number to set of filenames
-std::map<unsigned, std::set<std::string> > version_filename;
 
 extern void synchronize();
 extern void exit_simulation();
 
-static int load_static_globals( symbol_table *symtab, unsigned min_gaddr, unsigned max_gaddr, gpgpu_t *gpu );
-static int load_constants( symbol_table *symtab, addr_t min_gaddr, gpgpu_t *gpu );
-
-static kernel_info_t *gpgpu_cuda_ptx_sim_init_grid( const char *kernel_key, 
-		gpgpu_ptx_sim_arg_list_t args,
-		struct dim3 gridDim,
-		struct dim3 blockDim,
-		struct CUctx_st* context );
-
 /*DEVICE_BUILTIN*/
 struct cudaArray
 {
@@ -189,7 +176,7 @@ struct cudaArray
 
 cudaError_t g_last_cudaError = cudaSuccess;
 
-extern stream_manager *g_stream_manager;
+//extern stream_manager *g_stream_manager();
 
 void register_ptx_function( const char *name, function_info *impl )
 {
@@ -210,145 +197,10 @@ void register_ptx_function( const char *name, function_info *impl )
 # endif
 #endif
 
-struct _cuda_device_id {
-	_cuda_device_id(gpgpu_sim* gpu) {m_id = 0; m_next = NULL; m_gpgpu=gpu;}
-	struct _cuda_device_id *next() { return m_next; }
-	unsigned num_shader() const { return m_gpgpu->get_config().num_shader(); }
-	int num_devices() const {
-		if( m_next == NULL ) return 1;
-		else return 1 + m_next->num_devices();
-	}
-	struct _cuda_device_id *get_device( unsigned n )
-	{
-		assert( n < (unsigned)num_devices() );
-		struct _cuda_device_id *p=this;
-		for(unsigned i=0; i<n; i++)
-			p = p->m_next;
-		return p;
-	}
-	const struct cudaDeviceProp *get_prop() const
-	{
-		return m_gpgpu->get_prop();
-	}
-	unsigned get_id() const { return m_id; }
-
-	gpgpu_sim *get_gpgpu() { return m_gpgpu; }
-private:
-	unsigned m_id;
-	class gpgpu_sim *m_gpgpu;
-	struct _cuda_device_id *m_next;
-};
-
-struct CUctx_st {
-	CUctx_st( _cuda_device_id *gpu )
-	{
-		m_gpu = gpu;
-		m_binary_info.cmem = 0;
-		m_binary_info.gmem = 0;
-	}
-
-	_cuda_device_id *get_device() { return m_gpu; }
-
-	void add_binary( symbol_table *symtab, unsigned fat_cubin_handle )
-	{
-		m_code[fat_cubin_handle] = symtab;
-		m_last_fat_cubin_handle = fat_cubin_handle;
-	}
-
-	void add_ptxinfo( const char *deviceFun, const struct gpgpu_ptx_sim_info &info )
-	{
-		symbol *s = m_code[m_last_fat_cubin_handle]->lookup(deviceFun);
-		assert( s != NULL );
-		function_info *f = s->get_pc();
-		assert( f != NULL );
-		f->set_kernel_info(info);
-	}
-
-	void add_ptxinfo( const struct gpgpu_ptx_sim_info &info )
-	{
-		m_binary_info = info;
-	}
-
-	void register_function( unsigned fat_cubin_handle, const char *hostFun, const char *deviceFun )
-	{
-		if( m_code.find(fat_cubin_handle) != m_code.end() ) {
-			symbol *s = m_code[fat_cubin_handle]->lookup(deviceFun);
-			if(s != NULL) {
-				function_info *f = s->get_pc();
-				assert( f != NULL );
-				m_kernel_lookup[hostFun] = f;
-			}
-			else {
-				printf("Warning: cannot find deviceFun %s\n", deviceFun);
-				m_kernel_lookup[hostFun] = NULL;
-			}
-	//		assert( s != NULL );
-	//		function_info *f = s->get_pc();
-	//		assert( f != NULL );
-	//		m_kernel_lookup[hostFun] = f;
-		} else {
-			m_kernel_lookup[hostFun] = NULL;
-		}
-	}
-
-    void register_hostFun_function( const char*hostFun, function_info* f){
-        m_kernel_lookup[hostFun] = f;
-    }
-
-	function_info *get_kernel(const char *hostFun)
-	{
-		std::map<const void*,function_info*>::iterator i=m_kernel_lookup.find(hostFun);
-		assert( i != m_kernel_lookup.end() );
-		return i->second;
-	}
-
-private:
-	_cuda_device_id *m_gpu; // selected gpu
-	std::map<unsigned,symbol_table*> m_code; // fat binary handle => global symbol table
-	unsigned m_last_fat_cubin_handle;
-	std::map<const void*,function_info*> m_kernel_lookup; // unique id (CUDA app function address) => kernel entry point
-	struct gpgpu_ptx_sim_info m_binary_info;
-
-};
-
-class kernel_config {
-public:
-	kernel_config( dim3 GridDim, dim3 BlockDim, size_t sharedMem, struct CUstream_st *stream )
-	{
-		m_GridDim=GridDim;
-		m_BlockDim=BlockDim;
-		m_sharedMem=sharedMem;
-		m_stream = stream;
-	}
-	kernel_config()
-	{
-		m_GridDim=dim3(-1,-1,-1);
-		m_BlockDim=dim3(-1,-1,-1);
-		m_sharedMem=0;
-		m_stream =NULL;
-	}
-	void set_arg( const void *arg, size_t size, size_t offset )
-	{
-		m_args.push_front( gpgpu_ptx_sim_arg(arg,size,offset) );
-	}
-	dim3 grid_dim() const { return m_GridDim; }
-	dim3 block_dim() const { return m_BlockDim; }
-	void set_grid_dim(dim3 *d) { m_GridDim = *d; }
-	void set_block_dim(dim3 *d) { m_BlockDim = *d; }
-	gpgpu_ptx_sim_arg_list_t get_args() { return m_args; }
-	struct CUstream_st *get_stream() { return m_stream; }
-
-private:
-	dim3 m_GridDim;
-	dim3 m_BlockDim;
-	size_t m_sharedMem;
-	struct CUstream_st *m_stream;
-	gpgpu_ptx_sim_arg_list_t m_args;
-};
-
-struct _cuda_device_id *GPGPUSim_Init()
+struct _cuda_device_id *gpgpu_context::GPGPUSim_Init()
 {
-	static _cuda_device_id *the_device = NULL;
+	//static _cuda_device_id *the_device = NULL;
+	_cuda_device_id *the_device = GPGPUsim_ctx_ptr()->the_cude_device;
 	if( !the_device ) {
 		gpgpu_sim *the_gpu = gpgpu_ptx_sim_init_perf();
 
@@ -393,7 +245,8 @@ struct _cuda_device_id *GPGPUSim_Init()
 		prop->maxThreadsPerMultiProcessor = the_gpu->threads_per_core();
 #endif
 		the_gpu->set_prop(prop);
-		the_device = new _cuda_device_id(the_gpu);
+		GPGPUsim_ctx_ptr()->the_cude_device = new _cuda_device_id(the_gpu);
+		the_device = GPGPUsim_ctx_ptr()->the_cude_device;
 	}
 	start_sim_thread(1);
 	return the_device;
@@ -401,15 +254,27 @@ struct _cuda_device_id *GPGPUSim_Init()
 
 static CUctx_st* GPGPUSim_Context()
 {
-	static CUctx_st *the_context = NULL;
+	//static CUctx_st *the_context = NULL;
+	gpgpu_context *cur_ctx = GPGPU_Context();
+	CUctx_st *the_context = GPGPUsim_ctx_ptr()->the_context;
 	if( the_context == NULL ) {
-		_cuda_device_id *the_gpu = GPGPUSim_Init();
-		the_context = new CUctx_st(the_gpu);
+		_cuda_device_id *the_gpu = cur_ctx->GPGPUSim_Init();
+		GPGPUsim_ctx_ptr()->the_context = new CUctx_st(the_gpu);
+		the_context = GPGPUsim_ctx_ptr()->the_context;
 	}
 	return the_context;
 }
 
- void ptxinfo_addinfo()
+gpgpu_context* GPGPU_Context()
+{
+	static gpgpu_context *gpgpu_ctx = NULL;
+	if( gpgpu_ctx == NULL ) {
+	    gpgpu_ctx = new gpgpu_context();
+	}
+	return gpgpu_ctx;
+}
+
+ void ptxinfo_data::ptxinfo_addinfo()
 {
 	 if(!get_ptxinfo_kname()){
 		 /* This info is not per kernel (since CUDA 5.0 some info (e.g. gmem, and cmem) is added at the beginning for the whole binary ) */
@@ -480,29 +345,618 @@ typedef std::map<unsigned,CUevent_st*> event_tracker_t;
 
 int CUevent_st::m_next_event_uid;
 event_tracker_t g_timer_events;
-int g_active_device = 0; //active gpu that runs the code
-std::list<kernel_config> g_cuda_launch_stack;
 
-/*******************************************************************************
- *                                                                              *
- *                                                                              *
- *                                                                              *
- *******************************************************************************/
+extern int cuobjdump_lex_init(yyscan_t* scanner);
+extern void cuobjdump_set_in  (FILE * _in_str ,yyscan_t yyscanner );
+extern int cuobjdump_parse(yyscan_t scanner, struct cuobjdump_parser* parser, std::list<cuobjdumpSection*> &cuobjdumpSectionList);
+extern int cuobjdump_lex_destroy(yyscan_t scanner);
+
+enum cuobjdumpSectionType {
+	PTXSECTION=0,
+	ELFSECTION
+};
+
+
+// sectiontype: 0 for ptx, 1 for elf
+void addCuobjdumpSection(int sectiontype, std::list<cuobjdumpSection*> &cuobjdumpSectionList){
+	if (sectiontype)
+		cuobjdumpSectionList.push_front(new cuobjdumpELFSection());
+	else
+		cuobjdumpSectionList.push_front(new cuobjdumpPTXSection());
+	printf("## Adding new section %s\n", sectiontype?"ELF":"PTX");
+}
+
+void setCuobjdumparch(const char* arch, std::list<cuobjdumpSection*> &cuobjdumpSectionList){
+	unsigned archnum;
+	sscanf(arch, "sm_%u", &archnum);
+	assert (archnum && "cannot have sm_0");
+	printf("Adding arch: %s\n", arch);
+	cuobjdumpSectionList.front()->setArch(archnum);
+}
+
+void setCuobjdumpidentifier(const char* identifier, std::list<cuobjdumpSection*> &cuobjdumpSectionList){
+	printf("Adding identifier: %s\n", identifier);
+	cuobjdumpSectionList.front()->setIdentifier(identifier);
+}
+
+void setCuobjdumpptxfilename(const char* filename, std::list<cuobjdumpSection*> &cuobjdumpSectionList){
+	printf("Adding ptx filename: %s\n", filename);
+	cuobjdumpSection* x = cuobjdumpSectionList.front();
+	if (dynamic_cast<cuobjdumpPTXSection*>(x) == NULL){
+		assert (0 && "You shouldn't be trying to add a ptxfilename to an elf section");
+	}
+	(dynamic_cast<cuobjdumpPTXSection*>(x))->setPTXfilename(filename);
+}
+
+void setCuobjdumpelffilename(const char* filename, std::list<cuobjdumpSection*> &cuobjdumpSectionList){
+	if (dynamic_cast<cuobjdumpELFSection*>(cuobjdumpSectionList.front()) == NULL){
+		assert (0 && "You shouldn't be trying to add a elffilename to an ptx section");
+	}
+	(dynamic_cast<cuobjdumpELFSection*>(cuobjdumpSectionList.front()))->setELFfilename(filename);
+}
+
+void setCuobjdumpsassfilename(const char* filename, std::list<cuobjdumpSection*> &cuobjdumpSectionList){
+	if (dynamic_cast<cuobjdumpELFSection*>(cuobjdumpSectionList.front()) == NULL){
+		assert (0 && "You shouldn't be trying to add a sassfilename to an ptx section");
+	}
+	(dynamic_cast<cuobjdumpELFSection*>(cuobjdumpSectionList.front()))->setSASSfilename(filename);
+}
+
+//! Return the executable file of the process containing the PTX/SASS code 
+//!
+//! This Function returns the executable file ran by the process.  This
+//! executable is supposed to contain the PTX/SASS code.  It provides workaround
+//! for processes running on valgrind by dereferencing /proc/<pid>/exe within the
+//! GPGPU-Sim process before calling cuobjdump to extract PTX/SASS.  This is
+//! needed because valgrind uses x86 emulation to detect memory leak.  Other
+//! processes (e.g. cuobjdump) reading /proc/<pid>/exe will see the emulator
+//! executable instead of the application binary.  
+//! 
+std::string get_app_binary(){
+   char self_exe_path[1025];
+#ifdef __APPLE__
+   uint32_t size = sizeof(self_exe_path);
+   if( _NSGetExecutablePath(self_exe_path,&size) != 0 ) {
+	   printf("GPGPU-Sim ** ERROR: _NSGetExecutablePath input buffer too small\n");
+	   exit(1);
+   }
+#else
+   std::stringstream exec_link;
+   exec_link << "/proc/self/exe";
+
+   ssize_t path_length = readlink(exec_link.str().c_str(), self_exe_path, 1024); 
+   assert(path_length != -1); 
+   self_exe_path[path_length] = '\0'; 
+#endif
+
+   printf("self exe links to: %s\n", self_exe_path); 
+   return self_exe_path; 
+}
+
+//above func gives abs path whereas this give just the name of application.
+char* get_app_binary_name(std::string abs_path){
+   char *self_exe_path;
+#ifdef __APPLE__
+   //TODO: get apple device and check the result.
+   printf("WARNING: not tested for Apple-mac devices \n");
+   abort();
+#else
+   char* buf = strdup(abs_path.c_str());
+   char *token = strtok(buf, "/");
+   while(token !=NULL){
+	self_exe_path = token;
+	token = strtok(NULL,"/");
+   }
+#endif
+   self_exe_path = strtok(self_exe_path, ".");
+   printf("self exe links to: %s\n", self_exe_path);
+   return self_exe_path;
+}
+
+static int get_app_cuda_version() {
+    int app_cuda_version = 0;
+    char fname[1024];
+    snprintf(fname,1024,"_app_cuda_version_XXXXXX");
+    int fd=mkstemp(fname);
+    close(fd);
+    std::string app_cuda_version_command = "ldd " + get_app_binary() + " | grep libcudart.so | sed  's/.*libcudart.so.\\(.*\\) =>.*/\\1/' > " + fname;
+    system(app_cuda_version_command.c_str());
+    FILE * cmd = fopen(fname, "r");
+    char buf[256];
+    while (fgets(buf, sizeof(buf), cmd) != 0) {
+        std::cout << buf;
+        app_cuda_version = atoi(buf);
+    }
+    fclose(cmd);
+    if ( app_cuda_version == 0 ) {
+        printf( "Error - Cannot detect the app's CUDA version.\n" );
+        exit(1);
+    }
+    return app_cuda_version;
+}
+
+//! Keep track of the association between filename and cubin handle
+void cuda_runtime_api::cuobjdumpRegisterFatBinary(unsigned int handle, const char* filename, CUctx_st *context){
+	fatbinmap[handle] = filename;
+}
+
 
-extern "C" {
 
 /*******************************************************************************
- *                                                                              *
- *                                                                              *
- *                                                                              *
+ * Add internal cuda runtime API call to accept gpgpu_context                   *
  *******************************************************************************/
-cudaError_t cudaPeekAtLastError(void)
+cudaError_t cudaSetDeviceInternal(int device, gpgpu_context* gpgpu_ctx = NULL)
 {
-	return g_last_cudaError;
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+	//set the active device to run cuda
+	if ( device <= ctx->GPGPUSim_Init()->num_devices() ) {
+		ctx->api->g_active_device = device;
+		return g_last_cudaError = cudaSuccess;
+	} else {
+		return g_last_cudaError = cudaErrorInvalidDevice;
+	}
 }
 
-__host__ cudaError_t CUDARTAPI cudaMalloc(void **devPtr, size_t size) 
+cudaError_t cudaGetDeviceInternal(int *device, gpgpu_context* gpgpu_ctx = NULL)
+{
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+	*device = ctx->api->g_active_device;
+	return g_last_cudaError = cudaSuccess;
+}
+
+__host__ cudaError_t CUDARTAPI cudaDeviceGetLimitInternal( size_t* pValue, cudaLimit limit, gpgpu_context* gpgpu_ctx = NULL )
+{
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
+    if(g_debug_execution >= 3){
+	announce_call(__my_func__);
+    }
+    _cuda_device_id *dev = ctx->GPGPUSim_Init();
+    const struct cudaDeviceProp *prop = dev->get_prop();
+    const gpgpu_sim_config& config=dev->get_gpgpu()->get_config();
+    switch(limit) {
+	case 0:  // cudaLimitStackSize
+	    *pValue=config.stack_limit();
+	    break;
+	case 2:  // cudaLimitMallocHeapSize
+	    *pValue=config.heap_limit();
+	    break;
+#if (CUDART_VERSION > 5050)
+	case 3: // cudaLimitDevRuntimeSyncDepth
+	    if(prop->major > 2){
+		*pValue=config.sync_depth_limit();
+		break;
+	    }
+	    else{
+		printf("ERROR:Limit %s is not supported on this architecture \n",limit);
+		abort();
+	    }
+	case 4: // cudaLimitDevRuntimePendingLaunchCount
+	    if(prop->major > 2){
+		*pValue=config.pending_launch_count_limit();
+		break;
+	    }
+	    else{
+		printf("ERROR:Limit %s is not supported on this architecture \n",limit);
+		abort();
+	    }
+#endif
+	default:
+	    printf("ERROR:Limit %s unimplemented \n",limit);
+	    abort();
+    }
+    return g_last_cudaError = cudaSuccess;
+
+}
+
+
+void** cudaRegisterFatBinaryInternal( void *fatCubin, gpgpu_context* gpgpu_ctx = NULL)
 {
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+#if (CUDART_VERSION < 2010)
+	printf("GPGPU-Sim PTX: ERROR ** this version of GPGPU-Sim requires CUDA 2.1 or higher\n");
+	exit(1);
+#endif
+	CUctx_st *context = GPGPUSim_Context();
+	static unsigned next_fat_bin_handle = 1;
+	if(context->get_device()->get_gpgpu()->get_config().use_cuobjdump()) {
+		// The following workaround has only been verified on 64-bit systems. 
+		if (sizeof(void*) == 4) 
+			printf("GPGPU-Sim PTX: FatBin file name extraction has not been tested on 32-bit system.\n"); 
+
+        // This code will get the CUDA version the app was compiled with.
+        // We need this to determine how to handle the parsing of the binary.
+        // Making this a runtime variable based on the app, enables GPGPU-Sim compiled
+        // with a newer version of CUDA to run apps compiled with older versions of
+        // CUDA. This is especially useful for PTXPLUS execution.
+        //Skip cuda version check for pytorch application
+	std::string app_binary_path =  get_app_binary();
+        int pos = app_binary_path.find("python");
+        if (pos==std::string::npos){
+        	// Not pytorch app : checking cuda version
+		int app_cuda_version = get_app_cuda_version();
+        	assert( app_cuda_version == CUDART_VERSION / 1000  && "The app must be compiled with same major version as the simulator." );
+        }
+
+	//int app_cuda_version = get_app_cuda_version();
+        //assert( app_cuda_version == CUDART_VERSION / 1000  && "The app must be compiled with same major version as the simulator." );
+        const char* filename;
+#if CUDART_VERSION < 6000
+            // FatBin handle from the .fatbin.c file (one of the intermediate files generated by NVCC)
+            typedef struct {int m; int v; const unsigned long long* d; char* f;} __fatDeviceText __attribute__ ((aligned (8))); 
+            __fatDeviceText * fatDeviceText = (__fatDeviceText *) fatCubin;
+
+            // Extract the source code file name that generate the given FatBin. 
+            // - Obtains the pointer to the actual fatbin structure from the FatBin handle (fatCubin).
+            // - An integer inside the fatbin structure contains the relative offset to the source code file name.
+            // - This offset differs among different CUDA and GCC versions. 
+            char * pfatbin = (char*) fatDeviceText->d; 
+            int offset = *((int*)(pfatbin+48)); 
+            filename = (pfatbin+16+offset);
+#else
+            filename = "default";
+#endif
+
+		// The extracted file name is associated with a fat_cubin_handle passed
+		// into cudaLaunch().  Inside cudaLaunch(), the associated file name is
+		// used to find the PTX/SASS section from cuobjdump, which contains the
+		// PTX/SASS code for the launched kernel function.  
+		// This allows us to work around the fact that cuobjdump only outputs the
+		// file name associated with each section. 
+		unsigned long long fat_cubin_handle = next_fat_bin_handle;
+		next_fat_bin_handle++;
+		printf("GPGPU-Sim PTX: __cudaRegisterFatBinary, fat_cubin_handle = %llu, filename=%s\n", fat_cubin_handle, filename);
+		/*!
+		 * This function extracts all data from all files in first call
+		 * then for next calls, only returns the appropriate number
+		 */
+		assert(fat_cubin_handle >= 1);
+		if (fat_cubin_handle==1) ctx->api->cuobjdumpInit();
+		ctx->api->cuobjdumpRegisterFatBinary(fat_cubin_handle, filename, context);
+
+		return (void**)fat_cubin_handle;
+	} 
+#if (CUDART_VERSION < 8000)
+	else {
+		static unsigned source_num=1;
+		unsigned long long fat_cubin_handle = next_fat_bin_handle++;
+		__cudaFatCudaBinary *info =   (__cudaFatCudaBinary *)fatCubin;
+		assert( info->version >= 3 );
+		unsigned num_ptx_versions=0;
+		unsigned max_capability=0;
+		unsigned selected_capability=0;
+		bool found=false;
+		unsigned forced_max_capability = context->get_device()->get_gpgpu()->get_config().get_forced_max_capability();
+		if (!info->ptx){
+			printf("ERROR: Cannot find ptx code in cubin file\n"
+					"\tIf you are using CUDA 4.0 or higher, please enable -gpgpu_ptx_use_cuobjdump or downgrade to CUDA 3.1\n");
+			exit(1);
+		}
+		while( info->ptx[num_ptx_versions].gpuProfileName != NULL ) {
+			unsigned capability=0;
+			sscanf(info->ptx[num_ptx_versions].gpuProfileName,"compute_%u",&capability);
+			printf("GPGPU-Sim PTX: __cudaRegisterFatBinary found PTX versions for '%s', ", info->ident);
+			printf("capability = %s\n", info->ptx[num_ptx_versions].gpuProfileName );
+			if( forced_max_capability ) {
+				if( capability > max_capability && capability <= forced_max_capability ) {
+					found = true;
+					max_capability=capability;
+					selected_capability = num_ptx_versions;
+				}
+			} else {
+				if( capability > max_capability ) {
+					found = true;
+					max_capability=capability;
+					selected_capability = num_ptx_versions;
+				}
+			}
+			num_ptx_versions++;
+		}
+		if( found  ) {
+			printf("GPGPU-Sim PTX: Loading PTX for %s, capability = %s\n",
+					info->ident, info->ptx[selected_capability].gpuProfileName );
+			symbol_table *symtab;
+			const char *ptx = info->ptx[selected_capability].ptx;
+			if(context->get_device()->get_gpgpu()->get_config().convert_to_ptxplus() ) {
+				printf("GPGPU-Sim PTX: ERROR ** PTXPlus is only supported through cuobjdump\n"
+						"\tEither enable cuobjdump or disable PTXPlus in your configuration file\n");
+				exit(1);
+			} else {
+				symtab=ctx->gpgpu_ptx_sim_load_ptx_from_string(ptx,source_num);
+				context->add_binary(symtab,fat_cubin_handle);
+				ctx->gpgpu_ptxinfo_load_from_string( ptx, source_num, max_capability, context->no_of_ptx );
+			}
+			source_num++;
+			ctx->api->load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF,context->get_device()->get_gpgpu());
+			ctx->api->load_constants(symtab,STATIC_ALLOC_LIMIT,context->get_device()->get_gpgpu());
+		} else {
+			printf("GPGPU-Sim PTX: warning -- did not find an appropriate PTX in cubin\n");
+		}
+		return (void**)fat_cubin_handle;
+	}
+#else
+        else {
+		printf("ERROR **  __cudaRegisterFatBinary() needs to be updated\n");
+		abort();
+        }
+#endif
+}
+
+void cudaRegisterFunctionInternal(
+		void   **fatCubinHandle,
+		const char    *hostFun,
+		char    *deviceFun,
+		const char    *deviceName,
+		int      thread_limit,
+		uint3   *tid,
+		uint3   *bid,
+		dim3    *bDim,
+		dim3    *gDim,
+		gpgpu_context *gpgpu_ctx = NULL
+)
+{
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+	CUctx_st *context = GPGPUSim_Context();
+	unsigned fat_cubin_handle = (unsigned)(unsigned long long)fatCubinHandle;
+	printf("GPGPU-Sim PTX: __cudaRegisterFunction %s : hostFun 0x%p, fat_cubin_handle = %u\n",
+			deviceFun, hostFun, fat_cubin_handle);
+	if(context->get_device()->get_gpgpu()->get_config().use_cuobjdump())
+		ctx->cuobjdumpParseBinary(fat_cubin_handle);
+	context->register_function( fat_cubin_handle, hostFun, deviceFun );
+}
+
+void cudaRegisterVarInternal(
+		void **fatCubinHandle,
+		char *hostVar, //pointer to...something
+		char *deviceAddress, //name of variable
+		const char *deviceName, //name of variable (same as above)
+		int ext,
+		int size,
+		int constant,
+		int global,
+		gpgpu_context *gpgpu_ctx = NULL)
+{
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+	printf("GPGPU-Sim PTX: __cudaRegisterVar: hostVar = %p; deviceAddress = %s; deviceName = %s\n", hostVar, deviceAddress, deviceName);
+	printf("GPGPU-Sim PTX: __cudaRegisterVar: Registering const memory space of %d bytes\n", size);
+	if(GPGPUSim_Context()->get_device()->get_gpgpu()->get_config().use_cuobjdump())
+		ctx->cuobjdumpParseBinary((unsigned)(unsigned long long)fatCubinHandle);
+	fflush(stdout);
+	if ( constant && !global && !ext ) {
+		ctx->func_sim->gpgpu_ptx_sim_register_const_variable(hostVar,deviceName,size);
+	} else if ( !constant && !global && !ext ) {
+		ctx->func_sim->gpgpu_ptx_sim_register_global_variable(hostVar,deviceName,size);
+	} else cuda_not_implemented(__my_func__,__LINE__);
+}
+
+cudaError_t cudaConfigureCallInternal(dim3 gridDim, dim3 blockDim, size_t sharedMem, cudaStream_t stream, gpgpu_context* gpgpu_ctx = NULL)
+{
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+	struct CUstream_st *s = (struct CUstream_st *)stream;
+	ctx->api->g_cuda_launch_stack.push_back( kernel_config(gridDim,blockDim,sharedMem,s) );
+	return g_last_cudaError = cudaSuccess;
+}
+
+__host__ cudaError_t CUDARTAPI cudaGetDeviceCountInternal(int *count, gpgpu_context* gpgpu_ctx = NULL)
+{
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
+    if(g_debug_execution >= 3){
+	announce_call(__my_func__);
+    }
+    _cuda_device_id *dev = ctx->GPGPUSim_Init();
+    *count = dev->num_devices();
+    return g_last_cudaError = cudaSuccess;
+}
+
+__host__ cudaError_t CUDARTAPI cudaGetDevicePropertiesInternal(struct cudaDeviceProp *prop, int device, gpgpu_context* gpgpu_ctx = NULL)
+{
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
+    if(g_debug_execution >= 3){
+	announce_call(__my_func__);
+    }
+    _cuda_device_id *dev = ctx->GPGPUSim_Init();
+    if (device <= dev->num_devices() )  {
+	*prop= *dev->get_prop();
+	return g_last_cudaError = cudaSuccess;
+    } else {
+	return g_last_cudaError = cudaErrorInvalidDevice;
+    }
+}
+
+
+__host__ cudaError_t CUDARTAPI cudaChooseDeviceInternal(int *device, const struct cudaDeviceProp *prop, gpgpu_context* gpgpu_ctx = NULL)
+{
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+	_cuda_device_id *dev = ctx->GPGPUSim_Init();
+	*device = dev->get_id();
+	return g_last_cudaError = cudaSuccess;
+}
+
+cudaError_t cudaSetupArgumentInternal(const void *arg, size_t size, size_t offset, gpgpu_context* gpgpu_ctx = NULL)
+{
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+	gpgpusim_ptx_assert( !ctx->api->g_cuda_launch_stack.empty(), "empty launch stack" );
+	kernel_config &config = ctx->api->g_cuda_launch_stack.back();
+	config.set_arg(arg,size,offset);
+	printf("GPGPU-Sim PTX: Setting up arguments for %zu bytes starting at 0x%llx..\n",size, (unsigned long long) arg);
+
+	return g_last_cudaError = cudaSuccess;
+}
+
+cudaError_t cudaLaunchInternal( const char *hostFun, gpgpu_context* gpgpu_ctx = NULL )
+{
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+	CUctx_st* context = GPGPUSim_Context();
+	char *mode = getenv("PTX_SIM_MODE_FUNC");
+	if( 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();
+	{
+		dim3 gridDim = config.grid_dim();
+		dim3 blockDim = config.block_dim();
+		if (gridDim.x * gridDim.y * gridDim.z == 0 || blockDim.x * blockDim.y * blockDim.z == 0)
+		{
+			//can't launch
+			printf("can't launch a empty kernel\n");
+			ctx->api->g_cuda_launch_stack.pop_back();
+			return g_last_cudaError = cudaErrorInvalidConfiguration;
+		}
+	}
+	struct CUstream_st *stream = config.get_stream();
+
+	printf("\nGPGPU-Sim PTX: cudaLaunch for 0x%p (mode=%s) on stream %u\n", hostFun,
+			(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();
+	function_info *kernel_func_info = grid->entry();
+	if (kernel_func_info->is_pdom_set()) {
+    		printf("GPGPU-Sim PTX: PDOM analysis already done for %s \n", kname.c_str() );
+    	} else {
+    		printf("GPGPU-Sim PTX: finding reconvergence points for \'%s\'...\n", kname.c_str() );
+		kernel_func_info->do_pdom();
+		kernel_func_info->set_pdom();
+    	} 
+	dim3 gridDim = config.grid_dim();
+	dim3 blockDim = config.block_dim();
+		
+	gpgpu_t *gpu = context->get_device()->get_gpgpu();
+	checkpoint *g_checkpoint;
+	g_checkpoint = new checkpoint();
+	class memory_space *global_mem;
+	global_mem = gpu->get_global_memory();
+
+	if(gpu->resume_option ==1 && (grid->get_uid()==gpu->resume_kernel))
+	{
+		
+	    char f1name[2048];
+	    snprintf(f1name,2048,"checkpoint_files/global_mem_%d.txt", grid->get_uid());
+
+	    g_checkpoint->load_global_mem(global_mem, f1name);	
+		for (int i=0;i<gpu->resume_CTA;i++)
+			grid->increment_cta_id();
+	}
+	if(gpu->resume_option==1 && (grid->get_uid()<gpu->resume_kernel))
+	{
+		char f1name[2048];
+	    snprintf(f1name,2048,"checkpoint_files/global_mem_%d.txt", grid->get_uid());
+
+	    g_checkpoint->load_global_mem(global_mem, f1name);	
+		printf("Skipping kernel %d as resuming from kernel %d\n",grid->get_uid(),gpu->resume_kernel );
+		ctx->api->g_cuda_launch_stack.pop_back();
+		return g_last_cudaError = cudaSuccess;
+		
+	}
+	if(gpu->checkpoint_option==1 && (grid->get_uid()>gpu->checkpoint_kernel))
+	{
+		printf("Skipping kernel %d as checkpoint from kernel %d\n",grid->get_uid(),gpu->checkpoint_kernel );
+		ctx->api->g_cuda_launch_stack.pop_back();
+		return g_last_cudaError = cudaSuccess;
+		
+	}
+	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,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;
+}
+
+cudaError_t cudaMallocInternal(void **devPtr, size_t size, gpgpu_context* gpgpu_ctx = NULL)
+{
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
 	if(g_debug_execution >= 3){
 	    announce_call(__my_func__);
     }
@@ -510,7 +964,7 @@ __host__ cudaError_t CUDARTAPI cudaMalloc(void **devPtr, size_t size)
 	*devPtr = context->get_device()->get_gpgpu()->gpu_malloc(size);
 	if(g_debug_execution >= 3){
 		printf("GPGPU-Sim PTX: cudaMallocing %zu bytes starting at 0x%llx..\n",size, (unsigned long long) *devPtr);
-        g_mallocPtr_Size[(unsigned long long)*devPtr] = size;
+        ctx->api->g_mallocPtr_Size[(unsigned long long)*devPtr] = size;
     }
 	if ( *devPtr  ) {
 		return g_last_cudaError = cudaSuccess;
@@ -519,20 +973,270 @@ __host__ cudaError_t CUDARTAPI cudaMalloc(void **devPtr, size_t size)
 	}
 }
 
-__host__ cudaError_t CUDARTAPI cudaMallocHost(void **ptr, size_t size)
+cudaError_t cudaMallocHostInternal(void **ptr, size_t size, gpgpu_context* gpgpu_ctx = NULL)
 {
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
 	if(g_debug_execution >= 3){
 	    announce_call(__my_func__);
     }
-	GPGPUSim_Context();
 	*ptr = malloc(size);
 	if ( *ptr  ) {
 		//track pinned memory size allocated in the host so that same amount of memory is also allocated in GPU.
-		pinned_memory_size[*ptr]=size;
+		ctx->api->pinned_memory_size[*ptr]=size;
+		return g_last_cudaError = cudaSuccess;
+	} else {
+		return g_last_cudaError = cudaErrorMemoryAllocation;
+	}
+}
+
+cudaError_t cudaHostGetDevicePointerInternal(void **pDevice, void *pHost, unsigned int flags, gpgpu_context* gpgpu_ctx = NULL)
+{
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+	//only cpu memory allocation happens in cudaHostAlloc. Linking with device pointer to pinned memory happens here.
+	//TODO: once kernel is executed, the contents in global pointer of GPU must be copied back to CPU host pointer!
+	flags=0;
+	CUctx_st* context = GPGPUSim_Context();
+	gpgpu_t *gpu = context->get_device()->get_gpgpu();
+	std::map<void *, size_t>::const_iterator i = ctx->api->pinned_memory_size.find(pHost);
+	assert(i != ctx->api->pinned_memory_size.end());
+	size_t size = i->second;
+	*pDevice = gpu->gpu_malloc(size);
+	if(g_debug_execution >= 3){
+		printf("GPGPU-Sim PTX: cudaMallocing %zu bytes starting at 0x%llx..\n",size, (unsigned long long) *pDevice);
+        ctx->api->g_mallocPtr_Size[(unsigned long long)*pDevice] = size;
+    }
+	if ( *pDevice  ) {
+		ctx->api->pinned_memory[pHost]=pDevice;
+		//Copy contents in cpu to gpu
+		gpu->memcpy_to_gpu((size_t)*pDevice,pHost,size);
+		return g_last_cudaError = cudaSuccess;
+	} else {
+		return g_last_cudaError = cudaErrorMemoryAllocation;
+	}
+}
+
+cudaError_t cudaGLMapBufferObjectInternal(void** devPtr, GLuint bufferObj, gpgpu_context* gpgpu_ctx = NULL)
+{
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+#ifdef OPENGL_SUPPORT
+	GLint buffer_size=0;
+	CUctx_st* context = GPGPUSim_Context();
+
+	glbmap_entry_t *p = ctx->api->g_glbmap;
+	while ( p && p->m_bufferObj != bufferObj )
+		p = p->m_next;
+	if ( p == NULL ) {
+		glBindBuffer(GL_ARRAY_BUFFER,bufferObj);
+		glGetBufferParameteriv(GL_ARRAY_BUFFER,GL_BUFFER_SIZE,&buffer_size);
+		assert( buffer_size != 0 );
+		*devPtr = context->get_device()->get_gpgpu()->gpu_malloc(buffer_size);
+
+		// create entry and insert to front of list
+		glbmap_entry_t *n = (glbmap_entry_t *) calloc(1,sizeof(glbmap_entry_t));
+		n->m_next = ctx->api->g_glbmap;
+		ctx->api->g_glbmap = n;
+
+		// initialize entry
+		n->m_bufferObj = bufferObj;
+		n->m_devPtr = *devPtr;
+		n->m_size = buffer_size;
+
+		p = n;
+	} else {
+		buffer_size = p->m_size;
+		*devPtr = p->m_devPtr;
+	}
+
+	if ( *devPtr  ) {
+		char *data = (char *) calloc(p->m_size,1);
+		glGetBufferSubData(GL_ARRAY_BUFFER,0,buffer_size,data);
+		memcpy_to_gpu( (size_t) *devPtr, data, buffer_size );
+		free(data);
+		printf("GPGPU-Sim PTX: cudaGLMapBufferObject %zu bytes starting at 0x%llx..\n", (size_t)buffer_size,
+				(unsigned long long) *devPtr);
 		return g_last_cudaError = cudaSuccess;
 	} else {
 		return g_last_cudaError = cudaErrorMemoryAllocation;
 	}
+
+	return g_last_cudaError = cudaSuccess;
+#else
+	fflush(stdout);
+	fflush(stderr);
+	printf("GPGPU-Sim PTX: GPGPU-Sim support for OpenGL integration disabled -- exiting\n");
+	fflush(stdout);
+	exit(50);
+#endif
+}
+
+#if CUDART_VERSION >= 6050
+CUresult
+cuLinkAddFileInternal(CUlinkState state, CUjitInputType type, const char *path,
+    unsigned int numOptions, CUjit_option *options, void **optionValues, gpgpu_context* gpgpu_ctx = NULL)
+{
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+    static bool addedFile = false;
+    if (addedFile){
+        printf("GPGPU-Sim PTX: ERROR: cuLinkAddFile does not support multiple files\n");
+        abort();
+    }
+
+    //blocking
+    assert(type==CU_JIT_INPUT_PTX);
+	CUctx_st *context = GPGPUSim_Context();
+    char *file = getenv("PTX_JIT_PATH");
+    if(file==NULL){
+        printf("GPGPU-Sim PTX: ERROR: PTX_JIT_PATH has not been set\n");
+        abort();
+    }
+    strcat(file,"/");
+    strcat(file,path);
+	symbol_table *symtab = ctx->gpgpu_ptx_sim_load_ptx_from_filename( file );
+    std::string fname(path);
+    ctx->api->name_symtab[fname] = symtab;
+    context->add_binary(symtab, 1);
+    ctx->api->load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF,context->get_device()->get_gpgpu());
+    ctx->api->load_constants(symtab,STATIC_ALLOC_LIMIT,context->get_device()->get_gpgpu());
+    addedFile = true;
+	return CUDA_SUCCESS;
+}
+#endif
+
+#if (CUDART_VERSION >= 2010)
+
+cudaError_t cudaHostAllocInternal(void **pHost,  size_t bytes, unsigned int flags, gpgpu_context* gpgpu_ctx = NULL)
+{
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+	*pHost = malloc(bytes);
+	//need to track the size allocated so that cudaHostGetDevicePointer() can function properly.
+	//TODO: vary this function behavior based on flags value (following nvidia documentation)
+	ctx->api->pinned_memory_size[*pHost]=bytes;
+	if( *pHost )
+		return g_last_cudaError = cudaSuccess;
+	else
+		return g_last_cudaError = cudaErrorMemoryAllocation;
+}
+
+#endif
+
+size_t getMaxThreadsPerBlock(struct cudaFuncAttributes *attr, gpgpu_context *ctx) {
+  _cuda_device_id *dev = ctx->GPGPUSim_Init();
+  struct cudaDeviceProp prop;
+
+  prop = *dev->get_prop();
+
+  size_t max = prop.maxThreadsPerBlock;
+
+  if (attr->numRegs && (prop.regsPerBlock / attr->numRegs) < max)
+    max = prop.regsPerBlock / attr->numRegs;
+
+  if (attr->sharedSizeBytes && (prop.sharedMemPerBlock / attr->sharedSizeBytes) < max)
+      max = prop.sharedMemPerBlock / attr->sharedSizeBytes;
+
+  return max;
+}
+
+cudaError_t CUDARTAPI cudaFuncGetAttributesInternal(struct cudaFuncAttributes *attr, const char *hostFun, gpgpu_context* gpgpu_ctx = NULL )
+{
+    gpgpu_context *ctx;
+    if (gpgpu_ctx){
+	ctx = gpgpu_ctx;
+    } else {
+	ctx = GPGPU_Context();
+    }
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+	CUctx_st *context = GPGPUSim_Context();
+	function_info *entry = context->get_kernel(hostFun);
+	if( entry ) {
+		const struct gpgpu_ptx_sim_info *kinfo = entry->get_kernel_info();
+		attr->sharedSizeBytes = kinfo->smem;
+		attr->constSizeBytes  = kinfo->cmem;
+		attr->localSizeBytes  = kinfo->lmem;
+		attr->numRegs         = kinfo->regs;
+		if(kinfo->maxthreads > 0)
+		  attr->maxThreadsPerBlock = kinfo->maxthreads;
+		else
+		  attr->maxThreadsPerBlock = getMaxThreadsPerBlock(attr, ctx);
+#if CUDART_VERSION >= 3000
+		attr->ptxVersion      = kinfo->ptx_version;
+		attr->binaryVersion   = kinfo->sm_target;
+#endif
+	}
+	return g_last_cudaError = cudaSuccess;
+}
+
+
+/*******************************************************************************
+ *                                                                              *
+ *                                                                              *
+ *                                                                              *
+ *******************************************************************************/
+
+extern "C" {
+
+/*******************************************************************************
+ *                                                                              *
+ *                                                                              *
+ *                                                                              *
+ *******************************************************************************/
+cudaError_t cudaPeekAtLastError(void)
+{
+	return g_last_cudaError;
+}
+
+__host__ cudaError_t CUDARTAPI cudaMalloc(void **devPtr, size_t size) 
+{
+    return cudaMallocInternal(devPtr, size);
+}
+
+__host__ cudaError_t CUDARTAPI cudaMallocHost(void **ptr, size_t size)
+{
+    return cudaMallocHostInternal(ptr, size);
 }
 __host__ cudaError_t CUDARTAPI cudaMallocPitch(void **devPtr, size_t *pitch, size_t width, size_t height)
 {
@@ -617,21 +1321,21 @@ __host__ cudaError_t CUDARTAPI cudaMemcpy(void *dst, const void *src, size_t cou
 	if(g_debug_execution >= 3)
 		printf("GPGPU-Sim PTX: cudaMemcpy(): devPtr = %p\n", dst);
 	if( kind == cudaMemcpyHostToDevice )
-		g_stream_manager->push( stream_operation(src,(size_t)dst,count,0) );
+		g_stream_manager()->push( stream_operation(src,(size_t)dst,count,0) );
 	else if( kind == cudaMemcpyDeviceToHost )
-		g_stream_manager->push( stream_operation((size_t)src,dst,count,0) );
+		g_stream_manager()->push( stream_operation((size_t)src,dst,count,0) );
 	else if( kind == cudaMemcpyDeviceToDevice )
-		g_stream_manager->push( stream_operation((size_t)src,(size_t)dst,count,0) );
+		g_stream_manager()->push( stream_operation((size_t)src,(size_t)dst,count,0) );
 	else if ( kind == cudaMemcpyDefault ) {
 		if ((size_t)src >= GLOBAL_HEAP_START) {
 			if ((size_t)dst >= GLOBAL_HEAP_START)
-				g_stream_manager->push( stream_operation((size_t)src,(size_t)dst,count,0) ); // device to device
+				g_stream_manager()->push( stream_operation((size_t)src,(size_t)dst,count,0) ); // device to device
 			else
-				g_stream_manager->push( stream_operation((size_t)src,dst,count,0) ); // device to host
+				g_stream_manager()->push( stream_operation((size_t)src,dst,count,0) ); // device to host
 		}
 		else {
 			if ((size_t)dst >= GLOBAL_HEAP_START)
-				g_stream_manager->push( stream_operation(src,(size_t)dst,count,0) );
+				g_stream_manager()->push( stream_operation(src,(size_t)dst,count,0) );
 			else {
 				printf("GPGPU-Sim PTX: cudaMemcpy - ERROR : unsupported transfer: host to host\n");
 				abort();
@@ -773,7 +1477,7 @@ __host__ cudaError_t CUDARTAPI cudaMemcpyToSymbol(const char *symbol, const void
 	assert(kind == cudaMemcpyHostToDevice);
 	printf("GPGPU-Sim PTX: cudaMemcpyToSymbol: symbol = %p\n", symbol);
 	//stream_operation( const char *symbol, const void *src, size_t count, size_t offset )
-	g_stream_manager->push( stream_operation(src,symbol,count,offset,0) );
+	g_stream_manager()->push( stream_operation(src,symbol,count,offset,0) );
 	//gpgpu_ptx_sim_memcpy_symbol(symbol,src,count,offset,1,context->get_device()->get_gpgpu());
 	return g_last_cudaError = cudaSuccess;
 }
@@ -787,7 +1491,7 @@ __host__ cudaError_t CUDARTAPI cudaMemcpyFromSymbol(void *dst, const char *symbo
 	//CUctx_st *context = GPGPUSim_Context();
 	assert(kind == cudaMemcpyDeviceToHost);
 	printf("GPGPU-Sim PTX: cudaMemcpyFromSymbol: symbol = %p\n", symbol);
-	g_stream_manager->push( stream_operation(symbol,dst,count,offset,0) );
+	g_stream_manager()->push( stream_operation(symbol,dst,count,offset,0) );
 	//gpgpu_ptx_sim_memcpy_symbol(symbol,dst,count,offset,0,context->get_device()->get_gpgpu());
 	return g_last_cudaError = cudaSuccess;
 }
@@ -816,9 +1520,9 @@ __host__ cudaError_t CUDARTAPI cudaMemcpyAsync(void *dst, const void *src, size_
     }
 	struct CUstream_st *s = (struct CUstream_st *)stream;
 	switch( kind ) {
-	case cudaMemcpyHostToDevice: g_stream_manager->push( stream_operation(src,(size_t)dst,count,s) ); break;
-	case cudaMemcpyDeviceToHost: g_stream_manager->push( stream_operation((size_t)src,dst,count,s) ); break;
-	case cudaMemcpyDeviceToDevice: g_stream_manager->push( stream_operation((size_t)src,(size_t)dst,count,s) ); break;
+	case cudaMemcpyHostToDevice: g_stream_manager()->push( stream_operation(src,(size_t)dst,count,s) ); break;
+	case cudaMemcpyDeviceToHost: g_stream_manager()->push( stream_operation((size_t)src,dst,count,s) ); break;
+	case cudaMemcpyDeviceToDevice: g_stream_manager()->push( stream_operation((size_t)src,(size_t)dst,count,s) ); break;
 	default:
 		abort();
 	}
@@ -875,105 +1579,6 @@ __host__ cudaError_t CUDARTAPI cudaMemcpy2DFromArrayAsync(void *dst, size_t dpit
 	return g_last_cudaError = cudaErrorUnknown;
 }
 
-
-
-/*******************************************************************************
- *                                                                              *
- *                                                                              *
- *                                                                              *
- *******************************************************************************/
-
-__host__ cudaError_t CUDARTAPI cudaMemset(void *mem, int c, size_t count)
-{
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-	CUctx_st *context = GPGPUSim_Context();
-	gpgpu_t *gpu = context->get_device()->get_gpgpu();
-	gpu->gpu_memset((size_t)mem, c, count);
-	return g_last_cudaError = cudaSuccess;
-}
-
-//memset operation is done but i think its not async?
-__host__ cudaError_t CUDARTAPI cudaMemsetAsync(void *mem, int c, size_t count, 	cudaStream_t stream=0)
-{
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-	printf("GPGPU-Sim PTX: WARNING: Asynchronous memset not supported (%s)\n", __my_func__);
-	CUctx_st *context = GPGPUSim_Context();
-	gpgpu_t *gpu = context->get_device()->get_gpgpu();
-	gpu->gpu_memset((size_t)mem, c, count);
-	return g_last_cudaError = cudaSuccess;
-}
-
-__host__ cudaError_t CUDARTAPI cudaMemset2D(void *mem, size_t pitch, int c, size_t width, size_t height)
-{
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-	cuda_not_implemented(__my_func__,__LINE__);
-	return g_last_cudaError = cudaErrorUnknown;
-}
-
-
-
-/*******************************************************************************
- *                                                                              *
- *                                                                              *
- *                                                                              *
- *******************************************************************************/
-
-__host__ cudaError_t CUDARTAPI cudaGetSymbolAddress(void **devPtr, const char *symbol)
-{
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-	cuda_not_implemented(__my_func__,__LINE__);
-	return g_last_cudaError = cudaErrorUnknown;
-}
-
-
-__host__ cudaError_t CUDARTAPI cudaGetSymbolSize(size_t *size, const char *symbol)
-{
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-	cuda_not_implemented(__my_func__,__LINE__);
-	return g_last_cudaError = cudaErrorUnknown;
-}
-
-
-
-/*******************************************************************************
- *                                                                              *
- *                                                                              *
- *                                                                              *
- *******************************************************************************/
-__host__ cudaError_t CUDARTAPI cudaGetDeviceCount(int *count)
-{
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-	_cuda_device_id *dev = GPGPUSim_Init();
-	*count = dev->num_devices();
-	return g_last_cudaError = cudaSuccess;
-}
-
-__host__ cudaError_t CUDARTAPI cudaGetDeviceProperties(struct cudaDeviceProp *prop, int device)
-{
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-	_cuda_device_id *dev = GPGPUSim_Init();
-	if (device <= dev->num_devices() )  {
-		*prop= *dev->get_prop();
-		return g_last_cudaError = cudaSuccess;
-	} else {
-		return g_last_cudaError = cudaErrorInvalidDevice;
-	}
-}
-
 #if (CUDART_VERSION >= 8000)
 cudaError_t CUDARTAPI cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(int* numBlocks, const char *hostFunc, int blockSize, size_t dynamicSMemSize, unsigned int flags)
 {
@@ -1003,13 +1608,38 @@ cudaError_t CUDARTAPI cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(int
 #endif
 
 #if (CUDART_VERSION > 5000)
-__host__ cudaError_t CUDARTAPI cudaDeviceGetAttribute(int *value, enum cudaDeviceAttr attr, int device)
+
+/*******************************************************************************
+ *                                                                              *
+ *                                                                              *
+ *                                                                              *
+ *******************************************************************************/
+__host__ cudaError_t CUDARTAPI cudaMemset(void *mem, int c, size_t count)
 {
 	if(g_debug_execution >= 3){
 	    announce_call(__my_func__);
     }
+	CUctx_st *context = GPGPUSim_Context();
+	gpgpu_t *gpu = context->get_device()->get_gpgpu();
+	gpu->gpu_memset((size_t)mem, c, count);
+	return g_last_cudaError = cudaSuccess;
+}
+
+__host__ cudaError_t CUDARTAPI cudaDeviceGetAttributeInternal(int *value, enum cudaDeviceAttr attr, int device, gpgpu_context* gpgpu_ctx = NULL)
+{
+	    gpgpu_context *ctx;
+		if (gpgpu_ctx){
+		ctx = gpgpu_ctx;
+		} else {
+		ctx = GPGPU_Context();
+		}
+		if(g_debug_execution >= 3){
+			announce_call(__my_func__);
+		}
+
         const struct cudaDeviceProp *prop;
-        _cuda_device_id *dev = GPGPUSim_Init();
+        _cuda_device_id *dev = ctx->GPGPUSim_Init();
+
         if (device <= dev->num_devices() )  {
                 prop = dev->get_prop();
                 switch (attr) {
@@ -1193,82 +1823,98 @@ __host__ cudaError_t CUDARTAPI cudaDeviceGetAttribute(int *value, enum cudaDevic
                 return g_last_cudaError = cudaErrorInvalidDevice;
         }
 }
-#endif
 
-__host__ cudaError_t CUDARTAPI cudaChooseDevice(int *device, const struct cudaDeviceProp *prop)
+//memset operation is done but i think its not async?
+__host__ cudaError_t CUDARTAPI cudaMemsetAsync(void *mem, int c, size_t count, 	cudaStream_t stream=0)
 {
 	if(g_debug_execution >= 3){
 	    announce_call(__my_func__);
     }
-	_cuda_device_id *dev = GPGPUSim_Init();
-	*device = dev->get_id();
+	printf("GPGPU-Sim PTX: WARNING: Asynchronous memset not supported (%s)\n", __my_func__);
+	CUctx_st *context = GPGPUSim_Context();
+	gpgpu_t *gpu = context->get_device()->get_gpgpu();
+	gpu->gpu_memset((size_t)mem, c, count);
 	return g_last_cudaError = cudaSuccess;
 }
 
-__host__ cudaError_t CUDARTAPI cudaSetDevice(int device)
+__host__ cudaError_t CUDARTAPI cudaMemset2D(void *mem, size_t pitch, int c, size_t width, size_t height)
 {
 	if(g_debug_execution >= 3){
 	    announce_call(__my_func__);
     }
-	//set the active device to run cuda
-	if ( device <= GPGPUSim_Init()->num_devices() ) {
-		g_active_device = device;
-		return g_last_cudaError = cudaSuccess;
-	} else {
-		return g_last_cudaError = cudaErrorInvalidDevice;
-	}
+	cuda_not_implemented(__my_func__,__LINE__);
+	return g_last_cudaError = cudaErrorUnknown;
 }
 
-__host__ cudaError_t CUDARTAPI cudaGetDevice(int *device)
+
+
+/*******************************************************************************
+ *                                                                              *
+ *                                                                              *
+ *                                                                              *
+ *******************************************************************************/
+
+__host__ cudaError_t CUDARTAPI cudaGetSymbolAddress(void **devPtr, const char *symbol)
 {
 	if(g_debug_execution >= 3){
 	    announce_call(__my_func__);
     }
-	*device = g_active_device;
-	return g_last_cudaError = cudaSuccess;
+	cuda_not_implemented(__my_func__,__LINE__);
+	return g_last_cudaError = cudaErrorUnknown;
 }
 
-__host__ cudaError_t CUDARTAPI cudaDeviceGetLimit ( size_t* pValue, cudaLimit limit )
+
+__host__ cudaError_t CUDARTAPI cudaGetSymbolSize(size_t *size, const char *symbol)
 {
 	if(g_debug_execution >= 3){
-            announce_call(__my_func__);
-   	 }
-        _cuda_device_id *dev = GPGPUSim_Init();
-	const struct cudaDeviceProp *prop = dev->get_prop();
-	const gpgpu_sim_config& config=dev->get_gpgpu()->get_config();
-        switch(limit) {
-        case 0:  // cudaLimitStackSize
-        	*pValue=config.stack_limit();
-                break;
-        case 2:  // cudaLimitMallocHeapSize
-                *pValue=config.heap_limit();
-                break;
-#if (CUDART_VERSION > 5050)
-        case 3: // cudaLimitDevRuntimeSyncDepth
-		if(prop->major > 2){
-			*pValue=config.sync_depth_limit();
-	                break;
-		}
-		else{
-			printf("ERROR:Limit %s is not supported on this architecture \n",limit);
-			abort();
-		}
-        case 4: // cudaLimitDevRuntimePendingLaunchCount
-		if(prop->major > 2){
-     	        	*pValue=config.pending_launch_count_limit();
-	                break;
-		}
-		else{
-			printf("ERROR:Limit %s is not supported on this architecture \n",limit);
-			abort();
-		}
+	    announce_call(__my_func__);
+    }
+	cuda_not_implemented(__my_func__,__LINE__);
+	return g_last_cudaError = cudaErrorUnknown;
+}
+
+
+
+/*******************************************************************************
+ *                                                                              *
+ *                                                                              *
+ *                                                                              *
+ *******************************************************************************/
+__host__ cudaError_t CUDARTAPI cudaGetDeviceCount(int *count)
+{
+    return cudaGetDeviceCountInternal(count);
+}
+
+__host__ cudaError_t CUDARTAPI cudaGetDeviceProperties(struct cudaDeviceProp *prop, int device)
+{
+    return cudaGetDevicePropertiesInternal(prop, device);
+}
+
+#if (CUDART_VERSION > 5000)
+__host__ cudaError_t CUDARTAPI cudaDeviceGetAttribute(int *value, enum cudaDeviceAttr attr, int device)
+{
+    return cudaDeviceGetAttributeInternal(value, attr, device);
+}
 #endif
-        default:
-                        printf("ERROR:Limit %s unimplemented \n",limit);
-                        abort();
-        }
-	return g_last_cudaError = cudaSuccess;
 
+__host__ cudaError_t CUDARTAPI cudaChooseDevice(int *device, const struct cudaDeviceProp *prop)
+{
+    return cudaChooseDeviceInternal(device, prop);
+}
+
+__host__ cudaError_t CUDARTAPI cudaSetDevice(int device)
+{
+    return cudaSetDeviceInternal(device);
+}
+
+__host__ cudaError_t CUDARTAPI cudaGetDevice(int *device)
+{
+    return cudaGetDeviceInternal(device);
+}
+
+__host__ cudaError_t CUDARTAPI cudaDeviceGetLimit( size_t* pValue, cudaLimit limit )
+{
+    return cudaDeviceGetLimitInternal( pValue, limit );
 }
 
 __host__ cudaError_t CUDARTAPI cudaStreamGetPriority ( cudaStream_t hStream, int* priority )
@@ -1466,131 +2112,35 @@ __host__ const char* CUDARTAPI cudaGetErrorString(cudaError_t error)
 	return strdup(buf);
 }
 
-__host__ cudaError_t CUDARTAPI cudaConfigureCall(dim3 gridDim, dim3 blockDim, size_t sharedMem, cudaStream_t stream)
-{
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-	struct CUstream_st *s = (struct CUstream_st *)stream;
-	g_cuda_launch_stack.push_back( kernel_config(gridDim,blockDim,sharedMem,s) );
-	return g_last_cudaError = cudaSuccess;
-}
-
 __host__ cudaError_t CUDARTAPI cudaSetupArgument(const void *arg, size_t size, size_t offset)
 {
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-	gpgpusim_ptx_assert( !g_cuda_launch_stack.empty(), "empty launch stack" );
-	kernel_config &config = g_cuda_launch_stack.back();
-	config.set_arg(arg,size,offset);
-	printf("GPGPU-Sim PTX: Setting up arguments for %zu bytes starting at 0x%llx..\n",size, (unsigned long long) arg);
-
-	return g_last_cudaError = cudaSuccess;
+    return cudaSetupArgumentInternal(arg, size, offset);
 }
 
 
 __host__ cudaError_t CUDARTAPI cudaLaunch( const char *hostFun )
 {
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-	CUctx_st* context = GPGPUSim_Context();
-	char *mode = getenv("PTX_SIM_MODE_FUNC");
-	if( mode )
-		sscanf(mode,"%u", &g_ptx_sim_mode);
-	gpgpusim_ptx_assert( !g_cuda_launch_stack.empty(), "empty launch stack" );
-	kernel_config config = g_cuda_launch_stack.back();
-	{
-		dim3 gridDim = config.grid_dim();
-		dim3 blockDim = config.block_dim();
-		//if (gridDim.x * gridDim.y * gridDim.z == 0 || blockDim.x * blockDim.y * blockDim.z == 0)
-		//{
-			//can't launch
-		//	printf("can't launch a empty kernel\n");
-		//	g_cuda_launch_stack.pop_back();
-		//	return g_last_cudaError = cudaErrorInvalidConfiguration;
-		//}
-	}
-	struct CUstream_st *stream = config.get_stream();
-	if(g_stream_manager->is_blocking())
-		stream = 0;
-
-	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 );
-	kernel_info_t *grid = 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();
-	function_info *kernel_func_info = grid->entry();
-	if (kernel_func_info->is_pdom_set()) {
-    		printf("GPGPU-Sim PTX: PDOM analysis already done for %s \n", kname.c_str() );
-    	} else {
-    		printf("GPGPU-Sim PTX: finding reconvergence points for \'%s\'...\n", kname.c_str() );
-		kernel_func_info->do_pdom();
-		kernel_func_info->set_pdom();
-    	} 
-	dim3 gridDim = config.grid_dim();
-	dim3 blockDim = config.block_dim();
-		
-	gpgpu_t *gpu = context->get_device()->get_gpgpu();
-	checkpoint *g_checkpoint;
-	g_checkpoint = new checkpoint();
-	class memory_space *global_mem;
-	global_mem = gpu->get_global_memory();
-
-	if(gpu->resume_option ==1 && (grid->get_uid()==gpu->resume_kernel))
-	{
-		
-	    char f1name[2048];
-	    snprintf(f1name,2048,"checkpoint_files/global_mem_%d.txt", grid->get_uid());
-
-	    g_checkpoint->load_global_mem(global_mem, f1name);	
-		for (int i=0;i<gpu->resume_CTA;i++)
-			grid->increment_cta_id();
-	}
-	if(gpu->resume_option==1 && (grid->get_uid()<gpu->resume_kernel))
-	{
-		char f1name[2048];
-	    snprintf(f1name,2048,"checkpoint_files/global_mem_%d.txt", grid->get_uid());
-
-	    g_checkpoint->load_global_mem(global_mem, f1name);	
-		printf("Skipping kernel %d as resuming from kernel %d\n",grid->get_uid(),gpu->resume_kernel );
-		g_cuda_launch_stack.pop_back();
-		return g_last_cudaError = cudaSuccess;
-		
-	}
-	if(gpu->checkpoint_option==1 && (grid->get_uid()>gpu->checkpoint_kernel))
-	{
-		printf("Skipping kernel %d as checkpoint from kernel %d\n",grid->get_uid(),gpu->checkpoint_kernel );
-		g_cuda_launch_stack.pop_back();
-		return g_last_cudaError = cudaSuccess;
-		
-	}
-	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);
-	g_stream_manager->push(op);
-	g_cuda_launch_stack.pop_back();
-	return g_last_cudaError = cudaSuccess;
+    return cudaLaunchInternal( hostFun );
 }
 
 __host__ cudaError_t CUDARTAPI cudaLaunchKernel ( const char* hostFun, dim3 gridDim, dim3 blockDim, const void** args, size_t sharedMem, cudaStream_t stream )
 {
 
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    	}
-        CUctx_st *context = GPGPUSim_Context();
-        function_info *entry = context->get_kernel(hostFun);
-    
-	cudaConfigureCall(gridDim, blockDim, sharedMem, stream);
-    	for(unsigned i = 0; i < entry->num_args(); i++){
-        	std::pair<size_t, unsigned> p = entry->get_param_config(i);
-        	cudaSetupArgument(args[i], p.first, p.second);
-    	}  
+    if(g_debug_execution >= 3){
+        announce_call(__my_func__);
+    }
+    CUctx_st *context = GPGPUSim_Context();
+    function_info *entry = context->get_kernel(hostFun);
+#if CUDART_VERSION < 10000
+    cudaConfigureCallInternal(gridDim, blockDim, sharedMem, stream);
+#endif
+    for(unsigned i = 0; i < entry->num_args(); i++){
+        std::pair<size_t, unsigned> p = entry->get_param_config(i);
+        cudaSetupArgumentInternal(args[i], p.first, p.second);
+    }
 
-	cudaLaunch(hostFun);
-	return g_last_cudaError = cudaSuccess;
+    cudaLaunchInternal(hostFun);
+    return g_last_cudaError = cudaSuccess;
 }
 
 
@@ -1608,7 +2158,7 @@ __host__ cudaError_t CUDARTAPI cudaStreamCreate(cudaStream_t *stream)
 	printf("GPGPU-Sim PTX: cudaStreamCreate\n");
 #if (CUDART_VERSION >= 3000)
 	*stream = new struct CUstream_st();
-	g_stream_manager->add_stream(*stream);
+	g_stream_manager()->add_stream(*stream);
 #else
 	*stream = 0;
 	printf("GPGPU-Sim PTX: WARNING: Asynchronous kernel execution not supported (%s)\n", __my_func__);
@@ -1647,7 +2197,7 @@ __host__ cudaError_t CUDARTAPI cudaStreamDestroy(cudaStream_t stream)
 	//per-stream synchronization required for application using external libraries without explicit synchronization in the code to 
 	//avoid the stream_manager from spinning forever to destroy non-empty streams without making any forward progress. 
 	stream->synchronize();
-	g_stream_manager->destroy_stream(stream);
+	g_stream_manager()->destroy_stream(stream);
 #endif
 	return g_last_cudaError = cudaSuccess;
 }
@@ -1727,7 +2277,7 @@ __host__ cudaError_t CUDARTAPI cudaEventRecord(cudaEvent_t event, cudaStream_t s
 	if( !e ) return g_last_cudaError = cudaErrorUnknown;
 	struct CUstream_st *s = (struct CUstream_st *)stream;
 	stream_operation op(e,s);
-	g_stream_manager->push(op);
+	g_stream_manager()->push(op);
 	return g_last_cudaError = cudaSuccess;
 }
 
@@ -1743,11 +2293,11 @@ __host__ cudaError_t CUDARTAPI cudaStreamWaitEvent(cudaStream_t stream, cudaEven
       return g_last_cudaError = cudaSuccess;
    }
    if (!stream){
-      g_stream_manager->pushCudaStreamWaitEventToAllStreams(e, flags);
+      g_stream_manager()->pushCudaStreamWaitEventToAllStreams(e, flags);
    } else {
 	   struct CUstream_st *s = (struct CUstream_st *)stream;
 	   stream_operation op(s,e,flags);
-	   g_stream_manager->push(op);
+	   g_stream_manager()->push(op);
    }
 	return g_last_cudaError = cudaSuccess;
 }
@@ -1898,178 +2448,8 @@ __host__ cudaError_t CUDARTAPI cudaGetExportTable(const void **ppExportTable, co
 
 //#include "../../cuobjdump_to_ptxplus/cuobjdump_parser.h"
 
-enum cuobjdumpSectionType {
-	PTXSECTION=0,
-	ELFSECTION
-};
-
-
-class cuobjdumpSection {
-public:
-	//Constructor
-	cuobjdumpSection() {
-		arch = 0;
-		identifier = "";
-	}
-	virtual ~cuobjdumpSection() {}
-	unsigned getArch() {return arch;}
-	void setArch(unsigned a) {arch = a;}
-	std::string getIdentifier() {return identifier;}
-	void setIdentifier(std::string i) {identifier = i;}
-	virtual void print(){std::cout << "cuobjdump Section: unknown type" << std::endl;}
-private:
-	unsigned arch;
-	std::string identifier;
-};
-
-class cuobjdumpELFSection : public cuobjdumpSection
-{
-public:
-	cuobjdumpELFSection() {}
-	virtual ~cuobjdumpELFSection() {
-		elffilename = "";
-		sassfilename = "";
-	}
-	std::string getELFfilename() {return elffilename;}
-	void setELFfilename(std::string f) {elffilename = f;}
-	std::string getSASSfilename() {return sassfilename;}
-	void setSASSfilename(std::string f) {sassfilename = f;}
-	virtual void print() {
-		std::cout << "ELF Section:" << std::endl;
-		std::cout << "arch: sm_" << getArch() << std::endl;
-		std::cout << "identifier: " << getIdentifier() << std::endl;
-		std::cout << "elf filename: " << getELFfilename() << std::endl;
-		std::cout << "sass filename: " << getSASSfilename() << std::endl;
-		std::cout << std::endl;
-	}
-private:
-	std::string elffilename;
-	std::string sassfilename;
-};
-
-class cuobjdumpPTXSection : public cuobjdumpSection
-{
-public:
-	cuobjdumpPTXSection(){
-		ptxfilename = "";
-	}
-	std::string getPTXfilename() {return ptxfilename;}
-	void setPTXfilename(std::string f) {ptxfilename = f;}
-	virtual void print() {
-		std::cout << "PTX Section:" << std::endl;
-		std::cout << "arch: sm_" << getArch() << std::endl;
-		std::cout << "identifier: " << getIdentifier() << std::endl;
-		std::cout << "ptx filename: " << getPTXfilename() << std::endl;
-		std::cout << std::endl;
-	}
-private:
-	std::string ptxfilename;
-};
-
-std::list<cuobjdumpSection*> cuobjdumpSectionList;
-std::list<cuobjdumpSection*> libSectionList;
-
-// sectiontype: 0 for ptx, 1 for elf
-void addCuobjdumpSection(int sectiontype){
-	if (sectiontype)
-		cuobjdumpSectionList.push_front(new cuobjdumpELFSection());
-	else
-		cuobjdumpSectionList.push_front(new cuobjdumpPTXSection());
-	printf("## Adding new section %s\n", sectiontype?"ELF":"PTX");
-}
-
-void setCuobjdumparch(const char* arch){
-	unsigned archnum;
-	sscanf(arch, "sm_%u", &archnum);
-	assert (archnum && "cannot have sm_0");
-	printf("Adding arch: %s\n", arch);
-	cuobjdumpSectionList.front()->setArch(archnum);
-}
-
-void setCuobjdumpidentifier(const char* identifier){
-	printf("Adding identifier: %s\n", identifier);
-	cuobjdumpSectionList.front()->setIdentifier(identifier);
-}
-
-void setCuobjdumpptxfilename(const char* filename){
-	printf("Adding ptx filename: %s\n", filename);
-	cuobjdumpSection* x = cuobjdumpSectionList.front();
-	if (dynamic_cast<cuobjdumpPTXSection*>(x) == NULL){
-		assert (0 && "You shouldn't be trying to add a ptxfilename to an elf section");
-	}
-	(dynamic_cast<cuobjdumpPTXSection*>(x))->setPTXfilename(filename);
-}
-
-void setCuobjdumpelffilename(const char* filename){
-	if (dynamic_cast<cuobjdumpELFSection*>(cuobjdumpSectionList.front()) == NULL){
-		assert (0 && "You shouldn't be trying to add a elffilename to an ptx section");
-	}
-	(dynamic_cast<cuobjdumpELFSection*>(cuobjdumpSectionList.front()))->setELFfilename(filename);
-}
-
-void setCuobjdumpsassfilename(const char* filename){
-	if (dynamic_cast<cuobjdumpELFSection*>(cuobjdumpSectionList.front()) == NULL){
-		assert (0 && "You shouldn't be trying to add a sassfilename to an ptx section");
-	}
-	(dynamic_cast<cuobjdumpELFSection*>(cuobjdumpSectionList.front()))->setSASSfilename(filename);
-}
-extern int cuobjdump_parse();
-extern FILE *cuobjdump_in;
-
-//! Return the executable file of the process containing the PTX/SASS code 
-//!
-//! This Function returns the executable file ran by the process.  This
-//! executable is supposed to contain the PTX/SASS code.  It provides workaround
-//! for processes running on valgrind by dereferencing /proc/<pid>/exe within the
-//! GPGPU-Sim process before calling cuobjdump to extract PTX/SASS.  This is
-//! needed because valgrind uses x86 emulation to detect memory leak.  Other
-//! processes (e.g. cuobjdump) reading /proc/<pid>/exe will see the emulator
-//! executable instead of the application binary.  
-//! 
-std::string get_app_binary(){
-   char self_exe_path[1025];
-#ifdef __APPLE__
-   uint32_t size = sizeof(self_exe_path);
-   if( _NSGetExecutablePath(self_exe_path,&size) != 0 ) {
-	   printf("GPGPU-Sim ** ERROR: _NSGetExecutablePath input buffer too small\n");
-	   exit(1);
-   }
-#else
-   std::stringstream exec_link;
-   exec_link << "/proc/self/exe";
-
-   ssize_t path_length = readlink(exec_link.str().c_str(), self_exe_path, 1024); 
-   assert(path_length != -1); 
-   self_exe_path[path_length] = '\0'; 
-#endif
-
-   printf("self exe links to: %s\n", self_exe_path); 
-   return self_exe_path; 
-}
-
-//above func gives abs path whereas this give just the name of application.
-char* get_app_binary_name(std::string abs_path){
-   char *self_exe_path;
-#ifdef __APPLE__
-   //TODO: get apple device and check the result.
-   printf("WARNING: not tested for Apple-mac devices \n");
-   abort();
-#else
-   char* buf = strdup(abs_path.c_str());
-   char *token = strtok(buf, "/");
-   while(token !=NULL){
-	self_exe_path = token;
-	token = strtok(NULL,"/");
-   }
-#endif
-   self_exe_path = strtok(self_exe_path, ".");
-   printf("self exe links to: %s\n", self_exe_path);
-   return self_exe_path;
-}
-
 //extracts all ptx files from binary and dumps into prog_name.unique_no.sm_<>.ptx files
-void extract_ptx_files_using_cuobjdump(){
-    extern bool g_cdp_enabled;
+void cuda_runtime_api::extract_ptx_files_using_cuobjdump(CUctx_st *context){
     char command[1000];
     char *pytorch_bin = getenv("PYTORCH_BIN");
     std::string app_binary = get_app_binary(); 
@@ -2090,7 +2470,7 @@ void extract_ptx_files_using_cuobjdump(){
         printf("WARNING: Failed to execute cuobjdump to get list of ptx files \n");
         exit(0);
     }   
-    if(!g_cdp_enabled) {
+    if(!gpgpu_ctx->device_runtime->g_cdp_enabled) {
         //based on the list above, dump ptx files individually. Format of dumped ptx file is prog_name.unique_no.sm_<>.ptx
 
        std::ifstream infile(ptx_list_file_name);
@@ -2105,11 +2485,11 @@ void extract_ptx_files_using_cuobjdump(){
                 printf("ERROR: command: %s failed \n",command);
                 exit(0);
             }
-            no_of_ptx++;
+            context->no_of_ptx++;
        }
     }
 
-	 if(!no_of_ptx){
+	 if(!context->no_of_ptx){
 	 	printf("WARNING: Number of ptx in the executable file are 0. One of the reasons might be\n");
 	 	printf("\t1. CDP is enabled\n");
 	 	printf("\t2. When using PyTorch, PYTORCH_BIN is not set correctly\n");
@@ -2137,28 +2517,6 @@ void extract_ptx_files_using_cuobjdump(){
 
 }
 
-static int get_app_cuda_version() {
-    int app_cuda_version = 0;
-    char fname[1024];
-    snprintf(fname,1024,"_app_cuda_version_XXXXXX");
-    int fd=mkstemp(fname);
-    close(fd);
-    std::string app_cuda_version_command = "ldd " + get_app_binary() + " | grep libcudart.so | sed  's/.*libcudart.so.\\(.*\\) =>.*/\\1/' > " + fname;
-    system(app_cuda_version_command.c_str());
-    FILE * cmd = fopen(fname, "r");
-    char buf[256];
-    while (fgets(buf, sizeof(buf), cmd) != 0) {
-        std::cout << buf;
-        app_cuda_version = atoi(buf);
-    }
-    fclose(cmd);
-    if ( app_cuda_version == 0 ) {
-        printf( "Error - Cannot detect the app's CUDA version.\n" );
-        exit(1);
-    }
-    return app_cuda_version;
-}
-
 
 //! Call cuobjdump to extract everything (-elf -sass -ptx)
 /*!
@@ -2168,7 +2526,7 @@ static int get_app_cuda_version() {
  *	It is also responsible for extracting the libraries linked to the binary if the option is
  *	enabled
  * */
-void extract_code_using_cuobjdump(){
+void cuda_runtime_api::extract_code_using_cuobjdump(){
     CUctx_st *context = GPGPUSim_Context();
     unsigned forced_max_capability = context->get_device()->get_gpgpu()->get_config().get_forced_max_capability();
 
@@ -2185,12 +2543,11 @@ void extract_code_using_cuobjdump(){
     }
     // Running cuobjdump using dynamic link to current process
     // Needs the option '-all' to extract PTX from CDP-enabled binary 
-    extern bool g_cdp_enabled;
 
     //dump ptx for all individial ptx files into sepearte files which is later used by ptxas.
     int result=0;
 #if (CUDART_VERSION >= 6000)
-    extract_ptx_files_using_cuobjdump();
+    extract_ptx_files_using_cuobjdump(context);
     return;
 #endif
     //TODO: redundant to dump twice. how can it be prevented?
@@ -2200,7 +2557,7 @@ void extract_code_using_cuobjdump(){
 	snprintf(fname,1024,"_cuobjdump_complete_output_XXXXXX");
 	int fd=mkstemp(fname);
 	close(fd);
-	if(!g_cdp_enabled)
+	if(!gpgpu_ctx->device_runtime->g_cdp_enabled)
             snprintf(command,1000,"$CUDA_INSTALL_PATH/bin/cuobjdump -ptx -elf -sass %s > %s", app_binary.c_str(), fname);
 	else
             snprintf(command,1000,"$CUDA_INSTALL_PATH/bin/cuobjdump -ptx -elf -sass -all %s > %s", app_binary.c_str(), fname);
@@ -2223,9 +2580,16 @@ void extract_code_using_cuobjdump(){
 
         if (parse_output) {
             printf("Parsing file %s\n", fname);
+	    FILE *cuobjdump_in;
             cuobjdump_in = fopen(fname, "r");
 
-            cuobjdump_parse();
+	    struct cuobjdump_parser parser;
+	    parser.elfserial = 1;
+	    parser.ptxserial = 1;
+	    cuobjdump_lex_init(&(parser.scanner));
+	    cuobjdump_set_in(cuobjdump_in, (parser.scanner));
+	    cuobjdump_parse(parser.scanner, &parser, cuobjdumpSectionList);
+	    cuobjdump_lex_destroy(parser.scanner);
             fclose(cuobjdump_in);
             printf("Done parsing!!!\n");
         } else {
@@ -2273,8 +2637,15 @@ void extract_code_using_cuobjdump(){
                     std::cout << "Done" << std::endl;
 
                     std::cout << "Trying to parse " << libcodfn.str() << std::endl;
+		    FILE *cuobjdump_in;
                     cuobjdump_in = fopen(libcodfn.str().c_str(), "r");
-                    cuobjdump_parse();
+		    struct cuobjdump_parser parser;
+		    parser.elfserial = 1;
+		    parser.ptxserial = 1;
+		    cuobjdump_lex_init(&(parser.scanner));
+		    cuobjdump_set_in(cuobjdump_in, (parser.scanner));
+		    cuobjdump_parse(parser.scanner, &parser, cuobjdumpSectionList);
+		    cuobjdump_lex_destroy(parser.scanner);
                     fclose(cuobjdump_in);
                     std::getline(libsf, line);
             }
@@ -2324,7 +2695,7 @@ void printSectionList(std::list<cuobjdumpSection*> sl) {
 }
 
 //! Remove unecessary sm versions from the section list
-std::list<cuobjdumpSection*> pruneSectionList(std::list<cuobjdumpSection*> cuobjdumpSectionList, CUctx_st *context) {
+std::list<cuobjdumpSection*> cuda_runtime_api::pruneSectionList(CUctx_st *context) {
 	unsigned forced_max_capability = context->get_device()->get_gpgpu()->get_config().get_forced_max_capability();
 
 	//For ptxplus, force the max capability to 19 if it's higher or unspecified(0)
@@ -2377,7 +2748,7 @@ std::list<cuobjdumpSection*> pruneSectionList(std::list<cuobjdumpSection*> cuobj
 }
 
 //! Merge all PTX sections that have a specific identifier into one file
-std::list<cuobjdumpSection*> mergeMatchingSections(std::list<cuobjdumpSection*> cuobjdumpSectionList, std::string identifier){
+std::list<cuobjdumpSection*> cuda_runtime_api::mergeMatchingSections(std::string identifier){
 	const char *ptxcode = "";
 	std::list<cuobjdumpSection*>::iterator old_iter;
 	cuobjdumpPTXSection* old_ptxsection = NULL;
@@ -2420,7 +2791,7 @@ std::list<cuobjdumpSection*> mergeMatchingSections(std::list<cuobjdumpSection*>
 }
 
 //! Merge any PTX sections with matching identifiers
-std::list<cuobjdumpSection*> mergeSections(std::list<cuobjdumpSection*> cuobjdumpSectionList){
+std::list<cuobjdumpSection*> cuda_runtime_api::mergeSections(){
 	std::vector<std::string> identifier;
 	cuobjdumpPTXSection* ptxsection;
 
@@ -2442,7 +2813,7 @@ std::list<cuobjdumpSection*> mergeSections(std::list<cuobjdumpSection*> cuobjdum
 	for (	std::vector<std::string>::iterator iter = identifier.begin();
 			iter != identifier.end();
 			iter++) {
-		cuobjdumpSectionList = mergeMatchingSections(cuobjdumpSectionList, *iter);
+		cuobjdumpSectionList = mergeMatchingSections(*iter);
 	}
 
 	return cuobjdumpSectionList;
@@ -2467,7 +2838,7 @@ cuobjdumpELFSection* findELFSectionInList(std::list<cuobjdumpSection*> sectionli
 }
 
 //! Find an ELF section in all the known lists
-cuobjdumpELFSection* findELFSection(const std::string identifier){
+cuobjdumpELFSection* cuda_runtime_api::findELFSection(const std::string identifier){
 	cuobjdumpELFSection* sec = findELFSectionInList(cuobjdumpSectionList, identifier);
 	if (sec!=NULL)return sec;
 	sec = findELFSectionInList(libSectionList, identifier);
@@ -2478,7 +2849,7 @@ cuobjdumpELFSection* findELFSection(const std::string identifier){
 }
 
 //! Within the section list, find the PTX section corresponding to a given identifier
-cuobjdumpPTXSection* findPTXSectionInList(std::list<cuobjdumpSection*> sectionlist, const std::string identifier){
+cuobjdumpPTXSection* cuda_runtime_api::findPTXSectionInList(std::list<cuobjdumpSection*> &sectionlist, const std::string identifier){
 	std::list<cuobjdumpSection*>::iterator iter;
 	for (	iter = sectionlist.begin();
 			iter != sectionlist.end();
@@ -2489,8 +2860,7 @@ cuobjdumpPTXSection* findPTXSectionInList(std::list<cuobjdumpSection*> sectionli
 			if(ptxsection->getIdentifier() == identifier)
 				return ptxsection;
 			else {
-				extern bool g_cdp_enabled;
-				if(g_cdp_enabled) {
+				if(gpgpu_ctx->device_runtime->g_cdp_enabled) {
 					printf("Warning: __cudaRegisterFatBinary needs %s, but find PTX section with %s\n",
 						identifier.c_str(), ptxsection->getIdentifier().c_str());
 					return ptxsection;
@@ -2502,7 +2872,7 @@ cuobjdumpPTXSection* findPTXSectionInList(std::list<cuobjdumpSection*> sectionli
 }
 
 //! Find an PTX section in all the known lists
-cuobjdumpPTXSection* findPTXSection(const std::string identifier){
+cuobjdumpPTXSection* cuda_runtime_api::findPTXSection(const std::string identifier){
 	cuobjdumpPTXSection* sec = findPTXSectionInList(cuobjdumpSectionList, identifier);
 	if (sec!=NULL)return sec;
 	sec = findPTXSectionInList(libSectionList, identifier);
@@ -2515,35 +2885,27 @@ cuobjdumpPTXSection* findPTXSection(const std::string identifier){
 
 
 //! Extract the code using cuobjdump and remove unnecessary sections
-void cuobjdumpInit(){
+void cuda_runtime_api::cuobjdumpInit(){
 	CUctx_st *context = GPGPUSim_Context();
 	extract_code_using_cuobjdump(); //extract all the output of cuobjdump to _cuobjdump_*.*
 	const char* pre_load = getenv("CUOBJDUMP_SIM_FILE");
 	if (pre_load ==NULL || strlen(pre_load)==0){
-		cuobjdumpSectionList = pruneSectionList(cuobjdumpSectionList, context);
-		cuobjdumpSectionList = mergeSections(cuobjdumpSectionList);
+		cuobjdumpSectionList = pruneSectionList(context);
+		cuobjdumpSectionList = mergeSections();
 	}
 }
 
-std::map<int, std::string> fatbinmap;
-std::map<int, bool>fatbin_registered;
-std::map<std::string, symbol_table*> name_symtab;
-
-//! Keep track of the association between filename and cubin handle
-void cuobjdumpRegisterFatBinary(unsigned int handle, const char* filename){
-	fatbinmap[handle] = filename;
-}
 
 //! Either submit PTX for simulation or convert SASS to PTXPlus and submit it
-void cuobjdumpParseBinary(unsigned int handle){
+void gpgpu_context::cuobjdumpParseBinary(unsigned int handle){
 
-	if(fatbin_registered[handle]) return;
-	fatbin_registered[handle] = true;
 	CUctx_st *context = GPGPUSim_Context();
-	std::string fname = fatbinmap[handle];
+	if(api->fatbin_registered[handle]) return;
+	api->fatbin_registered[handle] = true;
+	std::string fname = api->fatbinmap[handle];
 
-	if (name_symtab.find(fname) != name_symtab.end()) {
-		symbol_table *symtab = name_symtab[fname];
+	if (api->name_symtab.find(fname) != api->name_symtab.end()) {
+		symbol_table *symtab = api->name_symtab[fname];
 		context->add_binary(symtab, handle);
 		return;
 	}
@@ -2552,7 +2914,7 @@ void cuobjdumpParseBinary(unsigned int handle){
 #if (CUDART_VERSION >= 6000)
    //loops through all ptx files from smallest sm version to largest
    std::map<unsigned,std::set<std::string> >::iterator itr_m;
-   for (itr_m = version_filename.begin(); itr_m!=version_filename.end(); itr_m++){
+   for (itr_m = api->version_filename.begin(); itr_m!=api->version_filename.end(); itr_m++){
       std::set<std::string>::iterator itr_s;
       for (itr_s = itr_m->second.begin(); itr_s!=itr_m->second.end(); itr_s++){
           std::string ptx_filename = *itr_s;
@@ -2560,11 +2922,11 @@ void cuobjdumpParseBinary(unsigned int handle){
           symtab = gpgpu_ptx_sim_load_ptx_from_filename( ptx_filename.c_str() );
       }
    }
-   name_symtab[fname] = symtab;
+   api->name_symtab[fname] = symtab;
    context->add_binary(symtab, handle);
-   load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF,context->get_device()->get_gpgpu());
-   load_constants(symtab,STATIC_ALLOC_LIMIT,context->get_device()->get_gpgpu());
-   for (itr_m = version_filename.begin(); itr_m!=version_filename.end(); itr_m++){
+   api->load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF,context->get_device()->get_gpgpu());
+   api->load_constants(symtab,STATIC_ALLOC_LIMIT,context->get_device()->get_gpgpu());
+   for (itr_m = api->version_filename.begin(); itr_m!=api->version_filename.end(); itr_m++){
       std::set<std::string>::iterator itr_s;
       for (itr_s = itr_m->second.begin(); itr_s!=itr_m->second.end(); itr_s++){
           std::string ptx_filename = *itr_s;
@@ -2576,8 +2938,8 @@ void cuobjdumpParseBinary(unsigned int handle){
 #endif
 
 	unsigned max_capability = 0;
-	for (	std::list<cuobjdumpSection*>::iterator iter = cuobjdumpSectionList.begin();
-			iter != cuobjdumpSectionList.end();
+	for (	std::list<cuobjdumpSection*>::iterator iter = api->cuobjdumpSectionList.begin();
+			iter != api->cuobjdumpSectionList.end();
 			iter++){
 		unsigned capability = (*iter)->getArch();
 		if (capability > max_capability) max_capability = capability;
@@ -2588,7 +2950,7 @@ void cuobjdumpParseBinary(unsigned int handle){
 	cuobjdumpPTXSection* ptx = NULL;
 	const char* pre_load = getenv("CUOBJDUMP_SIM_FILE");
 	if(pre_load==NULL || strlen(pre_load)==0)
-		ptx = findPTXSection(fname);
+		ptx = api->findPTXSection(fname);
 	char *ptxcode;
 	const char *override_ptx_name = getenv("PTX_SIM_KERNELFILE"); 
 	if (override_ptx_name == NULL or getenv("PTX_SIM_USE_PTX_FILE") == NULL or strlen(getenv("PTX_SIM_USE_PTX_FILE"))==0) {
@@ -2598,192 +2960,47 @@ void cuobjdumpParseBinary(unsigned int handle){
 		ptxcode = readfile(override_ptx_name);
 	}
 	if(context->get_device()->get_gpgpu()->get_config().convert_to_ptxplus() ) {
-		cuobjdumpELFSection* elfsection = findELFSection(ptx->getIdentifier());
+		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());
 		symtab=gpgpu_ptx_sim_load_ptx_from_string(ptxplus_str, handle);
 		printf("Adding %s with cubin handle %u\n", ptx->getPTXfilename().c_str(), handle);
 		context->add_binary(symtab, handle);
-		gpgpu_ptxinfo_load_from_string( ptxcode, handle, max_capability );
+		gpgpu_ptxinfo_load_from_string( ptxcode, handle, max_capability, context->no_of_ptx );
 		delete[] ptxplus_str;
 	} else {
 		symtab=gpgpu_ptx_sim_load_ptx_from_string(ptxcode, handle);
 		//if CUOBJDUMP_SIM_FILE is not set, ptx is NULL. So comment below.
 		//printf("Adding %s with cubin handle %u\n", ptx->getPTXfilename().c_str(), handle);
 		context->add_binary(symtab, handle);
-		gpgpu_ptxinfo_load_from_string( ptxcode, handle, max_capability );
+		gpgpu_ptxinfo_load_from_string( ptxcode, handle, max_capability, context->no_of_ptx );
 	}
-	load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF,context->get_device()->get_gpgpu());
-	load_constants(symtab,STATIC_ALLOC_LIMIT,context->get_device()->get_gpgpu());
-	name_symtab[fname] = symtab;
+	api->load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF,context->get_device()->get_gpgpu());
+	api->load_constants(symtab,STATIC_ALLOC_LIMIT,context->get_device()->get_gpgpu());
+	api->name_symtab[fname] = symtab;
 
 	//TODO: Remove temporarily files as per configurations
 }
+}
+
+extern "C" {
 
 void** CUDARTAPI __cudaRegisterFatBinary( void *fatCubin )
 {
 	if(g_debug_execution >= 3){
 	    announce_call(__my_func__);
     }
-#if (CUDART_VERSION < 2010)
-	printf("GPGPU-Sim PTX: ERROR ** this version of GPGPU-Sim requires CUDA 2.1 or higher\n");
-	exit(1);
-#endif
-	CUctx_st *context = GPGPUSim_Context();
-	static unsigned next_fat_bin_handle = 1;
-	if(context->get_device()->get_gpgpu()->get_config().use_cuobjdump()) {
-		// The following workaround has only been verified on 64-bit systems. 
-		if (sizeof(void*) == 4) 
-			printf("GPGPU-Sim PTX: FatBin file name extraction has not been tested on 32-bit system.\n"); 
-
-        // This code will get the CUDA version the app was compiled with.
-        // We need this to determine how to handle the parsing of the binary.
-        // Making this a runtime variable based on the app, enables GPGPU-Sim compiled
-        // with a newer version of CUDA to run apps compiled with older versions of
-        // CUDA. This is especially useful for PTXPLUS execution.
-        //Skip cuda version check for pytorch application
-	std::string app_binary_path =  get_app_binary();
-        int pos = app_binary_path.find("python");
-        if (pos==std::string::npos){
-        	// Not pytorch app : checking cuda version
-		int app_cuda_version = get_app_cuda_version();
-        	assert( app_cuda_version == CUDART_VERSION / 1000  && "The app must be compiled with same major version as the simulator." );
-        }
-
-	//int app_cuda_version = get_app_cuda_version();
-        //assert( app_cuda_version == CUDART_VERSION / 1000  && "The app must be compiled with same major version as the simulator." );
-        const char* filename;
-#if CUDART_VERSION < 6000
-            // FatBin handle from the .fatbin.c file (one of the intermediate files generated by NVCC)
-            typedef struct {int m; int v; const unsigned long long* d; char* f;} __fatDeviceText __attribute__ ((aligned (8))); 
-            __fatDeviceText * fatDeviceText = (__fatDeviceText *) fatCubin;
-
-            // Extract the source code file name that generate the given FatBin. 
-            // - Obtains the pointer to the actual fatbin structure from the FatBin handle (fatCubin).
-            // - An integer inside the fatbin structure contains the relative offset to the source code file name.
-            // - This offset differs among different CUDA and GCC versions. 
-            char * pfatbin = (char*) fatDeviceText->d; 
-            int offset = *((int*)(pfatbin+48)); 
-            filename = (pfatbin+16+offset);
-#else
-            filename = "default";
-#endif
-
-		// The extracted file name is associated with a fat_cubin_handle passed
-		// into cudaLaunch().  Inside cudaLaunch(), the associated file name is
-		// used to find the PTX/SASS section from cuobjdump, which contains the
-		// PTX/SASS code for the launched kernel function.  
-		// This allows us to work around the fact that cuobjdump only outputs the
-		// file name associated with each section. 
-		unsigned long long fat_cubin_handle = next_fat_bin_handle;
-		next_fat_bin_handle++;
-		printf("GPGPU-Sim PTX: __cudaRegisterFatBinary, fat_cubin_handle = %llu, filename=%s\n", fat_cubin_handle, filename);
-		/*!
-		 * This function extracts all data from all files in first call
-		 * then for next calls, only returns the appropriate number
-		 */
-		assert(fat_cubin_handle >= 1);
-		if (fat_cubin_handle==1) cuobjdumpInit();
-		cuobjdumpRegisterFatBinary(fat_cubin_handle, filename);
-
-		return (void**)fat_cubin_handle;
-	} 
-#if (CUDART_VERSION < 8000)
-	else {
-		static unsigned source_num=1;
-		unsigned long long fat_cubin_handle = next_fat_bin_handle++;
-		__cudaFatCudaBinary *info =   (__cudaFatCudaBinary *)fatCubin;
-		assert( info->version >= 3 );
-		unsigned num_ptx_versions=0;
-		unsigned max_capability=0;
-		unsigned selected_capability=0;
-		bool found=false;
-		unsigned forced_max_capability = context->get_device()->get_gpgpu()->get_config().get_forced_max_capability();
-		if (!info->ptx){
-			printf("ERROR: Cannot find ptx code in cubin file\n"
-					"\tIf you are using CUDA 4.0 or higher, please enable -gpgpu_ptx_use_cuobjdump or downgrade to CUDA 3.1\n");
-			exit(1);
-		}
-		while( info->ptx[num_ptx_versions].gpuProfileName != NULL ) {
-			unsigned capability=0;
-			sscanf(info->ptx[num_ptx_versions].gpuProfileName,"compute_%u",&capability);
-			printf("GPGPU-Sim PTX: __cudaRegisterFatBinary found PTX versions for '%s', ", info->ident);
-			printf("capability = %s\n", info->ptx[num_ptx_versions].gpuProfileName );
-			if( forced_max_capability ) {
-				if( capability > max_capability && capability <= forced_max_capability ) {
-					found = true;
-					max_capability=capability;
-					selected_capability = num_ptx_versions;
-				}
-			} else {
-				if( capability > max_capability ) {
-					found = true;
-					max_capability=capability;
-					selected_capability = num_ptx_versions;
-				}
-			}
-			num_ptx_versions++;
-		}
-		if( found  ) {
-			printf("GPGPU-Sim PTX: Loading PTX for %s, capability = %s\n",
-					info->ident, info->ptx[selected_capability].gpuProfileName );
-			symbol_table *symtab;
-			const char *ptx = info->ptx[selected_capability].ptx;
-			if(context->get_device()->get_gpgpu()->get_config().convert_to_ptxplus() ) {
-				printf("GPGPU-Sim PTX: ERROR ** PTXPlus is only supported through cuobjdump\n"
-						"\tEither enable cuobjdump or disable PTXPlus in your configuration file\n");
-				exit(1);
-			} else {
-				symtab=gpgpu_ptx_sim_load_ptx_from_string(ptx,source_num);
-				context->add_binary(symtab,fat_cubin_handle);
-				gpgpu_ptxinfo_load_from_string( ptx, source_num, max_capability );
-			}
-			source_num++;
-			load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF,context->get_device()->get_gpgpu());
-			load_constants(symtab,STATIC_ALLOC_LIMIT,context->get_device()->get_gpgpu());
-		} else {
-			printf("GPGPU-Sim PTX: warning -- did not find an appropriate PTX in cubin\n");
-		}
-		return (void**)fat_cubin_handle;
-	}
-#else
-        else {
-		printf("ERROR **  __cudaRegisterFatBinary() needs to be updated\n");
-		abort();
-        }
-#endif
+    return cudaRegisterFatBinaryInternal(fatCubin);
 }
 
-void __cudaUnregisterFatBinary(void **fatCubinHandle)
+void CUDARTAPI __cudaRegisterFatBinaryEnd( void **fatCubinHandle )
 {
 	if(g_debug_execution >= 3){
 	    announce_call(__my_func__);
     }
-	;
-}
-
-cudaError_t cudaDeviceReset ( void ) {
-	// Should reset the simulated GPU
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-	return g_last_cudaError = cudaSuccess;
-}
-cudaError_t CUDARTAPI cudaDeviceSynchronize(void){
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-	//Blocks until the device has completed all preceding requested tasks
-	synchronize();
-	return g_last_cudaError = cudaSuccess;
-}
-
-void CUDARTAPI __cudaRegisterFatBinaryEnd( void **fatCubinHandle )
-{
-
 }
 
 unsigned CUDARTAPI __cudaPushCallConfiguration(dim3 gridDim,
@@ -2791,7 +3008,10 @@ unsigned CUDARTAPI __cudaPushCallConfiguration(dim3 gridDim,
                                       size_t sharedMem = 0, 
                                       struct CUstream_st *stream = 0)
 {
-
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+    cudaConfigureCallInternal(gridDim, blockDim, sharedMem, stream);
 }
 
 cudaError_t CUDARTAPI __cudaPopCallConfiguration(
@@ -2801,9 +3021,13 @@ cudaError_t CUDARTAPI __cudaPopCallConfiguration(
   void         *stream
 )
 {
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
     return g_last_cudaError = cudaSuccess;
 }
 
+
 void CUDARTAPI __cudaRegisterFunction(
 		void   **fatCubinHandle,
 		const char    *hostFun,
@@ -2814,18 +3038,19 @@ void CUDARTAPI __cudaRegisterFunction(
 		uint3   *bid,
 		dim3    *bDim,
 		dim3    *gDim
-)
-{
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-	CUctx_st *context = GPGPUSim_Context();
-	unsigned fat_cubin_handle = (unsigned)(unsigned long long)fatCubinHandle;
-	printf("GPGPU-Sim PTX: __cudaRegisterFunction %s : hostFun 0x%p, fat_cubin_handle = %u\n",
-			deviceFun, hostFun, fat_cubin_handle);
-	if(context->get_device()->get_gpgpu()->get_config().use_cuobjdump())
-		cuobjdumpParseBinary(fat_cubin_handle);
-	context->register_function( fat_cubin_handle, hostFun, deviceFun );
+) {
+    cudaRegisterFunctionInternal(
+	    fatCubinHandle,
+	    hostFun,
+	    deviceFun,
+	    deviceName,
+	    thread_limit,
+	    tid,
+	    bid,
+	    bDim,
+	    gDim
+	    );
+
 }
 
 extern void __cudaRegisterVar(
@@ -2838,19 +3063,43 @@ extern void __cudaRegisterVar(
 		int constant,
 		int global )
 {
+    cudaRegisterVarInternal(
+	    fatCubinHandle,
+	    hostVar,
+	    deviceAddress,
+	    deviceName,
+	    ext,
+	    size,
+	    constant,
+	    global );
+}
+
+__host__ cudaError_t CUDARTAPI cudaConfigureCall(dim3 gridDim, dim3 blockDim, size_t sharedMem, cudaStream_t stream)
+{
+    return cudaConfigureCallInternal(gridDim, blockDim, sharedMem, stream);
+}
+
+void __cudaUnregisterFatBinary(void **fatCubinHandle)
+{
 	if(g_debug_execution >= 3){
 	    announce_call(__my_func__);
     }
-	printf("GPGPU-Sim PTX: __cudaRegisterVar: hostVar = %p; deviceAddress = %s; deviceName = %s\n", hostVar, deviceAddress, deviceName);
-	printf("GPGPU-Sim PTX: __cudaRegisterVar: Registering const memory space of %d bytes\n", size);
-	if(GPGPUSim_Context()->get_device()->get_gpgpu()->get_config().use_cuobjdump())
-		cuobjdumpParseBinary((unsigned)(unsigned long long)fatCubinHandle);
-	fflush(stdout);
-	if ( constant && !global && !ext ) {
-		gpgpu_ptx_sim_register_const_variable(hostVar,deviceName,size);
-	} else if ( !constant && !global && !ext ) {
-		gpgpu_ptx_sim_register_global_variable(hostVar,deviceName,size);
-	} else cuda_not_implemented(__my_func__,__LINE__);
+}
+
+cudaError_t cudaDeviceReset ( void ) {
+	// Should reset the simulated GPU
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+	return g_last_cudaError = cudaSuccess;
+}
+cudaError_t CUDARTAPI cudaDeviceSynchronize(void){
+	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+	//Blocks until the device has completed all preceding requested tasks
+	synchronize();
+	return g_last_cudaError = cudaSuccess;
 }
 
 
@@ -2922,10 +3171,6 @@ char __cudaInitModule(
 }
 
 
-#ifndef OPENGL_SUPPORT
-typedef unsigned long GLuint;
-#endif
-
 cudaError_t cudaGLRegisterBufferObject(GLuint bufferObj)
 {
 	if(g_debug_execution >= 3){
@@ -2935,70 +3180,9 @@ cudaError_t cudaGLRegisterBufferObject(GLuint bufferObj)
 	return g_last_cudaError = cudaSuccess;
 }
 
-struct glbmap_entry {
-	GLuint m_bufferObj;
-	void *m_devPtr;
-	size_t m_size;
-	struct glbmap_entry *m_next;
-};
-typedef struct glbmap_entry glbmap_entry_t;
-
-glbmap_entry_t* g_glbmap = NULL;
-
 cudaError_t cudaGLMapBufferObject(void** devPtr, GLuint bufferObj) 
 {
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-#ifdef OPENGL_SUPPORT
-	GLint buffer_size=0;
-	CUctx_st* ctx = GPGPUSim_Context();
-
-	glbmap_entry_t *p = g_glbmap;
-	while ( p && p->m_bufferObj != bufferObj )
-		p = p->m_next;
-	if ( p == NULL ) {
-		glBindBuffer(GL_ARRAY_BUFFER,bufferObj);
-		glGetBufferParameteriv(GL_ARRAY_BUFFER,GL_BUFFER_SIZE,&buffer_size);
-		assert( buffer_size != 0 );
-		*devPtr = ctx->get_device()->get_gpgpu()->gpu_malloc(buffer_size);
-
-		// create entry and insert to front of list
-		glbmap_entry_t *n = (glbmap_entry_t *) calloc(1,sizeof(glbmap_entry_t));
-		n->m_next = g_glbmap;
-		g_glbmap = n;
-
-		// initialize entry
-		n->m_bufferObj = bufferObj;
-		n->m_devPtr = *devPtr;
-		n->m_size = buffer_size;
-
-		p = n;
-	} else {
-		buffer_size = p->m_size;
-		*devPtr = p->m_devPtr;
-	}
-
-	if ( *devPtr  ) {
-		char *data = (char *) calloc(p->m_size,1);
-		glGetBufferSubData(GL_ARRAY_BUFFER,0,buffer_size,data);
-		memcpy_to_gpu( (size_t) *devPtr, data, buffer_size );
-		free(data);
-		printf("GPGPU-Sim PTX: cudaGLMapBufferObject %zu bytes starting at 0x%llx..\n", (size_t)buffer_size,
-				(unsigned long long) *devPtr);
-		return g_last_cudaError = cudaSuccess;
-	} else {
-		return g_last_cudaError = cudaErrorMemoryAllocation;
-	}
-
-	return g_last_cudaError = cudaSuccess;
-#else
-	fflush(stdout);
-	fflush(stderr);
-	printf("GPGPU-Sim PTX: GPGPU-Sim support for OpenGL integration disabled -- exiting\n");
-	fflush(stdout);
-	exit(50);
-#endif
+    return cudaGLMapBufferObjectInternal(devPtr, bufferObj);
 }
 
 cudaError_t cudaGLUnmapBufferObject(GLuint bufferObj)
@@ -3007,7 +3191,8 @@ cudaError_t cudaGLUnmapBufferObject(GLuint bufferObj)
 	    announce_call(__my_func__);
     }
 #ifdef OPENGL_SUPPORT
-	glbmap_entry_t *p = g_glbmap;
+	CUctx_st* ctx = GPGPUSim_Context();
+	glbmap_entry_t *p = ctx->api->g_glbmap;
 	while ( p && p->m_bufferObj != bufferObj )
 		p = p->m_next;
 	if ( p == NULL )
@@ -3041,45 +3226,12 @@ cudaError_t cudaGLUnregisterBufferObject(GLuint bufferObj)
 
 cudaError_t CUDARTAPI cudaHostAlloc(void **pHost,  size_t bytes, unsigned int flags)
 {
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-	*pHost = malloc(bytes);
-	//need to track the size allocated so that cudaHostGetDevicePointer() can function properly.
-	//TODO: vary this function behavior based on flags value (following nvidia documentation)
-	pinned_memory_size[*pHost]=bytes;
-	if( *pHost )
-		return g_last_cudaError = cudaSuccess;
-	else
-		return g_last_cudaError = cudaErrorMemoryAllocation;
+    return cudaHostAllocInternal(pHost, bytes, flags);
 }
 
 cudaError_t CUDARTAPI cudaHostGetDevicePointer(void **pDevice, void *pHost, unsigned int flags)
 {
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-	//only cpu memory allocation happens in cudaHostAlloc. Linking with device pointer to pinned memory happens here.
-	//TODO: once kernel is executed, the contents in global pointer of GPU must be copied back to CPU host pointer!
-	flags=0;
-	CUctx_st* context = GPGPUSim_Context();
-	gpgpu_t *gpu = context->get_device()->get_gpgpu();
-	std::map<void *, size_t>::const_iterator i = pinned_memory_size.find(pHost);
-	assert(i != pinned_memory_size.end());
-	size_t size = i->second;
-	*pDevice = gpu->gpu_malloc(size);
-	if(g_debug_execution >= 3){
-		printf("GPGPU-Sim PTX: cudaMallocing %zu bytes starting at 0x%llx..\n",size, (unsigned long long) *pDevice);
-        g_mallocPtr_Size[(unsigned long long)*pDevice] = size;
-    }
-	if ( *pDevice  ) {
-		pinned_memory[pHost]=pDevice;
-		//Copy contents in cpu to gpu
-		gpu->memcpy_to_gpu((size_t)*pDevice,pHost,size);
-		return g_last_cudaError = cudaSuccess;
-	} else {
-		return g_last_cudaError = cudaErrorMemoryAllocation;
-	}
+    return cudaHostGetDevicePointerInternal(pDevice, pHost, flags);
 }
 
 __host__ cudaError_t CUDARTAPI cudaPointerGetAttributes(
@@ -3143,61 +3295,12 @@ cudaError_t CUDARTAPI cudaSetDeviceFlags( int flags )
     }
 }
 
-size_t getMaxThreadsPerBlock(struct cudaFuncAttributes *attr) {
-  _cuda_device_id *dev = GPGPUSim_Init();
-  struct cudaDeviceProp prop;
-
-  prop = *dev->get_prop();
-
-  size_t max = prop.maxThreadsPerBlock;
-
-  if (attr->numRegs && (prop.regsPerBlock / attr->numRegs) < max)
-    max = prop.regsPerBlock / attr->numRegs;
 
-  if (attr->sharedSizeBytes && (prop.sharedMemPerBlock / attr->sharedSizeBytes) < max)
-      max = prop.sharedMemPerBlock / attr->sharedSizeBytes;
 
-  return max;
-}
 
 cudaError_t CUDARTAPI cudaFuncGetAttributes(struct cudaFuncAttributes *attr, const char *hostFun )
 {
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-	CUctx_st *context = GPGPUSim_Context();
-	function_info *entry = context->get_kernel(hostFun);
-	if( entry ) {
-		const struct gpgpu_ptx_sim_info *kinfo = entry->get_kernel_info();
-		attr->sharedSizeBytes = kinfo->smem;
-		attr->constSizeBytes  = kinfo->cmem;
-		attr->localSizeBytes  = kinfo->lmem;
-		attr->numRegs         = kinfo->regs;
-		if(kinfo->maxthreads > 0)
-		  attr->maxThreadsPerBlock = kinfo->maxthreads;
-		else
-		  attr->maxThreadsPerBlock = getMaxThreadsPerBlock(attr);
-#if CUDART_VERSION >= 3000
-		attr->ptxVersion      = kinfo->ptx_version;
-		attr->binaryVersion   = kinfo->sm_target;
-#endif
-	}
-	return g_last_cudaError = cudaSuccess;
-}
-
-cudaError_t CUDARTAPI cudaEventCreateWithFlags(cudaEvent_t *event, int flags)
-{
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-	CUevent_st *e = new CUevent_st(flags==cudaEventBlockingSync);
-	g_timer_events[e->get_uid()] = e;
-#if CUDART_VERSION >= 3000
-	*event = e;
-#else
-	*event = e->get_uid();
-#endif
-	return g_last_cudaError = cudaSuccess;
+    return cudaFuncGetAttributesInternal(attr, hostFun );
 }
 
 cudaError_t CUDARTAPI cudaDriverGetVersion(int *driverVersion)
@@ -3237,26 +3340,64 @@ __host__ cudaError_t CUDARTAPI cudaDeviceSetLimit(enum cudaLimit limit, size_t v
     return g_last_cudaError = cudaSuccess;
 }
 
-#if CUDART_VERSION >= 9000
-__host__  cudaError_t cudaFuncSetAttribute ( const void* func, enum cudaFuncAttribute attr, int value ) {
+//#if CUDART_VERSION >= 9000
+//__host__  cudaError_t cudaFuncSetAttribute ( const void* func, enum cudaFuncAttribute attr, int value ) {
 
 	//ignore this Attribute for now, and the default is that carveout = cudaSharedmemCarveoutDefault;   //  (-1)
-	return g_last_cudaError = cudaSuccess;
-}
+//	return g_last_cudaError = cudaSuccess;
+//}
 
-#if CUDART_VERSION >= 9020
-__host__ __device__ unsigned __cudaPushCallConfiguration(dim3 gridDim, dim3 blockDim,
-                                                size_t sharedMem = 0,
-                                                void *stream = 0) {
 
-	return 1;
-}
 #endif
 
 #endif
 
 #endif
 
+
+#if CUDART_VERSION >= 9000
+/**
+ * \brief Set attributes for a given function
+ *
+ * This function sets the attributes of a function specified via \p entry.
+ * The parameter \p entry must be a pointer to a function that executes
+ * on the device. The parameter specified by \p entry must be declared as a \p __global__
+ * function. The enumeration defined by \p attr is set to the value defined by \p value
+ * If the specified function does not exist, then ::cudaErrorInvalidDeviceFunction is returned.
+ * If the specified attribute cannot be written, or if the value is incorrect, 
+ * then ::cudaErrorInvalidValue is returned.
+ *
+ * Valid values for \p attr are:
+ * ::cuFuncAttrMaxDynamicSharedMem - Maximum size of dynamic shared memory per block
+ * ::cudaFuncAttributePreferredSharedMemoryCarveout - Preferred shared memory-L1 cache split ratio
+ *
+ * \param entry - Function to get attributes of
+ * \param attr  - Attribute to set
+ * \param value - Value to set
+ *
+ * \return
+ * ::cudaSuccess,
+ * ::cudaErrorInitializationError,
+ * ::cudaErrorInvalidDeviceFunction,
+ * ::cudaErrorInvalidValue
+ * \notefnerr
+ *
+ * \ref ::cudaLaunchKernel(const T *func, dim3 gridDim, dim3 blockDim, void **args, size_t sharedMem, cudaStream_t stream) "cudaLaunchKernel (C++ API)",
+ * \ref ::cudaFuncSetCacheConfig(T*, enum cudaFuncCache) "cudaFuncSetCacheConfig (C++ API)",
+ * \ref ::cudaFuncGetAttributes(struct cudaFuncAttributes*, const void*) "cudaFuncGetAttributes (C API)",
+ * ::cudaSetDoubleForDevice,
+ * ::cudaSetDoubleForHost,
+ * \ref ::cudaSetupArgument(T, size_t) "cudaSetupArgument (C++ API)"
+ */
+cudaError_t CUDARTAPI cudaFuncSetAttribute(const void *func, enum cudaFuncAttribute attr, int value)
+{
+   	if(g_debug_execution >= 3){
+	    announce_call(__my_func__);
+    }
+	printf("GPGPU-Sim PTX: Execution warning: ignoring call to \"%s ( func=%p, attr=%d, value=%d )\"\n",
+        __my_func__, func, attr, value );
+    return g_last_cudaError = cudaSuccess;
+}
 #endif
 
 cudaError_t CUDARTAPI cudaGLSetGLDevice(int device)
@@ -3328,17 +3469,9 @@ int CUDARTAPI __cudaSynchronizeThreads(void**, void*)
 
 ////////
 
-extern int ptx_parse();
-extern int ptx__scan_string(const char*);
-extern FILE *ptx_in;
-
-extern int ptxinfo_parse();
-extern int ptxinfo_debug;
-extern FILE *ptxinfo_in;
-
 /// static functions
 
-static int load_static_globals( symbol_table *symtab, unsigned min_gaddr, unsigned max_gaddr, gpgpu_t *gpu ) 
+int cuda_runtime_api::load_static_globals( symbol_table *symtab, unsigned min_gaddr, unsigned max_gaddr, gpgpu_t *gpu ) 
 {
 	if(g_debug_execution >= 3){
 	    announce_call(__my_func__);
@@ -3377,7 +3510,7 @@ static int load_static_globals( symbol_table *symtab, unsigned min_gaddr, unsign
 	return ng_bytes;
 }
 
-static int load_constants( symbol_table *symtab, addr_t min_gaddr, gpgpu_t *gpu ) 
+int cuda_runtime_api::load_constants( symbol_table *symtab, addr_t min_gaddr, gpgpu_t *gpu ) 
 {
 	if(g_debug_execution >= 3){
 	    announce_call(__my_func__);
@@ -3423,7 +3556,7 @@ static int load_constants( symbol_table *symtab, addr_t min_gaddr, gpgpu_t *gpu
 	return nc_bytes;
 }
 
-kernel_info_t *gpgpu_cuda_ptx_sim_init_grid( const char *hostFun, 
+kernel_info_t * cuda_runtime_api::gpgpu_cuda_ptx_sim_init_grid( const char *hostFun,
 		gpgpu_ptx_sim_arg_list_t args,
 		struct dim3 gridDim,
 		struct dim3 blockDim,
@@ -3451,7 +3584,7 @@ kernel_info_t *gpgpu_cuda_ptx_sim_init_grid( const char *hostFun,
 	}
 
 	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){
@@ -3919,33 +4052,8 @@ CUresult CUDAAPI
 cuLinkAddFile(CUlinkState state, CUjitInputType type, const char *path,
     unsigned int numOptions, CUjit_option *options, void **optionValues)
 {
-	if(g_debug_execution >= 3){
-	    announce_call(__my_func__);
-    }
-    static bool addedFile = false;
-    if (addedFile){
-        printf("GPGPU-Sim PTX: ERROR: cuLinkAddFile does not support multiple files\n");
-        abort();
-    }
-
-    //blocking
-    assert(type==CU_JIT_INPUT_PTX);
-	CUctx_st *context = GPGPUSim_Context();
-    char *file = getenv("PTX_JIT_PATH");
-    if(file==NULL){
-        printf("GPGPU-Sim PTX: ERROR: PTX_JIT_PATH has not been set\n");
-        abort();
-    }
-    strcat(file,"/");
-    strcat(file,path);
-	symbol_table *symtab = gpgpu_ptx_sim_load_ptx_from_filename( file );
-    std::string fname(path);
-    name_symtab[fname] = symtab;
-    context->add_binary(symtab, 1);
-    load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF,context->get_device()->get_gpgpu());
-    load_constants(symtab,STATIC_ALLOC_LIMIT,context->get_device()->get_gpgpu());
-    addedFile = true;
-	return CUDA_SUCCESS;
+    return cuLinkAddFileInternal(state, type, path,
+	    numOptions, options, optionValues);
 }
 #endif
 
@@ -4939,12 +5047,12 @@ CUresult CUDAAPI cuLaunchKernel(CUfunction f,
     const char *hostFun = (const char*) f;
 	CUctx_st *context = GPGPUSim_Context();
 	function_info *entry = context->get_kernel(hostFun);
-    cudaConfigureCall(dim3(gridDimX, gridDimY, gridDimZ), dim3(blockDimX, blockDimY, blockDimZ), sharedMemBytes, (cudaStream_t) hStream);
+    cudaConfigureCallInternal(dim3(gridDimX, gridDimY, gridDimZ), dim3(blockDimX, blockDimY, blockDimZ), sharedMemBytes, (cudaStream_t) hStream);
     for(unsigned i = 0; i < entry->num_args(); i++){
         std::pair<size_t, unsigned> p = entry->get_param_config(i);
         cudaSetupArgument(kernelParams[i], p.first, p.second);
     }
-    cudaLaunch(hostFun);
+    cudaLaunchInternal(hostFun);
 	return CUDA_SUCCESS;
 }
 #endif /* CUDART_VERSION >= 4000 */
diff --git a/libcuda/cuobjdump.h b/libcuda/cuobjdump.h
new file mode 100644
index 0000000..6ab6778
--- /dev/null
+++ b/libcuda/cuobjdump.h
@@ -0,0 +1,80 @@
+#ifndef __cuobjdump_h__
+#define __cuobjdump_h__
+#include <string>
+#include <list>
+#include <iostream>
+
+typedef void * yyscan_t;
+struct cuobjdump_parser {
+    yyscan_t scanner;
+    int elfserial;
+    int ptxserial;
+    FILE *ptxfile;
+    FILE *elffile;
+    FILE *sassfile;
+    char filename [1024];
+};
+
+class cuobjdumpSection {
+public:
+	//Constructor
+	cuobjdumpSection() {
+		arch = 0;
+		identifier = "";
+	}
+	virtual ~cuobjdumpSection() {}
+	unsigned getArch() {return arch;}
+	void setArch(unsigned a) {arch = a;}
+	std::string getIdentifier() {return identifier;}
+	void setIdentifier(std::string i) {identifier = i;}
+	virtual void print(){std::cout << "cuobjdump Section: unknown type" << std::endl;}
+private:
+	unsigned arch;
+	std::string identifier;
+};
+
+class cuobjdumpELFSection : public cuobjdumpSection
+{
+public:
+	cuobjdumpELFSection() {}
+	virtual ~cuobjdumpELFSection() {
+		elffilename = "";
+		sassfilename = "";
+	}
+	std::string getELFfilename() {return elffilename;}
+	void setELFfilename(std::string f) {elffilename = f;}
+	std::string getSASSfilename() {return sassfilename;}
+	void setSASSfilename(std::string f) {sassfilename = f;}
+	virtual void print() {
+		std::cout << "ELF Section:" << std::endl;
+		std::cout << "arch: sm_" << getArch() << std::endl;
+		std::cout << "identifier: " << getIdentifier() << std::endl;
+		std::cout << "elf filename: " << getELFfilename() << std::endl;
+		std::cout << "sass filename: " << getSASSfilename() << std::endl;
+		std::cout << std::endl;
+	}
+private:
+	std::string elffilename;
+	std::string sassfilename;
+};
+
+class cuobjdumpPTXSection : public cuobjdumpSection
+{
+public:
+	cuobjdumpPTXSection(){
+		ptxfilename = "";
+	}
+	std::string getPTXfilename() {return ptxfilename;}
+	void setPTXfilename(std::string f) {ptxfilename = f;}
+	virtual void print() {
+		std::cout << "PTX Section:" << std::endl;
+		std::cout << "arch: sm_" << getArch() << std::endl;
+		std::cout << "identifier: " << getIdentifier() << std::endl;
+		std::cout << "ptx filename: " << getPTXfilename() << std::endl;
+		std::cout << std::endl;
+	}
+private:
+	std::string ptxfilename;
+};
+
+#endif /* __cuobjdump_h__ */
diff --git a/libcuda/cuobjdump.l b/libcuda/cuobjdump.l
index 2b0dac8..5a19d65 100644
--- a/libcuda/cuobjdump.l
+++ b/libcuda/cuobjdump.l
@@ -30,6 +30,7 @@
 %{
 #include <stdio.h>
 #include <string.h>
+#include "cuobjdump.h"
 #include "cuobjdump_parser.h"
 
 #define YY_NEVER_INTERACTIVE 1
@@ -38,9 +39,9 @@
 
 #define YYDEBUG 1
 
-#define yylval cuobjdump_lval
-
-void cuobjdump_error(const char*);
+void cuobjdump_error(yyscan_t yyscanner, struct cuobjdump_parser* parser, std::list<cuobjdumpSection*> &cuobjdumpSectionList, const char* msg);
+#define YY_DECL int cuobjdump_lex \
+	       (YYSTYPE * yylval_param , yyscan_t yyscanner, struct cuobjdump_parser* parser, std::list<cuobjdumpSection*> &cuobjdumpSectionList)
 %}
 
 %option stack
@@ -48,6 +49,8 @@ void cuobjdump_error(const char*);
 %option yylineno
 %option nounput
 
+%option bison-bridge
+%option reentrant
 
 %s ptxcode
 %s sasscode
@@ -69,54 +72,54 @@ newlines	{newline}+
 
 
 "ptxasOptions"{notnewline}*{newline}
-[1-9]{numeric}* yylval.string_value = strdup(yytext); return DECIMAL;
+[1-9]{numeric}* yylval->string_value = strdup(yytext); return DECIMAL;
 
 "has debug info"{newline}
 
 
 "Fatbin ptx code:"{newline}	{
-	yy_push_state(ptxcode);
-	yy_push_state(identifier);
-	yy_push_state(ptxheader);
-	yylval.string_value = strdup(yytext);
+	yy_push_state(ptxcode, yyscanner);
+	yy_push_state(identifier, yyscanner);
+	yy_push_state(ptxheader, yyscanner);
+	yylval->string_value = strdup(yytext);
 	return PTXHEADER;
 }
 "Fatbin elf code:"{newline}	{
-	yy_push_state(elfcode);
-	yy_push_state(identifier);
-	yy_push_state(elfheader);
-	yylval.string_value = strdup(yytext);
+	yy_push_state(elfcode, yyscanner);
+	yy_push_state(identifier, yyscanner);
+	yy_push_state(elfheader, yyscanner);
+	yylval->string_value = strdup(yytext);
 	return ELFHEADER;
 }
 	/*PTX code tokens*/
-<ptxcode>{notnewline}*{newline}	yylval.string_value = strdup(yytext); return PTXLINE;
+<ptxcode>{notnewline}*{newline}	yylval->string_value = strdup(yytext); return PTXLINE;
 
 	/*ELF code tokens*/
 <elfcode>{whitespace}*"code for sm_"{numeric}+{newline}	{
 	BEGIN(sasscode); 
-	yylval.string_value = strdup(yytext); 
+	yylval->string_value = strdup(yytext);
 	return SASSLINE;
 }
-<elfcode>{notnewline}*{newline} yylval.string_value = strdup(yytext); return ELFLINE;
+<elfcode>{notnewline}*{newline} yylval->string_value = strdup(yytext); return ELFLINE;
 
 	/*SASS code tokens*/
-<sasscode>{notnewline}*{newline} yylval.string_value = strdup(yytext); return SASSLINE;
+<sasscode>{notnewline}*{newline} yylval->string_value = strdup(yytext); return SASSLINE;
 
 <identifier>{newline}"compressed"{newline}	BEGIN(conidentifier); return H_COMPRESSED;
 <identifier>{newline}"identifier = "	BEGIN(endidentifier); return H_IDENTIFIER;
-<identifier>{newline}{newline} yy_pop_state();
+<identifier>{newline}{newline} yy_pop_state(yyscanner);
 
 
 <conidentifier>"identifier = "	BEGIN(endidentifier); return H_IDENTIFIER;
-<conidentifier>{newline} yy_pop_state();
+<conidentifier>{newline} yy_pop_state(yyscanner);
 
 
-<endidentifier>{notnewline}+	yylval.string_value = strdup(yytext); yy_pop_state(); return FILENAME;
+<endidentifier>{notnewline}+	yylval->string_value = strdup(yytext); yy_pop_state(yyscanner); return FILENAME;
 
 
 
 	/*Header tokens*/
-<elfheader>[[:alnum:]_]+ yylval.string_value = strdup(yytext); return IDENTIFIER;
+<elfheader>[[:alnum:]_]+ yylval->string_value = strdup(yytext); return IDENTIFIER;
 <elfheader>"================"	return H_SEPARATOR;
 <elfheader>"arch = "		return H_ARCH;
 <elfheader>"code version = "	return H_CODEVERSION;
@@ -128,7 +131,7 @@ newlines	{newline}+
 
 
 	/*Header tokens*/
-<ptxheader>[[:alnum:]_]+ yylval.string_value = strdup(yytext); return IDENTIFIER;
+<ptxheader>[[:alnum:]_]+ yylval->string_value = strdup(yytext); return IDENTIFIER;
 <ptxheader>"================"	return H_SEPARATOR;
 <ptxheader>"arch = "		return H_ARCH;
 <ptxheader>"code version = "	return H_CODEVERSION;
@@ -143,7 +146,7 @@ newlines	{newline}+
 
 
 	/* Looking for the identifier (filename) then the header is done */
-<endheader>{notnewline}+	yylval.string_value = strdup(yytext); yy_pop_state(); return IDENTIFIER;
+<endheader>{notnewline}+	yylval->string_value = strdup(yytext); yy_pop_state(yyscanner); return IDENTIFIER;
 
 
 
@@ -153,11 +156,12 @@ newlines	{newline}+
 <<EOF>> BEGIN(INITIAL);return 0;
 
 	/*No other rule matched. Throw an error*/
-.	cuobjdump_error("Invalid token");
+.	cuobjdump_error(yyscanner, parser, cuobjdumpSectionList, "Invalid token");
 
 %%
-void cuobjdump_error(const char* message)
+void cuobjdump_error(yyscan_t yyscanner, struct cuobjdump_parser* parser, std::list<cuobjdumpSection*> &cuobjdumpSectionList, const char* msg)
 {
-	printf("   %s near \"%s\"",message, yytext);
-	printf(" on line %i\n",yylineno); 
+    struct yyguts_t * yyg = (struct yyguts_t*)yyscanner;
+	printf("   %s near \"%s\"",msg, yytext);
+	printf(" on line %i\n",yylineno);
 }
diff --git a/libcuda/cuobjdump.y b/libcuda/cuobjdump.y
index 31760f7..8d1bca6 100644
--- a/libcuda/cuobjdump.y
+++ b/libcuda/cuobjdump.y
@@ -29,24 +29,31 @@
 %{
 #include <stdio.h>
 
-int yylex(void);
-void yyerror(const char*);
-extern void addCuobjdumpSection(int sectiontype);
-void setCuobjdumparch(const char* arch);
-void setCuobjdumpidentifier(const char* identifier);
-void setCuobjdumpptxfilename(const char* filename);
-void setCuobjdumpelffilename(const char* filename);
-void setCuobjdumpsassfilename(const char* filename);
-int elfserial = 1;
-int ptxserial = 1;
-FILE *ptxfile;
-FILE *elffile;
-FILE *sassfile;
-char filename [1024];
+typedef void * yyscan_t;
+#include "cuobjdump.h"
+
+extern void addCuobjdumpSection(int sectiontype, std::list<cuobjdumpSection*> &cuobjdumpSectionList);
+void setCuobjdumparch(const char* arch, std::list<cuobjdumpSection*> &cuobjdumpSectionList);
+void setCuobjdumpidentifier(const char* identifier, std::list<cuobjdumpSection*> &cuobjdumpSectionList);
+void setCuobjdumpptxfilename(const char* filename, std::list<cuobjdumpSection*> &cuobjdumpSectionList);
+void setCuobjdumpelffilename(const char* filename, std::list<cuobjdumpSection*> &cuobjdumpSectionList);
+void setCuobjdumpsassfilename(const char* filename, std::list<cuobjdumpSection*> &cuobjdumpSectionList);
 %}
+%define api.pure full
+%parse-param {yyscan_t scanner}
+%parse-param {struct cuobjdump_parser* parser}
+%parse-param {std::list<cuobjdumpSection*> &cuobjdumpSectionList}
+%lex-param {yyscan_t scanner}
+%lex-param {struct cuobjdump_parser* parser}
+%lex-param {std::list<cuobjdumpSection*> &cuobjdumpSectionList}
+
 %union {
 	char* string_value;
 }
+%{
+int yylex(YYSTYPE * yylval_param, yyscan_t yyscanner, struct cuobjdump_parser* parser, std::list<cuobjdumpSection*> &cuobjdumpSectionList);
+void yyerror(yyscan_t yyscanner, struct cuobjdump_parser* parser, std::list<cuobjdumpSection*> &cuobjdumpSectionList, const char* msg);
+%}
 %token <string_value> H_SEPARATOR H_ARCH H_CODEVERSION H_PRODUCER H_HOST H_COMPILESIZE H_IDENTIFIER H_UNKNOWN H_COMPRESSED
 %token <string_value> CODEVERSION
 %token <string_value> STRING
@@ -70,25 +77,25 @@ emptylines	:	emptylines NEWLINE
 			|	;
 
 section :	PTXHEADER {
-				addCuobjdumpSection(0);
-				snprintf(filename, 1024, "_cuobjdump_%d.ptx", ptxserial++);
-				ptxfile = fopen(filename, "w");
-				setCuobjdumpptxfilename(filename);
+				addCuobjdumpSection(0, cuobjdumpSectionList);
+				snprintf(parser->filename, 1024, "_cuobjdump_%d.ptx", parser->ptxserial++);
+				parser->ptxfile = fopen(parser->filename, "w");
+				setCuobjdumpptxfilename(parser->filename, cuobjdumpSectionList);
 			} headerinfo compressedkeyword identifier ptxcode {
-				fclose(ptxfile);
+				fclose(parser->ptxfile);
 			}
 		|	ELFHEADER {
-				addCuobjdumpSection(1);
-				snprintf(filename, 1024, "_cuobjdump_%d.elf", elfserial);
-				elffile = fopen(filename, "w");
-				setCuobjdumpelffilename(filename);
+				addCuobjdumpSection(1, cuobjdumpSectionList);
+				snprintf(parser->filename, 1024, "_cuobjdump_%d.elf", parser->elfserial);
+				parser->elffile = fopen(parser->filename, "w");
+				setCuobjdumpelffilename(parser->filename, cuobjdumpSectionList);
 			} headerinfo compressedkeyword identifier elfcode {
-				fclose(elffile);
-				snprintf(filename, 1024, "_cuobjdump_%d.sass", elfserial++);
-				sassfile = fopen(filename, "w");
-				setCuobjdumpsassfilename(filename);
+				fclose(parser->elffile);
+				snprintf(parser->filename, 1024, "_cuobjdump_%d.sass", parser->elfserial++);
+				parser->sassfile = fopen(parser->filename, "w");
+				setCuobjdumpsassfilename(parser->filename, cuobjdumpSectionList);
 			} sasscode { 
-				fclose(sassfile);
+				fclose(parser->sassfile);
 			};
 
 headerinfo :	H_SEPARATOR NEWLINE
@@ -96,27 +103,27 @@ headerinfo :	H_SEPARATOR NEWLINE
 				H_CODEVERSION CODEVERSION NEWLINE
 				H_PRODUCER H_UNKNOWN NEWLINE
 				H_HOST IDENTIFIER NEWLINE
-				H_COMPILESIZE IDENTIFIER  {setCuobjdumparch($4);};
+				H_COMPILESIZE IDENTIFIER  {setCuobjdumparch($4, cuobjdumpSectionList);};
 			|   H_SEPARATOR NEWLINE
 				H_ARCH IDENTIFIER NEWLINE
 				H_CODEVERSION CODEVERSION NEWLINE
 				H_PRODUCER IDENTIFIER NEWLINE
 				H_HOST IDENTIFIER NEWLINE
-				H_COMPILESIZE IDENTIFIER {setCuobjdumparch($4);};
+				H_COMPILESIZE IDENTIFIER {setCuobjdumparch($4, cuobjdumpSectionList);};
 
-identifier : H_IDENTIFIER FILENAME emptylines {setCuobjdumpidentifier($2);}
-			 |	{setCuobjdumpidentifier("default");};
+identifier : H_IDENTIFIER FILENAME emptylines {setCuobjdumpidentifier($2, cuobjdumpSectionList);}
+			 |	{setCuobjdumpidentifier("default", cuobjdumpSectionList);};
 
 compressedkeyword : H_COMPRESSED emptylines
                     | ;
 
-ptxcode :	ptxcode PTXLINE {fprintf(ptxfile, "%s", $2);}
+ptxcode :	ptxcode PTXLINE {fprintf(parser->ptxfile, "%s", $2);}
 		|	;
 
-elfcode :	elfcode ELFLINE {fprintf(elffile, "%s", $2);}
+elfcode :	elfcode ELFLINE {fprintf(parser->elffile, "%s", $2);}
 		|	;
 
-sasscode :	sasscode SASSLINE {fprintf(sassfile, "%s", $2);}
+sasscode :	sasscode SASSLINE {fprintf(parser->sassfile, "%s", $2);}
 		 |	;
 
 
diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h
new file mode 100644
index 0000000..45c5cdd
--- /dev/null
+++ b/libcuda/gpgpu_context.h
@@ -0,0 +1,71 @@
+#ifndef __gpgpu_context_h__
+#define __gpgpu_context_h__
+#include "cuda_api_object.h"
+#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"
+#include "../src/cuda-sim/ptx-stats.h"
+
+class gpgpu_context {
+    public:
+	gpgpu_context() {
+	    g_global_allfiles_symbol_table = NULL;
+	    sm_next_access_uid=0;
+	    warp_inst_sm_next_uid=0;
+	    operand_info_sm_next_uid = 1;
+	    kernel_info_m_next_uid = 1;
+	    g_num_ptx_inst_uid = 0;
+	    g_ptx_cta_info_uid = 1;
+	    symbol_sm_next_uid = 1;
+	    function_info_sm_next_uid = 1;
+	    debug_tensorcore = 0;
+	    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);
+	    device_runtime = new cuda_device_runtime(this);
+	    stats = new ptx_stats(this);
+	}
+	// global list
+	symbol_table *g_global_allfiles_symbol_table;
+	const char *g_filename;
+	unsigned sm_next_access_uid;
+	unsigned warp_inst_sm_next_uid;
+	unsigned operand_info_sm_next_uid;//uid for operand_info
+	unsigned kernel_info_m_next_uid;//uid for kernel_info_t
+	unsigned g_num_ptx_inst_uid; //uid for ptx inst inside ptx_instruction
+	unsigned long long g_ptx_cta_info_uid;
+	unsigned symbol_sm_next_uid; //uid for symbol
+	unsigned function_info_sm_next_uid;
+	std::vector<ptx_instruction*> s_g_pc_to_insn; // a direct mapping from PC to instruction
+	bool debug_tensorcore;
+
+	// objects pointers for each file
+	cuda_runtime_api* api;
+	ptxinfo_data* ptxinfo;
+	ptx_recognizer* ptx_parser;
+	GPGPUsim_ctx* the_gpgpusim;
+	cuda_sim* func_sim;
+	cuda_device_runtime* device_runtime;
+	ptx_stats* stats;
+	// member function list
+	void cuobjdumpParseBinary(unsigned int handle);
+	class symbol_table *gpgpu_ptx_sim_load_ptx_from_string( const char *p, unsigned source_num );
+	class symbol_table *gpgpu_ptx_sim_load_ptx_from_filename( const char *filename );
+	void gpgpu_ptx_info_load_from_filename( const char *filename, unsigned sm_version);
+	void gpgpu_ptxinfo_load_from_string( const char *p_for_info, unsigned source_num, unsigned sm_version=20, int no_of_ptx=0 );
+	void print_ptx_file( const char *p, unsigned source_num, const char *filename );
+	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);
+	const ptx_instruction* pc_to_instruction(unsigned pc); 
+	const warp_inst_t *ptx_fetch_inst( address_type pc );
+	unsigned translate_pc_to_ptxlineno(unsigned pc);
+};
+gpgpu_context* GPGPU_Context();
+
+#endif /* __gpgpu_context_h__ */