Merging

//depot/gpgpu_sim_research/fermi-test/cuobjdump_to_ptxplus/...

to //depot/gpgpu_sim_research/fermi/distribution/cuobjdump_to_ptxplus/...

[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 12022]

13 files changed, 4058 insertions, 0 deletions

diff --git a/cuobjdump_to_ptxplus/Makefile b/cuobjdump_to_ptxplus/Makefile
new file mode 100644
index 0000000..f58a0fd
--- /dev/null
+++ b/cuobjdump_to_ptxplus/Makefile
@@ -0,0 +1,145 @@
+BOOST_ROOT ?= /usr/include/boost
+
+LEX			= flex
+LEXFLAGS	= -B 
+YACC		= bison
+YFLAGS		= -t -d -v --report=all
+CC			= gcc
+CCFLAGS		= -ggdb -I${BOOST_ROOT} -lboost_regex -lboost_filesystem -fPIC
+CXX			= g++
+CXXFLAGS	= ${CCFLAGS}
+LD			= g++
+LDFLAGS		= ${CCFLAGS}
+
+
+CUOBJDUMP_PARSER_OBJECTS = cuobjdump_parser.o cuobjdump_lexer.o 
+SASS_PARSER_OBJECTS = sass_lexer.o sass_parser.o 
+ELF_PARSER_OBJECTS = elf_lexer.o elf_parser.o
+HEADER_PARSER_OBJECTS = headerParser.o headerLexer.o
+PTX_PARSER_OBJECTS = ptx.tab.o lex.ptx_.o
+
+all: cuobjdump_to_ptxplus
+
+cuobjdump_to_ptxplus: stringList.o decudaInst.o decudaInstList.o cuobjdumpInst.o cuobjdumpInstList.o cuobjdump_to_ptxplus.o $(PTX_PARSER_OBJECTS) $(SASS_PARSER_OBJECTS) $(ELF_PARSER_OBJECTS) $(HEADER_PARSER_OBJECTS) $(CUOBJDUMP_PARSER_OBJECTS) 
+	${LD} ${LDFLAGS} -ocuobjdump_to_ptxplus stringList.o decudaInst.o decudaInstList.o headerParser.o headerLexer.o cuobjdumpInst.o cuobjdumpInstList.o cuobjdump_to_ptxplus.o lex.ptx_.o ptx.tab.o $(SASS_PARSER_OBJECTS) $(ELF_PARSER_OBJECTS) 
+
+stringList.o: ../decuda_to_ptxplus/stringList.cc ../decuda_to_ptxplus/stringList.h
+	${CXX} ${CXXFLAGS} -c -ostringList.o ../decuda_to_ptxplus/stringList.cc
+
+decudaInst.o: ../decuda_to_ptxplus/decudaInst.cc ../decuda_to_ptxplus/decudaInst.h ../decuda_to_ptxplus/stringList.h
+	${CXX} ${CXXFLAGS} -c -odecudaInst.o ../decuda_to_ptxplus/decudaInst.cc
+
+decudaInstList.o: ../decuda_to_ptxplus/decudaInstList.cc ../decuda_to_ptxplus/decudaInstList.h ../decuda_to_ptxplus/decudaInst.h
+	${CXX} ${CXXFLAGS} -c -odecudaInstList.o ../decuda_to_ptxplus/decudaInstList.cc
+
+cuobjdumpInst.o: cuobjdumpInst.cc cuobjdumpInst.h ../decuda_to_ptxplus/decudaInst.h ./decudaInstList.h
+	${CXX} ${CXXFLAGS} -c -ocuobjdumpInst.o cuobjdumpInst.cc
+
+cuobjdumpInstList.o: cuobjdumpInstList.cc cuobjdumpInstList.h ../decuda_to_ptxplus/decudaInstList.h
+	${CXX} ${CXXFLAGS} -c -ocuobjdumpInstList.o cuobjdumpInstList.cc
+
+#sassParser.cpp: sassParser.y ../distribution/decuda_to_ptxplus/decudaInstList.h cuobjdumpInstList.h
+#	${YACC} ${YFLAGS} -osassParser.cpp sassParser.y
+
+#sassParser.hpp: sassParser.cpp
+
+#sassLexer.cpp: sassLexer.l sassParser.hpp
+#	${LEX} ${LEXFLAGS} -osassLexer.cpp sassLexer.l
+
+headerParser.cpp : ../decuda_to_ptxplus/header.y ../decuda_to_ptxplus/decudaInstList.h
+	${YACC} ${YFLAGS} --name-prefix=header_ -oheaderParser.cpp ../decuda_to_ptxplus/header.y
+
+headerParser.hpp : headerParser.cpp
+
+headerLexer.cpp : ../decuda_to_ptxplus/header.l headerParser.hpp
+	${LEX} ${LEXFLAGS} -oheaderLexer.cpp ../decuda_to_ptxplus/header.l
+
+#sassLexer.o: sassLexer.cpp
+#	${CXX} ${CXXFLAGS} -c -Igenerated -osassLexer.o sassLexer.cpp
+#
+#sassParser.o: sassParser.cpp ../distribution/decuda_to_ptxplus/stringList.h ../distribution/decuda_to_ptxplus/decudaInst.h ../distribution/decuda_to_ptxplus/decudaInstList.h cuobjdumpInstList.h
+#	${CXX} ${CXXFLAGS} -c -Igenerated -osassParser.o sassParser.cpp
+
+headerLexer.o : headerLexer.cpp
+	${CXX} ${CXXFLAGS} -c -Igenerated -oheaderLexer.o headerLexer.cpp
+
+headerParser.o : headerParser.cpp ../decuda_to_ptxplus/stringList.h ../decuda_to_ptxplus/decudaInst.h ./decudaInstList.h
+	${CXX} ${CXXFLAGS} -c -Igenerated -oheaderParser.o headerParser.cpp
+
+decudaInstList.h : ../decuda_to_ptxplus/decudaInstList.h ./decudaInst.h
+	rm -f decudaInstList.h
+	ln -s ../decuda_to_ptxplus/decudaInstList.h decudaInstList.h
+
+decudaInst.h : ../decuda_to_ptxplus/decudaInst.h ./stringList.h
+	rm -f decudaInst.h
+	ln -s ../decuda_to_ptxplus/decudaInst.h decudaInst.h
+
+stringList.h : ../decuda_to_ptxplus/stringList.h
+	rm -f stringList.h
+	ln -s ../decuda_to_ptxplus/stringList.h stringList.h
+
+cuobjdump_to_ptxplus.o: cuobjdump_to_ptxplus.cc ../decuda_to_ptxplus/decudaInstList.h cuobjdumpInstList.h
+	${CXX} ${CXXFLAGS} -c -ocuobjdump_to_ptxplus.o cuobjdump_to_ptxplus.cc
+
+lex.ptx_.c : ../src/cuda-sim/ptx.l
+	${LEX} ${LEXFLAGS} ../src/cuda-sim/ptx.l
+
+ptx.tab.c : ../src/cuda-sim/ptx.y
+	${YACC} ${YFLAGS} --name-prefix=ptx_ -v ../src/cuda-sim/ptx.y
+
+ptx.tab.h : ptx.tab.c
+
+lex.ptx_.o : lex.ptx_.c ptx.tab.h
+	rm -f opcodes.def opcodes.h
+	ln -s ../src/cuda-sim/opcodes.def opcodes.def
+	ln -s ../src/cuda-sim/opcodes.h opcodes.h
+	${CXX} ${CXXFLAGS} -c lex.ptx_.c
+
+ptx.tab.o : ptx.tab.c ptx_parser.h
+	${CXX} ${CXXFLAGS} -c ptx.tab.c
+
+.PRECIOUS: %.c %.h %_lexer.c %_parser.c %_parser.h %_parser.cpp %_parser.hpp %_lexer.cpp
+
+%_lexer.c: %.l %_parser.h
+	$(LEX) $(LEXFLAGS) -P $*_ -o$@ $<
+
+%_lexer.cpp: %.l %_parser.hpp
+	$(LEX) $(LEXFLAGS) -P $*_ -o$@ $<
+
+%_parser.c: %.y
+	$(YACC) $(YFLAGS) -p $*_ -o$@ $<
+
+%_parser.cpp: %.y
+	$(YACC) $(YFLAGS) -p $*_ -o$@ $<
+
+# The ':' is needed otherwise make gets confused and doesn't understand that it knows how to make %Parser.hpp from %Parser.cpp.
+# : in bash is equivalent to NOP (no operation)
+%_parser.h: %_parser.c
+	:
+
+%_parser.hpp: %_parser.cpp
+	:
+
+cuobjdump_parser.o: cuobjdump_parser.c
+	${CC} ${CCFLAGS} -c -o $@ $<
+
+cuobjdump_lexer.o: cuobjdump_lexer.c
+	${CC} ${CCFLAGS} -c -o $@ $<
+
+%.o: %.cpp
+	$(CXX) ${CXXFLAGS} -c -o $@ $<
+
+clean:
+	rm -f ./opcodes.def ./opcodes.h
+	rm -f ./lex.ptx_.c ./ptx.tab.c ./ptx.tab.h ./lex.ptx_.o ./ptx.tab.o ptx.output
+	rm -f ./decudaInstList.h ./decudaInst.h ./stringList.h
+	rm -f ./headerParser.o ./headerLexer.o ./headerParser.cpp ./headerParser.hpp ./headerLexer.cpp ./headerParser.output
+	rm -f ./stringList.o ./decudaInst.o ./decudaInstList.o
+	rm -f ./cuobjdumpInstList.o ./cuobjdumpInst.o
+	rm -f ./sass_lexer.o ./sass_parser.o ./cuobjdump_to_ptxplus.o
+	rm -f ./sass_parser.hpp ./sass_parser.output
+	rm -f ./sass_parser.cpp ./sass_lexer.cpp
+	rm -f ./cuobjdump_to_ptxplus
+	rm -f ./cuobjdump_parser.o ./cuobjdump_parser.c ./cuobjdump_parser.output ./cuobjdump_parser.h
+	rm -f ./cuobjdump_lexer.o ./cuobjdump_lexer.c
+	rm -f ./elf_lexer.o ./elf_parser.o ./elf_lexer.cpp ./elf_parser.cpp ./elf_parser.hpp ./elf_parser.output
diff --git a/cuobjdump_to_ptxplus/cuobjdump.l b/cuobjdump_to_ptxplus/cuobjdump.l
new file mode 100644
index 0000000..fcd43f8
--- /dev/null
+++ b/cuobjdump_to_ptxplus/cuobjdump.l
@@ -0,0 +1,96 @@
+	/*Lex file for output of cuobjdump*/
+%{
+#include <stdio.h>
+#include <string.h>
+#include "cuobjdump_parser.h"
+
+#define YY_NEVER_INTERACTIVE 1
+
+#define YYDEBUG 1
+
+#define yylval cuobjdump_lval
+
+void cuobjdump_error(const char*);
+%}
+
+%option stack
+%option noyywrap
+%option yylineno
+
+%s ptxcode
+%s sasscode
+%s elfcode
+%s header
+%s endheader
+
+alpha		[a-zA-Z]
+numeric		[0-9]
+whitespace	[ \t]
+newline		[\n]
+notnewline	[^\n]
+newlines	{newline}+
+%%
+
+
+"ptxasOptions"{notnewline}*{newline}
+[1-9]{numeric}* yylval.string_value = strdup(yytext); return DECIMAL;
+
+
+"Fatbin ptx code:"{newline}	{
+	yy_push_state(ptxcode);
+	yy_push_state(header);
+	yylval.string_value = strdup(yytext);
+	return PTXHEADER;
+}
+"Fatbin elf code:"{newline}	{
+	yy_push_state(elfcode);
+	yy_push_state(header);
+	yylval.string_value = strdup(yytext);
+	return ELFHEADER;
+}
+	/*PTX code tokens*/
+<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); 
+	return SASSLINE;
+}
+<elfcode>{notnewline}*{newline} yylval.string_value = strdup(yytext); return ELFLINE;
+
+	/*SASS code tokens*/
+<sasscode>{notnewline}*{newline} yylval.string_value = strdup(yytext); return SASSLINE;
+
+	/*Header tokens*/
+<header>[[:alnum:]_]+ yylval.string_value = strdup(yytext); return IDENTIFIER;
+<header>"================"	return H_SEPARATOR;
+<header>"arch = "		return H_ARCH;
+<header>"code version = "	return H_CODEVERSION;
+<header>\[{numeric},{numeric}\]	return CODEVERSION;
+<header>"producer = "		return H_PRODUCER;
+<header>"host = "			return H_HOST;
+<header>"compile_size = "	return H_COMPILESIZE;
+<header>"identifier = "		BEGIN(endheader); return H_IDENTIFIER;
+<header>"has debug info"{newline}
+
+	/* Looking for the identifier (filename) then the header is done */
+<endheader>[[:alnum:]_\.]+{newline}	yylval.string_value = strdup(yytext); yy_pop_state(); return FILENAME;
+
+
+{newline}	return NEWLINE;
+
+	/*Reached end of file*/
+<<EOF>> return 0;
+
+	/*No other rule matched. Throw an error*/
+.	cuobjdump_error("Invalid token");
+
+%%
+void cuobjdump_error(const char* message)
+{
+	printf("   "); printf(message); printf(" near \""); printf(yytext); printf("\"");
+	printf(" on line "); 
+	char line[5]; sprintf(line, "%i", yylineno); printf(line);
+	printf("\n");
+}
diff --git a/cuobjdump_to_ptxplus/cuobjdump.y b/cuobjdump_to_ptxplus/cuobjdump.y
new file mode 100644
index 0000000..0b20a83
--- /dev/null
+++ b/cuobjdump_to_ptxplus/cuobjdump.y
@@ -0,0 +1,83 @@
+	/*Yacc file for output of cuobjdump*/
+%{
+#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];
+%}
+%union {
+	char* string_value;
+}
+%token <string_value> H_SEPARATOR H_ARCH H_CODEVERSION H_PRODUCER H_HOST H_COMPILESIZE H_IDENTIFIER
+%token <string_value> CODEVERSION
+%token <string_value> STRING
+%token <string_value> FILENAME
+%token <string_value> DECIMAL
+%token <string_value> PTXHEADER ELFHEADER
+%token <string_value> PTXLINE
+%token <string_value> ELFLINE
+%token <string_value> SASSLINE
+%token <string_value> IDENTIFIER
+%token <string_value> NEWLINE
+
+%%
+
+program :	emptylines section
+		|	program section;
+
+emptylines	:	emptylines NEWLINE
+			|	;
+
+section :	PTXHEADER {
+				addCuobjdumpSection(0);
+				snprintf(filename, 1024, "_cuobjdump_%d.ptx", ptxserial++);
+				ptxfile = fopen(filename, "w");
+				setCuobjdumpptxfilename(filename);
+			} headerinfo ptxcode {
+				fclose(ptxfile);
+			}
+		|	ELFHEADER {
+				addCuobjdumpSection(1);
+				snprintf(filename, 1024, "_cuobjdump_%d.elf", elfserial);
+				elffile = fopen(filename, "w");
+				setCuobjdumpelffilename(filename);
+			} headerinfo elfcode { 
+				fclose(elffile);
+				snprintf(filename, 1024, "_cuobjdump_%d.sass", elfserial++);
+				sassfile = fopen(filename, "w");
+				setCuobjdumpsassfilename(filename);
+			} sasscode { 
+				fclose(sassfile);
+			};
+
+headerinfo :	H_SEPARATOR NEWLINE
+				H_ARCH IDENTIFIER NEWLINE
+				H_CODEVERSION CODEVERSION NEWLINE
+				H_PRODUCER IDENTIFIER NEWLINE
+				H_HOST IDENTIFIER NEWLINE
+				H_COMPILESIZE IDENTIFIER NEWLINE
+				H_IDENTIFIER FILENAME {setCuobjdumparch($4); setCuobjdumpidentifier($19);};
+
+ptxcode :	ptxcode PTXLINE {fprintf(ptxfile, "%s", $2);}
+		|	;
+
+elfcode :	elfcode ELFLINE {fprintf(elffile, "%s", $2);}
+		|	;
+
+sasscode :	sasscode SASSLINE {fprintf(sassfile, "%s", $2)}
+		 |	;
+
+
+%%
diff --git a/cuobjdump_to_ptxplus/cuobjdumpInst.cc b/cuobjdump_to_ptxplus/cuobjdumpInst.cc
new file mode 100644
index 0000000..b014888
--- /dev/null
+++ b/cuobjdump_to_ptxplus/cuobjdumpInst.cc
@@ -0,0 +1,1690 @@
+#include "cuobjdumpInst.h"
+#include <sstream>
+#include <iostream>
+
+extern void output(const char * text);
+
+//Constructor
+cuobjdumpInst::cuobjdumpInst()
+{
+	//initilize everything to empty
+}
+
+bool cuobjdumpInst::checkCubojdumpLabel(std::list<std::string> labelList, std::string label)
+{
+	if(labelList.empty())
+		return false;
+
+	std::list<std::string>::iterator labelIterator;
+
+	for( labelIterator=labelList.begin(); labelIterator!=labelList.end(); labelIterator++ )
+	{
+		if(label.compare(*labelIterator) == 0)
+			return true;
+	}
+
+	return false;
+}
+
+void cuobjdumpInst::printCuobjdumpLabel(std::list<std::string> labelList)
+{
+	if((strcmp(m_label, "")!=0)&&(checkCubojdumpLabel(labelList, m_label))) {
+		output(m_label);
+		output(": ");
+	}
+}
+
+void cuobjdumpInst::printCuobjdumpPredicate()
+{
+	stringListPiece* currentPiece = m_predicate->getListStart();
+	if(currentPiece!=NULL)
+	{
+		output("@$p");
+		char modString[2];
+		modString[0]=currentPiece->stringText[1];
+		modString[1]='\0';
+		output(modString);
+
+		stringListPiece* currentPiece2 = m_predicateModifiers->getListStart();
+		for(int i=0; (i<m_predicateModifiers->getSize())&&(currentPiece2!=NULL); i++)
+		{
+			const char* modString2 = currentPiece2->stringText;
+			char* modString3 = new char[strlen(currentPiece2->stringText)+1];
+
+			for(int i=0; i<strlen(modString2); i++)
+			{
+				if(modString2[i] >= 'A' && modString2[i] <= 'Z')
+					modString3[i] = modString2[i] + 32;
+				else
+					modString3[i] = modString2[i];
+			}
+			modString3[strlen(modString2)]='\0';
+			if( strcmp(modString3, ".not_sign")==0)
+			{
+				output(".nsf");
+			} else if( strcmp(modString3, ".sign")==0)
+			{
+				output(".sf");
+			} else if( strcmp(modString3, ".carry")==0)
+			{
+				output(".cf");
+			}
+			else
+			{
+				output(modString3);
+			}
+			currentPiece = currentPiece2->nextString;
+		}
+		output(" ");
+	}
+}
+
+void cuobjdumpInst::printCuobjdumpTypeModifiers()
+{
+	stringListPiece* currentPiece = m_typeModifiers->getListStart();
+	for(int i=0; (i<m_typeModifiers->getSize())&&(currentPiece!=NULL); i++)
+	{
+		const char* modString = currentPiece->stringText;
+		if( strcmp(modString, ".F16")==0 )
+			output(".f16");
+		else if( strcmp(modString, ".F32")==0 )
+			output(".f32");
+		else if( strcmp(modString, ".F64")==0 )
+			output(".ff64");
+		else if( strcmp(modString, ".S8")==0 )
+			output(".s8");
+		else if( strcmp(modString, ".S16")==0 )
+			output(".s16");
+		else if( strcmp(modString, ".S32")==0 )
+			output(".s32");
+		else if( strcmp(modString, ".S64")==0 )
+			output(".bb64"); //TODO: might have to change to .ss64 in the future.
+		else if( strcmp(modString, ".S128")==0 )
+			output(".bb128"); //TODO: might have to change to .ss64 in the future.
+		else if( strcmp(modString, ".U8")==0 )
+			output(".u8");
+		else if( strcmp(modString, ".U16")==0 )
+			output(".u16");
+		else if( strcmp(modString, ".U32")==0 )
+			output(".u32");
+		else if( strcmp(modString, ".U64")==0 )
+			output(".bb64"); //TODO: might have to change to .ss64 in the future.
+		else if( strcmp(modString, ".HI")==0 )
+			output(".hi");
+		else
+		{
+			printf("Unknown Type: ");
+			printf(modString);
+			printf("\n");
+			output("Unknown Type: "); output(modString);
+			assert(0);
+		}
+
+		currentPiece = currentPiece->nextString;
+	}
+}
+
+void cuobjdumpInst::printCuobjdumpBaseModifiers()
+{
+	stringListPiece* currentPiece = m_baseModifiers->getListStart();
+	for(int i=0; (i<m_baseModifiers->getSize())&&(currentPiece!=NULL); i++)
+	{
+		const char* modString = currentPiece->stringText;
+		if( strcmp(modString, "EQ")==0 )
+			output(".eq");
+		else if( strcmp(modString, "EQU")==0 )
+			output(".equ");
+		else if( strcmp(modString, "GE")==0 )
+			output(".ge");
+		else if( strcmp(modString, "GEU")==0 )
+			output(".geu");
+		else if( strcmp(modString, "GT")==0 )
+			output(".gt");
+		else if( strcmp(modString, "GTU")==0 )
+			output(".gtu");
+		else if( strcmp(modString, "LE")==0 )
+			output(".le");
+		else if( strcmp(modString, "LEU")==0 )
+			output(".leu");
+		else if( strcmp(modString, "LT")==0 )
+			output(".lt");
+		else if( strcmp(modString, "LTU")==0 )
+			output(".ltu");
+		else if( strcmp(modString, "NE")==0 )
+			output(".ne");
+		else if( strcmp(modString, "NEU")==0 )
+			output(".neu");
+		else if( strcmp(modString, ".exit")==0 )
+			output(modString);
+		else if( strcmp(modString, ".abs")==0 )
+		{
+			if((strcmp(m_base, "F2F")!=0) && (strcmp(m_base, "I2I")!=0))
+			{
+				output(modString);
+			}
+		}
+		else if( strcmp(modString, "ex2")==0 )
+			output(modString);
+		else if( strcmp(modString, "sin")==0 )
+			output(modString);
+		else if( strcmp(modString, "cos")==0 )
+			output(modString);
+		else if( strcmp(modString, ".rz")==0 )
+			output(modString);
+		else if( strcmp(modString, ".bext")==0 )
+		{
+			//".bext" is a modifier that indicated u8 to u16 type conversion, I think
+		}
+		else if( strcmp(modString, ".s")==0 )
+		{
+			//".s" is the same as ".join" in decuda.
+		}
+		else if( strcmp(modString, ".sfu")==0 )
+		{
+			//".sfu" is an unknown base modifier, TODO: find out what it is
+		}
+		else if( strcmp(modString, ".x")==0 )
+		{
+			//".x" is an unknown base modifier, TODO: find out what it is
+			output(modString);
+		}
+		else if( strcmp(modString, ".e")==0 )
+		{
+			//".e" is an unknown base modifier, TODO: find out what it is
+			output(modString);
+		}
+		else if( strcmp(modString, ".ir")==0 )
+		{
+			//".ir" is an unknown base modifier, TODO: find out what it is
+			output(modString);
+		}
+		else if(strcmp(modString, "IADD")==0 ||
+				strcmp(modString, "IMIN")==0 ||
+				strcmp(modString, "IMAX")==0
+				//strcmp(modString, "INC")==0 ||
+				//strcmp(modString, "DEC")==0 ||
+				//strcmp(modString, "IAND")==0 ||
+				//strcmp(modString, "IOR")==0 ||
+				//strcmp(modString, "IXOR")==0
+				)
+		{
+			/*
+			 * This is the case of a GRED or GATOM operation
+			 */
+			output(".");
+			std::string modstr = modString;
+			modstr = modstr.substr(1);
+			for (int i=0; i<modstr.length(); i++){
+				modstr[i] = tolower(modstr[i]);
+			}
+			output(modstr.c_str());
+		}
+		else
+		{
+			printf("Unknown Base Mod: ");
+			printf(modString);
+			printf("\n");
+			output("Unknown Base Mod: ");
+			output(modString);
+			assert(0);
+		}
+
+		currentPiece = currentPiece->nextString;
+	}
+}
+
+/*
+ * Remove the trailing 'l' or 'h' and output the operand followed by ".lo" or ".hi" respectively
+ */
+void cuobjdumpInst::printCuobjdumpOperandlohi(std::string op) {
+	if (op.substr(op.length()-1) == "l") {
+		output(op.substr(0, op.length()-1).c_str());
+		output(".lo");
+	} else if (op.substr(op.length()-1) == "h"){
+		output(op.substr(0, op.length()-1).c_str());
+		output(".hi");
+	} else {
+		output(op.c_str());
+	}
+}
+
+void cuobjdumpInst::printCuobjdumpOperand(stringListPiece* currentPiece, std::string operandDelimiter, const char* base)
+{
+
+	output(operandDelimiter.c_str());
+	output(" ");
+
+	//Current piece
+	std::string currp = currentPiece->stringText;
+	std::string mod;
+
+	if(currp[0] == '-') {
+		mod = currp.substr(1);
+
+		if(mod.substr(0,2) == "0x") {
+			unsigned int immValue;
+			std::stringstream hexStringConvert2;
+			hexStringConvert2 << std::hex << mod;
+			hexStringConvert2 >> immValue;
+
+			if(immValue){
+				immValue = ~immValue + 1;
+			}
+
+			std::stringstream outputhexstream;
+			outputhexstream << std::hex << immValue;
+			std::string outputhex;
+			outputhex = outputhexstream.str();
+			output("0x");
+			for(int i=8; i > outputhex.length(); i--)
+			{
+				output("0");
+			}
+			output(outputhex.c_str());
+			return;
+		}
+
+		output("-");
+	}
+	else
+	{
+		mod = std::string(currp);
+	}
+
+	// Make it lower case
+	if(mod.substr(0,9)!= "constant1")
+	for(int i=0; i<mod.length(); i++)
+	{
+		mod[i] = tolower(mod[i]);
+	}
+
+	//double destination
+	if(mod[0]=='.') //double destination
+	{
+		std::string temp = mod.substr(2,2);
+		output("$p");
+		output(mod.substr(2,2).c_str());
+
+		mod = mod.substr(4);
+	}
+
+
+	if(mod == "g [0x1].u16") { //handling special register case: %ntid.x
+		output("%%ntid.x");
+	} else if(mod == "g [0x2].u16") { //handling special register case: %ntid.y
+		output("%%ntid.y");
+	} else if(mod == "g [0x3].u16") { //handling special register case: %ntid.x
+		output("%%ntid.z");
+	} else if(mod =="g [0x4].u16") { //handling special register case: %nctaid.x
+		output("%%nctaid.x");
+	} else if(mod =="g [0x5].u16") { //handling special register case: %nctaid.y
+		output("%%nctaid.y");
+	} else if(mod == "g [0x6].u16") {//handling special register case: %ctaid.x
+		output("%%ctaid.x");
+	} else if(mod == "g [0x7].u16") {//handling special register case: %ctaid.y
+		output("%%ctaid.y");
+	} else if(mod[0]=='r') { //basic register
+		if(	(strcmp(m_base, "DADD")==0) ||
+			(strcmp(m_base, "DMUL")==0) ||
+			(strcmp(m_base, "DFMA")==0) ||
+			(	(m_typeModifiers->getSize()==1) &&
+				(strcmp((m_typeModifiers->getListStart()->stringText), ".S64")==0) &&
+				(	(strcmp(m_base, "G2R")==0) ||
+					(strcmp(m_base, "R2G")==0) ||
+					(strcmp(m_base, "GLD")==0) ||
+					(strcmp(m_base, "GST")==0) ||
+					(strcmp(m_base, "LST")==0) ))) {
+			std::string modsub = mod.substr(1);
+			int regNumInt = atoi(modsub.c_str());
+			std::stringstream temp;
+			temp << "{$r" << (regNumInt) << ",$r"<< (regNumInt+1) << "}";
+			output(temp.str().c_str());
+		} else if(	(m_typeModifiers->getSize()==1) &&
+					(strcmp((m_typeModifiers->getListStart()->stringText), ".S128")==0)) {
+			std::string modsub = mod.substr(1);
+			int regNumInt = atoi(modsub.c_str());
+			std::stringstream temp;
+			temp << "{$r" << (regNumInt);
+			temp << ",$r" << (regNumInt+1);
+			temp << ",$r" << (regNumInt+2);
+			temp << ",$r" << (regNumInt+3) << "}";
+			output(temp.str().c_str());
+		} else {
+			output("$");
+			printCuobjdumpOperandlohi(mod);
+		}
+	} else if(mod[0] == 'c' && mod.length() == 2) { //predicate register (conditional code)
+		output("$p");
+		output(mod.substr(1,1).c_str());
+	} else if(mod[0]=='a') {//offset register
+		output("$ofs");
+		mod = mod.substr(1);
+		printCuobjdumpOperandlohi(mod);
+	} else if(mod[0]=='o') {//output register
+		output("$o127");
+	} else if (	mod[0]=='g' ||
+				mod.substr(0,2) == "lo" ||
+				mod.substr(0,8) == "constant") {
+		//memory operands, global14 = global, g = shared, c = constant
+		std::string modsub;
+		std::string modsub2;
+		std::string modsub3;
+		modsub = mod.c_str();
+		int localFlag = 0;
+
+		if(mod.find("global14") != std::string::npos) {
+			//Those instructions don't need the dereferencing done by g [*]
+			if(	strcmp(base,"GRED")==0 ||
+				strcmp(base, "GATOM")==0 ||
+				strcmp(base, "GST")==0 ||
+				strcmp(base, "GLD")==0)
+				output("[");
+			else
+				output("g[");
+		} else if(mod[0]=='g') {
+			//Shared memory
+			output("s[");
+		} else if(mod.find("local") !=  std::string::npos) {
+			modsub3 = modsub.substr(4, modsub.length()-10);
+			output(modsub3.c_str());
+			output("[");
+			localFlag = 1;
+		} else if(mod.substr(0,9) == "constant1") {
+			output(modsub.substr(0, modsub.find_first_of("[]")+1).c_str());
+		} else if(mod.substr(0,9)=="constant0"){
+			output("constant0[");
+		} else {
+			printf("Unidentified modifier: %s\n", mod.c_str());
+			assert(0);
+		}
+
+		modsub = modsub.substr(modsub.find_first_of("[]")+1);
+		modsub = modsub.substr(0, modsub.length()-1);
+		//Here we handle whatever was inside the []
+		int plusequalFlag = 0;
+		if(modsub.find("+++") !=  std::string::npos) {
+			plusequalFlag = 1;
+		}
+
+		if(modsub.find("+") !=  std::string::npos)
+		{
+			//Handle a1+++0x1 or a1+0x1
+			modsub2 = modsub.substr(modsub.find("+")+1);
+
+			output("$ofs");
+			output(modsub.substr(1,1).c_str());
+
+			if(plusequalFlag == 1) {
+				output("+=");
+			} else {
+				output("+");
+			}
+
+
+			if(modsub2[0]=='r') {
+				output("$");
+				printCuobjdumpOperandlohi(modsub2);
+			} else {
+				modsub2 = modsub2.substr(2);
+				unsigned int addrValue;
+				std::stringstream hexStringConvert;
+				hexStringConvert << std::hex << modsub2;
+				hexStringConvert >> addrValue;
+				unsigned chunksize = 4;
+				if (	this->m_typeModifiers->getSize()>0 &&
+						(	strcmp((m_typeModifiers->getListStart()->stringText), ".S16")==0 ||
+							strcmp((m_typeModifiers->getListStart()->stringText), ".U16")==0 )) chunksize = 2;
+				if (	this->m_typeModifiers->getSize()>0 &&
+						(	strcmp((m_typeModifiers->getListEnd()->stringText), ".U8")==0 ||
+							strcmp((m_typeModifiers->getListEnd()->stringText), ".S8")==0 )) chunksize = 1;
+				addrValue = addrValue*chunksize;
+				char outputHex[10];
+				sprintf(outputHex, "%x", addrValue);
+				std::stringstream outputhex;
+				outputhex << std::hex << addrValue;
+				output("0x");
+				for(int i=4; i > outputhex.str().length(); i--)
+				{
+					output("0");
+				}
+				output(outputhex.str().c_str());
+			}
+		} else {
+			if(modsub[0]=='r') {
+				//Register
+				output("$");
+				printCuobjdumpOperandlohi(modsub);
+			} else {
+				// Immediate offset
+				mod = mod.substr(2);
+
+				unsigned int addrValue;
+				std::stringstream hexStringConvert;
+				hexStringConvert << std::hex << modsub;
+				hexStringConvert >> addrValue;
+
+				if(localFlag == 0)
+				{
+					unsigned chunksize = 4;
+					if (	this->m_typeModifiers->getSize()>0 &&
+							(strcmp((m_typeModifiers->getListStart()->stringText), ".S16")==0 ||
+									strcmp((m_typeModifiers->getListStart()->stringText), ".U16")==0 )) chunksize = 2;
+					addrValue = addrValue*chunksize;
+				}
+				std::stringstream outputhex;
+				outputhex << std::hex << addrValue;
+				output("0x");
+				for(int i=4; i > outputhex.str().length(); i--) {
+					output("0");
+				}
+				output(outputhex.str().c_str());
+			}
+		}
+		output("]");
+	} else if(mod.substr(0,2) == "0x") { //immediate value
+		output("0x");
+		std::string outputhex;
+		outputhex = mod.substr(2);
+
+		for(int i=8; i > outputhex.length(); i--) {
+			output("0");
+		}
+		output(outputhex.c_str());
+	} else if(mod.substr(0,3) == "l0x") { //label
+		output(mod.c_str());
+	} else {//variable name
+		output(mod.c_str());
+	}
+}
+
+void cuobjdumpInst::printCuobjdumpOperands()
+{
+	stringListPiece* currentPiece = m_operands->getListStart();
+	std::string delimiter = "";
+	for(int i=0; (i<m_operands->getSize())&&(currentPiece!=NULL); i++)
+	{
+		if((strcmp(m_base, "LOP.PASS_B")==0 || strcmp(m_base, "LOP.S.PASS_B")==0) && (i==1))
+		{
+			currentPiece = currentPiece->nextString;
+			continue;
+		}
+		printCuobjdumpOperand(currentPiece, delimiter, m_base);
+		currentPiece = currentPiece->nextString;
+		delimiter = ",";
+	}
+}
+
+char * int_default_mod () { return ".u32" ;}
+
+
+std::string breaktarget;
+
+void cuobjdumpInst::printCuobjdumpPtxPlus(std::list<std::string> labelList, std::list<std::string> texList)
+{
+	stringListPiece* currentPiece;
+
+	printCuobjdumpLabel(labelList);
+
+	if(strcmp(m_base, "")==0)
+	{
+	}
+	else if(strcmp(m_base, ".entry")==0)
+	{
+		/*do nothing here*/
+	}
+	else if(strcmp(m_base, "BAR.ARV.WAIT b0, 0xfff")==0)
+	{
+		printCuobjdumpPredicate();
+		output("bar.sync 0x00000000;");
+	}
+	else if(strcmp(m_base, "ADA")==0)
+	{
+		printCuobjdumpPredicate();
+		output("add");
+		printCuobjdumpBaseModifiers();
+
+		if(m_typeModifiers->getSize() == 0)
+			output(int_default_mod()); //TODO: setting default type modifier but I'm not sure if this is right.
+		else
+			printCuobjdumpTypeModifiers();
+
+		printCuobjdumpOperands();
+		output(";");
+	}
+	else if(strcmp(m_base, "BRA")==0)
+	{
+		printCuobjdumpPredicate();
+		output("bra");
+		printCuobjdumpBaseModifiers();
+		printCuobjdumpTypeModifiers();
+		printCuobjdumpOperands();
+		output(";");
+	}
+	else if(strcmp(m_base, "CAL")==0)
+	{
+		printCuobjdumpPredicate();
+		output("callp");
+		printCuobjdumpBaseModifiers();
+		printCuobjdumpTypeModifiers();
+		printCuobjdumpOperands();
+		output(";");
+	}
+	else if(strcmp(m_base, "COS")==0)
+	{
+		printCuobjdumpPredicate();
+		//output("nop;");
+		//output(" //cos");
+		output("cos");
+		printCuobjdumpBaseModifiers();
+
+		if(m_typeModifiers->getSize() == 0)
+			output(".f32"); //TODO: setting default type modifier but I'm not sure if this is right.
+		else
+			printCuobjdumpTypeModifiers();
+
+		printCuobjdumpOperands();
+	}
+	else if(strcmp(m_base, "DADD")==0)
+	{
+		printCuobjdumpPredicate();
+		output("add");
+		printCuobjdumpBaseModifiers();
+
+		if(m_typeModifiers->getSize() == 0)
+			output(".ff64"); //TODO: setting default type modifier but I'm not sure if this is right.
+		else
+			printCuobjdumpTypeModifiers();
+
+		printCuobjdumpOperands();
+		output(";");
+	}
+	else if(strcmp(m_base, "DFMA")==0)
+	{
+		printCuobjdumpPredicate();
+		output("mad");
+		printCuobjdumpBaseModifiers();
+
+		if(m_typeModifiers->getSize() == 0)
+			output(".ff64"); //TODO: setting default type modifier but I'm not sure if this is right.
+		else
+			printCuobjdumpTypeModifiers();
+
+		printCuobjdumpOperands();
+		output(";");
+	}
+	else if(strcmp(m_base, "DMUL")==0)
+	{
+		printCuobjdumpPredicate();
+		output("mul");
+		printCuobjdumpBaseModifiers();
+
+		if(m_typeModifiers->getSize() == 0)
+			output(".ff64"); //TODO: setting default type modifier but I'm not sure if this is right.
+		else
+			printCuobjdumpTypeModifiers();
+
+		printCuobjdumpOperands();
+		output(";");
+	}
+	else if(strcmp(m_base, "EX2")==0)
+	{
+		printCuobjdumpPredicate();
+		//output("nop;");
+		//output(" //ex2");
+		output("ex2");
+		printCuobjdumpBaseModifiers();
+
+		if(m_typeModifiers->getSize() == 0)
+			output(".f32"); //TODO: setting default type modifier but I'm not sure if this is right.
+		else
+			printCuobjdumpTypeModifiers();
+
+		printCuobjdumpOperands();
+		output(";");
+	}
+	else if(strcmp(m_base, "F2F")==0)
+	{
+		printCuobjdumpPredicate();
+
+		int absFlag = 0;
+		stringListPiece* currentPiece = m_baseModifiers->getListStart();
+
+		while(currentPiece != NULL)
+		{
+			if(strcmp(currentPiece->stringText, ".abs")==0)
+			{
+				output("abs");
+				absFlag = 1;
+				break;
+			}
+			currentPiece = currentPiece->nextString;
+		}
+		if(absFlag == 0)
+		{
+			output("cvt");
+		}
+
+		printCuobjdumpBaseModifiers();
+
+		if(absFlag == 0)
+		{
+			printCuobjdumpTypeModifiers();
+		}
+		else
+		{
+			output(".f32"); //TODO: setting default type modifier but I'm not sure if this is right.
+		}
+
+		printCuobjdumpOperands();
+		output(";");
+	}
+	else if(strcmp(m_base, "F2I")==0)
+	{
+		printCuobjdumpPredicate();
+		output("cvt");
+		printCuobjdumpBaseModifiers();
+		printCuobjdumpTypeModifiers();
+		printCuobjdumpOperands();
+		output(";");
+	}
+	else if(strcmp(m_base, "FADD")==0 || strcmp(m_base, "FADD32I")==0)
+	{
+		printCuobjdumpPredicate();
+		output("add");
+		printCuobjdumpBaseModifiers();
+
+		if(m_typeModifiers->getSize() == 0)
+			output(".f32"); //TODO: setting default type modifier but I'm not sure if this is right.
+		else
+			printCuobjdumpTypeModifiers();
+
+		printCuobjdumpOperands();
+		output(";");
+	}
+	else if(strcmp(m_base, "FADD32")==0)
+	{
+		printCuobjdumpPredicate();
+		output("add.half");
+		printCuobjdumpBaseModifiers();
+
+		if(m_typeModifiers->getSize() == 0)
+			output(".f32"); //TODO: setting default type modifier but I'm not sure if this is right.
+		else
+			printCuobjdumpTypeModifiers();
+
+		printCuobjdumpOperands();
+		output(";");
+	}
+	else if(strcmp(m_base, "FMAD")==0 || strcmp(m_base, "FMAD32I")==0)
+	{
+		printCuobjdumpPredicate();
+		output("mad");
+		printCuobjdumpBaseModifiers();
+
+		if(m_typeModifiers->getSize() == 0)
+			output(".f32"); //TODO: setting default type modifier but I'm not sure if this is right.
+		else
+			printCuobjdumpTypeModifiers();
+
+		printCuobjdumpOperands();
+		output(";");
+	}
+	else if(strcmp(m_base, "FMUL")==0 || strcmp(m_base, "FMUL32I")==0)
+	{
+		printCuobjdumpPredicate();
+		output("mul");
+		printCuobjdumpBaseModifiers();
+
+		if(m_typeModifiers->getSize() == 0)
+			output(".f32"); //TODO: setting default type modifier but I'm not sure if this is right.
+		else
+			printCuobjdumpTypeModifiers();
+
+		printCuobjdumpOperands();
+		output(";");
+	}
+	else if(strcmp(m_base, "FMUL32")==0)
+	{
+		printCuobjdumpPredicate();
+		output("mul.half");
+		printCuobjdumpBaseModifiers();
+
+		if(m_typeModifiers->getSize() == 0)
+			output(".f32"); //TODO: setting default type modifier but I'm not sure if this is right.
+		else
+			printCuobjdumpTypeModifiers();
+
+		printCuobjdumpOperands();
+		output(";");
+	}
+	else if(strcmp(m_base, "FSET")==0)
+	{
+		printCuobjdumpPredicate();
+		output("set");
+		printCuobjdumpBaseModifiers();
+
+		if(m_typeModifiers->getSize() == 0)
+		{
+			output(".f32"); output(".f32");
+		}
+		else
+			printCuobjdumpTypeModifiers();
+
+		printCuobjdumpOperands();
+		output(";");
+	}
+	else if(strcmp(m_base, "G2R")==0)
+	{
+		printCuobjdumpPredicate();
+		output("mov");
+		printCuobjdumpBaseModifiers();
+
+		if( m_typeModifiers->getSize() == 2 ) {
+			std::string type1, type2, type;
+			int type1Size, type2Size;
+			stringListPiece* currentPiece = m_typeModifiers->getListStart();
+			type1 = currentPiece->stringText;
+			type2 = currentPiece->nextString->stringText;
+
+			type1Size = atoi(type1.substr(2, type1.size()-2).c_str());
+			type2Size = atoi(type2.substr(2, type2.size()-2).c_str());
+			type = (type1Size < type2Size) ? type1 : type2;
+
+			if( strcmp(type.c_str(), ".F16")==0 )
+				output(".f16");
+			else if( strcmp(type.c_str(), ".F32")==0 )
+				output(".f32");
+			else if( strcmp(type.c_str(), ".F64")==0 )
+				output(".ff64");
+			else if( strcmp(type.c_str(), ".S8")==0 )
+				output(".s8");
+			else if( strcmp(type.c_str(), ".S16")==0 )
+				output(".s16");
+			else if( strcmp(type.c_str(), ".S32")==0 )
+				output(".s32");
+			else if( strcmp(type.c_str(), ".S64")==0 )
+				output(".bb64");
+			else if( strcmp(type.c_str(), ".S128")==0 )
+				output(".bb128");
+			else if( strcmp(type.c_str(), ".U8")==0 )
+				output(".u8");
+			else if( strcmp(type.c_str(), ".U16")==0 )
+				output(".u16");
+			else if( strcmp(type.c_str(), ".U32")==0 )
+				output(".u32");
+			else if( strcmp(type.c_str(), ".U64")==0 )
+				output(".bb64");
+			else
+				output(type.c_str());
+		} else if( m_typeModifiers->getSize() == 1 ) {
+			printCuobjdumpTypeModifiers();
+		} else {
+			output("Error: unsupported number of type modifiers. ");
+                }
+		printCuobjdumpOperands();
+		output(";");
+	}
+	else if(strcmp(m_base, "GLD")==0)
+	{
+		printCuobjdumpPredicate();
+		//output("mov");
+		output("ld.global");
+		printCuobjdumpBaseModifiers();
+		if (strcmp(m_typeModifiers->getListStart()->stringText, ".S128")==0) {
+			output(".v4.u32");
+		} else if (strcmp(m_typeModifiers->getListStart()->stringText, ".S64")==0) {
+			output(".v2.u32");
+		} else {
+			printCuobjdumpTypeModifiers();
+		}
+		printCuobjdumpOperands();
+		output(";");
+	}
+	else if(strcmp(m_base, "GST")==0)
+	{
+		printCuobjdumpPredicate();
+		//output("mov");
+		output("st.global");
+		printCuobjdumpBaseModifiers();
+		if (strcmp(m_typeModifiers->getListStart()->stringText, ".S128")==0) {
+			output(".v4.u32");
+		} else if (strcmp(m_typeModifiers->getListStart()->stringText, ".S64")==0) {
+			output(".v2.u32");
+		} else {
+			printCuobjdumpTypeModifiers();
+		}
+		printCuobjdumpOperands();
+		output(";");
+	}
+	else if(strcmp(m_base, "I2F")==0)
+	{
+		printCuobjdumpPredicate();
+		output("cvt");
+		printCuobjdumpBaseModifiers();
+		printCuobjdumpTypeModifiers();
+		printCuobjdumpOperands();
+		output(";");
+	}
+	else if(strcmp(m_base, "I2I")==0)
+	{
+		printCuobjdumpPredicate();
+
+		int absFlag = 0;
+		int bextFlag = 0;
+		stringListPiece* currentPiece = m_baseModifiers->getListStart();
+
+		while(currentPiece != NULL)
+		{
+			if(strcmp(currentPiece->stringText, ".abs")==0)
+			{
+				output("abs");
+				absFlag = 1;
+				break;
+			}
+			if(strcmp(currentPiece->stringText, ".bext")==0)
+			{
+				bextFlag = 1;
+				break;
+			}
+			currentPiece = currentPiece->nextString;
+		}
+		if(absFlag == 0)
+		{
+			output("cvt");
+		}
+
+		printCuobjdumpBaseModifiers();
+
+		currentPiece = this->m_typeModifiers->getListStart();
+		unsigned int maxlength=16;
+		unsigned int currlength = 16;
+		bool issigned = false;
+		std::string tmpstr;
+		while(currentPiece != NULL)
+		{
+			tmpstr = currentPiece->stringText;
+			if(tmpstr[1] == 'S')
+			{
+				issigned = true;
+			}
+			if(tmpstr.substr(2,2) ==  "32") {currlength=32;}
+			if (currlength > maxlength) {maxlength = currlength;}
+			currentPiece = currentPiece->nextString;
+		}
+
+		if(absFlag == 0)
+		{
+			if(bextFlag == 0)
+				printCuobjdumpTypeModifiers();
+			else{
+				output(".");
+				if (issigned) {output("s");}
+				else {output("u");}
+				std::stringstream tmp;
+				tmp << maxlength;
+				output(tmp.str().c_str());
+				output(".");
+				if (issigned) {output("s");}
+				else {output("u");}
+				output("8");
+				//output(".u16.u8");
+			}
+		}
+		else
+		{
+			output(int_default_mod()); //TODO: setting default type modifier but I'm not sure if this is right.
+		}
+
+		printCuobjdumpOperands();
+		output(";");
+	}
+	else if(strcmp(m_base, "IADD")==0)
+	{
+		printCuobjdumpPredicate();
+		output("add");
+		printCuobjdumpBaseModifiers();
+
+		if(m_typeModifiers->getSize() == 0)
+			output(int_default_mod()); //TODO: setting default type modifier but I'm not sure if this is right.
+		else
+			printCuobjdumpTypeModifiers();
+
+		printCuobjdumpOperands();
+		output(";");
+	}
+	else if(strcmp(m_base, "IADD32")==0  || strcmp(m_base, "IADD32I")==0 )
+	{
+		printCuobjdumpPredicate();
+		output("add.half");
+		printCuobjdumpBaseModifiers();
+
+		if(m_typeModifiers->getSize() == 0)
+			output(".u32"); //TODO: setting default type modifier but I'm not sure if this is right.
+		else
+			printCuobjdumpTypeModifiers();
+		printCuobjdumpOperands();
+		output(";");
+	}
+	else if(strcmp(m_base, "IMAD")==0 || strcmp(m_base, "IMAD32I")==0  || strcmp(m_base, "IMAD32")==0)
+	{
+		printCuobjdumpPredicate();
+		output("mad.lo");
+		printCuobjdumpBaseModifiers();
+
+		if(m_typeModifiers->getSize() == 0)
+			output(int_default_mod()); //TODO: setting default type modifier but I'm not sure if this is right.
+		else
+			printCuobjdumpTypeModifiers();
+		printCuobjdumpOperands();
+		output(";");
+        }
+	else if(strcmp(m_base, "IMAD.U24")==0)
+	{
+		printCuobjdumpPredicate();
+		output("mad24.lo");
+		printCuobjdumpBaseModifiers();
+
+		if(m_typeModifiers->getSize() == 0)
+			output(".u32"); //TODO: setting default type modifier but I'm not sure if this is right.
+		else
+			printCuobjdumpTypeModifiers();
+		printCuobjdumpOperands();
+		output(";");
+        }
+	else if(strcmp(m_base, "IMAD.S24")==0)
+	{
+		printCuobjdumpPredicate();
+		output("mad24.lo");
+		printCuobjdumpBaseModifiers();
+
+		if(m_typeModifiers->getSize() == 0)
+			output(".s32"); //TODO: setting default type modifier but I'm not sure if this is right.
+		else
+			printCuobjdumpTypeModifiers();
+		printCuobjdumpOperands();
+		output(";");
+	}
+	else if(strcmp(m_base, "IMUL32I.U24.U24")==0 ||
+			strcmp(m_base, "IMUL32I.S24.S24")==0 ||
+			strcmp(m_base, "IMUL32.U24.U24")==0 ||
+			strcmp(m_base, "IMUL32.S24.S24")==0 ||
+			strcmp(m_base, "IMUL.U24.U24")==0 ||
+			strcmp(m_base, "IMUL.S24.S24")==0)
+	{
+		printCuobjdumpPredicate();
+		output("mul24.lo");
+		printCuobjdumpBaseModifiers();
+
+		if(m_typeModifiers->getSize() == 0)
+		{
+			output(".u32"); //TODO: setting default type modifier but I'm not sure if this is right.
+		}
+		else if(m_typeModifiers->getSize() == 2)
+		{
+			std::string type1, type2, type;
+			int type1Size, type2Size;
+			stringListPiece* currentPiece = m_typeModifiers->getListStart();
+			char tempString[5];
+
+			type1 = currentPiece->stringText;
+			type2 = currentPiece->nextString->stringText;
+
+			type1Size = atoi(type1.substr(2, type1.size()-2).c_str());
+			type2Size = atoi(type2.substr(2, type2.size()-2).c_str());
+			type = (type1Size > type2Size) ? type1 : type2;
+			strcpy(tempString, type.c_str());
+			/*if(type1Size==16 && type2Size==16)
+				output(".lo");*/
+			if(tempString[1] >= 'A' && tempString[1] <= 'Z')
+				tempString[1] += 32;			
+			output(tempString);
+		}	
+		else
+		{
+			printCuobjdumpTypeModifiers();
+		}
+		printCuobjdumpOperands();
+		output(";");
+	}
+	else if(strcmp(m_base, "IMUL")==0 || strcmp(m_base, "IMUL32I")==0)
+	{
+		printCuobjdumpPredicate();
+		output("mul.lo");
+		printCuobjdumpBaseModifiers();
+
+		if(m_typeModifiers->getSize() == 0)
+		{
+			output(int_default_mod()); //TODO: setting default type modifier but I'm not sure if this is right.
+		}
+		else if(m_typeModifiers->getSize() == 2)
+		{
+			std::string type1, type2, type;
+			int type1Size, type2Size;
+			stringListPiece* currentPiece = m_typeModifiers->getListStart();
+			char tempString[5];
+
+			type1 = currentPiece->stringText;
+			type2 = currentPiece->nextString->stringText;
+
+			type1Size = atoi(type1.substr(2, type1.size()-2).c_str());
+			type2Size = atoi(type2.substr(2, type2.size()-2).c_str());
+			type = (type1Size > type2Size) ? type1 : type2;
+			strcpy(tempString, type.c_str());
+			/*if(type1Size==16 && type2Size==16)
+					output(".lo");*/
+			if(tempString[1] >= 'A' && tempString[1] <= 'Z')
+				tempString[1] += 32;			
+			output(tempString);
+		}	
+		else
+		{
+			printCuobjdumpTypeModifiers();
+		}
+		printCuobjdumpOperands();
+		output(";");
+	}
+	else if(strcmp(m_base, "IMUL32")==0)
+	{
+		printCuobjdumpPredicate();
+		output("mul.half.lo");
+		printCuobjdumpBaseModifiers();
+
+		if(m_typeModifiers->getSize() == 2)
+		{
+			std::string type1, type2, type;
+			int type1Size, type2Size;
+			stringListPiece* currentPiece = m_typeModifiers->getListStart();
+			char tempString[5];
+
+			type1 = currentPiece->stringText;
+			type2 = currentPiece->nextString->stringText;
+
+			type1Size = atoi(type1.substr(2, type1.size()-2).c_str());
+			type2Size = atoi(type2.substr(2, type2.size()-2).c_str());
+			type = (type1Size > type2Size) ? type1 : type2;
+			strcpy(tempString, type.c_str());
+			/*if(type1Size==16 && type2Size==16)
+				output(".lo");*/
+			if(tempString[1] >= 'A' && tempString[1] <= 'Z')
+				tempString[1] += 32;			
+			output(tempString);
+
+		}	
+		else
+			printCuobjdumpTypeModifiers();
+
+		printCuobjdumpOperands();
+		output(";");
+	}
+	else if(strcmp(m_base, "IMUL32.S24.S24")==0)
+	{
+		printCuobjdumpPredicate();
+		output("mul24.half.lo");
+		printCuobjdumpBaseModifiers();
+
+		if(m_typeModifiers->getSize() == 0)
+		{
+			output(".u32"); //TODO: setting default type modifier but I'm not sure if this is right.
+		}
+		else if(m_typeModifiers->getSize() == 2)
+		{
+			std::string type1, type2, type;
+			int type1Size, type2Size;
+			stringListPiece* currentPiece = m_typeModifiers->getListStart();
+			char tempString[5];
+
+			type1 = currentPiece->stringText;
+			type2 = currentPiece->nextString->stringText;
+
+			type1Size = atoi(type1.substr(2, type1.size()-2).c_str());
+			type2Size = atoi(type2.substr(2, type2.size()-2).c_str());
+			type = (type1Size > type2Size) ? type1 : type2;
+			strcpy(tempString, type.c_str());
+			/*if(type1Size==16 && type2Size==16)
+				output(".lo");*/
+			if(tempString[1] >= 'A' && tempString[1] <= 'Z')
+				tempString[1] += 32;			
+			output(tempString);
+
+		}	
+		else
+			printCuobjdumpTypeModifiers();
+
+		printCuobjdumpOperands();
+		output(";");
+	}
+	else if(strcmp(m_base, "ISET")==0)
+	{
+		printCuobjdumpPredicate();
+		output("set");
+		printCuobjdumpBaseModifiers();
+
+		if(m_typeModifiers->getSize() == 0)
+			output(".u32.u32"); //TODO: setting default type modifier but I'm not sure if this is right.
+		else if(m_typeModifiers->getSize() == 1)
+		{
+			printCuobjdumpTypeModifiers(); printCuobjdumpTypeModifiers();
+		}
+		else
+			printCuobjdumpTypeModifiers();
+		printCuobjdumpOperands();
+		output(";");
+	}
+	else if(strcmp(m_base, "LG2")==0)
+	{
+		printCuobjdumpPredicate();
+		output("lg2");
+		printCuobjdumpBaseModifiers();
+
+		if(m_typeModifiers->getSize() == 0)
+			output(".f32"); //TODO: setting default type modifier but I'm not sure if this is right.
+		else
+			printCuobjdumpTypeModifiers();
+
+		printCuobjdumpOperands();
+		output(";");
+	}
+	else if(strcmp(m_base, "LST")==0)
+	{
+		printCuobjdumpPredicate();
+		output("mov");
+		printCuobjdumpBaseModifiers();
+		printCuobjdumpTypeModifiers();
+		printCuobjdumpOperands();
+		output(";");
+	}
+	else if(strcmp(m_base, "MVC")==0) {
+		printCuobjdumpPredicate();
+		output("cvt");
+		printCuobjdumpBaseModifiers();
+
+		if(m_typeModifiers->getSize() == 0)
+			output(".u32.u32"); //TODO: setting default type modifier but I'm not sure if this is right.
+		else
+			printCuobjdumpTypeModifiers();
+		printCuobjdumpOperands();
+		output(";");
+	}
+
+	else if(strcmp(m_base, "MOV")==0 || strcmp(m_base, "MVI")==0 || strcmp(m_base, "MVC")==0)
+	{
+		printCuobjdumpPredicate();
+		output("mov");
+		printCuobjdumpBaseModifiers();
+
+		if(m_typeModifiers->getSize() == 0)
+			output(int_default_mod()); //TODO: setting default type modifier but I'm not sure if this is right.
+		else
+			printCuobjdumpTypeModifiers();
+		printCuobjdumpOperands();
+		output(";");
+	}
+	else if(strcmp(m_base, "MOV32")==0)
+	{
+		printCuobjdumpPredicate();
+		output("mov.half");
+		printCuobjdumpBaseModifiers();
+
+		if(m_typeModifiers->getSize() == 0)
+			output(".u32"); //TODO: setting default type modifier but I'm not sure if this is right.
+		else
+			printCuobjdumpTypeModifiers();
+		printCuobjdumpOperands();
+		output(";");
+	}
+	else if(strcmp(m_base, "NOP")==0 || strcmp(m_base, "SSY")==0)
+	{
+		printCuobjdumpPredicate();
+		output("nop");
+		printCuobjdumpBaseModifiers();
+		output(";");
+	}
+	else if(strcmp(m_base, "LLD")==0)
+	{
+		printCuobjdumpPredicate();
+		output("mov");
+		printCuobjdumpBaseModifiers();
+		printCuobjdumpTypeModifiers();
+		printCuobjdumpOperands();
+		output(";");
+	}
+	else if(strcmp(m_base, "LOP.AND")==0 || strcmp(m_base, "LOP.S.AND")==0)
+	{
+		printCuobjdumpPredicate();
+		output("and");
+		printCuobjdumpBaseModifiers();
+
+		if(m_typeModifiers->getSize() == 0)
+			output(".b32"); //TODO: setting default type modifier but I'm not sure if this is right.
+		else
+			printCuobjdumpTypeModifiers();
+
+		printCuobjdumpOperands();
+		output(";");
+	}
+	else if(strcmp(m_base, "LOP.OR")==0 || strcmp(m_base, "LOP.S.OR")==0)
+	{
+		printCuobjdumpPredicate();
+		output("or");
+		printCuobjdumpBaseModifiers();
+
+		if(m_typeModifiers->getSize() == 0)
+			output(".b32"); //TODO: setting default type modifier but I'm not sure if this is right.
+		else
+			printCuobjdumpTypeModifiers();
+
+		printCuobjdumpOperands();
+		output(";");
+	}
+	else if(strcmp(m_base, "LOP.PASS_B")==0 || strcmp(m_base, "LOP.S.PASS_B")==0)
+	{
+		printCuobjdumpPredicate();
+		output("not");
+		printCuobjdumpBaseModifiers();
+
+		if(m_typeModifiers->getSize() == 0)
+			output(".b32"); //TODO: setting default type modifier but I'm not sure if this is right.
+		else
+			printCuobjdumpTypeModifiers();
+
+		printCuobjdumpOperands();
+		output(";");
+	}
+	else if(strcmp(m_base, "LOP.XOR")==0 || strcmp(m_base, "LOP.S.XOR")==0)
+	{
+		printCuobjdumpPredicate();
+		output("xor");
+		printCuobjdumpBaseModifiers();
+
+		if(m_typeModifiers->getSize() == 0)
+			output(".b32"); //TODO: setting default type modifier but I'm not sure if this is right.
+		else
+			printCuobjdumpTypeModifiers();
+
+		printCuobjdumpOperands();
+		output(";");
+	}
+	else if(strcmp(m_base, "R2A")==0)
+	{
+		printCuobjdumpPredicate();
+		output("shl");
+		printCuobjdumpBaseModifiers();
+
+		if(m_typeModifiers->getSize() == 0)
+			output(".b32"); //TODO: setting default type modifier but I'm not sure if this is right.
+		else
+			printCuobjdumpTypeModifiers();
+		printCuobjdumpOperands();
+		if(m_operands->getSize() == 2)
+			output(", 0x0");
+		output(";");
+	}
+	else if(strcmp(m_base, "R2G.U16.U8")==0){
+		/*
+		 * This code handles a cuobjdump bug that causes the wrong register number to be printed
+		 */
+		printCuobjdumpPredicate();
+		output("mov.u8");
+		stringListPiece* currentPiece = m_operands->getListStart();
+		std::string delimiter = "";
+		printCuobjdumpOperand(currentPiece, delimiter, m_base);
+		output(", ");
+		std::string curr = currentPiece->nextString->stringText;
+		curr = curr.substr(1);
+		int regnum;
+		std::istringstream(curr)>>regnum;
+		output("$r");
+		std::stringstream finalregnum;
+		finalregnum << (regnum/2);
+		output(finalregnum.str().c_str());
+		output( regnum%2==0? ".lo": ".hi");
+		output(";");
+	}
+	else if(strcmp(m_base, "R2G")==0)
+	{
+		printCuobjdumpPredicate();
+		output("mov");
+		printCuobjdumpBaseModifiers();
+
+		if( m_typeModifiers->getSize() == 2 ) {
+			std::string type1, type2, type;
+			int type1Size, type2Size;
+			stringListPiece* currentPiece = m_typeModifiers->getListStart();
+			type1 = currentPiece->stringText;
+			type2 = currentPiece->nextString->stringText;
+
+			type1Size = atoi(type1.substr(2, type1.size()-2).c_str());
+			type2Size = atoi(type2.substr(2, type2.size()-2).c_str());
+			type = (type1Size < type2Size) ? type1 : type2;
+
+			if( strcmp(type.c_str(), ".F16")==0 )
+				output(".f16");
+			else if( strcmp(type.c_str(), ".F32")==0 )
+				output(".f32");
+			else if( strcmp(type.c_str(), ".F64")==0 )
+				output(".ff64");
+			else if( strcmp(type.c_str(), ".S8")==0 )
+				output(".s8");
+			else if( strcmp(type.c_str(), ".S16")==0 )
+				output(".s16");
+			else if( strcmp(type.c_str(), ".S32")==0 )
+				output(".s32");
+			else if( strcmp(type.c_str(), ".S64")==0 )
+				output(".bb64");
+			else if( strcmp(type.c_str(), ".S128")==0 )
+				output(".bb128");
+			else if( strcmp(type.c_str(), ".U8")==0 )
+				output(".u8");
+			else if( strcmp(type.c_str(), ".U16")==0 )
+				output(".u16");
+			else if( strcmp(type.c_str(), ".U32")==0 )
+				output(".u32");
+			else if( strcmp(type.c_str(), ".U64")==0 )
+				output(".bb64");
+			else
+				output(type.c_str());
+		} else if( m_typeModifiers->getSize() == 1 ) {
+			printCuobjdumpTypeModifiers();
+		} else {
+			output("Error: unsupported number of type modifiers. ");
+                }
+
+		printCuobjdumpOperands();
+		output(";");
+	}
+	else if(strcmp(m_base, "RCP")==0)
+	{
+		printCuobjdumpPredicate();
+		output("rcp");
+		printCuobjdumpBaseModifiers();
+
+		if(m_typeModifiers->getSize() == 0)
+			output(".f32"); //TODO: setting default type modifier but I'm not sure if this is right.
+		else
+			printCuobjdumpTypeModifiers();
+
+		printCuobjdumpOperands();
+		output(";");
+	}
+	else if(strcmp(m_base, "RCP32")==0)
+	{
+		printCuobjdumpPredicate();
+		output("rcp.half");
+		printCuobjdumpBaseModifiers();
+
+		if(m_typeModifiers->getSize() == 0)
+			output(".f32");
+		else
+			printCuobjdumpTypeModifiers();
+		printCuobjdumpOperands();
+		output(";");
+	}
+	else if(strcmp(m_base, "RET")==0)
+	{
+		printCuobjdumpPredicate();
+		output("retp;");
+	}
+	else if(strcmp(m_base, "RRO")==0)
+	{
+		output("nop; //");
+		printCuobjdumpPredicate();
+		printCuobjdumpBaseModifiers();
+
+		if(m_typeModifiers->getSize() == 0)
+			output(".f32"); //TODO: setting default type modifier but I'm not sure if this is right.
+		else
+			printCuobjdumpTypeModifiers();
+
+		printCuobjdumpOperands();
+		output(";");
+	}
+	else if(strcmp(m_base, "RSQ")==0)
+	{
+		printCuobjdumpPredicate();
+		output("rsqrt");
+		printCuobjdumpBaseModifiers();
+
+		if(m_typeModifiers->getSize() == 0)
+			output(".f32"); //TODO: setting default type modifier but I'm not sure if this is right.
+		else
+			printCuobjdumpTypeModifiers();
+
+		printCuobjdumpOperands();
+		output(";");
+	}
+	else if(strcmp(m_base, "SHL")==0)
+	{
+		printCuobjdumpPredicate();
+		output("shl");
+		printCuobjdumpBaseModifiers();
+
+		if(m_typeModifiers->getSize() == 0)
+			output(int_default_mod()); //TODO: setting default type modifier but I'm not sure if this is right.
+		else
+			printCuobjdumpTypeModifiers();
+
+		printCuobjdumpOperands();
+		output(";");
+	}
+	else if(strcmp(m_base, "SHR")==0)
+	{
+		printCuobjdumpPredicate();
+		output("shr");
+		printCuobjdumpBaseModifiers();
+
+		if(m_typeModifiers->getSize() == 0)
+			output(int_default_mod()); //TODO: setting default type modifier but I'm not sure if this is right.
+		else
+			printCuobjdumpTypeModifiers();
+
+		printCuobjdumpOperands();
+		output(";");
+	}
+	else if(strcmp(m_base, "SIN")==0)
+	{
+		printCuobjdumpPredicate();
+		output("sin");
+		printCuobjdumpBaseModifiers();
+
+		if(m_typeModifiers->getSize() == 0)
+			output(".f32"); //TODO: setting default type modifier but I'm not sure if this is right.
+		else
+			printCuobjdumpTypeModifiers();
+
+		printCuobjdumpOperands();
+	}
+	else if(strcmp(m_base, "S2R")==0)
+	{
+		printCuobjdumpPredicate();
+		output("cvt");
+		printCuobjdumpBaseModifiers();
+
+		if(m_typeModifiers->getSize() == 0)
+			output(".u32");
+		else
+			printCuobjdumpTypeModifiers();
+		printCuobjdumpOperands();
+	}
+	else if(strcmp(m_base, "LD")==0)
+	{
+		printCuobjdumpPredicate();
+		output("mov");
+		printCuobjdumpBaseModifiers();
+		printCuobjdumpTypeModifiers();
+		printCuobjdumpOperands();
+		output(";");
+	}
+	else if(strcmp(m_base, "STS")==0)
+	{
+		printCuobjdumpPredicate();
+		output("mov");
+		printCuobjdumpBaseModifiers();
+		printCuobjdumpTypeModifiers();
+		printCuobjdumpOperands();
+		output(";");
+	}
+	else if(strcmp(m_base, "BAR")==0)
+	{
+		printCuobjdumpPredicate();
+		output("bar.sync 0;");
+	}
+	else if(strcmp(m_base, "LDS")==0)
+	{
+		// If there is not global address space that includes shared memory, then fix this.
+		// Same for STS
+		printCuobjdumpPredicate();
+		output("mov");
+		printCuobjdumpBaseModifiers();
+		printCuobjdumpTypeModifiers();
+		printCuobjdumpOperands();
+		output(";");
+	}
+	else if(strcmp(m_base, "ST")==0)
+	{
+		printCuobjdumpPredicate();
+		output("mov");
+		printCuobjdumpBaseModifiers();
+		printCuobjdumpTypeModifiers();
+		printCuobjdumpOperands();
+		output(";");
+	} else if(strcmp(m_base, "IMIN")==0) {
+		printCuobjdumpPredicate();
+		output("min");
+		printCuobjdumpBaseModifiers();
+		if(m_typeModifiers->getSize() == 0)
+			output(".s32"); //TODO: setting default type modifier but I'm not sure if this is right.
+		else
+			printCuobjdumpTypeModifiers();
+		printCuobjdumpOperands();
+		output(";");
+	} else if(strcmp(m_base, "IMAX")==0) {
+		printCuobjdumpPredicate();
+		output("max");
+		printCuobjdumpBaseModifiers();
+		if(m_typeModifiers->getSize() == 0)
+			output(".s32"); //TODO: setting default type modifier but I'm not sure if this is right.
+		else
+			printCuobjdumpTypeModifiers();
+		printCuobjdumpOperands();
+		output(";");
+	} else if(strcmp(m_base, "FMIN")==0) {
+			printCuobjdumpPredicate();
+			output("min");
+			printCuobjdumpBaseModifiers();
+			if(m_typeModifiers->getSize() == 0)
+				output(".f32"); //TODO: setting default type modifier but I'm not sure if this is right.
+			else
+				printCuobjdumpTypeModifiers();
+			printCuobjdumpOperands();
+			output(";");
+	}
+	else if(strcmp(m_base, "FMAX")==0) {
+			printCuobjdumpPredicate();
+			output("max");
+			printCuobjdumpBaseModifiers();
+			if(m_typeModifiers->getSize() == 0)
+				output(".f32"); //TODO: setting default type modifier but I'm not sure if this is right.
+			else
+				printCuobjdumpTypeModifiers();
+			printCuobjdumpOperands();
+			output(";");
+	}
+	else if(strcmp(m_base, "A2R")==0) {
+		printCuobjdumpPredicate();
+		output("mov");
+		printCuobjdumpBaseModifiers();
+		if(m_typeModifiers->getSize() == 0)
+			output(int_default_mod()); //TODO: setting default type modifier but I'm not sure if this is right.
+		else
+			printCuobjdumpTypeModifiers();
+		printCuobjdumpOperands();
+		output(";");
+	}
+	else if((strcmp(m_base, "TEX")==0) ||
+			(strcmp(m_base, "TEX32")==0)) {
+		printCuobjdumpPredicate();
+		output("tex.1d.v4.f32.s32 ");
+		stringListPiece *currPiece;
+		stringListPiece reg;
+		std::string addrReg, tex_id;
+		currPiece = m_operands->getListStart();
+		output("{");
+		printCuobjdumpOperand(currPiece, "", "");
+		output(",_,_,_} , ");
+		reg = *currPiece;
+		currPiece = currPiece->nextString;
+		tex_id = currPiece->stringText;
+		unsigned int tex_id_int;
+		std::stringstream ss;
+		ss << std::hex << tex_id;
+		ss >> tex_id_int;
+		std::list<std::string>::iterator texIter = texList.begin();
+		for (int i=0; i<tex_id_int; i++) {
+			assert (texIter!=texList.end());
+			texIter++;
+		}
+		output((*texIter).c_str());
+		output(",{");
+		printCuobjdumpOperand(&reg, "", "");
+		output(",_,_,_};");
+	}
+	else if(strcmp(m_base, "EXIT")==0) {
+		printCuobjdumpPredicate();
+		output("exit");
+		output(";");
+	}
+	else if(strcmp(m_base, "GRED")==0) {
+		printCuobjdumpPredicate();
+		// ptx instruction atom can be used to perform reduction using destination register '_'
+		output("atom.global");
+		printCuobjdumpBaseModifiers();
+			printCuobjdumpTypeModifiers();
+		output(" _, ");
+		printCuobjdumpOperands();
+		output(";");
+	} else if(strcmp(m_base, "GATOM")==0) {
+		printCuobjdumpPredicate();
+		output("atom.global");
+		printCuobjdumpBaseModifiers();
+		printCuobjdumpTypeModifiers();
+
+		printCuobjdumpOperands();
+		output(";");
+	}else if(strcmp(m_base, "PBK")==0) {
+		//PKB specifies the target of later BRK (break) instructions
+		// Here we convert it to nop and store the target
+		output("nop;");
+		breaktarget = m_operands->getListStart()->stringText;
+	} else if(strcmp(m_base, "BRK")==0) {
+		printCuobjdumpPredicate();
+		/*
+		 * Convert it **at compile time** to a branch to the break target saved earlier
+		 * Correct operation not guaranteed
+		 */
+
+		output("bra ");
+		output(breaktarget.c_str());
+		output(";");
+	} else if(strcmp(m_base, "C2R")==0) {
+		printCuobjdumpPredicate();
+		output("mov.u32");
+		printCuobjdumpOperands();
+		output(";");
+	} else if(strcmp(m_base, "R2C")==0) {
+		printCuobjdumpPredicate();
+		output("mov.pred");
+		printCuobjdumpOperands();
+		output(";");
+	}
+	else
+	{
+		printf("Unknown Instruction: ");
+		printf(m_base);
+		printf("\n");
+		output("Unknown Instruction: ");
+		output(m_base);
+		assert(0);
+	}
+
+}
+
diff --git a/cuobjdump_to_ptxplus/cuobjdumpInst.h b/cuobjdump_to_ptxplus/cuobjdumpInst.h
new file mode 100644
index 0000000..20d31be
--- /dev/null
+++ b/cuobjdump_to_ptxplus/cuobjdumpInst.h
@@ -0,0 +1,22 @@
+#include "decudaInstList.h"
+
+class cuobjdumpInst: public decudaInst
+{
+
+public:
+	//Constructor
+	cuobjdumpInst();
+
+	bool checkCubojdumpLabel(std::list<std::string> labelList, std::string label);
+
+	void printCuobjdumpLabel(std::list<std::string> labelList);
+	void printCuobjdumpPredicate();
+	void printCuobjdumpTypeModifiers();
+	void printCuobjdumpBaseModifiers();
+	void printCuobjdumpOperand(stringListPiece* currentPiece, std::string operandDelimiter, const char* base);
+	void printCuobjdumpOperandlohi(std::string op);
+	void printCuobjdumpOperands();
+
+	void printCuobjdumpPtxPlus(std::list<std::string> labelList, std::list<std::string> texList);
+};
+
diff --git a/cuobjdump_to_ptxplus/cuobjdumpInstList.cc b/cuobjdump_to_ptxplus/cuobjdumpInstList.cc
new file mode 100644
index 0000000..5eaf67f
--- /dev/null
+++ b/cuobjdump_to_ptxplus/cuobjdumpInstList.cc
@@ -0,0 +1,553 @@
+#include <iostream>
+#include "cuobjdumpInstList.h"
+
+extern void output(const char * text);
+
+//Constructor
+cuobjdumpInstList::cuobjdumpInstList()
+{
+        //initilize everything to empty
+}
+
+
+
+
+//TODO: Some register processing work is supposed to be done here.
+void cuobjdumpInstList::addCuobjdumpRegister(std::string reg, bool lo)
+{
+	int vectorFlag = 0;
+	char * regString;
+	regString = new char [reg.size()+1];
+
+	stringList* typeModifiers = getListEnd().getTypeModifiers();
+	const char* baseInst = getListEnd().getBase();
+
+	//TODO: support for 64bit vectors and 128bit vectors
+	if((strcmp(baseInst, "DADD")==0) || (strcmp(baseInst, "DMUL")==0) || (strcmp(baseInst, "DFMA")==0) ||
+		((typeModifiers->getSize()==1) &&
+		(strcmp((typeModifiers->getListStart()->stringText), ".S64")==0) &&
+		((strcmp(baseInst, "G2R")==0)||(strcmp(baseInst, "R2G")==0)||
+		(strcmp(baseInst, "GLD")==0)||(strcmp(baseInst, "GST")==0)||
+		(strcmp(baseInst, "LST")==0))))
+	{
+		vectorFlag = 64;
+	}
+	else if((typeModifiers->getSize()==1) && (strcmp((typeModifiers->getListStart()->stringText), ".S128")==0))
+	{
+		vectorFlag = 128;
+	}
+
+	//TODO: does the vector flag ever need to be set?
+	std::string parsedReg = parseCuobjdumpRegister(reg, lo, vectorFlag);
+	
+	strcpy(regString, parsedReg.c_str());
+
+	getListEnd().addOperand(regString);
+}
+
+// add memory operand
+// memType: 0=constant, 1=shared, 2=global, 3=local
+int currconstmem =1;
+void cuobjdumpInstList::addCuobjdumpMemoryOperand(std::string mem, int memType) {
+	std::string origMem = mem;
+	bool neg = false;
+
+	// If constant memory type, add prefix for entry specific constant memory
+	if(memType == 0) {
+		// Global memory c14
+		// Replace this with the actual global memory name
+		if(mem.substr(0,1) == "-") {
+			//Remove minus sign if exists
+			mem = mem.substr(1, mem.size()-1);
+			neg = true;
+		}
+
+		if(mem.substr(0, 7) == "c [0xe]") {
+			// Find the global memory identifier based on the offset provided
+			int offset;
+			sscanf(mem.substr(9,mem.size()-10).c_str(), "%x", &offset);
+			// Find memory
+			bool found = false;
+			std::list<globalMemory>::iterator g;
+			for(g=m_globalMemoryList.begin(); g!=m_globalMemoryList.end(); ++g) {
+				if(g->offset == offset) {
+					mem = g->name;
+					found = true;
+					break;
+				}
+			}
+
+			if(!found) {
+				printf("Could not find a global memory with this offset in: %s\n", mem.c_str());
+				output("Could not find a global memory with this offset.\n");
+				assert(0);
+			}
+
+		}
+		else if(mem.substr(0, 7) == "c [0x0]"){
+			mem = "constant0" + mem.substr(7, mem.length());
+		}
+		else if(mem.substr(0, 5) == "c [0x"){
+			if(currconstmem != m_entryList.size()) currconstmem++;
+			std::string out;
+			out = "constant1" + m_entryList.back().m_entryName + mem.substr(8);
+			mem = out.c_str();
+		}
+		else {
+			output("Unrecognized memory type:");
+			output(mem.c_str());
+			output("\n");
+			assert(0);
+		}
+
+		if (neg) {
+			mem = "-"+mem;
+		}
+	}
+
+	// Local memory
+	if(memType == 3) {
+		std::stringstream out;
+		printf("Trying to find lmem for: %s\n", m_entryList.back().m_entryName.c_str());
+		printf("Original memory: %s\n", mem.c_str());
+		assert(kernellmemmap[m_entryList.back().m_entryName] !=0 );
+		out << "l" << kernellmemmap[m_entryList.back().m_entryName];// << mem;
+		mem = out.str();
+	}
+
+	// Add the memory operand to instruction operand list
+	char* memName = new char [strlen(mem.c_str())+1];
+	strcpy(memName, mem.c_str());
+	getListEnd().addOperand(memName);
+}
+
+// increment register list and parse register
+std::string cuobjdumpInstList::parseCuobjdumpRegister(std::string reg, bool lo, int vectorFlag)
+{
+	std::string origReg = reg;
+	// Make sure entry list is not empty
+	if(m_entryList.size() == 0) {
+		output("ERROR: Adding a register before adding an entry.\n");
+		assert(0);
+	}
+
+	// remove minus sign if exists
+	if(reg.substr(0,1) == "-")
+		reg = reg.substr(1, reg.size()-1);
+
+	// if lo or hi register, get register name only (remove 'H' or 'L')
+	if(lo)
+		reg = reg.substr(0, reg.size()-1);
+
+	// Increase register number if needed
+	// Two types of registers, R# or A#
+	if(reg.substr(0, 1) == "R") {
+		reg = reg.substr(1, reg.size()-1);
+		int regNum = atoi(reg.c_str());
+
+		// Remove register overlap at 64
+		// TODO: is this still needed?
+		/*if(regNum > 63 && regNum < 124) {
+			regNum -= 64;
+			// Fix the origReg string
+			std::stringstream out;
+			out << ((origReg.substr(0,1)=="-") ? "-" : "")
+			    << "$r" << regNum
+			    << (lo ? origReg.substr(origReg.size()-3, 3) : "");
+			origReg = out.str();
+		}*/
+
+		if(vectorFlag==64)
+			regNum += 1;
+		if(vectorFlag==128)
+			regNum += 3;
+
+		if( m_entryList.back().m_largestRegIndex < regNum && regNum < 124 )
+			m_entryList.back().m_largestRegIndex = regNum;
+		else if( regNum == 124 )
+			m_entryList.back().m_reg124 = true;
+	} else if(reg.substr(0, 1) == "A") {
+		reg = reg.substr(1, reg.size()-1);
+		int regNum = atoi(reg.c_str());
+
+		if( m_entryList.back().m_largestOfsRegIndex < regNum && regNum < 124 )
+			m_entryList.back().m_largestOfsRegIndex = regNum;
+	} else if(reg == "o [0x7f]") {
+		m_entryList.back().m_oreg127 = true;
+	} else if (reg.substr(0,3) == "SR_") {
+		if(reg.substr(3,3)=="Tid") {
+			origReg = "%%tid";
+			if(reg.substr(7,1)=="X") {
+				origReg += ".x";
+			}
+		}
+	} else {
+		output("ERROR: unknown register type.\n");
+		printf("\nERROR: unknown register type: ");
+		printf(reg.c_str());
+		printf("\n");
+		assert(0);
+	}
+	return origReg;
+}
+
+// pred|reg double operand
+void cuobjdumpInstList::addCuobjdumpDoublePredReg(std::string pred, std::string reg, bool lo)
+{
+	std::string parsedPred = parseCuobjdumpPredicate(pred);
+	std::string parsedReg = parseCuobjdumpRegister(reg, lo, 0);
+
+	std::string doublePredReg;
+	if(
+		strcmp(getListEnd().getBase(), "DSET")==0 ||
+		strcmp(getListEnd().getBase(), "FSET")==0 ||
+		strcmp(getListEnd().getBase(), "ISET")==0
+	)
+		doublePredReg = parsedPred + "/" + parsedReg;
+	else
+		doublePredReg = parsedPred + "|" + parsedReg;
+
+	char* doublePredRegName = new char [strlen(doublePredReg.c_str())];
+	strcpy(doublePredRegName, doublePredReg.c_str());
+	getListEnd().addOperand(doublePredRegName);
+}
+
+std::string cuobjdumpInstList::parseCuobjdumpPredicate(std::string pred)
+{
+	std::string origPred = pred;
+
+	// Make sure entry list is not empty
+	if(m_entryList.size() == 0) {
+		output("ERROR: Adding a predicate before adding an entry.\n");
+		assert(0);
+	}
+
+	// increase predicate numbers if needed
+	pred = pred.substr(2, pred.size()-2);
+	int predNum = atoi(pred.c_str());
+	if( m_entryList.back().m_largestPredIndex < predNum )
+		m_entryList.back().m_largestPredIndex = predNum;
+
+	return origPred;
+}
+
+void cuobjdumpInstList::addCubojdumpLabel(std::string label)
+{
+
+	if(!(m_entryList.back().m_labelList.empty()))
+	{
+		std::list<std::string>::iterator labelIterator;
+
+		for( labelIterator=m_entryList.back().m_labelList.begin(); labelIterator!=m_entryList.back().m_labelList.end(); labelIterator++ )
+		{
+			if(label.compare(*labelIterator) == 0)
+				return;
+		}
+	}
+
+	m_entryList.back().m_labelList.push_back(label);
+}
+
+void cuobjdumpInstList::setConstMemoryMap(const char* kernelname, int index){
+	std::string kernel = kernelname;
+	kernel = kernel.substr(14, kernel.length()-1);
+	kernel = kernel.substr(0, kernel.find("\t"));
+	printf("Setting kernelcmemmap[%s]=%d\n", kernel.c_str(), index);
+	kernelcmemmap[kernel] = index;
+}
+
+void cuobjdumpInstList::setLocalMemoryMap(const char* kernelname, int index){
+	std::string kernel = kernelname;
+	kernel = kernel.substr(10, kernel.length()-1);
+	kernel = kernel.substr(0, kernel.find("\t"));
+	printf("Setting kernellmemmap[%s]=%d\n", kernel.c_str(), index);
+	kernellmemmap[kernel] = index;
+}
+
+
+//NOT USED
+void cuobjdumpInstList::reverseConstMemory() {
+	int total = kernelcmemmap.size();
+	for (	std::map<std::string,int>::iterator iter = kernelcmemmap.begin();
+			iter != kernelcmemmap.end();
+			iter++){
+		(*iter).second = total - (*iter).second;
+	}
+}
+
+
+// create new entry specific constant memory "bank"
+void cuobjdumpInstList::addEntryConstMemory(int index, int entryIndex)
+{
+	constMemory newConstMem;
+	newConstMem.index = index;
+	newConstMem.entryIndex = entryIndex;
+	m_constMemoryList.push_back(newConstMem);
+}
+
+void cuobjdumpInstList::addEntryConstMemory2(char* kernelname)
+{
+	std::string kernel = kernelname;
+	kernel = kernel.substr(14, kernel.length()-1);
+	kernel = kernel.substr(0, kernel.find("\t"));
+	constMemory2 newConstMem2;
+	newConstMem2.kernel = strdup(kernel.c_str());
+	m_constMemoryList2.push_back(newConstMem2);
+}
+
+void cuobjdumpInstList::addEntryLocalMemory(int value, int entryIndex)
+{
+	localMemory newLocalMem;
+	newLocalMem.value = value;
+	newLocalMem.entryIndex = entryIndex;
+	m_localMemoryList.push_back(newLocalMem);
+}
+
+
+// Read in constant memory from bin file
+// Two cases of constant memory have been noticed so far
+// 1 - All the constant memory is initialized in original ptx file. The assembler combines all this memory into c0
+// 2 - Constant memory is declared in ptx, but not initialized (initialized by host). The assembler still calls this c0
+void cuobjdumpInstList::readConstMemoryFromElfFile(std::string elf)
+{
+	unsigned k=1;
+	printf("Trying to find constant memory in elf file:\n");
+
+	// Get each constant segment
+	const boost::regex constPattern("^\\.nv\\.constant1\\.[^\n]+\n[ x0-9a-f\t]+$");
+	// Parse each constseg
+	const boost::sregex_token_iterator end;
+	for (
+		boost::sregex_token_iterator i(elf.begin(),elf.end(), constPattern);
+		i != end;
+		++i
+		)
+	{
+		std::string memseg = *i;
+		boost::smatch memResult;
+
+		const boost::regex memValuePattern("(0x[A-Fa-f0-9]{8,8})");
+
+		bool memExists = boost::regex_search(memseg, memResult, memValuePattern);
+
+
+		std::list<std::string> c1;
+		std::list<std::string>::iterator it = c1.begin();
+
+		const boost::sregex_token_iterator end2;
+		for (
+			boost::sregex_token_iterator j(memseg.begin(),memseg.end(), memValuePattern);
+			j != end2;
+			++j ){
+			c1.insert(it, *j);
+		}
+
+		addEntryConstMemory(1, k);
+		setConstMemoryType(".u32");
+
+		std::list<std::string>::iterator c;
+		if(c1.size() > 0) {
+			for(c=c1.begin(); c!=c1.end(); ++c) {
+				std::string a = *c;
+				//printf("%s ", a.c_str());
+				addConstMemoryValue(a);
+			}
+		}
+
+
+		printf("Found constant memory\n");
+		printf(memseg.c_str());
+		printf("\n");
+		k++;
+	}
+	m_kernelCount = k-1;
+}
+void cuobjdumpInstList::setKernelCount(int k){
+	m_kernelCount = k;
+}
+void cuobjdumpInstList::readOtherConstMemoryFromBinFile(std::string binString)
+{
+	// Initialize a list to store memory values
+	// std::list<std::string> c0;
+
+	// Get each code segment
+	//const boost::regex codePattern("(code \\{[^\\{\\}]*(const \\{[^\\{\\}]*(mem \\{[^\\{\\}]*\\}[^\\{\\}]*)+\\}[^\\{\\}]*)+bincode \\{[^\\{\\}]*\\}[^\\{\\}]*\\})");
+	const boost::regex codePattern("(code \\{[^\\{\\}]*(const \\{[^\\{\\}]*(mem \\{[^\\{\\}]*\\}[^\\{\\}]*)+\\}[^\\{\\}]*)*bincode \\{[^\\{\\}]*\\}[^\\{\\}]*\\})");
+
+	int k=1;
+
+	// Parse each codeseg
+	const boost::sregex_token_iterator end;
+	for(
+		boost::sregex_token_iterator i(binString.begin(),binString.end(), codePattern);
+		i != end;
+		++i
+	)
+	{
+		std::list<std::string> c1;
+
+		// For each code segment, get the seg numbers and memory values string
+		std::string codeSeg_s = *i;
+		std::string segnum_s, lmem_s, mem;
+		int segnum;
+		int lmem;
+
+		boost::smatch segnumResult;
+		boost::smatch lmemResult;
+		boost::smatch memResult;
+
+		const boost::regex segnumPattern("segnum\\s*=\\s(\\d*)");
+		const boost::regex lmemPattern("lmem\\s*=\\s(\\d*)");
+		const boost::regex memPattern("mem \\{([^\\}]*)\\}");
+
+		boost::regex_search(codeSeg_s, segnumResult, segnumPattern);
+		boost::regex_search(codeSeg_s, lmemResult, lmemPattern);
+		bool memExists = boost::regex_search(codeSeg_s, memResult, memPattern);
+
+		lmem_s = lmemResult[1];
+		lmem = atoi(lmem_s.c_str());
+
+		addEntryLocalMemory(lmem, k);
+
+		if(memExists)
+		{
+		segnum_s = segnumResult[1];
+		segnum = atoi(segnum_s.c_str());
+
+		mem = memResult[1];
+		const boost::regex memValuePattern("(0x[A-Fa-f0-9]{8,8})");
+
+		std::list<std::string>::iterator it = c1.begin();
+
+		const boost::sregex_token_iterator end2;
+		for (
+			boost::sregex_token_iterator j(mem.begin(),mem.end(), memValuePattern);
+			j != end2;
+			++j
+			)
+		{
+			c1.insert(it, *j);
+		}
+
+		addEntryConstMemory(segnum, k);
+		setConstMemoryType(".u32");
+
+		std::list<std::string>::iterator c;
+		if(c1.size() > 0) {
+			for(c=c1.begin(); c!=c1.end(); ++c) {
+				std::string a = *c;
+				//printf("%s ", a.c_str());
+				addConstMemoryValue(a);
+			}
+		}
+		}
+		k++;
+	}
+	m_kernelCount = k-1;
+}
+
+void cuobjdumpInstList::printCuobjdumpInstList()
+{
+	// Each entry
+	std::list<decudaEntry>::iterator e;
+	for(e=m_entryList.begin(); e!=m_entryList.end(); ++e) {
+
+
+
+
+		for(	std::list<decudaInst>::iterator currentInst=e->m_instList.begin();
+				currentInst!=e->m_instList.end();
+				++currentInst) {
+			// Output the instruction
+			output("\t");
+			currentInst->printDecudaInst();
+			output("\n");
+		}
+	}
+}
+
+void cuobjdumpInstList::printCuobjdumpLocalMemory()
+{
+	for(	std::list<localMemory>::iterator i=m_localMemoryList.begin();
+			i!=m_localMemoryList.end();
+			++i) {
+		char line[40];
+		//if(i->value > 0)
+		{
+			sprintf(line, ".local .b8 l%d[%d];\n", i->entryIndex, i->value);
+			output(line);
+		}
+	}
+}
+
+void cuobjdumpInstList::printCuobjdumpPtxPlusList(cuobjdumpInstList* headerInfo)
+{
+	output("\n");
+	printMemory();
+	printCuobjdumpLocalMemory();
+	// Each entry
+	std::list<decudaEntry>::reverse_iterator e;
+	for(e=m_entryList.rbegin(); e!=m_entryList.rend(); ++e) {
+
+		output("\n");
+
+		// Output the header information for this entry using headerInfo
+		// First, find the matching entry in headerInfo
+		decudaEntry headerEntry;
+
+		if( headerInfo->findEntry(e->m_entryName, headerEntry) ) {
+			// Entry for current header found, print it out
+			std::list<decudaInst>::iterator headerInstIter;
+			for(headerInstIter=headerEntry.m_instList.begin();
+				headerInstIter!=headerEntry.m_instList.end();
+				++headerInstIter) {
+				if(headerInstIter!=headerEntry.m_instList.begin()) {
+					output("\t");
+				}
+				headerInstIter->printHeaderPtx();
+				output("\n");
+			}
+			output("{\n");
+		} else {
+			// Couldn't find this entry in ptx file
+			// Check if it is a dummy entry
+			if(e->m_entryName == "__cuda_dummy_entry__") {
+				output(".entry ");
+				output("__cuda_dummy_entry__");
+				output("\n");
+				output("{\n");
+			} else {
+				output("Mismatch in entry names between decuda output and original ptx file.\n");
+				assert(0);
+			}
+		}
+		assert( &*e != NULL);
+		printRegNames(*e);
+		printPredNames(*e);
+		printOutOfBoundRegisters(*e);
+		output("\n");
+
+		for(std::list<decudaInst>::iterator currentInst=e->m_instList.begin(); currentInst!=e->m_instList.end(); ++currentInst){
+			// Output the instruction
+			//cuobjdumpInst* outputInst = &*currentInst;
+			cuobjdumpInst* outputInst = static_cast<cuobjdumpInst*>(&*currentInst);
+			output("\t");
+			//outputInst->printCuobjdumpPtxPlus(m_entryList.back().m_labelList);
+			outputInst->printCuobjdumpPtxPlus(e->m_labelList, this->m_realTexList);
+			output("\n");
+		}
+		output("\n\tl_exit: exit;\n");
+		output("}\n");
+	}
+}
+
+void cuobjdumpInstList::addConstMemoryPtr(const char* offset, const char* size, const char* name){
+	constMemoryPtr ptr;
+	ptr.offset = atoi(offset);
+	ptr.bytes = atoi(size);
+	ptr.name = name;
+	ptr.destination = "constant0";
+	m_constMemoryPtrList.push_back(ptr);
+	printf("\naddConstMemoryPtr: %s, size: %d, offset: %d\n", ptr.name.c_str(), ptr.bytes, ptr.offset);
+}
+
diff --git a/cuobjdump_to_ptxplus/cuobjdumpInstList.h b/cuobjdump_to_ptxplus/cuobjdumpInstList.h
new file mode 100644
index 0000000..1bc8f10
--- /dev/null
+++ b/cuobjdump_to_ptxplus/cuobjdumpInstList.h
@@ -0,0 +1,47 @@
+#include "cuobjdumpInst.h"
+
+// Used for local memory segments
+struct localMemory
+{
+	int value;
+	int entryIndex;
+};
+
+class cuobjdumpInstList: public decudaInstList
+{
+protected:
+	std::string parseCuobjdumpPredicate(std::string pred);
+	int m_kernelCount;
+	std::map<std::string,int>kernelcmemmap;
+	std::map<std::string,int>kernellmemmap;
+	std::list<localMemory> m_localMemoryList;
+
+public:
+	//Constructor
+	cuobjdumpInstList();
+
+	void setKernelCount(int k);
+	void readConstMemoryFromElfFile(std::string elf);
+
+	void addCuobjdumpRegister(std::string reg, bool lo=false); //add register
+	void addCuobjdumpMemoryOperand(std::string mem, int memType);
+	std::string parseCuobjdumpRegister(std::string reg, bool lo, int vectorFlag);
+	void addCuobjdumpDoublePredReg(std::string pred, std::string reg, bool lo=false);
+
+	void addCubojdumpLabel(std::string label);
+
+	void addEntryConstMemory(int index, int entryIndex);
+	void addEntryConstMemory2(char* kernel);
+	void addConstMemoryPtr(const char* bytes, const char* offset, const char* name);
+	void setConstMemoryMap(const char* kernelname, int index);
+	void setLocalMemoryMap(const char* kernelname, int index);
+	void reverseConstMemory();
+	void addEntryLocalMemory(int value, int entryIndex);
+	void readOtherConstMemoryFromBinFile(std::string binString); // read in constant memory from bin file
+
+	void printCuobjdumpLocalMemory();
+
+	void printCuobjdumpInstList();
+	void printCuobjdumpPtxPlusList(cuobjdumpInstList* headerInfo);
+};
+
diff --git a/cuobjdump_to_ptxplus/cuobjdump_to_ptxplus.cc b/cuobjdump_to_ptxplus/cuobjdump_to_ptxplus.cc
new file mode 100644
index 0000000..e8d45ef
--- /dev/null
+++ b/cuobjdump_to_ptxplus/cuobjdump_to_ptxplus.cc
@@ -0,0 +1,162 @@
+#include <iostream>
+#include "cuobjdumpInstList.h"
+#include <stdio.h>
+#include <fstream>
+
+using namespace std;
+
+cuobjdumpInstList *g_instList = new cuobjdumpInstList();
+cuobjdumpInstList *g_headerList = new cuobjdumpInstList();
+
+int sass_parse();
+extern "C" FILE *sass_in;
+
+int ptx_parse();
+extern "C" FILE *ptx_in;
+
+int elf_parse();
+extern "C" FILE *elf_in;
+
+extern int g_error_detected;
+
+FILE *bin_in;
+FILE *ptxplus_out;
+
+void output(const char * text)
+{
+	printf(text);
+	fprintf(ptxplus_out, text);
+}
+
+std::string fileToString(const char * fileName) {
+	ifstream fileStream(fileName, ios::in);
+	string text, line;
+	while(getline(fileStream,line)) {
+		text += (line + "\n");
+	}
+	fileStream.close();
+	return text;
+}
+
+std::string extractFilename( const std::string& path )
+{
+	return path.substr( path.find_last_of( '/' ) +1 );
+}
+int main(int argc, char* argv[])
+{
+	//TODO: Output to file not yet supported.
+	if(argc != 5)
+	{
+		cout << "Usage: cuobjdump_to_ptxplus ptxfile sassfile elffile ptxplusfile(output)\n";
+		return 0;
+	}
+	int result;
+	string ptxfile = argv[1];
+	string sassfile = argv[2];
+	string elffile = argv[3];
+	string ptxplusfile = argv[4];
+
+	/*
+	char commandline[1024];
+
+	snprintf(commandline,1024,"cuobjdump -ptx %s > %s.ptx", exefile.c_str(), basefilename.c_str());
+	fflush(stdout);
+	result = system(commandline);
+	if (result) {printf("ERROR: could not execute %s\n", commandline); exit(1);}
+
+
+	snprintf(commandline,1024,"sed '/arch\\ =\\ sm_20/,/1a$/d' %s.ptx > %s.stripped2.ptx", basefilename.c_str(), basefilename.c_str());
+	fflush(stdout);
+	result = system(commandline);
+	if (result) {printf("ERROR: could not execute %s\n", commandline); exit(1);}
+
+	snprintf(commandline,1024,"grep -v \"^arch\\ =\\|Fatbin\\ \\|^code\\ version\\|===========\\|producer\\ =\\|host\\ =\\|compile_size\\ =\\|identifier\\ =\" %s.stripped2.ptx > %s.stripped.ptx", basefilename.c_str(), basefilename.c_str());
+	fflush(stdout);
+	result = system(commandline);
+	if (result) {printf("ERROR: could not execute %s\n", commandline); exit(1);}
+
+	snprintf(commandline,1024,"cuobjdump -sass %s > %s.sass", exefile.c_str(), basefilename.c_str());
+	fflush(stdout);
+	result = system(commandline);
+	if (result) {printf("ERROR: could not execute %s\n", commandline); exit(1);}
+
+	snprintf(commandline,1024,"sed '/arch\\ =\\ sm_20/,/1a$/d' %s.sass > %s.stripped.sass", basefilename.c_str(), basefilename.c_str());
+	fflush(stdout);
+	result = system(commandline);
+	if (result) {printf("ERROR: could not execute %s\n", commandline); exit(1);}
+
+	snprintf(commandline,1024,"cuobjdump -elf %s > %s.elf", exefile.c_str(), basefilename.c_str());
+	fflush(stdout);
+	result = system(commandline);
+	if (result) {printf("ERROR: could not execute %s\n", commandline); exit(1);}
+
+	snprintf(commandline,1024,"sed '/arch\\ =\\ sm_20/,/1a$/d' %s.elf > %s.stripped.elf", basefilename.c_str(), basefilename.c_str());
+	fflush(stdout);
+	result = system(commandline);
+	if (result) {printf("ERROR: could not execute %s\n", commandline); exit(1);}
+	*/
+
+	sass_in = fopen(sassfile.c_str(), "r" );
+	ptx_in = fopen(ptxfile.c_str(), "r" );
+	elf_in = fopen(elffile.c_str(), "r");
+	ptxplus_out = fopen(ptxplusfile.c_str(), "w" );
+
+
+	std::string elf = fileToString(elffile.c_str());
+
+	printf("RUNNING cuobjdump_to_ptxplus ...\n");
+
+
+	printf("Parsing .elf file %s\n", elffile.c_str());
+	elf_parse();
+	printf("Finished parsing .elf file %s\n", elffile.c_str());
+
+	//Parse original ptx
+	printf("Parsing .ptx file %s\n", ptxfile.c_str());
+	ptx_parse();
+	if (g_error_detected){
+		assert(0 && "ptx parsing failed");
+	}
+	printf("Finished parsing .ptx file %s\n", ptxfile.c_str());
+
+	// Copy real tex list from ptx to ptxplus instruction list
+	g_instList->setRealTexList(g_headerList->getRealTexList());
+	// Insert constant memory from bin file
+//	g_instList->readConstMemoryFromElfFile(elf);
+//	g_instList->readOtherConstMemoryFromBinFile(fileToString(binFilename));
+
+	// Insert global memory from bin file
+//	g_instList->readGlobalMemoryFromBinFile(fileToString(binFilename));
+
+	// Parse cuobjdump output
+	printf("Parsing .sass file %s\n", sassfile.c_str());
+	sass_parse();
+	printf("Finished parsing .sass file %s\n", sassfile.c_str());
+	/*
+	printf("################################################## Instruction List dump\n");
+	g_instList->printCuobjdumpInstList();
+	printf("################################################## END Instruction List dump\n");
+	printf("################################################## Header Instructions Dump\n");
+	g_headerList->printCuobjdumpInstList();
+	printf("################################################## END Header Instructions Dump\n");
+	*/
+	// Print ptxplus
+	output("//HEADER\n");
+	g_headerList->printHeaderInstList();
+	output("//END HEADER\n\n\n");
+	output("//INSTRUCTIONS\n");
+	g_instList->printCuobjdumpPtxPlusList(g_headerList);
+	output("//END INSTRUCTIONS\n");
+	// TODO: remove this. Prints recorded cuobjdump output.
+//	g_instList->printCuobjdumpInstList();
+
+	fclose(sass_in);
+	fclose(ptx_in);
+//	fclose(bin_in);
+	fclose(ptxplus_out);
+
+	printf("DONE. \n");
+
+	return 0;
+}
+
diff --git a/cuobjdump_to_ptxplus/elf.l b/cuobjdump_to_ptxplus/elf.l
new file mode 100644
index 0000000..3457a7a
--- /dev/null
+++ b/cuobjdump_to_ptxplus/elf.l
@@ -0,0 +1,114 @@
+	/*Lex file for elf files output by cuobjdump*/
+%{
+#include <stdio.h>
+#include <string.h>
+#include "elf_parser.hpp"
+
+#define yylval elf_lval
+
+void elf_error(const char*);
+%}
+
+%option stack
+%option noyywrap
+%option yylineno
+
+%x cmem
+%x symtab
+
+whitespace	[ \t]
+newline		[\n]
+notnewline	[^\n]
+newlines	{newline}+
+ident		[[:alnum:]_\.]
+hexdigit	[[:digit:]abcdef]
+%%
+
+<cmem,INITIAL>\.nv\.constant1\.{notnewline}*{newline}	{
+	BEGIN(cmem);
+	yylval.string_value = strdup(yytext); 
+	return C1BEGIN;
+}
+
+<cmem,INITIAL>\.nv\.local{notnewline}*{newline}	{
+	BEGIN(INITIAL);
+	yylval.string_value = strdup(yytext); 
+	return LOCALMEM;
+}
+
+<cmem,INITIAL>\.nv\.constant0{notnewline}*{newline}	{
+	BEGIN(cmem);
+	return C0BEGIN;
+}
+
+
+\.section" "\.symtab\tSYMTAB{newline}	{
+	BEGIN(symtab);
+	return STBEGIN;
+}
+
+\.section{notnewline}*{newline}	{
+	//These are the lines before the symtab
+}
+
+<cmem,INITIAL>\.{notnewline}*{newline}	{
+	BEGIN(INITIAL);
+}
+
+{notnewline}*{newline} {
+	//printf("Discarding line %s", yytext);
+}
+
+<cmem>"  "{newline}?	{
+	//printf("#SPACE#");
+	return SPACE2;
+}
+
+<cmem>0x[0-9a-fA-F]{8}	{
+	//printf("CMEM VAL: %s \n", yytext);
+	yylval.string_value = strdup(yytext); 
+	return CMEMVAL;
+}
+
+<symtab>{whitespace}*[[:alpha:]]+{whitespace}+[[:alpha:]]+{whitespace}+[[:alpha:]]+{whitespace}+[[:alpha:]]+{whitespace}+[[:alpha:]]+{whitespace}+[[:alpha:]]+{whitespace}+[[:alpha:]]+{whitespace}+{newline} {
+	return STHEADER;
+}
+
+<symtab>{whitespace}+ {}
+
+<symtab>{newline} {}
+
+<symtab>{hexdigit}+ {
+	yylval.string_value = strdup(yytext);
+	return NUMBER;
+}
+
+<symtab>{ident}+ {
+	yylval.string_value = strdup(yytext);
+	return IDENTIFIER;
+}
+
+<symtab>{newline}{whitespace}*{newline}{whitespace}* {
+	BEGIN(INITIAL);
+}
+	/*Reached end of file*/
+<<EOF>> return 0;
+
+	/*No other rule matched. Throw an error*/
+.	elf_error("Invalid token");
+
+%%
+void elf_error(const char* message)
+{
+	printf("   ");
+	printf(message);
+	printf(" near \"");
+	printf(yytext);
+	printf("\"");
+	printf(" on line "); 
+	char line[5];
+	sprintf(line, "%i", yylineno);
+	printf(line);
+	printf("\n");
+	abort();
+}
diff --git a/cuobjdump_to_ptxplus/elf.y b/cuobjdump_to_ptxplus/elf.y
new file mode 100644
index 0000000..8f5163e
--- /dev/null
+++ b/cuobjdump_to_ptxplus/elf.y
@@ -0,0 +1,81 @@
+	/*Yacc file for elf files output by cuobjdump*/
+%{
+#include <stdio.h>
+#include "cuobjdumpInstList.h"
+
+int yylex(void);
+void yyerror(const char*);
+//void addEntryConstMemory(int, int);
+//void setConstMemoryType(const char*);
+//void addConstMemoryValue(const char*);
+extern cuobjdumpInstList *g_instList;
+int cmemcount=1;
+int lmemcount=1;
+bool lastcmem = false;// false = constrant0, true = constant1
+%}
+%union {
+	char* string_value;
+}
+%token <string_value> C1BEGIN CMEMVAL SPACE2 C0BEGIN STBEGIN STHEADER
+%token <string_value> NUMBER HEXNUMBER IDENTIFIER LOCALMEM
+
+%%
+
+elffile	:	symtab program
+		;
+
+symtab	:	STBEGIN STHEADER stcontent
+		;
+
+stcontent	:	stcontent stline
+			|	stline;
+
+stline	:	NUMBER NUMBER NUMBER NUMBER NUMBER NUMBER IDENTIFIER {
+				if (strcmp($4, "11")==0) {
+					g_instList->addConstMemoryPtr($2, $3, $7);
+				}
+			}
+		|	NUMBER NUMBER NUMBER NUMBER NUMBER NUMBER {}
+		;
+
+program	:	program cmemsection
+		|	program localmemsection
+		|	{
+				g_instList->setKernelCount(cmemcount-1);
+				//g_instList->reverseConstMemory();
+			};
+
+localmemsection	:	LOCALMEM {
+						printf("Found LocalMem section number %d\n", lmemcount);
+						g_instList->addEntryLocalMemory(0, lmemcount);
+						g_instList->setLocalMemoryMap($1, lmemcount);
+						lmemcount++;
+					};
+
+cmemsection	:	C1BEGIN {
+					printf("Found ConstMem section number %d\n", cmemcount);
+					//g_instList->addEntryConstMemory(1, cmemcount);
+					g_instList->addEntryConstMemory2($1);
+					g_instList->setConstMemoryType2(".u32");
+					//g_instList->setConstMemoryType(".u32");
+					//g_instList->setConstMemoryMap($1,cmemcount);
+					cmemcount++;
+					lastcmem = true;
+				} cmemvals
+			|	C0BEGIN {
+					printf("Found ConstMem c0 section\n");
+					g_instList->addConstMemory(0);
+					g_instList->setConstMemoryType(".u32");
+					lastcmem = false;
+				} cmemvals;
+
+cmemvals	:	cmemvals CMEMVAL SPACE2 {
+					//printf("Found ConstMem value\n");
+					printf("addConstMemoryValue %s\n", $3);
+					if (lastcmem)
+						g_instList->addConstMemoryValue2($3);
+					else
+						g_instList->addConstMemoryValue($3);
+				}
+			|	;
+%%
diff --git a/cuobjdump_to_ptxplus/ptx_parser.h b/cuobjdump_to_ptxplus/ptx_parser.h
new file mode 100644
index 0000000..5111743
--- /dev/null
+++ b/cuobjdump_to_ptxplus/ptx_parser.h
@@ -0,0 +1,336 @@
+#include <cstdlib>
+#include <cstring>
+#include <string>
+#include <list>
+#include <map>
+#include <vector>
+#include <assert.h>
+#include <iostream>
+#include <sstream>
+#include "decudaInstList.h"
+#include <string.h>
+
+#define NON_ARRAY_IDENTIFIER 1
+#define ARRAY_IDENTIFIER_NO_DIM 2
+#define ARRAY_IDENTIFIER 3
+#define P_DEBUG 1
+#define DPRINTF(...) \
+   if(P_DEBUG) { \
+      printf("(%s:%s:%u) ", __FILE__, __FUNCTION__, __LINE__); \
+      printf(__VA_ARGS__); \
+      printf("\n"); \
+      fflush(stdout); \
+   }
+
+
+enum _memory_space_t {
+   undefined_space=0,
+   reg_space,
+   local_space,
+   shared_space,
+   param_space_unclassified,
+   param_space_kernel,  /* global to all threads in a kernel : read-only */
+   param_space_local,   /* local to a thread : read-writable */
+   const_space,
+   tex_space,
+   surf_space,
+   global_space,
+   generic_space
+};
+
+int g_error_detected;
+const char *g_filename = "";
+int g_func_decl;
+
+void set_symtab( void* a ) {DPRINTF("");}
+void end_function() {DPRINTF("");}
+void add_directive() {DPRINTF("");}
+void add_function_arg() {DPRINTF("");}
+void add_instruction() {DPRINTF("");}
+void add_file( unsigned a, const char *b ) {DPRINTF("");}
+void add_variables() {DPRINTF("");}
+void set_variable_type() {DPRINTF("");}
+void add_option(int a ) {DPRINTF("");}
+void add_array_initializer() {DPRINTF("");}
+void add_label( const char *a ) {DPRINTF("");}
+void set_return() {DPRINTF("");}
+void add_opcode( int a ) {DPRINTF("");}
+void add_pred( const char *a, int b, int c ) {DPRINTF("");}
+void add_scalar_operand( const char *a ) {DPRINTF("%s", a);}
+void add_neg_pred_operand( const char *a ) {DPRINTF("");}
+void add_address_operand( const char *a, int b ) {DPRINTF("%s", a);}
+void change_operand_lohi( int a ) {DPRINTF("");}
+void change_double_operand_type( int a ) {DPRINTF("");}
+void change_operand_neg( ) {DPRINTF("");}
+void add_double_operand( const char *a, const char *b ) {DPRINTF("");}
+void add_2vector_operand( const char *a, const char *b ) {DPRINTF("");}
+void add_3vector_operand( const char *a, const char *b, const char *c ) {DPRINTF("");}
+void add_4vector_operand( const char *a, const char *b, const char *c, const char *d ) {DPRINTF("");}
+void add_builtin_operand( int a, int b ) {DPRINTF("");}
+void add_memory_operand() {DPRINTF("");}
+void change_memory_addr_space( const char *a ) {DPRINTF("");}
+void add_literal_int( int a ) {DPRINTF("");}
+void add_literal_float( float a ) {DPRINTF("");}
+void add_literal_double( double a ) {DPRINTF("");}
+void add_extern_spec() {DPRINTF("");}
+void add_alignment_spec( int ) {DPRINTF("");}
+void add_pragma( const char *a ) {DPRINTF("");}
+void add_constptr(const char* identifier1, const char* identifier2, int offset) {DPRINTF("");}
+
+/*non-dummy stuff below this point*/
+
+
+
+extern decudaInstList *g_headerList;
+
+// Global variable to track if we are currently inside a entry directive
+bool inEntryDirective = false;
+// Global variable to track is we are currently inside the parameter definitions for an entry
+bool inParamDirective = false;
+
+bool inConstDirective = false;
+
+// Global variable to track if we are currently inside a tex directive
+bool inTexDirective = false;
+
+
+void add_identifier( const char *a, int b, unsigned c ) {
+	DPRINTF("name=%s", a);
+	if(inConstDirective){
+		//g_headerList->getListEnd()
+	}
+}
+
+void add_function_name( const char *headerInput )
+{
+	DPRINTF("name=%s", headerInput);
+	char* headerInfo = (char*) headerInput;
+	const char* compareString = g_headerList->getListEnd().getBase();
+
+	if((strcmp(compareString, ".entry")==0)||(strcmp(compareString, ".func")==0))
+	{
+		g_headerList->setLastEntryName(headerInfo);
+		g_headerList->getListEnd().addOperand(headerInfo);
+	}
+}
+
+//void add_space_spec(int headerInput)
+void add_space_spec( enum _memory_space_t spec, int value )
+{
+	DPRINTF("spec=%u", spec);
+	decudaInst *instEntry;
+	static int constmemindex=1;
+	switch(spec)
+	{
+		case param_space_unclassified:
+			if(inEntryDirective && inParamDirective) {
+				instEntry = new decudaInst();
+				instEntry->setBase(".param");
+				g_headerList->add(instEntry);
+			}
+			break;
+		case tex_space:
+			inTexDirective = true;
+			/*
+			instEntry = new decudaInst();
+			instEntry->setBase(".tex");
+			g_headerList->add(instEntry);
+			*/
+			break;
+		case const_space:
+			if(!inEntryDirective) {
+				/*
+				inConstDirective = true;
+				instEntry = new decudaInst();
+				instEntry->setBase(".const");
+				g_headerList->add(instEntry);
+				*/
+				//g_headerList->addConstMemory(constmemindex++);
+			}
+			break;
+	}
+}
+
+void add_scalar_type_spec( int headerInput )
+{
+	DPRINTF("");
+	const char* compareString = g_headerList->getListEnd().getBase();
+
+	if( (inEntryDirective && inParamDirective) || inTexDirective || inConstDirective)
+	{
+		switch(headerInput)
+		{
+			case S8_TYPE:
+				g_headerList->getListEnd().addBaseModifier(".s8");
+				break;
+			case S16_TYPE:
+				g_headerList->getListEnd().addBaseModifier(".s16");
+				break;
+			case S32_TYPE:
+				g_headerList->getListEnd().addBaseModifier(".s32");
+				break;
+			case S64_TYPE:
+				g_headerList->getListEnd().addBaseModifier(".s64");
+				break;
+			case U8_TYPE:
+				g_headerList->getListEnd().addBaseModifier(".u8");
+				break;
+			case U16_TYPE:
+				g_headerList->getListEnd().addBaseModifier(".u16");
+				break;
+			case U32_TYPE:
+				g_headerList->getListEnd().addBaseModifier(".u32");
+				break;
+			case U64_TYPE:
+				g_headerList->getListEnd().addBaseModifier(".u64");
+				break;
+			case F16_TYPE:
+				g_headerList->getListEnd().addBaseModifier(".f16");
+				break;
+			case F32_TYPE:
+				g_headerList->getListEnd().addBaseModifier(".f32");
+				break;
+			case F64_TYPE:
+				g_headerList->getListEnd().addBaseModifier(".f64");
+				break;
+			case B8_TYPE:
+				g_headerList->getListEnd().addBaseModifier(".b8");
+				break;
+			case B16_TYPE:
+				g_headerList->getListEnd().addBaseModifier(".b16");
+				break;
+			case B32_TYPE:
+				g_headerList->getListEnd().addBaseModifier(".b32");
+				break;
+			case B64_TYPE:
+				g_headerList->getListEnd().addBaseModifier(".b64");
+				break;
+			case PRED_TYPE:
+				g_headerList->getListEnd().addBaseModifier(".pred");
+				break;
+			default:
+				std::cout << "Unknown type spec" << "\n";
+				break;
+		}
+	}
+}
+
+//void version_header(double versionNumber)
+void add_version_info( float versionNumber, unsigned ext)
+{
+	DPRINTF("");
+	decudaInst *instEntry = new decudaInst();
+	instEntry->setBase(".version");
+	g_headerList->add(instEntry);
+
+
+	//convert double to char*
+	std::ostringstream strs;
+	strs << versionNumber;
+	char *versionNumber2 = strdup(strs.str().c_str());
+
+	g_headerList->getListEnd().addOperand(versionNumber2);
+	//g_headerList->getListEnd().addOperand("1.4");
+}
+
+void target_header(char* firstTarget)
+{
+	DPRINTF("%s", firstTarget);
+	decudaInst *instEntry = new decudaInst();
+	instEntry->setBase(".target");
+	g_headerList->add(instEntry);	
+
+	g_headerList->getListEnd().addOperand(firstTarget);
+}
+
+void target_header2(char* firstTarget, char* secondTarget)
+{
+	DPRINTF("%s, %s", firstTarget, secondTarget);
+	decudaInst *instEntry = new decudaInst();
+	instEntry->setBase(".target");
+	g_headerList->add(instEntry);	
+
+	g_headerList->getListEnd().addOperand(firstTarget);
+
+	g_headerList->getListEnd().addOperand(secondTarget);
+}
+
+void target_header3(char* firstTarget, char* secondTarget, char* thirdTarget)
+{
+	DPRINTF("%s, %s, %s", firstTarget, secondTarget, thirdTarget);
+	decudaInst *instEntry = new decudaInst();
+	instEntry->setBase(".target");
+	g_headerList->add(instEntry);
+
+	g_headerList->getListEnd().addOperand(firstTarget);
+
+	g_headerList->getListEnd().addOperand(secondTarget);
+
+	g_headerList->getListEnd().addOperand(thirdTarget);
+}
+
+void start_function( int a )
+{
+	DPRINTF("");
+	inEntryDirective = true;
+}
+
+void* reset_symtab()
+{
+	DPRINTF("");
+	inEntryDirective = false;
+	void* a;
+	return a;
+}
+
+void func_header(const char* headerBase)
+{
+	DPRINTF("%s", headerBase);
+	// If start of an entry
+	if((strcmp(headerBase, ".entry")==0)||(strcmp(headerBase, ".func")==0)) {
+		inEntryDirective = true;
+		g_headerList->addEntry("");
+		decudaInst *instEntry = new decudaInst();
+		instEntry->setBase(headerBase);
+		g_headerList->add(instEntry);
+
+	}
+}
+
+void func_header_info(const char* headerInfo)
+{
+	DPRINTF("%s", headerInfo);
+	const char* compareString = g_headerList->getListEnd().getBase();
+
+	if(inEntryDirective && !inTexDirective) {
+		g_headerList->getListEnd().addOperand(headerInfo);
+		// If start of parameters
+		if(strcmp(headerInfo,"(")==0)
+			inParamDirective = true;
+
+		// If end of parameters
+		if(strcmp(headerInfo,")")==0) {
+			inParamDirective = false;
+		}
+	} else if(inTexDirective) {
+		g_headerList->addTex(headerInfo);
+		inTexDirective = false;
+	} else if(inConstDirective){
+
+	} else {
+		printf("Unkown header info: #%s#\n", headerInfo);
+	}
+
+}
+
+void func_header_info_int(const char* s, int i)
+{
+	DPRINTF("%s %d", s, i);
+	if(inEntryDirective && !inTexDirective) {
+		g_headerList->getListEnd().addOperand(s);
+		char *buff = (char*) malloc(30*sizeof(char));
+		sprintf(buff, "%d", i);
+		assert (i>=0);
+		g_headerList->getListEnd().addOperand(buff);
+	}
+}
diff --git a/cuobjdump_to_ptxplus/sass.l b/cuobjdump_to_ptxplus/sass.l
new file mode 100644
index 0000000..a4beeb1
--- /dev/null
+++ b/cuobjdump_to_ptxplus/sass.l
@@ -0,0 +1,324 @@
+%{
+#include <stdio.h>
+#include "sass_parser.hpp"
+#include <string.h>
+
+
+/* Not sure if this is needed*/
+#define	YY_NEVER_INTERACTIVE	1
+
+#define yylval sass_lval
+
+void sass_error(const char*);
+%}
+%option noyywrap
+%option yylineno
+
+/*Regular expresions go here*/
+alpha		[a-zA-Z]
+numeric		[0-9]
+hexalpha	[A-Fa-f]
+alphanum	({alpha}|{numeric})
+hexdigit	({numeric}|{hexalpha})
+whitespace	[ \t]
+newline		[\n]
+anycharbutnewline [^\n]
+anychar		({newline}|{anycharbutnewline})
+pipe		\|
+
+identifier	({alpha}({alphanum}|["_$"])*)|([_$%]({alphanum}|["_$"])+)
+hexliteral	0[xX]{hexdigit}+U?
+octliteral	0[0-7]+U?
+binliteral	0[bB][01]+U?
+decliteral	([1-9][0-9]*U?)|0U?
+fltliteral	[0-9]+\.[0-9]+
+
+%%
+	/*Translation rules*/
+
+	/*Compiler directives*/
+
+	/*constant memory directives*/
+
+	/*guard predicate*/
+
+	/*label*/
+
+"[" yylval.string_value = strdup(yytext); return OSQBRACKET;
+"]" yylval.string_value = strdup(yytext); return CSQBRACKET; 
+
+	/*base instructions*/
+"BAR.ARV.WAIT b0, 0xfff"	yylval.string_value = strdup(yytext); return BAR;
+ADA	yylval.string_value = strdup(yytext); return ADA;
+BRA	yylval.string_value = strdup(yytext); return BRA;
+CAL	yylval.string_value = strdup(yytext); return CAL;
+COS	yylval.string_value = strdup(yytext); return COS;
+DADD	yylval.string_value = strdup(yytext); return DADD;
+DFMA	yylval.string_value = strdup(yytext); return DFMA;
+DMUL	yylval.string_value = strdup(yytext); return DMUL;
+EX2	yylval.string_value = strdup(yytext); return EX2;
+F2F	yylval.string_value = strdup(yytext); return F2F;
+F2I	yylval.string_value = strdup(yytext); return F2I;
+FADD	yylval.string_value = strdup(yytext); return FADD;
+FADD32	yylval.string_value = strdup(yytext); return FADD32;
+FADD32I	yylval.string_value = strdup(yytext); return FADD32I;
+FMAD	yylval.string_value = strdup(yytext); return FMAD;
+FMAD32I	yylval.string_value = strdup(yytext); return FMAD32I;
+FMUL	yylval.string_value = strdup(yytext); return FMUL;
+FMUL32	yylval.string_value = strdup(yytext); return FMUL32;
+FMUL32I	yylval.string_value = strdup(yytext); return FMUL32I;
+FSET	yylval.string_value = strdup(yytext); return FSET;
+G2R	yylval.string_value = strdup(yytext); return G2R;
+GLD	yylval.string_value = strdup(yytext); return GLD;
+GST	yylval.string_value = strdup(yytext); return GST;
+I2F	yylval.string_value = strdup(yytext); return I2F;
+I2I	yylval.string_value = strdup(yytext); return I2I;
+S2R	yylval.string_value = strdup(yytext); return S2R;
+STS	yylval.string_value = strdup(yytext); return STS;
+LD	yylval.string_value = strdup(yytext); return LD;
+ST	yylval.string_value = strdup(yytext); return ST;
+IADD	yylval.string_value = strdup(yytext); return IADD;
+IADD32	yylval.string_value = strdup(yytext); return IADD32;
+IADD32I	yylval.string_value = strdup(yytext); return IADD32I;
+IMAD	yylval.string_value = strdup(yytext); return IMAD;
+IMAD\.U24	yylval.string_value = strdup(yytext); return IMAD24;
+IMAD\.S24	yylval.string_value = strdup(yytext); return IMAD24;
+IMAD32	yylval.string_value = strdup(yytext); return IMAD32;
+IMAD32I	yylval.string_value = strdup(yytext); return IMAD32I;
+IMUL	yylval.string_value = strdup(yytext); return IMUL;
+IMUL\.U24\.U24	yylval.string_value = strdup(yytext); return IMUL24;
+IMUL\.S24\.S24	yylval.string_value = strdup(yytext); return IMULS24;
+IMUL32	yylval.string_value = strdup(yytext); return IMUL32;
+IMUL32\.S24\.S24	yylval.string_value = strdup(yytext); return IMUL32S24;
+IMUL32\.U24\.U24	yylval.string_value = strdup(yytext); return IMUL32U24;
+IMUL32I\.S24\.S24	yylval.string_value = strdup(yytext); return IMUL32IS24;
+IMUL32I	yylval.string_value = strdup(yytext); return IMUL32I;
+IMUL32I\.U24\.U24	yylval.string_value = strdup(yytext); return IMUL32I24;
+ISET	yylval.string_value = strdup(yytext); return ISET;
+LG2	yylval.string_value = strdup(yytext); return LG2;
+LLD	yylval.string_value = strdup(yytext); return LLD;
+LOP\.AND	yylval.string_value = strdup(yytext); return AND;
+LOP\.OR	yylval.string_value = strdup(yytext); return OR;
+LOP\.PASS_B	yylval.string_value = strdup(yytext); return NOT;
+LOP\.XOR	yylval.string_value = strdup(yytext); return XOR;
+LOP\.S\.AND	yylval.string_value = strdup(yytext); return ANDS;
+LOP\.S\.OR	yylval.string_value = strdup(yytext); return ORS;
+LOP\.S\.PASS_B	yylval.string_value = strdup(yytext); return NOTS;
+LOP\.S\.XOR	yylval.string_value = strdup(yytext); return XORS;
+LST	yylval.string_value = strdup(yytext); return LST;
+MOV	yylval.string_value = strdup(yytext); return MOV;
+MOV32	yylval.string_value = strdup(yytext); return MOV32;
+MVC	yylval.string_value = strdup(yytext); return MVC;
+MVI	yylval.string_value = strdup(yytext); return MVI;
+NOP	yylval.string_value = strdup(yytext); return NOP;
+R2A	yylval.string_value = strdup(yytext); return R2A;
+R2G	yylval.string_value = strdup(yytext); return R2G;
+R2G.U16.U8	yylval.string_value = strdup(yytext); return R2GU16U8;
+RCP	yylval.string_value = strdup(yytext); return RCP;
+RCP32	yylval.string_value = strdup(yytext); return RCP32;
+RET	yylval.string_value = strdup(yytext); return RET;
+RRO	yylval.string_value = strdup(yytext); return RRO;
+RSQ	yylval.string_value = strdup(yytext); return RSQ;
+SIN	yylval.string_value = strdup(yytext); return SIN;
+SHL	yylval.string_value = strdup(yytext); return SHL;
+SHR	yylval.string_value = strdup(yytext); return SHR;
+SSY	yylval.string_value = strdup(yytext); return SSY;
+BAR	yylval.string_value = strdup(yytext); return BAR;
+LDS	yylval.string_value = strdup(yytext); return LDS;
+IMIN	yylval.string_value = strdup(yytext); return IMIN;
+IMAX	yylval.string_value = strdup(yytext); return IMAX;
+FMAX	yylval.string_value = strdup(yytext); return FMAX;
+FMIN	yylval.string_value = strdup(yytext); return FMIN;
+C2R	yylval.string_value = strdup(yytext); return C2R;
+R2C	yylval.string_value = strdup(yytext); return R2C;
+A2R	yylval.string_value = strdup(yytext); return A2R;
+TEX	yylval.string_value = strdup(yytext); return TEX;
+TEX32	yylval.string_value = strdup(yytext); return TEX32;
+EXIT	yylval.string_value = strdup(yytext); return EXIT;
+GRED	yylval.string_value = strdup(yytext); return GRED;
+GATOM	yylval.string_value = strdup(yytext); return GATOM;
+PBK	yylval.string_value = strdup(yytext); return PBK;
+BRK	yylval.string_value = strdup(yytext); return BRK;
+
+	/*instruction modifiers*/
+EQ	yylval.string_value = strdup(yytext); return EQ;
+EQU	yylval.string_value = strdup(yytext); return EQU;
+GE	yylval.string_value = strdup(yytext); return GE;
+GEU	yylval.string_value = strdup(yytext); return GEU;
+GT	yylval.string_value = strdup(yytext); return GT;
+GTU	yylval.string_value = strdup(yytext); return GTU;
+LE	yylval.string_value = strdup(yytext); return LE;
+LEU	yylval.string_value = strdup(yytext); return LEU;
+LT	yylval.string_value = strdup(yytext); return LT;
+LTU	yylval.string_value = strdup(yytext); return LTU;
+NE	yylval.string_value = strdup(yytext); return NE;
+NEU	yylval.string_value = strdup(yytext); return NEU;
+\.BEXT	yylval.string_value = strdup(yytext); return DOTBEXT;
+\.S	yylval.string_value = strdup(yytext); return DOTS;
+\.SFU	yylval.string_value = strdup(yytext); return DOTSFU;
+\.TRUNC	yylval.string_value = strdup(yytext); return DOTTRUNC;
+\.IR	yylval.string_value = strdup(yytext); /*return DOTIR;*/
+\.UN	yylval.string_value = strdup(yytext); return DOTUN;
+\.NODEP	yylval.string_value = strdup(yytext); return DOTNODEP;
+\.CARRY0	yylval.string_value = strdup(yytext); return DOTCARRY0;
+\.SAT	yylval.string_value = strdup(yytext);
+
+
+	/* comparison operators */
+\.EQ	yylval.string_value = strdup(yytext); return DOTEQ;
+\.EQU	yylval.string_value = strdup(yytext); return DOTEQU;
+\.GE	yylval.string_value = strdup(yytext); return DOTGE;
+\.GEU	yylval.string_value = strdup(yytext); return DOTGEU;
+\.GT	yylval.string_value = strdup(yytext); return DOTGT;
+\.GTU	yylval.string_value = strdup(yytext); return DOTGTU;
+\.LE	yylval.string_value = strdup(yytext); return DOTLE;
+\.LEU	yylval.string_value = strdup(yytext); return DOTLEU;
+\.LT	yylval.string_value = strdup(yytext); return DOTLT;
+\.LTU	yylval.string_value = strdup(yytext); return DOTLTU;
+\.NE	yylval.string_value = strdup(yytext); return DOTNE;
+\.NEU	yylval.string_value = strdup(yytext); return DOTNEU;
+\.NOT_SIGN	yylval.string_value = strdup(yytext); return DOTNSF;
+\.SIGN	yylval.string_value = strdup(yytext); return DOTSF;
+\.CARRY	yylval.string_value = strdup(yytext); return DOTCARRY;
+\.CC	yylval.string_value = strdup(yytext); return DOTCC;
+\.X		yylval.string_value = strdup(yytext); return DOTX;
+\.E		yylval.string_value = strdup(yytext); return DOTE;
+\.RED	yylval.string_value = strdup(yytext); return DOTRED;
+\.POPC	yylval.string_value = strdup(yytext); return DOTPOPC;
+
+	/*operand types*/
+\.F16	yylval.string_value = strdup(yytext); return DOTF16;
+\.F32	yylval.string_value = strdup(yytext); return DOTF32;
+\.F64	yylval.string_value = strdup(yytext); return DOTF64;
+\.S8	yylval.string_value = strdup(yytext); return DOTS8;
+\.S16	yylval.string_value = strdup(yytext); return DOTS16;
+\.S32	yylval.string_value = strdup(yytext); return DOTS32;
+\.S64	yylval.string_value = strdup(yytext); return DOTS64;
+\.S128	yylval.string_value = strdup(yytext); return DOTS128;
+\.U8	yylval.string_value = strdup(yytext); return DOTU8;
+\.U16	yylval.string_value = strdup(yytext); return DOTU16;
+\.U32	yylval.string_value = strdup(yytext); return DOTU32;
+\.U24	yylval.string_value = strdup(yytext); return DOTU24;
+\.U64	yylval.string_value = strdup(yytext); return DOTU64;
+\.HI	yylval.string_value = strdup(yytext); return DOTHI;
+\.NOINC	yylval.string_value = strdup(yytext); return DOTNOINC;
+
+	/*geometries*/
+
+	/*register operands*/
+[\-]?R{numeric}+	yylval.string_value = strdup(yytext); return REGISTER;
+[\-]?R{numeric}+L	yylval.string_value = strdup(yytext); return REGISTERLO;
+[\-]?R{numeric}+H    yylval.string_value = strdup(yytext); return REGISTERHI;
+A{numeric}+	yylval.string_value = strdup(yytext); return OFFSETREGISTER;
+\.C{numeric}	yylval.string_value = strdup(yytext); return PREDREGISTER;
+"o ["{hexliteral}\]	yylval.string_value = strdup(yytext); return PREDREGISTER2;
+C{numeric}	 yylval.string_value = strdup(yytext); return PREDREGISTER3;
+"SR_"[^\.; \t]+	 yylval.string_value = strdup(yytext); return SREGISTER;
+
+	/*tex operand*/
+
+	/*vector operands*/
+
+	/*special registers*/
+
+	/* cuobjdump header */
+"Fatbin ptx code:" yylval.string_value = strdup(yytext); return PTXHEADER;
+"Fatbin elf code:" yylval.string_value = strdup(yytext); return ELFHEADER;
+"================"
+"arch = " yylval.string_value = strdup(yytext); return INFOARCHVERSION;
+"code version = " yylval.string_value = strdup(yytext); return INFOCODEVERSION_HEADER;
+\[{numeric},{numeric}\] yylval.string_value = strdup(yytext); return INFOCODEVERSION;
+"producer = " yylval.string_value = strdup(yytext); return INFOPRODUCER;
+"host = " yylval.string_value = strdup(yytext); return INFOHOST;
+{numeric}{numeric}"bit" yylval.string_value = strdup(yytext); return INFOCOMPILESIZE;
+"compile_size = " yylval.string_value = strdup(yytext); return INFOCOMPILESIZE_HEADER;
+"identifier = " yylval.string_value = strdup(yytext); return INFOIDENTIFIER;
+"code for "	yylval.string_value = strdup(yytext); return VERSIONHEADER;
+"Function : "	yylval.string_value = strdup(yytext); return FUNCTIONHEADER;
+"/*"{hexliteral}{whitespace}*"*/" 	yylval.string_value = strdup(yytext); return INSTHEX;
+"." yylval.string_value = strdup(yytext); return DOT;
+
+	/*memory location operands*/
+[\-]?"|g| ["{hexliteral}"]"".U16"?	yylval.string_value = strdup(yytext); return ABSSMEMLOCATION;
+[\-]?"|g| [R"{numeric}+\]".U16"?	yylval.string_value = strdup(yytext); return ABSSMEMLOCATION;
+[\-]?"|g| [A"{numeric}+\+{hexliteral}\]".U16"?	yylval.string_value = strdup(yytext); return ABSSMEMLOCATION;
+[\-]?"|g| [A"{numeric}+\+\+\+{hexliteral}\]".U16"?	yylval.string_value = strdup(yytext); return ABSSMEMLOCATION;
+
+[\-]?"g ["{hexliteral}"]"".U16"?	yylval.string_value = strdup(yytext); return SMEMLOCATION;
+[\-]?"g [R"{numeric}+\]".U16"?	yylval.string_value = strdup(yytext); return SMEMLOCATION;
+[\-]?"g [A"{numeric}+\+{hexliteral}\]".U16"?	yylval.string_value = strdup(yytext); return SMEMLOCATION;
+[\-]?"g [A"{numeric}+\+\+\+{hexliteral}\]".U16"?	yylval.string_value = strdup(yytext); return SMEMLOCATION;
+
+[\-]?"global14 ["{hexliteral}"]"".U16"?	yylval.string_value = strdup(yytext); return GMEMLOCATION;
+[\-]?"global14 [R"{numeric}+\]".U16"?	yylval.string_value = strdup(yytext); return GMEMLOCATION;
+[\-]?"global14 [A"{numeric}+\+{hexliteral}\]".U16"?	yylval.string_value = strdup(yytext); return GMEMLOCATION;
+[\-]?"global14 [A"{numeric}+\+\+\+{hexliteral}\]".U16"?	yylval.string_value = strdup(yytext); return GMEMLOCATION;
+
+[\-]?"c ["{hexliteral}"] ["{hexliteral}"]"".U16"?	yylval.string_value = strdup(yytext); return CMEMLOCATION;
+[\-]?"c ["{hexliteral}"] [R"{numeric}+\]".U16"?	yylval.string_value = strdup(yytext); return CMEMLOCATION;
+[\-]?"c ["{hexliteral}"] [A"{numeric}+\+{hexliteral}\]".U16"?	yylval.string_value = strdup(yytext); return CMEMLOCATION;
+[\-]?"c ["{hexliteral}"] [A"{numeric}+\+\+\+{hexliteral}\]".U16"?	yylval.string_value = strdup(yytext); return CMEMLOCATION;
+
+[\-]?"local ["{hexliteral}"]"".U16"?	yylval.string_value = strdup(yytext); return LMEMLOCATION;
+[\-]?"local [R"{numeric}+\]".U16"?	yylval.string_value = strdup(yytext); return LMEMLOCATION;
+[\-]?"local [A"{numeric}+\+{hexliteral}\]".U16"?	yylval.string_value = strdup(yytext); return LMEMLOCATION;
+[\-]?"local [A"{numeric}+\+\+\+{hexliteral}\]".U16"?	yylval.string_value = strdup(yytext); return LMEMLOCATION;
+
+	/*identifier + literals*/
+	
+{identifier}	yylval.string_value = strdup(yytext); return IDENTIFIER;
+[\-]?{hexliteral}	yylval.string_value = strdup(yytext); return HEXLITERAL; //change to int later?
+
+	/*label ends*/
+"/*"	yylval.string_value = strdup(yytext); return LABELSTART;
+"*/"	yylval.string_value = strdup(yytext); return LABELEND;
+{hexdigit}{hexdigit}{hexdigit}{hexdigit} yylval.string_value = strdup(yytext); return LABEL;
+
+	/*braces*/
+
+	/*brackets*/
+\(	yylval.string_value = strdup(yytext); return LEFTBRACKET;
+\)	yylval.string_value = strdup(yytext); return RIGHTBRACKET;
+
+	/* pound sign*/
+
+	/* comma*/
+","	/*yylval.string_value = strdup(yytext); return COMMA;*/
+
+	/* underscore sign*/
+
+	/* pipe*/
+"|"	yylval.string_value = strdup(yytext); return PIPE;
+"~"	yylval.string_value = strdup(yytext); return TILDE;
+
+	/*Need to create to newline tokens*/
+{newline}+	return NEWLINE;
+
+	/* semicolon */
+;	return SEMICOLON;
+
+	/*discard whitespace*/
+{whitespace}+
+	/*discard comments*/
+"//"{anycharbutnewline}*
+	/*discard dividers*/
+"-----"[\-]+
+"....."[\.]+
+
+
+	/*reached end of file*/
+<<EOF>>	return 0;
+
+	/*None of the rules matched. Throw error*/
+.			sass_error("Invalid token");
+
+%%
+
+void sass_error(const char* message)
+{
+	printf("   "); printf(message); printf(" near \""); printf(yytext); printf("\"");
+	printf(" on line "); 
+	char line[5]; sprintf(line, "%i", yylineno); printf(line);
+	printf("\n");
+	exit(1);
+}
diff --git a/cuobjdump_to_ptxplus/sass.y b/cuobjdump_to_ptxplus/sass.y
new file mode 100644
index 0000000..06d2750
--- /dev/null
+++ b/cuobjdump_to_ptxplus/sass.y
@@ -0,0 +1,405 @@
+%{
+#include <stdio.h>
+#include "cuobjdumpInstList.h"
+
+int yylex(void);
+void yyerror(const char*);
+
+extern cuobjdumpInstList *g_instList;
+
+cuobjdumpInst *instEntry;
+%}
+
+
+%union {
+  double double_value;
+  float  float_value;
+  int    int_value;
+  char * string_value;
+  void * ptr_value;
+}
+
+%token <string_value> BAR
+%token <string_value> ADA AND ANDS BRA CAL COS DADD DFMA DMUL EX2 F2F F2I FADD
+%token <string_value> FADD32 FADD32I FMAD FMAD32I FMUL FMUL32 FMUL32I FSET G2R
+%token <string_value> GLD GST I2F I2I IADD IADD32 IADD32I IMAD IMAD24 IMAD32I IMAD32
+%token <string_value> IMUL IMUL24 IMULS24 IMUL32 IMUL32S24 IMUL32U24 IMUL32I IMUL32I24 IMUL32IS24
+%token <string_value> ISET LG2 LLD LST MOV MOV32 MVC MVI NOP NOT NOTS OR ORS
+%token <string_value> R2A R2G R2GU16U8 RCP RCP32 RET RRO RSQ SIN SHL SHR SSY XOR XORS
+%token <string_value> S2R LD STS LDS ST IMIN IMAX A2R FMAX FMIN TEX TEX32 C2R EXIT
+%token <string_value> GRED PBK BRK R2C GATOM
+
+%token <string_value> EQ EQU GE GEU GT GTU LE LEU LT LTU NE NEU
+%token <string_value> DOTBEXT DOTS DOTSFU
+%token <string_value> DOTTRUNC DOTIR DOTUN DOTNODEP DOTCARRY0 DOTSAT
+%token <string_value> DOTF16 DOTF32 DOTF64 DOTS8 DOTS16 DOTS32 DOTS64 DOTS128 DOTU8 DOTU16 DOTU32 DOTU24 DOTU64
+%token <string_value> DOTHI DOTNOINC
+%token <string_value> DOTEQ DOTEQU DOTGE DOTGEU DOTGT DOTGTU DOTLE DOTLEU DOTLT DOTLTU DOTNE DOTNEU DOTNSF DOTSF DOTCARRY
+%token <string_value> DOTCC DOTX DOTE DOTRED DOTPOPC
+%token <string_value> REGISTER REGISTERLO REGISTERHI OFFSETREGISTER
+%token <string_value> PREDREGISTER PREDREGISTER2 PREDREGISTER3 SREGISTER
+%token <string_value> VERSIONHEADER FUNCTIONHEADER
+%token <string_value> SMEMLOCATION ABSSMEMLOCATION GMEMLOCATION CMEMLOCATION LMEMLOCATION
+%token <string_value> IDENTIFIER
+%token <string_value> HEXLITERAL
+%token <string_value> LEFTBRACKET RIGHTBRACKET
+%token <string_value> PIPE TILDE
+%token <string_value> NEWLINE SEMICOLON /*COMMA*/
+%token <string_value> LABEL LABELSTART LABELEND
+%token <string_value> PTXHEADER ELFHEADER
+%token <string_value> INFOARCHVERSION
+%token <string_value> INFOCODEVERSION_HEADER INFOCODEVERSION
+%token <string_value> INFOPRODUCER
+%token <string_value> INFOHOST
+%token <string_value> INFOCOMPILESIZE_HEADER INFOCOMPILESIZE
+%token <string_value> INFOIDENTIFIER DOT
+%token <string_value> INSTHEX
+%token <string_value> OSQBRACKET CSQBRACKET
+	/* set types for rules */
+%type<string_value> simpleInstructions
+%type<string_value> predicateModifier
+%type<string_value> opTypes
+
+%%
+
+	/*translation rules*/
+program		: NEWLINE sassCode
+			| program sassCode
+			| ;
+
+sassCode	: VERSIONHEADER IDENTIFIER functionList			{ printf($1); printf($2); printf(" No parsing errors\n\n");  }
+		| newlines VERSIONHEADER IDENTIFIER functionList	{ printf($2); printf($3); printf(" No parsing errors\n\n");  }
+		;
+
+functionList	: functionList FUNCTIONHEADER IDENTIFIER {
+						printf($2); 
+						printf($3);
+						printf("\n");
+						g_instList->addEntry($3);
+						instEntry = new cuobjdumpInst();
+						instEntry->setBase(".entry");
+						g_instList->add(instEntry);
+						g_instList->getListEnd().addOperand($3);} statementList { /*g_instList->getListEnd().addBaseModifier(".exit");*/ }
+/*		| FUNCTIONHEADER IDENTIFIER { printf($1); printf($2);
+						g_instList->addEntry($2);
+						instEntry = new cuobjdumpInst();
+						instEntry->setBase(".entry");
+						g_instList->add(instEntry);
+						g_instList->getListEnd().addOperand($2);} statementList { g_instList->getListEnd().addBaseModifier(".exit"); }*/
+				| functionList newlines
+				| 
+				;
+
+statementList	: statementList statement newlines	{ printf("\n"); }
+		| statement newlines			{ printf("\n"); }
+		| statementList statement SEMICOLON newlines	{ printf(";\n"); }
+		| statement SEMICOLON newlines			{ printf(";\n"); }
+		| newlines				{}
+		;
+
+newlines	: NEWLINE
+			| newlines NEWLINE
+			;
+
+statement	: { instEntry = new cuobjdumpInst(); } instructionLabel instructionHex assemblyInstruction
+			;
+
+instructionHex	: INSTHEX
+		;
+
+instructionLabel	: LABELSTART LABEL LABELEND	{ char* tempInput = $2;
+							  char* tempLabel = new char[12];
+							  tempLabel[0] = 'l';
+							  tempLabel[1] = '0';
+							  tempLabel[2] = 'x';
+							  for(int i=0; i<(8-strlen(tempInput)); i++)
+							  {
+								tempLabel[3+i] = '0';
+							  }
+							  for(int i=(11-strlen(tempInput)); i<11; i++)
+							  {
+								tempLabel[i] = tempInput[i-(11-strlen(tempInput))];
+							  }
+							  tempLabel[11] = '\0';
+							  instEntry->setLabel(tempLabel); }
+			;
+
+assemblyInstruction	: baseInstruction modifierList operandList	{ }
+					/*| baseInstruction operandList			{ }*/
+					/*| baseInstruction modifierList			{ }*/
+					/*| baseInstruction				{ }*/
+					;
+
+baseInstruction : simpleInstructions	{ printf($1); instEntry->setBase($1); g_instList->add(instEntry);}
+		| branchInstructions
+		| GRED DOT simpleInstructions	{ printf($1); instEntry->setBase($1); g_instList->add(instEntry); g_instList->getListEnd().addBaseModifier($3);}
+		| GATOM DOT simpleInstructions	{ printf($1); instEntry->setBase($1); g_instList->add(instEntry); g_instList->getListEnd().addBaseModifier($3);}
+		| pbkInstruction
+		;
+
+simpleInstructions	: ADA | AND | ANDS | COS | DADD | DFMA | DMUL | EX2 | F2F 
+					| F2I | FADD | FADD32 | FADD32I | FMAD | FMAD32I | FMUL 
+					| FMUL32 | FMUL32I | FSET | G2R | GLD | GST | I2F | I2I 
+					| IADD | IADD32 | IADD32I | IMAD | IMAD24 | IMAD32I | IMAD32 | IMUL 
+					| IMUL24 | IMULS24 | IMUL32 | IMUL32S24 | IMUL32I | IMUL32I24 | IMUL32IS24
+					| IMUL32U24
+					| ISET | LG2 | LLD | LST | MOV | MOV32 | MVC | MVI | NOP 
+					| NOT | NOTS | OR | ORS | R2A | R2G | R2GU16U8 | RCP | RCP32 | RET | RRO 
+					| RSQ | SHL | SHR | SIN | SSY | XOR | XORS | S2R | LD | STS 
+					| LDS | ST | EXIT | BAR | IMIN | IMAX | A2R | FMAX | FMIN 
+					| TEX | TEX32 | C2R | BRK | R2C
+					;
+
+pbkInstruction	:	PBK {
+						printf($1); instEntry->setBase($1); g_instList->add(instEntry);
+					} HEXLITERAL {
+						char* tempInput = $3;
+						char* tempLabel = new char[12];
+						tempLabel[0] = 'l';
+						tempLabel[1] = '0';
+						tempLabel[2] = 'x';
+						for(int i=0; i<(10-strlen(tempInput)); i++)
+						{
+							tempLabel[3+i] = '0';
+						}
+						for(int i=(13-strlen(tempInput)); i<11; i++)
+						{
+							tempLabel[i] = tempInput[i-(11-strlen(tempInput))];
+						}
+						tempLabel[11] = '\0';
+						g_instList->getListEnd().addOperand(tempLabel);
+						g_instList->addCubojdumpLabel(tempLabel);
+					}
+				;
+
+branchInstructions	: BRA {printf($1); instEntry->setBase($1); g_instList->add(instEntry);} instructionPredicate HEXLITERAL
+				{ printf($4);
+				  char* tempInput = $4;
+				  char* tempLabel = new char[12];
+				  tempLabel[0] = 'l';
+				  tempLabel[1] = '0';
+				  tempLabel[2] = 'x';
+				  for(int i=0; i<(10-strlen(tempInput)); i++)
+				  {
+					tempLabel[3+i] = '0';
+				  }
+				  for(int i=(13-strlen(tempInput)); i<11; i++)
+				  {
+					tempLabel[i] = tempInput[i-(11-strlen(tempInput))];
+				  }
+				  tempLabel[11] = '\0';
+				  g_instList->getListEnd().addOperand(tempLabel);
+				  g_instList->addCubojdumpLabel(tempLabel);}
+			| BRA {printf($1); instEntry->setBase($1); g_instList->add(instEntry);} HEXLITERAL
+				{ printf($3);
+				  char* tempInput = $3;
+				  char* tempLabel = new char[12];
+				  tempLabel[0] = 'l';
+				  tempLabel[1] = '0';
+				  tempLabel[2] = 'x';
+				  for(int i=0; i<(10-strlen(tempInput)); i++)
+				  {
+					tempLabel[3+i] = '0';
+				  }
+				  for(int i=(13-strlen(tempInput)); i<11; i++)
+				  {
+					tempLabel[i] = tempInput[i-(11-strlen(tempInput))];
+				  }
+				  tempLabel[11] = '\0';
+				  g_instList->getListEnd().addOperand(tempLabel);
+				  g_instList->addCubojdumpLabel(tempLabel);}
+			| CAL {printf($1); instEntry->setBase($1); g_instList->add(instEntry);} HEXLITERAL
+				{ printf($3);
+				  char* tempInput = $3;
+				  char* tempLabel = new char[12];
+				  tempLabel[0] = 'l';
+				  tempLabel[1] = '0';
+				  tempLabel[2] = 'x';
+				  for(int i=0; i<(10-strlen(tempInput)); i++)
+				  {
+					tempLabel[3+i] = '0';
+				  }
+				  for(int i=(13-strlen(tempInput)); i<11; i++)
+				  {
+					tempLabel[i] = tempInput[i-(11-strlen(tempInput))];
+				  }
+				  tempLabel[11] = '\0';
+				  g_instList->getListEnd().addOperand(tempLabel);
+				  g_instList->addCubojdumpLabel(tempLabel);}
+			
+			| CAL {printf($1); instEntry->setBase($1); g_instList->add(instEntry);} DOTNOINC HEXLITERAL
+				{ printf($4);
+				  char* tempInput = $4;
+				  char* tempLabel = new char[12];
+				  tempLabel[0] = 'l';
+				  tempLabel[1] = '0';
+				  tempLabel[2] = 'x';
+				  for(int i=0; i<(10-strlen(tempInput)); i++)
+				  {
+					tempLabel[3+i] = '0';
+				  }
+				  for(int i=(13-strlen(tempInput)); i<11; i++)
+				  {
+					tempLabel[i] = tempInput[i-(11-strlen(tempInput))];
+				  }
+				  tempLabel[11] = '\0';
+				  g_instList->getListEnd().addOperand(tempLabel);
+				  g_instList->addCubojdumpLabel(tempLabel);}
+
+			;
+
+modifierList	: modifier modifierList
+				/*| modifier */
+				|
+				;
+
+modifier	: opTypes	{ printf($1); g_instList->getListEnd().addTypeModifier($1);}
+		| DOTBEXT		{ g_instList->getListEnd().addBaseModifier(".bext"); }
+		| DOTS			{ g_instList->getListEnd().addBaseModifier(".s"); }
+		| DOTSFU		{ g_instList->getListEnd().addBaseModifier(".sfu"); }
+		| DOTTRUNC		{ g_instList->getListEnd().addBaseModifier(".rz"); }
+		| DOTX			{ g_instList->getListEnd().addBaseModifier(".x"); }
+		| DOTE			{ g_instList->getListEnd().addBaseModifier(".e"); }
+		| DOTRED		{ g_instList->getListEnd().addBaseModifier(".red"); }
+		| DOTPOPC		{ g_instList->getListEnd().addBaseModifier(".popc"); }
+		| DOTIR			{ g_instList->getListEnd().addBaseModifier(".ir"); }
+		| DOTUN			{ /*g_instList->getListEnd().addBaseModifier(".un"); */}
+		| DOTNODEP			{ /*g_instList->getListEnd().addBaseModifier(".nodep"); */}
+		| DOTCARRY0			{ /*g_instList->getListEnd().addBaseModifier(".carry0"); */}
+		;
+
+opTypes		: DOTF16	//{ printf($1); g_instList->getListEnd().addTypeModifier($1);}
+		| DOTF32	//{ printf($1); g_instList->getListEnd().addTypeModifier($1);}
+		| DOTF64	//{ printf($1); g_instList->getListEnd().addTypeModifier($1);}
+		| DOTS8		//{ printf($1); g_instList->getListEnd().addTypeModifier($1);}
+		| DOTS16	//{ printf($1); g_instList->getListEnd().addTypeModifier($1);}
+		| DOTS32	//{ printf($1); g_instList->getListEnd().addTypeModifier($1);}
+		| DOTS64	//{ printf($1); g_instList->getListEnd().addTypeModifier($1);}
+		| DOTS128	//{ printf($1); g_instList->getListEnd().addTypeModifier($1);}
+		| DOTU8		//{ printf($1); g_instList->getListEnd().addTypeModifier($1);}
+		| DOTU16	//{ printf($1); g_instList->getListEnd().addTypeModifier($1);}
+		| DOTU32	//{ printf($1); g_instList->getListEnd().addTypeModifier($1);}
+		| DOTU24	//{ printf($1); g_instList->getListEnd().addTypeModifier($1);}
+		| DOTU64	//{ printf($1); g_instList->getListEnd().addTypeModifier($1);}
+		| DOTHI		//{ printf($1); g_instList->getListEnd().addTypeModifier($1);}
+		;
+
+operandList	: operandList { printf(" "); } /*COMMA*/ operand	{}
+			/*| { printf(" "); } operand		{}*/
+			|
+			;
+
+operand		: registerlocation
+		| PIPE registerlocation PIPE	{ g_instList->getListEnd().addBaseModifier(".abs"); }
+		| TILDE registerlocation
+		| LEFTBRACKET instructionPredicate RIGHTBRACKET
+		| memorylocation opTypes { printf($2); g_instList->getListEnd().addTypeModifier($2);}
+		| memorylocation
+		| immediateValue
+		| extraModifier
+		| operandPredicate
+		| preOperand
+		;
+/* Register of the format [R0] will be converted to R0 */
+/* regMod will be also ignored */
+registerlocation	: REGISTER regMod	{ printf($1); g_instList->addCuobjdumpRegister($1);}
+			| OSQBRACKET REGISTER CSQBRACKET	{ printf($1); printf($2); printf($3); g_instList->addCuobjdumpRegister($2);}
+			| REGISTERLO	{ printf($1); g_instList->addCuobjdumpRegister($1,true);}
+			| REGISTERHI	{ printf($1); g_instList->addCuobjdumpRegister($1,true);}
+			| SREGISTER		{ printf($1); g_instList->addCuobjdumpRegister($1,false);}
+			| OFFSETREGISTER	{ printf($1); g_instList->addCuobjdumpRegister($1);}
+			| PREDREGISTER PREDREGISTER2	{ printf($1); printf(" "); printf($2); g_instList->addCuobjdumpDoublePredReg($1, $2);}
+			| PREDREGISTER REGISTER	{ printf($1); printf(" "); printf($2); g_instList->addCuobjdumpDoublePredReg($1, $2);}
+			/*| REGISTER PREDREGISTER3 { printf($1); printf(" "); printf($2); g_instList->addCuobjdumpRegister($1); printf("WEIRD CASE\n");}*/
+			;
+
+regMod		: DOTCC
+			|
+			;
+
+
+memorylocation	: SMEMLOCATION	{ printf($1); g_instList->addCuobjdumpMemoryOperand($1,1);}
+		|	ABSSMEMLOCATION {
+				printf($1);
+				char* input = $1;
+				char* temp = new char[99];
+				temp[0] = input[1];
+				int i=1;
+				while (i < strlen(input)-2) {
+					temp[i] = input[i+2];
+					i++;
+				}
+				g_instList->addCuobjdumpMemoryOperand(temp,1);
+				g_instList->getListEnd().addBaseModifier(".abs");
+			}
+		| GMEMLOCATION	{ printf($1); g_instList->addCuobjdumpMemoryOperand($1,2);}
+		| CMEMLOCATION	{ printf($1); g_instList->addCuobjdumpMemoryOperand($1,0);}
+		| LMEMLOCATION	{ printf($1); g_instList->addCuobjdumpMemoryOperand($1,3);}
+		;
+
+immediateValue	: IDENTIFIER { printf($1); g_instList->getListEnd().addOperand($1);}
+		| HEXLITERAL { printf($1); g_instList->getListEnd().addOperand($1);}
+		;
+
+extraModifier	: EQ	{ printf($1); g_instList->getListEnd().addBaseModifier($1);} 
+		| EQU	{ printf($1); g_instList->getListEnd().addBaseModifier($1);}
+		| GE	{ printf($1); g_instList->getListEnd().addBaseModifier($1);}
+		| GEU	{ printf($1); g_instList->getListEnd().addBaseModifier($1);}
+		| GT	{ printf($1); g_instList->getListEnd().addBaseModifier($1);}
+		| GTU	{ printf($1); g_instList->getListEnd().addBaseModifier($1);}
+		| LE	{ printf($1); g_instList->getListEnd().addBaseModifier($1);}
+		| LEU	{ printf($1); g_instList->getListEnd().addBaseModifier($1);}
+		| LT	{ printf($1); g_instList->getListEnd().addBaseModifier($1);}
+		| LTU	{ printf($1); g_instList->getListEnd().addBaseModifier($1);}
+		| NE	{ printf($1); g_instList->getListEnd().addBaseModifier($1);}
+		| NEU	{ printf($1); g_instList->getListEnd().addBaseModifier($1);}
+		;
+
+instructionPredicate	: PREDREGISTER3	predicateModifier {printf($1); printf($2);
+								g_instList->getListEnd().setPredicate($1);
+								g_instList->getListEnd().addPredicateModifier($2);}
+						| PREDREGISTER3 {printf($1); g_instList->getListEnd().setPredicate($1);}
+			;
+
+operandPredicate	:	PREDREGISTER3	predicateModifier {
+							printf($1); 
+							printf($2);
+							//g_instList->getListEnd().addOperand($1);
+							g_instList->getListEnd().setPredicate($1);
+							g_instList->getListEnd().addPredicateModifier($2);
+							/*May be the modifier needs to be added too*/
+						}
+					|	PREDREGISTER3 {
+							printf("HELLO: "); 
+							printf($1); 
+							g_instList->getListEnd().addOperand($1);
+						}
+					;
+
+
+preOperand	: EX2	{ printf($1); g_instList->getListEnd().addBaseModifier("ex2");}
+		| SIN	{ printf($1); g_instList->getListEnd().addBaseModifier("sin");}
+		| COS	{ printf($1); g_instList->getListEnd().addBaseModifier("cos");}
+		;
+
+predicateModifier	: DOTEQ	{ }
+			| DOTEQU	{ }
+			| DOTGE	{ }
+			| DOTGEU	{ }
+			| DOTGT	{ }
+			| DOTGTU	{ }
+			| DOTLE	{ }
+			| DOTLEU	{ }
+			| DOTLT	{ }
+			| DOTLTU	{ }
+			| DOTNE	{ }
+			| DOTNEU	{ }
+			| DOTNSF	{ }
+			| DOTSF	{ }
+			| DOTCARRY	{ }
+			;
+
+%%
+
+/*support c++ functions go here*/
+