Merge pull request #12 from echoedit/dev

Dev

7 files changed, 138 insertions, 126 deletions

diff --git a/libcuda/cuda_api_object.h b/libcuda/cuda_api_object.h
index 2001f91..41337c6 100644
--- a/libcuda/cuda_api_object.h
+++ b/libcuda/cuda_api_object.h
@@ -1,14 +1,68 @@
 #ifndef __cuda_api_object_h__
 #define __cuda_api_object_h__
+
+#include <list>
+#include <map>
+#include <set>
+#include <string>
+
 class cuobjdumpSection;
+class cuobjdumpELFSection;
+class cuobjdumpPTXSection;
+class symbol_table;
+class gpgpu_ptx_sim_arg;
 class kernel_config;
+class kernel_info_t;
+
+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;
 
 
 class cuda_runtime_api {
     public:
+	cuda_runtime_api() {
+	    g_glbmap = NULL;
+	}
 	// 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;
 	// member function list
+	void cuobjdumpInit();
+	void extract_code_using_cuobjdump();
+	void extract_ptx_files_using_cuobjdump(CUctx_st *context);
+	void cuobjdumpParseBinary(unsigned int handle);
+	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);
+	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 );
 };
 #endif /* __cuda_api_object_h__ */
diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index 500b5f2..71340b3 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -626,7 +626,7 @@ static int get_app_cuda_version() {
 }
 
 //! Keep track of the association between filename and cubin handle
-void gpgpu_context::cuobjdumpRegisterFatBinary(unsigned int handle, const char* filename, CUctx_st *context){
+void cuda_runtime_api::cuobjdumpRegisterFatBinary(unsigned int handle, const char* filename, CUctx_st *context){
 	fatbinmap[handle] = filename;
 }
 
@@ -705,8 +705,8 @@ void** cudaRegisterFatBinaryInternal( void *fatCubin, gpgpu_context* gpgpu_ctx =
 		 * then for next calls, only returns the appropriate number
 		 */
 		assert(fat_cubin_handle >= 1);
-		if (fat_cubin_handle==1) ctx->cuobjdumpInit();
-		ctx->cuobjdumpRegisterFatBinary(fat_cubin_handle, filename, context);
+		if (fat_cubin_handle==1) ctx->api->cuobjdumpInit();
+		ctx->api->cuobjdumpRegisterFatBinary(fat_cubin_handle, filename, context);
 
 		return (void**)fat_cubin_handle;
 	} 
@@ -803,7 +803,7 @@ void cudaRegisterFunctionInternal(
 	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);
+		ctx->api->cuobjdumpParseBinary(fat_cubin_handle);
 	context->register_function( fat_cubin_handle, hostFun, deviceFun );
 }
 
@@ -830,7 +830,7 @@ void cudaRegisterVarInternal(
 	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);
+		ctx->api->cuobjdumpParseBinary((unsigned)(unsigned long long)fatCubinHandle);
 	fflush(stdout);
 	if ( constant && !global && !ext ) {
 		gpgpu_ptx_sim_register_const_variable(hostVar,deviceName,size);
@@ -894,7 +894,7 @@ cudaError_t cudaLaunchInternal( const char *hostFun, gpgpu_context* gpgpu_ctx =
 	struct CUstream_st *stream = config.get_stream();
 	printf("\nGPGPU-Sim PTX: cudaLaunch for 0x%p (mode=%s) on stream %u\n", hostFun,
 			g_ptx_sim_mode?"functional simulation":"performance simulation", stream?stream->get_uid():0 );
-	kernel_info_t *grid = ctx->gpgpu_cuda_ptx_sim_init_grid(hostFun,config.get_args(),config.grid_dim(),config.block_dim(),context);
+	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();
@@ -965,7 +965,7 @@ cudaError_t cudaMallocInternal(void **devPtr, size_t size, gpgpu_context* gpgpu_
 	*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);
-        ctx->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;
@@ -988,7 +988,7 @@ cudaError_t cudaMallocHostInternal(void **ptr, size_t size, gpgpu_context* gpgpu
 	*ptr = malloc(size);
 	if ( *ptr  ) {
 		//track pinned memory size allocated in the host so that same amount of memory is also allocated in GPU.
-		ctx->pinned_memory_size[*ptr]=size;
+		ctx->api->pinned_memory_size[*ptr]=size;
 		return g_last_cudaError = cudaSuccess;
 	} else {
 		return g_last_cudaError = cudaErrorMemoryAllocation;
@@ -1014,16 +1014,16 @@ cudaError_t cudaHostGetDevicePointerInternal(void **pDevice, void *pHost, unsign
 	flags=0;
 	CUctx_st* context = GPGPUSim_Context();
 	gpgpu_t *gpu = context->get_device()->get_gpgpu();
-	std::map<void *, size_t>::const_iterator i = ctx->pinned_memory_size.find(pHost);
-	assert(i != ctx->pinned_memory_size.end());
+	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->g_mallocPtr_Size[(unsigned long long)*pDevice] = size;
+        ctx->api->g_mallocPtr_Size[(unsigned long long)*pDevice] = size;
     }
 	if ( *pDevice  ) {
-		ctx->pinned_memory[pHost]=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;
@@ -1050,7 +1050,7 @@ cudaError_t cudaGLMapBufferObjectInternal(void** devPtr, GLuint bufferObj, gpgpu
 	GLint buffer_size=0;
 	CUctx_st* context = GPGPUSim_Context();
 
-	glbmap_entry_t *p = ctx->g_glbmap;
+	glbmap_entry_t *p = ctx->api->g_glbmap;
 	while ( p && p->m_bufferObj != bufferObj )
 		p = p->m_next;
 	if ( p == NULL ) {
@@ -1061,8 +1061,8 @@ cudaError_t cudaGLMapBufferObjectInternal(void** devPtr, GLuint bufferObj, gpgpu
 
 		// 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->g_glbmap;
-		ctx->g_glbmap = n;
+		n->m_next = ctx->api->g_glbmap;
+		ctx->api->g_glbmap = n;
 
 		// initialize entry
 		n->m_bufferObj = bufferObj;
@@ -1129,7 +1129,7 @@ cuLinkAddFileInternal(CUlinkState state, CUjitInputType type, const char *path,
     strcat(file,path);
 	symbol_table *symtab = gpgpu_ptx_sim_load_ptx_from_filename( file );
     std::string fname(path);
-    ctx->name_symtab[fname] = symtab;
+    ctx->api->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());
@@ -1154,7 +1154,7 @@ cudaError_t cudaHostAllocInternal(void **pHost,  size_t bytes, unsigned int flag
 	*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->pinned_memory_size[*pHost]=bytes;
+	ctx->api->pinned_memory_size[*pHost]=bytes;
 	if( *pHost )
 		return g_last_cudaError = cudaSuccess;
 	else
@@ -2431,7 +2431,7 @@ __host__ cudaError_t CUDARTAPI cudaGetExportTable(const void **ppExportTable, co
 //#include "../../cuobjdump_to_ptxplus/cuobjdump_parser.h"
 
 //extracts all ptx files from binary and dumps into prog_name.unique_no.sm_<>.ptx files
-void gpgpu_context::extract_ptx_files_using_cuobjdump(CUctx_st *context){
+void cuda_runtime_api::extract_ptx_files_using_cuobjdump(CUctx_st *context){
     extern bool g_cdp_enabled;
     char command[1000];
     char *pytorch_bin = getenv("PYTORCH_BIN");
@@ -2509,7 +2509,7 @@ void gpgpu_context::extract_ptx_files_using_cuobjdump(CUctx_st *context){
  *	It is also responsible for extracting the libraries linked to the binary if the option is
  *	enabled
  * */
-void gpgpu_context::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();
 
@@ -2633,7 +2633,7 @@ void gpgpu_context::extract_code_using_cuobjdump(){
                     fclose(cuobjdump_in);
                     std::getline(libsf, line);
             }
-            api->libSectionList = cuobjdumpSectionList;
+            libSectionList = cuobjdumpSectionList;
 
             //Restore the original section list
             cuobjdumpSectionList = tmpsl;
@@ -2679,7 +2679,7 @@ void printSectionList(std::list<cuobjdumpSection*> sl) {
 }
 
 //! Remove unecessary sm versions from the section list
-std::list<cuobjdumpSection*> gpgpu_context::pruneSectionList(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)
@@ -2732,7 +2732,7 @@ std::list<cuobjdumpSection*> gpgpu_context::pruneSectionList(CUctx_st *context)
 }
 
 //! Merge all PTX sections that have a specific identifier into one file
-std::list<cuobjdumpSection*> gpgpu_context::mergeMatchingSections(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;
@@ -2775,7 +2775,7 @@ std::list<cuobjdumpSection*> gpgpu_context::mergeMatchingSections(std::string id
 }
 
 //! Merge any PTX sections with matching identifiers
-std::list<cuobjdumpSection*> gpgpu_context::mergeSections(){
+std::list<cuobjdumpSection*> cuda_runtime_api::mergeSections(){
 	std::vector<std::string> identifier;
 	cuobjdumpPTXSection* ptxsection;
 
@@ -2822,10 +2822,10 @@ cuobjdumpELFSection* findELFSectionInList(std::list<cuobjdumpSection*> sectionli
 }
 
 //! Find an ELF section in all the known lists
-cuobjdumpELFSection* gpgpu_context::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(api->libSectionList, identifier);
+	sec = findELFSectionInList(libSectionList, identifier);
 	if (sec!=NULL)return sec;
 	std::cout << "Could not find " << identifier << std::endl;
 	assert(0 && "Could not find the required ELF section");
@@ -2857,10 +2857,10 @@ cuobjdumpPTXSection* findPTXSectionInList(std::list<cuobjdumpSection*> &sectionl
 }
 
 //! Find an PTX section in all the known lists
-cuobjdumpPTXSection* gpgpu_context::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(api->libSectionList, identifier);
+	sec = findPTXSectionInList(libSectionList, identifier);
 	if (sec!=NULL)return sec;
 	std::cout << "Could not find " << identifier << std::endl;
 	assert(0 && "Could not find the required PTX section");
@@ -2870,7 +2870,7 @@ cuobjdumpPTXSection* gpgpu_context::findPTXSection(const std::string identifier)
 
 
 //! Extract the code using cuobjdump and remove unnecessary sections
-void gpgpu_context::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");
@@ -2882,7 +2882,7 @@ void gpgpu_context::cuobjdumpInit(){
 
 
 //! Either submit PTX for simulation or convert SASS to PTXPlus and submit it
-void gpgpu_context::cuobjdumpParseBinary(unsigned int handle){
+void cuda_runtime_api::cuobjdumpParseBinary(unsigned int handle){
 
 	CUctx_st *context = GPGPUSim_Context();
 	if(fatbin_registered[handle]) return;
@@ -3176,7 +3176,7 @@ cudaError_t cudaGLUnmapBufferObject(GLuint bufferObj)
     }
 #ifdef OPENGL_SUPPORT
 	CUctx_st* ctx = GPGPUSim_Context();
-	glbmap_entry_t *p = ctx->g_glbmap;
+	glbmap_entry_t *p = ctx->api->g_glbmap;
 	while ( p && p->m_bufferObj != bufferObj )
 		p = p->m_next;
 	if ( p == NULL )
@@ -3531,7 +3531,7 @@ static int load_constants( symbol_table *symtab, addr_t min_gaddr, gpgpu_t *gpu
 	return nc_bytes;
 }
 
-kernel_info_t * gpgpu_context::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,
diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h
index 7569ea6..4622f00 100644
--- a/libcuda/gpgpu_context.h
+++ b/libcuda/gpgpu_context.h
@@ -1,68 +1,15 @@
 #ifndef __gpgpu_context_h__
 #define __gpgpu_context_h__
-#include <list>
-#include <map>
-#include <set>
-#include <string>
 #include "cuda_api_object.h"
 
-class cuobjdumpSection;
-class cuobjdumpELFSection;
-class cuobjdumpPTXSection;
-class symbol_table;
-class gpgpu_ptx_sim_arg;
-class kernel_info_t;
-
-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;
-
 class gpgpu_context {
     public:
 	gpgpu_context() {
 	    api = new cuda_runtime_api();
-	    g_glbmap = NULL;
 	}
 	// global list
-	std::list<cuobjdumpSection*> cuobjdumpSectionList;
-	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;
 	// objects pointers for each file
 	cuda_runtime_api* api;
 	// member function list
-	void cuobjdumpInit();
-	void cuobjdumpParseBinary(unsigned int handle);
-	void extract_code_using_cuobjdump();
-	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);
-	void extract_ptx_files_using_cuobjdump(CUctx_st *context);
-	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 );
-
 };
 #endif /* __gpgpu_context_h__ */
diff --git a/src/cuda-sim/Makefile b/src/cuda-sim/Makefile
index 999dad7..85d1c8c 100644
--- a/src/cuda-sim/Makefile
+++ b/src/cuda-sim/Makefile
@@ -79,16 +79,16 @@ libgpgpu_ptx_sim.a: $(OBJS)
 	ar rcs $(OUTPUT_DIR)/libgpgpu_ptx_sim.a $(OUTPUT_DIR)/ptx.tab.o $(OUTPUT_DIR)/lex.ptx_.o $(OUTPUT_DIR)/ptxinfo.tab.o $(OUTPUT_DIR)/lex.ptxinfo_.o $(OBJS)
 
 $(OUTPUT_DIR)/ptx.tab.o: $(OUTPUT_DIR)/ptx.tab.c
-	$(CPP) -c $(OPT) -DYYDEBUG $(OUTPUT_DIR)/ptx.tab.c -o $(OUTPUT_DIR)/ptx.tab.o
+	$(CPP) -c $(CXX_OPT) -DYYDEBUG $(OUTPUT_DIR)/ptx.tab.c -o $(OUTPUT_DIR)/ptx.tab.o
 
 $(OUTPUT_DIR)/lex.ptx_.o: $(OUTPUT_DIR)/lex.ptx_.c 
-	$(CPP) -c $(OPT) $(OUTPUT_DIR)/lex.ptx_.c -o $(OUTPUT_DIR)/lex.ptx_.o
+	$(CPP) -c $(CXX_OPT) $(OUTPUT_DIR)/lex.ptx_.c -o $(OUTPUT_DIR)/lex.ptx_.o
 
 $(OUTPUT_DIR)/ptxinfo.tab.o: $(OUTPUT_DIR)/ptxinfo.tab.c
-	$(CPP) -c $(OPT) -DYYDEBUG $(OUTPUT_DIR)/ptxinfo.tab.c -o $(OUTPUT_DIR)/ptxinfo.tab.o
+	$(CPP) -c $(CXX_OPT) -DYYDEBUG $(OUTPUT_DIR)/ptxinfo.tab.c -o $(OUTPUT_DIR)/ptxinfo.tab.o
 
 $(OUTPUT_DIR)/lex.ptxinfo_.o: $(OUTPUT_DIR)/lex.ptxinfo_.c $(OUTPUT_DIR)/ptxinfo.tab.c
-	$(CPP) -c $(OPT) $(OUTPUT_DIR)/lex.ptxinfo_.c -o $(OUTPUT_DIR)/lex.ptxinfo_.o
+	$(CPP) -c $(CXX_OPT) $(OUTPUT_DIR)/lex.ptxinfo_.c -o $(OUTPUT_DIR)/lex.ptxinfo_.o
 
 $(OUTPUT_DIR)/ptx.tab.c: ptx.y
 	bison --name-prefix=ptx_ -v -d ptx.y --file-prefix=$(OUTPUT_DIR)/ptx
diff --git a/src/cuda-sim/ptx.y b/src/cuda-sim/ptx.y
index 8447693..837fbe9 100644
--- a/src/cuda-sim/ptx.y
+++ b/src/cuda-sim/ptx.y
@@ -229,7 +229,6 @@ class ptx_recognizer;
 	#include <string.h>
 	#include <math.h>
 	void syntax_not_implemented(yyscan_t yyscanner, ptx_recognizer* recognizer);
-	extern int g_func_decl;
 	int ptx_lex(YYSTYPE * yylval_param, yyscan_t yyscanner, ptx_recognizer* recognizer);
 	int ptx_error( yyscan_t yyscanner, ptx_recognizer* recognizer, const char *s );
 %}
@@ -260,21 +259,21 @@ block_spec_list: block_spec
 
 function_decl: function_decl_header LEFT_PAREN { recognizer->start_function($1); recognizer->func_header_info("(");} param_entry RIGHT_PAREN {recognizer->func_header_info(")");} function_ident_param { $$ = recognizer->reset_symtab(); }
 	| function_decl_header { recognizer->start_function($1); } function_ident_param { $$ = recognizer->reset_symtab(); }
-	| function_decl_header { recognizer->start_function($1); recognizer->add_function_name(""); g_func_decl=0; $$ = recognizer->reset_symtab(); }
+	| function_decl_header { recognizer->start_function($1); recognizer->add_function_name(""); recognizer->g_func_decl=0; $$ = recognizer->reset_symtab(); }
 	;
 
-function_ident_param: IDENTIFIER { recognizer->add_function_name($1); } LEFT_PAREN {recognizer->func_header_info("(");} param_list RIGHT_PAREN { g_func_decl=0; recognizer->func_header_info(")"); }
-	| IDENTIFIER { recognizer->add_function_name($1); g_func_decl=0; }
+function_ident_param: IDENTIFIER { recognizer->add_function_name($1); } LEFT_PAREN {recognizer->func_header_info("(");} param_list RIGHT_PAREN { recognizer->g_func_decl=0; recognizer->func_header_info(")"); }
+	| IDENTIFIER { recognizer->add_function_name($1); recognizer->g_func_decl=0; }
 	;
 
-function_decl_header: ENTRY_DIRECTIVE { $$ = 1; g_func_decl=1; recognizer->func_header(".entry"); }
-	| VISIBLE_DIRECTIVE ENTRY_DIRECTIVE { $$ = 1; g_func_decl=1; recognizer->func_header(".entry"); }
-	| WEAK_DIRECTIVE ENTRY_DIRECTIVE { $$ = 1; g_func_decl=1; recognizer->func_header(".entry"); }
-	| FUNC_DIRECTIVE { $$ = 0; g_func_decl=1; recognizer->func_header(".func"); }
-	| VISIBLE_DIRECTIVE FUNC_DIRECTIVE { $$ = 0; g_func_decl=1; recognizer->func_header(".func"); }
-	| WEAK_DIRECTIVE FUNC_DIRECTIVE { $$ = 0; g_func_decl=1; recognizer->func_header(".func"); }
-	| EXTERN_DIRECTIVE FUNC_DIRECTIVE { $$ = 2; g_func_decl=1; recognizer->func_header(".func"); }
-	| WEAK_DIRECTIVE FUNC_DIRECTIVE { $$ = 0; g_func_decl=1; recognizer->func_header(".func"); }
+function_decl_header: ENTRY_DIRECTIVE { $$ = 1; recognizer->g_func_decl=1; recognizer->func_header(".entry"); }
+	| VISIBLE_DIRECTIVE ENTRY_DIRECTIVE { $$ = 1; recognizer->g_func_decl=1; recognizer->func_header(".entry"); }
+	| WEAK_DIRECTIVE ENTRY_DIRECTIVE { $$ = 1; recognizer->g_func_decl=1; recognizer->func_header(".entry"); }
+	| FUNC_DIRECTIVE { $$ = 0; recognizer->g_func_decl=1; recognizer->func_header(".func"); }
+	| VISIBLE_DIRECTIVE FUNC_DIRECTIVE { $$ = 0; recognizer->g_func_decl=1; recognizer->func_header(".func"); }
+	| WEAK_DIRECTIVE FUNC_DIRECTIVE { $$ = 0; recognizer->g_func_decl=1; recognizer->func_header(".func"); }
+	| EXTERN_DIRECTIVE FUNC_DIRECTIVE { $$ = 2; recognizer->g_func_decl=1; recognizer->func_header(".func"); }
+	| WEAK_DIRECTIVE FUNC_DIRECTIVE { $$ = 0; recognizer->g_func_decl=1; recognizer->func_header(".func"); }
 	;
 
 param_list: /*empty*/
diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc
index a61dc16..897beaf 100644
--- a/src/cuda-sim/ptx_parser.cc
+++ b/src/cuda-sim/ptx_parser.cc
@@ -49,7 +49,6 @@ void set_ptx_warp_size(const struct core_config * warp_size)
 
 static bool g_debug_ir_generation=false;
 const char *g_filename;
-unsigned g_max_regs_per_thread = 0;
 
 // the program intermediate representation...
 static symbol_table *g_global_allfiles_symbol_table = NULL;
@@ -61,27 +60,6 @@ static symbol *g_last_symbol = NULL;
 
 int g_error_detected = 0;
 
-// type specifier stuff:
-memory_space_t g_space_spec = undefined_space;
-memory_space_t g_ptr_spec = undefined_space;
-int g_scalar_type_spec = -1;
-int g_vector_spec = -1;
-int g_alignment_spec = -1;
-int g_extern_spec = 0;
-
-// variable declaration stuff:
-type_info *g_var_type = NULL;
-
-// instruction definition stuff:
-const symbol *g_pred;
-int g_neg_pred;
-int g_pred_mod;
-symbol *g_label;
-int g_opcode = -1;
-std::list<operand_info> g_operands;
-std::list<int> g_options;
-std::list<int> g_wmma_options;
-std::list<int> g_scalar_type;
 
 #define PTX_PARSE_DPRINTF(...) \
    if( g_debug_ir_generation ) { \
@@ -362,10 +340,6 @@ bool check_for_duplicates( const char *identifier )
 extern std::set<std::string>   g_globals;
 extern std::set<std::string>   g_constants;
 
-int g_func_decl = 0;
-int g_ident_add_uid = 0;
-unsigned g_const_alloc = 1;
-
 // Returns padding that needs to be inserted ahead of address to make it aligned to min(size, maxalign)
 /*
  * @param address the address in bytes
diff --git a/src/cuda-sim/ptx_parser.h b/src/cuda-sim/ptx_parser.h
index fb1b7f2..c52fe20 100644
--- a/src/cuda-sim/ptx_parser.h
+++ b/src/cuda-sim/ptx_parser.h
@@ -29,6 +29,7 @@
 #define ptx_parser_INCLUDED
 
 #include "../abstract_hardware_model.h"
+#include "ptx_ir.h"
 extern const char *g_filename;
 extern int g_error_detected;
 
@@ -44,7 +45,20 @@ class ptx_recognizer {
 	    scanner = NULL;
 	    g_size = -1;
 	    g_add_identifier_cached__identifier = NULL;
+	    g_alignment_spec = -1;
+	    g_var_type = NULL;
+	    g_opcode = -1;
+	    g_space_spec = undefined_space;
+	    g_ptr_spec = undefined_space;
+	    g_scalar_type_spec = -1;
+	    g_vector_spec = -1;
+	    g_extern_spec = 0;
+	    g_func_decl = 0;
+	    g_ident_add_uid = 0;
+	    g_const_alloc = 1;
+	    g_max_regs_per_thread = 0;
 	}
+	// global list
 	yyscan_t scanner;
 #define LINEBUF_SIZE (4*1024)
 	char linebuf[LINEBUF_SIZE];
@@ -53,7 +67,31 @@ class ptx_recognizer {
 	char *g_add_identifier_cached__identifier;
 	int g_add_identifier_cached__array_dim;
 	int g_add_identifier_cached__array_ident;
+	int g_alignment_spec;
+	// variable declaration stuff:
+	type_info *g_var_type;
+	// instruction definition stuff:
+	const symbol *g_pred;
+	int g_neg_pred;
+	int g_pred_mod;
+	symbol *g_label;
+	int g_opcode;
+	std::list<operand_info> g_operands;
+	std::list<int> g_options;
+	std::list<int> g_wmma_options;
+	std::list<int> g_scalar_type;
+	// type specifier stuff:
+	memory_space_t g_space_spec;
+	memory_space_t g_ptr_spec;
+	int g_scalar_type_spec;
+	int g_vector_spec;
+	int g_extern_spec;
+	int g_func_decl;
+	int g_ident_add_uid;
+	unsigned g_const_alloc;
+	unsigned g_max_regs_per_thread;
 
+	// member function list
 	void init_directive_state();
 	void init_instruction_state();
 	void start_function( int entry_point );